第2章 CUDA编程基础

第2章：CUDA编程基础

好，咱们开始聊CUDA编程。说实话，很多做深度学习推理的同学，一开始接触CUDA都会有点懵。我当年也是，看着那一堆Grid、Block、Thread，心里直犯嘀咕——这玩意儿到底怎么映射到GPU上的？别急，咱们一步步来。

2.1 CUDA架构简介

CUDA，全称是Compute Unified Device Architecture。说白了，就是NVIDIA给GPU开的一个"后门"，让我们能用C语言去指挥GPU干活。

GPU和CPU最大的区别是什么？CPU是"全能选手"，什么活都能干，但一次只能干几件事。GPU是"专业选手"，擅长干大量重复的简单计算。你想想看，深度学习里的矩阵乘法，不就是成千上万个乘加操作吗？这活交给GPU再合适不过了。

我记得第一次用CUDA做矩阵乘法加速时，看到几百倍的加速比，确实被震撼到了。但后来也踩过不少坑——比如没注意内存带宽限制，结果性能还不如CPU。嗯，这些坑后面我会一一提到。

核心概念：CUDA将GPU视为一个可以并行执行大量线程的协处理器。CPU称为Host，GPU称为Device。

2.2 线程层次结构

这是CUDA最核心的概念。我建议你把它想象成一个军队的组织结构：

Grid（网格）：相当于整个军队
Block（线程块）：相当于一个连队
Thread（线程）：相当于一个士兵

每个Grid包含多个Block，每个Block包含多个Thread。为什么搞这么复杂？因为GPU的硬件就是这样组织的——SM（流多处理器）执行Block，而SM里的CUDA Core执行Thread。

我的经验：Block的大小很关键。我一般习惯用256或512个线程一个Block。太小了SM利用率低，太大了寄存器不够用。具体选多少，得看你的kernel用了多少寄存器。

来看一个简单的例子，感受一下线程是怎么组织的：

// 定义Grid和Block的维度
dim3 gridDim(2, 2);      // 2x2的Grid，共4个Block
dim3 blockDim(4, 4);     // 4x4的Block，共16个Thread

// 启动kernel
myKernel<<<gridDim, blockDim>>>(...);

这段代码启动了4×16=64个线程。每个线程可以通过内置变量知道自己的位置：

threadIdx.x：线程在Block内的x索引（0-3）
blockIdx.x：Block在Grid内的x索引（0-1）
blockDim.x：Block的x维度大小（4）
gridDim.x：Grid的x维度大小（2）

要计算全局唯一的线程ID，公式是：

int tid = blockIdx.x * blockDim.x + threadIdx.x;

这个公式我建议你背下来。写CUDA kernel时，几乎每个程序都要用到它。

2.3 内存层次结构

搞懂了线程怎么组织，接下来就是数据怎么存放了。CUDA的内存层次，说白了就是"越靠近计算单元，速度越快，容量越小"。

内存类型	位置	访问速度	作用域	生命周期
Global Memory	显存	慢（~200周期）	所有线程	程序运行期间
Shared Memory	SM内部	快（~5周期）	同一个Block	Block生命周期
Local Memory	显存（但每个线程私有）	慢	单个线程	线程生命周期
Registers	SM内部	最快（1周期）	单个线程	线程生命周期

这里我要特别强调一下Shared Memory。我在做深度学习推理优化时，80%的优化技巧都和Shared Memory有关。为什么？因为它快啊！而且同一个Block内的线程可以共享数据。

曾经踩过的坑：有一次我写一个矩阵乘法的kernel，把所有数据都放在Global Memory里，结果性能惨不忍睹。后来把分块数据搬到Shared Memory里，速度直接提升了10倍。记住：尽量减少Global Memory的访问次数。

2.4 简单的Kernel编写与启动

好了，理论说完了，咱们动手写一个最简单的kernel。就做向量加法吧——把两个数组的对应元素相加。

// kernel定义
__global__ void vecAdd(float* A, float* B, float* C, int N) {
    // 计算全局线程ID
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 边界检查
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

// 主函数中的调用
int main() {
    // ... 分配内存、拷贝数据等操作 ...
    
    // 启动kernel
    int N = 1000000;
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    
    // ... 拷贝结果回Host ...
}

注意几个关键点：

__global__修饰符表示这是一个在Device上执行的kernel函数
<<<...>>>是CUDA特有的语法，用来指定Grid和Block的维度
一定要做边界检查！因为线程总数可能超过数组大小

我个人习惯在写kernel时，先用一个简单的测试用例验证正确性。比如N=1000，手动算几个值对比一下。别嫌麻烦，这能省下后面debug的大量时间。

小技巧：计算blocksPerGrid时，用 (N + threadsPerBlock - 1) / threadsPerBlock 这个公式。它保证即使N不能被threadsPerBlock整除，也能覆盖所有元素。

嗯，到这里，CUDA编程的基础框架就搭起来了。你可能会问：就这么简单？是的，入门确实不难。但真正用好它，让程序跑出接近理论峰值的性能，那又是另一回事了。后面的章节，我会带你一步步深入。

下一章，咱们聊聊内存访问模式——为什么同样的计算，换个数据排列方式，速度能差好几倍。

📚 课程目录