第二讲：CUDA线程层次结构——Grid、Block、Thread的关系与索引计算
各位同学好，今天我们聊一个非常核心的话题——CUDA的线程层次结构。
说实话，我刚接触CUDA时，最头疼的就是搞不清Grid、Block、Thread这三者的关系。明明都是线程，为什么还要分三层？直接开一堆线程跑不就完了？嗯，后来我在一个图像处理项目里吃了大亏，才真正理解了这种设计的精妙之处。
一、为什么需要三层结构？
先问大家一个问题：如果你的GPU有几千个核心，你要处理一张4K图片（约800万个像素），你会怎么分配线程？
直接开800万个线程？理论上可以。但实际跑起来，你会发现性能惨不忍睹。为什么？因为GPU的硬件资源是有限的——寄存器、共享内存、线程调度器都有上限。
CUDA的设计者想了个办法：把线程分组管理。
- Thread（线程）：最基础的计算单元，每个线程执行一个任务
- Block（线程块）：一组线程的集合，共享一块共享内存
- Grid（线程网格）：所有Block的集合，对应一个kernel启动
说白了，这就是一个「分级管理」的思路。Grid是总指挥部，Block是各个作战小组，Thread就是每个士兵。
核心要点：一个Grid包含多个Block，一个Block包含多个Thread。所有Thread执行相同的kernel函数，但处理不同的数据。
二、硬件映射关系
这里我要强调一个关键点：Block是映射到SM（流多处理器）的。
我在调试一个深度学习推理引擎时，发现同样的kernel在不同GPU上性能差异巨大。后来才意识到，不同GPU的SM数量不同，每个SM能同时容纳的Block数量也不同。
| 硬件单元 | 对应软件概念 | 说明 |
|---|---|---|
| SM（流多处理器） | Block | 一个SM可以同时运行多个Block |
| CUDA Core | Thread | 每个CUDA Core执行一个线程 |
| 全局内存 | Grid | 所有Block共享全局内存 |
你想想看，如果Block数量太少，SM就闲着，利用率低。如果Block太多，每个Block分到的资源又不够。这就是调优的难点所在。
三、线程索引计算——核心公式
好，现在进入最实用的部分：怎么在kernel里知道「我是谁」？
CUDA提供了几个内置变量：
threadIdx.x：当前线程在Block内的索引blockIdx.x：当前Block在Grid内的索引blockDim.x：每个Block包含的线程数gridDim.x：每个Grid包含的Block数
那么，全局线程ID怎么算？
// 一维情况
int tid = blockIdx.x * blockDim.x + threadIdx.x;
// 二维情况
int tid_x = blockIdx.x * blockDim.x + threadIdx.x;
int tid_y = blockIdx.y * blockDim.y + threadIdx.y;
int global_id = tid_y * gridDim.x * blockDim.x + tid_x;
// 三维情况
int tid_x = blockIdx.x * blockDim.x + threadIdx.x;
int tid_y = blockIdx.y * blockDim.y + threadIdx.y;
int tid_z = blockIdx.z * blockDim.z + threadIdx.z;
我的习惯：在写kernel时，我通常会先算好全局ID，然后直接用它来访问数组。这样代码更清晰，也更容易调试。
四、实际代码示例
来看一个完整的例子：向量加法。
__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 N = 1000000;
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
注意那个边界检查 if (i < N)。我曾经在一个项目中忘了加这个判断，结果数组越界，数据全乱了。排查了整整一个下午……嗯，从那以后我再也不敢省略边界检查了。
五、多维线程索引
处理图像或矩阵时，用二维索引更直观。
__global__ void imageProcess(float *input, float *output,
int width, int height) {
// 二维线程索引
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
// 边界检查
if (col < width && row < height) {
int idx = row * width + col;
output[idx] = input[idx] * 2.0f;
}
}
// 调用方式
dim3 threadsPerBlock(16, 16);
dim3 blocksPerGrid((width + 15) / 16, (height + 15) / 16);
imageProcess<<<blocksPerGrid, threadsPerBlock>>>(d_in, d_out, width, height);
避坑指南：我曾经在二维索引计算时，把row * width + col写成了row * height + col。结果图像显示出来全是乱的。记住：行索引乘以宽度，不是高度！
六、Block和Thread的配置建议
这里我总结了一些经验值，供大家参考：
| GPU架构 | 推荐Threads/Block | 最大Threads/Block | 最大Block/Grid |
|---|---|---|---|
| Fermi (SM 2.x) | 128-256 | 1024 | 65535 |
| Kepler (SM 3.x) | 256-512 | 1024 | 2^31-1 |
| Maxwell/Pascal (SM 5.x/6.x) | 256-512 | 1024 | 2^31-1 |
| Turing/Ampere (SM 7.x/8.x) | 256-512 | 1024 | 2^31-1 |
我个人习惯：Threads/Block取256或512，然后根据数据量计算Block数量。这样既保证了足够的并行度，又不会让每个Block的资源太少。
七、常见问题与调试技巧
最后分享几个我在实际项目中踩过的坑：
- 索引越界：总是加边界检查，哪怕你觉得「肯定不会越界」
- Block数量太少：如果数据量小，可以适当减少Threads/Block，增加Block数量
- 忘记同步：同一个Block内的线程需要同步时，用
__syncthreads() - 调试方法：先用小数据量（比如N=32）跑，打印每个线程的索引，确认计算正确
总结一下：Grid、Block、Thread的三层结构，本质上是为GPU硬件量身定做的。理解索引计算是CUDA编程的第一步，也是最关键的一步。多写、多调试、多思考，慢慢就会形成直觉。
下一讲我们会深入探讨共享内存的使用，到时候会用到今天讲的Block内线程协作。大家先把索引计算练熟，后面就轻松了。
公众号:蓝海资料掘金营，微信deep3321