CUDA编程模型入门：GPU架构基础

说实话，我第一次接触CUDA的时候，也被那一堆术语搞懵了。SM是什么？Warp又是什么鬼？别急，咱们一步步来拆解。

GPU架构基础：SM、Warp、Thread

先说说GPU的核心计算单元——SM（Streaming Multiprocessor，流式多处理器）。你可以把SM想象成一个迷你CPU，但它跟CPU完全不同。CPU有十几个大核，每个核都很强；GPU有几十上百个SM，每个SM里又藏着很多小核心。

我当年在优化一个图像处理算法时，死活跑不满GPU利用率。后来才发现，是我对SM的理解太浅了。每个SM内部有自己的指令调度器、寄存器文件、共享内存。说白了，SM才是真正干活的地方。

那Warp是什么呢？Warp是GPU执行指令的最小单位。一个Warp包含32个线程，它们同时执行同一条指令。嗯，这里要注意——这32个线程必须走同一条执行路径。如果遇到if-else分支，那就麻烦了，一部分线程走if，另一部分走else，性能直接腰斩。

我曾经在一个项目里，因为没注意分支发散问题，性能从80%利用率掉到了30%。排查了两天才找到原因，教训深刻啊。

CUDA核函数（Kernel）的编写与调用

Kernel就是跑在GPU上的函数。写法上跟C函数差不多，但前面要加个__global__修饰符。我习惯把Kernel叫做「GPU入口函数」——CPU调用它，GPU执行它。

来看个最简单的例子：

// 定义一个Kernel函数
__global__ void vecAdd(float* A, float* B, float* C, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

// 调用Kernel
int main() {
    // ... 分配内存、拷贝数据 ...
    vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);
    // ... 拷贝结果回CPU ...
}

注意那个尖括号语法<<<gridSize, blockSize>>>，这是CUDA特有的调用方式。第一个参数是网格大小（多少个Block），第二个是块大小（每个Block多少个线程）。

我个人习惯把blockSize设成256或512，这是经验值。太小了SM吃不饱，太大了寄存器不够用。你想想看，如果每个线程用太多寄存器，SM能同时跑的Warp就少了，这叫「occupancy（占用率）」问题。

线程层次结构：Grid、Block、Thread

CUDA的线程组织方式分三层：Grid（网格）→ Block（线程块）→ Thread（线程）。

你可以这样理解：

• Grid：整个任务，包含所有Block
• Block：一组线程，可以协作（通过共享内存）
• Thread：单个执行单元

每个线程都有唯一的ID，通过threadIdx和blockIdx来定位。我刚开始写CUDA时，经常搞混一维、二维、三维索引。后来总结了个口诀：「blockIdx是块号，threadIdx是块内编号，blockDim是块大小」。

举个例子，如果你要处理一张1024x1024的图片：

dim3 blockSize(16, 16);   // 每个Block 16x16 = 256个线程
dim3 gridSize(64, 64);    // 64x64 = 4096个Block
// 总线程数 = 4096 * 256 = 1,048,576

__global__ void processImage(int* img, int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < width && y < height) {
        // 处理像素 (x, y)
    }
}

为什么用二维？因为图片本身就是二维的，用二维索引更直观。三维同理，适合体数据或视频帧处理。

内存层次结构：Global、Shared、Local、Register

GPU的内存体系，说白了就是「速度越快的，容量越小」。我画了张图帮你理解：

这张图我画了好几次才满意。从上到下，速度越来越慢，容量越来越大。

寄存器（Register）：每个线程私有的，最快。但数量有限，每个SM的寄存器总数是固定的（比如65536个）。如果每个线程用太多寄存器，SM能同时跑的线程就少了。

局部内存（Local Memory）：当寄存器不够用时，编译器会把变量「溢出」到局部内存。其实局部内存物理上在全局内存里，所以速度慢。我建议你尽量少用局部变量，尤其是大数组。

共享内存（Shared Memory）：这是CUDA编程的精华所在。同一个Block内的线程可以共享数据，延迟只有几个周期。我在做矩阵乘法时，就是靠共享内存把性能提升了10倍。

全局内存（Global Memory）：所有线程都能访问，但延迟高达400-800周期。访问全局内存时，要尽量做到「合并访问」——相邻线程访问相邻地址，这样能一次加载多个数据。

举个例子，如果你要计算两个大数组的和：

// 不好的写法：非合并访问
__global__ void badAdd(int* A, int* B, int* C) {
    int i = threadIdx.x;  // 所有线程都访问同一个地址？
    C[i] = A[i] + B[i];   // 不对，这样写有问题
}

// 好的写法：合并访问
__global__ void goodAdd(int* A, int* B, int* C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        C[i] = A[i] + B[i];  // 相邻线程访问相邻地址
    }
}

你看，好的写法里，线程0访问地址0，线程1访问地址1，这就是合并访问。硬件一次能加载128字节，正好是32个int。效率自然高。

好了，这一章的内容就到这里。CUDA编程模型的核心就是理解这些层次结构——线程怎么组织，数据怎么流动。把这些基础打牢，后面学起来就顺了。

---