3、CUDA编程模型：Kernel启动、Grid-Block-Thread映射、线程束执行

好，咱们今天聊聊CUDA编程模型里最核心的三个概念：Kernel怎么启动、线程怎么映射、以及线程束怎么执行。这三个东西搞明白了，你写CUDA代码就有底了。

我个人习惯把CUDA编程模型想象成一个「三层指挥体系」。你想想看，CPU是总司令，GPU是前线部队。总司令（CPU）下命令，前线部队（GPU）去执行。这个命令，就是Kernel。

3.1 Kernel启动：CPU给GPU下的一道命令

Kernel说白了就是一个在GPU上执行的函数。写法上跟C函数差不多，但前面要加个 __global__ 修饰符。调用的时候，用三括号语法 <<<grid, block>>> 指定启动配置。

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

// CPU端调用
int main() {
    // ... 数据准备 ...
    vecAdd<<<256, 256>>>(d_A, d_B, d_C, N);
    // ... 同步、结果回传 ...
}

这里有个坑，我曾经踩过。三括号里的参数顺序是 <<<gridDim, blockDim, sharedMem, stream>>>。很多人只写前两个，第三个共享内存大小默认是0，第四个流默认是0。但如果你要用动态共享内存，第三个参数必须显式指定。我刚开始写的时候忘了这茬，结果数据对不上，排查了半天。

注意：Kernel启动是异步的。CPU发完命令就继续往下跑了，不会等GPU算完。如果你紧接着就要用结果，记得加 cudaDeviceSynchronize() 同步一下。

3.2 Grid-Block-Thread映射：谁在哪儿干活

好，现在命令发出去了。GPU怎么分配人手？这就涉及到Grid、Block、Thread三层结构。

你可以这样理解：

  • Grid：整个任务，相当于一个工程队
  • Block：工程队里的班组，每个班组有自己的一亩三分地
  • Thread：班组里的工人，每个人干具体的活

每个线程都有一个唯一的ID，通过 threadIdx 和 blockIdx 组合计算出来。公式很简单：

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

如果是二维或三维的，稍微复杂一点：

// 二维Grid，二维Block
int blockId = blockIdx.x + blockIdx.y * gridDim.x;
int threadId = blockId * (blockDim.x * blockDim.y) 
             + threadIdx.y * blockDim.x + threadIdx.x;

嗯，这里要注意。映射方式直接影响访存效率。我个人习惯尽量用一维映射，除非处理图像或矩阵这种天然二维的数据。为什么？因为一维映射的地址计算简单，不容易出错，而且对编译器优化更友好。

小技巧：启动配置时，blockDim.x 建议取32的倍数。为什么是32？因为线程束（Warp）就是32个线程一组。取32的倍数能保证线程束满负荷工作，不会出现「半空」的线程束。

3.3 线程束执行：GPU真正的执行单元

说到线程束，这才是GPU执行的最小单位。你写代码时感觉是每个线程独立执行，但实际上，GPU是以32个线程为一组（一个Warp）来调度和执行的。

为什么会这样？因为GPU的设计哲学是「大量线程，简单控制」。它把32个线程打包成一个Warp，共享同一个指令单元。也就是说，一个Warp里的32个线程，在同一时刻执行同一条指令。

这就引出了一个经典问题：线程束发散（Warp Divergence）。

// 不好的写法：线程束发散
if (threadIdx.x % 2 == 0) {
    // 偶数线程走这里
    C[idx] = A[idx] + B[idx];
} else {
    // 奇数线程走这里
    C[idx] = A[idx] * B[idx];
}

// 好的写法：避免发散
C[idx] = (threadIdx.x % 2 == 0) ? (A[idx] + B[idx]) : (A[idx] * B[idx]);

第一个例子，同一个Warp里的线程，有的走if分支，有的走else分支。GPU没办法同时执行两条不同的指令，只能先执行if分支（偶数线程干活，奇数线程闲着），再执行else分支（奇数线程干活，偶数线程闲着）。性能直接砍半。

第二个例子用了三元运算符，虽然逻辑一样，但编译器有可能优化成无分支的代码。当然，最理想的情况是让同一个Warp里的线程走同一个分支。

核心原则：同一个Warp内的线程，尽量走相同的执行路径。如果实在避免不了，让分支粒度大于Warp大小（32），这样至少能保证一个Warp内没有发散。

我记得有一次优化一个归约（Reduction）算子，性能死活上不去。后来用nvprof一分析，发现线程束发散率高达40%。改成分块归约，每个Warp独立处理一块数据，发散率降到5%以下，性能提升了3倍多。

3.4 一张图看懂CUDA执行模型

下面这张SVG图，把Kernel启动、线程映射、线程束执行串起来了。你可以对照着看，理解会更直观。

CUDA执行模型：从Kernel启动到线程束执行 CPU（总司令） 调用 <<<G,B>>> Kernel启动 Grid（任务） gridDim.x × gridDim.y Block（班组） blockDim.x × blockDim.y × blockDim.z Thread（工人） threadIdx.x + blockIdx.x × blockDim.x Warp（线程束） 32个线程一组，共享指令单元

3.5 几个实战建议

讲完理论，给几个我实际项目中总结出来的建议：

  1. Block大小怎么选？ 我个人习惯用128或256。太小了（比如32）会导致SM利用率低；太大了（比如1024）又受限于寄存器资源。256是个比较平衡的值。
  2. Grid大小怎么算？ 保证总线程数大于等于数据量。比如处理100万个元素，block用256，那grid就是 ceil(1e6 / 256) = 3907。
  3. 别忘了边界检查。 总线程数可能比数据量大，多出来的线程要直接return，别让他们访问越界内存。
  4. 用Occupancy API辅助决策。 cudaOccupancyMaxPotentialBlockSize 这个函数能帮你算出当前硬件下最优的block大小，省得自己瞎猜。
避坑指南：我曾经在一个项目里，把blockDim设成了31。结果性能比设成32差了将近一倍。原因就是线程束对齐问题——31个线程占不满一个Warp，导致大量计算资源浪费。从那以后，我所有blockDim都是32的倍数。

好了，这一章的内容就这些。Kernel启动、线程映射、线程束执行，这三个概念是CUDA编程的基石。搞懂了它们，后面的优化技巧才能用得上。

想要更多资料，联系可可

← 上一章 📖 返回目录 下一章 →