线程层次结构：Grid、Block、Thread的三层映射
说实话，刚接触CUDA的时候，我也被Grid、Block、Thread这三层结构搞得有点晕。你想想看，CPU编程里我们最多就考虑多核多线程，到了GPU这里突然冒出三层来，确实需要一点时间消化。
但别担心，这三层结构其实很直观。我习惯把它想象成一个军队的编制：Grid是军团，Block是连队，Thread就是每个士兵。每个士兵（Thread）执行相同的指令，但处理不同的数据——这就是SIMT（单指令多线程）的精髓。
为什么需要三层映射？
GPU有成千上万个核心，但你不能让所有线程都乱跑。需要一套组织体系来管理它们。Grid和Block就是这套管理体系。
- Grid：整个任务，对应一个核函数调用
- Block：任务的一个子块，可以独立执行
- Thread：最小的执行单元，每个线程执行核函数中的代码
我在项目中遇到过一个问题：如果只用一维的线程索引，处理二维图像时特别别扭。后来发现CUDA支持一维、二维、三维的索引，这才舒服多了。
内置变量：你的导航工具
CUDA提供了几个内置变量，帮你定位当前线程在全局中的位置。这些变量在核函数里可以直接用，不需要声明。
| 变量名 | 类型 | 含义 |
|---|---|---|
| threadIdx | dim3 | 当前线程在Block内的索引 |
| blockIdx | dim3 | 当前Block在Grid内的索引 |
| blockDim | dim3 | Block的维度大小（每个Block有多少线程） |
| gridDim | dim3 | Grid的维度大小（有多少个Block） |
嗯，这里要注意：threadIdx 和 blockIdx 都是从0开始的。我第一次写的时候忘了这一点，结果数组越界了，排查了半天。
一维到三维索引计算
这是最核心的部分。你要学会如何把三维的线程索引映射到一维的数据上。说白了，就是算一个全局ID。
一维情况
最简单的情况。每个Block是一维的，Grid也是一维的。
int tid = blockIdx.x * blockDim.x + threadIdx.x;这个公式我闭着眼睛都能写出来。它把Block的偏移量和Block内的偏移量加起来，得到全局唯一的线程ID。
二维情况
处理图像时最常用。每个Block是二维的（比如16x16），Grid也是二维的。
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int tid = y * width + x; // 映射到一维数组我个人习惯用 row-major 的存储方式，也就是先存行再存列。这样 y * width + x 就是正确的线性索引。
三维情况
处理体数据（比如CT扫描）时会用到。每个Block是三维的，Grid也是三维的。
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 tid = z * width * height + y * width + x;核心公式：全局线程ID = blockIdx * blockDim + threadIdx
这个公式适用于任意维度，只是每个维度单独计算。
一个完整的例子
来看一个向量加法的核函数。假设我们要把两个长度为N的向量相加。
__global__ void vecAdd(float* A, float* B, float* C, int N) {
// 计算全局线程ID
int tid = blockIdx.x * blockDim.x + threadIdx.x;
// 边界检查，防止越界
if (tid < N) {
C[tid] = A[tid] + B[tid];
}
}调用时，我们需要指定Grid和Block的大小：
int N = 1000000;
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);小技巧：Block的大小通常是32的倍数（因为warp大小是32）。我一般用128或256，性能比较稳定。
避坑指南
我曾经踩过的坑：
- 忘记做边界检查：当N不能被blockDim整除时，最后一个Block会有多余的线程。不做边界检查就会访问非法内存。
- 混淆了blockDim和gridDim：blockDim是每个Block的大小，gridDim是Block的数量。我刚开始总搞反。
- 三维索引计算错误：处理三维数据时，一定要搞清楚哪个维度是x，哪个是y，哪个是z。CUDA的维度顺序是(x, y, z)。
用SVG图来理解三层结构
下面这张图展示了Grid、Block、Thread的层次关系，以及索引的计算方式。我建议你仔细看看，理解了这个图，三层映射就掌握了八成。
这张图里，Grid有3x2=6个Block，每个Block有4x3=12个Thread。总共72个线程。每个线程的全局ID可以通过 blockIdx * blockDim + threadIdx 计算出来。
实际开发中的选择建议
我做了几年CUDA开发，总结了一些经验：
- Block大小：建议128-512之间。太小了硬件利用率低，太大了可能超过资源限制（每个Block的共享内存和寄存器是有限的）。
- Grid大小：尽量让Block数量远大于SM数量（通常是几十倍），这样硬件可以自动做负载均衡。
- 维度选择：处理一维数据用一维索引，处理二维图像用二维索引，处理三维体数据用三维索引。不要为了炫技强行用三维。
说白了，线程层次结构就是CUDA的骨架。理解了它，你就知道怎么把数据切分给成千上万个线程去并行处理。我刚开始学的时候，花了一周才真正搞明白。但一旦理解了，后面的学习就顺畅多了。
记住那个核心公式：全局ID = blockIdx * blockDim + threadIdx。这是你写所有CUDA核函数的基础。
练习建议：写一个简单的核函数，打印出每个线程的blockIdx、threadIdx和计算出的全局ID。观察输出结果，验证你的理解是否正确。
想要更多资料，联系可可