线程层次结构：Grid、Block、Thread的关系
各位同学，今天我们来聊聊CUDA里最基础、也最核心的概念——线程层次结构。说实话，我刚开始接触CUDA时，也被Grid、Block、Thread这三层关系搞得有点晕。但后来在项目中调试一个图像处理算法时，才真正体会到：不理解这个层次结构，你连bug都找不到在哪。
一、三层结构：从大到小看
CUDA的线程组织方式，说白了就是三级管理：
- Grid（网格）：最顶层，包含多个Block
- Block（线程块）：中间层，包含多个Thread
- Thread（线程）：最底层，实际执行计算的最小单元
你可以想象成一个公司：Grid是总公司，Block是各个部门，Thread就是每个员工。我习惯把Grid看作一个任务的总包，Block是任务的分片，Thread是实际干活的。
关键点：一个Grid里的所有Block执行相同的核函数，但处理不同的数据。这就是SIMT（单指令多线程）的精髓。
二、线程索引计算：怎么找到自己的位置？
每个线程在运行时，都需要知道自己是「谁」。CUDA提供了几个内置变量：
| 变量名 | 含义 | 取值范围 |
|---|---|---|
| threadIdx | 线程在Block内的索引 | 0 到 blockDim-1 |
| blockIdx | Block在Grid内的索引 | 0 到 gridDim-1 |
| blockDim | Block的维度（包含多少线程） | 由程序员指定 |
| gridDim | Grid的维度（包含多少Block） | 由程序员指定 |
计算全局线程ID的公式，我建议你直接记住：
// 一维情况
int tid = blockIdx.x * blockDim.x + threadIdx.x;
// 二维情况（常用）
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int idy = blockIdx.y * blockDim.y + threadIdx.y;
int global_id = idy * width + idx; // 线性化访问嗯，这里要注意：千万别把blockIdx和threadIdx搞混。我在项目中遇到过一位同事，写了个矩阵乘法，结果把blockIdx当threadIdx用，算出来的结果全是乱的。调试了一下午才发现是索引算错了。
三、线程束（Warp）：真正的执行单元
讲完了逻辑结构，咱们聊聊硬件层面的东西。线程束（Warp）是GPU真正执行指令的单位。一个Warp包含32个连续的线程。
为什么会这样？因为GPU的SM（流多处理器）一次取指令，会同时发给32个线程执行。你想想看，如果Block里有128个线程，那它会被分成4个Warp来执行。
个人经验：我建议你在设计核函数时，尽量让Block的线程数是32的倍数。比如128、256、512。这样能保证Warp被填满，不会浪费计算资源。我曾经为了省显存，设了个96线程的Block，结果性能反而下降了——因为最后32个线程凑不满一个Warp，白白浪费了计算单元。
四、Warp发散：性能杀手
这里有个坑，我必须提醒你：Warp发散（Warp Divergence）。
当一个Warp里的线程走不同的分支时（比如if-else），GPU会串行执行所有分支。什么意思呢？就是本来32个线程可以并行干活，现在得分成两拨，先执行if的线程，再执行else的线程。性能直接砍半。
// 不好的写法：Warp发散
if (threadIdx.x % 2 == 0) {
// 偶数线程执行
data[threadIdx.x] *= 2;
} else {
// 奇数线程执行
data[threadIdx.x] += 1;
}
// 好的写法：避免Warp内分支
int mask = threadIdx.x & 1; // 用位运算代替分支
data[threadIdx.x] += mask;
data[threadIdx.x] += (1 - mask) * data[threadIdx.x];我曾经在一个图像滤波项目里，因为没注意Warp发散，性能比预期慢了3倍。后来改成用位运算和三元表达式，速度一下就上来了。
五、Block与Grid的维度选择
选多大Block？选多少Grid？这是每个CUDA程序员都会遇到的问题。我的建议是：
- Block大小：128到512线程之间比较合适。太小了SM利用率低，太大了寄存器压力大。
- Grid大小：至少是SM数量的2倍以上，保证有足够的Block来隐藏访存延迟。
- 维度选择：处理图像用2D，处理体数据用3D，处理数组用1D。别为了炫技用3D处理一维数组，没意义。
注意：每个Block的线程数不能超过1024（对于现代GPU）。这是硬件限制，硬编码的。我曾经在调试一个老项目时，发现有人设了2048线程的Block，结果核函数根本跑不起来，报错说启动失败。查了半天才发现是Block太大。
六、知识体系图
下面这张图，我画了Grid、Block、Thread和Warp之间的关系，你一看就明白了：
七、总结一下
线程层次结构，说白了就是三件事：
- Grid → Block → Thread：逻辑上的三级组织，每个线程通过blockIdx和threadIdx定位自己。
- Warp：硬件上的执行单元，32个线程一组。设计时尽量让Block大小是32的倍数。
- 避免Warp发散：同一个Warp里的线程，尽量走相同的代码路径。
我记得刚学CUDA那会儿，总觉得这些概念太抽象。后来在项目里亲手调过几次性能，踩过几个坑，才真正理解。你写代码的时候，多想想「我这个线程在哪个Block里？这个Block在Grid的什么位置？Warp里的线程会不会走分支？」——想清楚了，性能自然就上去了。
想要更多资料，联系可可