第2章 CUDA线程模型

第2章：CUDA线程模型：线程层次结构与执行模型

好，咱们今天来聊聊CUDA里最核心的概念之一——线程模型。说实话，我见过不少新手，写CUDA代码时把线程索引算错了，结果数据对不上，排查半天。嗯，这章咱们就把这个基础打牢。

2.1 线程层次结构：grid、block、thread

CUDA的线程组织方式，说白了就是三层结构：线程（thread）→ 线程块（block）→ 线程网格（grid）。你想想看，这就像军队编制：一个士兵是thread，一个班是block，一个团是grid。

thread：最基础的执行单元，每个thread执行相同的kernel函数
block：一组thread的集合，block内的thread可以共享内存（shared memory）并同步
grid：一组block的集合，一个kernel启动时对应一个grid

我个人习惯把grid想象成「任务大包」，block是「任务分包」，thread就是「干活的工人」。每个工人干的事一样，但处理的数据不同。

关键点：block内的thread可以通过__syncthreads()同步，但不同block之间的thread无法直接同步。这是硬件设计决定的，别想着跨block做同步，那得用全局内存加原子操作，性能会打折扣。

2.2 线程索引计算：一维、二维、三维

线程索引计算是写CUDA的必修课。我刚开始学的时候，经常把threadIdx和blockIdx搞混。后来总结了一个口诀：先算block内的偏移，再算block间的偏移。

2.2.1 一维网格与一维块

最简单的情况，grid和block都是一维的。全局线程ID的计算公式：

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

这个公式我闭着眼睛都能写出来。每个block有blockDim.x个线程，当前是第blockIdx.x个block，所以前面有blockIdx.x * blockDim.x个线程，再加上block内的偏移threadIdx.x。

2.2.2 二维网格与二维块

处理图像时，二维索引更直观。假设图像宽W、高H，每个像素对应一个线程：

int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
int idx = row * W + col;  // 全局线性索引

我在项目中遇到过一个问题：图像边界处理。如果线程数不是图像尺寸的整数倍，多余的线程会访问越界。所以一定要加边界检查：

if (col < W && row < H) {
    // 安全地处理像素
}

避坑指南：我曾经因为忘记加边界检查，导致程序在特定尺寸的图像上崩溃。排查了整整一下午，最后发现是线程数比图像像素多，访问了非法内存。从那以后，我写任何kernel都会先加边界判断。

2.2.3 三维网格与三维块

三维索引用于体数据（volume data）处理，比如CT扫描、流体模拟。计算方式类似：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
int idx = z * W * H + y * W + x;

说实话，三维用得不多。大部分场景一维或二维就够了。但做科学计算的朋友可能会用到。

2.3 线程块大小与网格大小的选择

block大小怎么选？这是个经典问题。我一般遵循以下原则：

因素	建议	原因
block内线程数	128~512	太少则占用率低，太多则寄存器压力大
block维度	尽量用一维或二维	三维索引计算开销稍大
block大小	32的倍数	warp大小是32，非32倍数会浪费线程
grid大小	至少能填满所有SM	保证足够的并行度

小技巧：我习惯用cudaOccupancyMaxPotentialBlockSize这个API来获取最优block大小。它会根据kernel的寄存器使用量和共享内存使用量，自动算出最高占用率的配置。省心省力。

2.4 Warp与SIMT执行模型

终于聊到warp了。这是CUDA最精髓的部分，也是很多性能问题的根源。

2.4.1 什么是Warp

warp是GPU调度和执行的基本单位，大小为32个线程。一个block被分成若干个warp，每个warp内的线程执行相同的指令（SIMT模型）。

你想想看，GPU有几千个核心，但调度器是以warp为单位发指令的。也就是说，同一warp内的32个线程，在同一时刻执行同一条指令，只是处理的数据不同。

2.4.2 SIMT vs SIMD

很多人问SIMT和SIMD有什么区别。我打个比方：

SIMD（单指令多数据）：像军训队列，所有人必须做完全一样的动作，步调一致
SIMT（单指令多线程）：像自由体操，虽然教练喊的是同一个口令，但每个运动员可以有自己的节奏（通过分支、掩码等）

SIMT更灵活。每个线程有自己的程序计数器、寄存器和执行状态。但代价是：如果warp内出现分支分歧（divergence），性能会下降。

2.4.3 Warp Divergence（分支分歧）

这是性能杀手。当warp内的线程走不同的分支时，GPU会串行执行所有分支路径，然后通过掩码屏蔽掉不执行的线程。

// 不好的写法：warp内分支分歧
if (threadIdx.x % 2 == 0) {
    // 偶数线程执行
} else {
    // 奇数线程执行
}

// 好的写法：避免warp内分歧
// 可以通过数据重排或使用三元运算符

我在项目中遇到过这样一个案例：一个图像处理kernel，因为根据像素值做了条件判断，导致warp分歧严重，性能只有理论值的60%。后来我通过将数据按条件排序，让同一warp内的线程处理相同类型的像素，性能直接翻倍。

核心原则：尽量让同一warp内的线程执行相同的指令路径。如果无法避免，至少保证分支条件是基于warpId（即threadIdx.x / 32）的，而不是基于threadIdx.x的奇偶性。

2.4.4 Warp调度与延迟隐藏

GPU通过快速切换warp来隐藏内存访问延迟。当一个warp在等待数据时，调度器会切换到另一个就绪的warp执行。这就是为什么需要足够的warp来「填满」流水线。

我建议：block内的warp数量至少是4~8个，这样调度器有足够的warp来切换。如果block太小，warp数量不足，GPU就会空闲等待，利用率上不去。

2.5 实战经验总结

嗯，讲了这么多，最后总结几条我踩过的坑：

线程索引计算一定要验证：写个简单的打印kernel，输出每个线程的全局ID，确认和预期一致
block大小选256最稳妥：大多数情况下，256线程的block能获得不错的占用率
警惕warp分歧：用nvprof或Nsight Compute查看分支效率，低于90%就要优化
grid大小要足够大：至少是SM数量的2~4倍，保证负载均衡

最后一个小建议：刚开始学CUDA时，别急着优化。先把线程索引算对，把边界检查写好，保证功能正确。性能优化是第二步的事。我见过太多人一开始就追求极致性能，结果代码跑出来全是错的。

下一章我们会聊共享内存和同步机制，那是另一个容易踩坑的地方。咱们到时候见。

📚 课程目录