4、线程层次结构：Grid、Block、Thread的三层关系
好，咱们今天聊点硬核的。线程层次结构，说白了就是GPU怎么组织它那成千上万个干活的小工。你想想看，CPU里顶多十几个线程，管理起来简单。但GPU动辄几万个线程同时跑，没有一套清晰的层级关系，那不乱套了？
我个人习惯把GPU的线程组织想象成一个军队：Grid是集团军，Block是连队，Thread就是每个士兵。这个比喻虽然简单，但很管用。
4.1 三层结构：从宏观到微观
先看一张表，把这三层的关系理清楚：
| 层级 | 英文名 | 我的理解 | 特点 |
|---|---|---|---|
| 网格 | Grid | 整个任务 | 一个kernel启动对应一个Grid |
| 线程块 | Block | 任务小组 | 块内线程可以协作（共享内存、同步） |
| 线程 | Thread | 最小执行单元 | 每个线程执行相同的kernel代码 |
嗯，这里要注意：不同Block之间的线程是独立的，它们不能直接通信，也不能互相同步。这是GPU设计上的一个关键约束。我在项目中遇到过有人试图跨Block共享数据，结果折腾了半天才发现这条路走不通——最后只能用全局内存加原子操作来曲线救国。
4.2 内置变量：你的定位工具
CUDA给了我们几个内置变量，用来在代码里定位当前线程是谁、在哪儿。说白了，就是每个线程的「身份证」和「地址簿」。
threadIdx：线程在块内的编号
这个变量告诉你，在当前Block里，你是第几个线程。它是一个三维向量：threadIdx.x、threadIdx.y、threadIdx.z。为什么是三维？因为GPU处理图像、体数据这些三维问题很常见，三维索引用起来更直观。
blockIdx：块在网格内的编号
这个变量告诉你，当前Block在整个Grid里是第几个。同样也是三维的：blockIdx.x、blockIdx.y、blockIdx.z。
blockDim：块的维度大小
每个Block里有多少个线程？blockDim告诉你答案。比如blockDim.x = 256，意味着每个Block在x方向上有256个线程。
gridDim：网格的维度大小
整个Grid里有多少个Block？gridDim就是干这个的。比如gridDim.x = 8，表示Grid在x方向上有8个Block。
核心公式：要计算一个线程在全局中的唯一ID，通常这样写：
int tid = blockIdx.x * blockDim.x + threadIdx.x;这个公式我几乎每个kernel里都会用到，已经形成肌肉记忆了。
4.3 实战：一维到三维的映射
光讲理论不过瘾，咱们直接上代码。假设我们要处理一个长度为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];
}
}
// 调用方式
int N = 1000000;
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);这段代码里，blocksPerGrid的计算方式我特别想强调一下：(N + threadsPerBlock - 1) / threadsPerBlock。这个技巧叫「向上取整除法」，保证即使N不能被threadsPerBlock整除，也能覆盖所有数据。我曾经因为忘了加这个-1，导致数组末尾几个元素没被处理，调试了整整一个下午。
4.4 二维和三维：处理图像数据
处理图像时，二维索引更自然。比如一张1024x768的图片：
__global__ void imageProcess(unsigned char *img, int width, int height) {
// 计算二维坐标
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
// 边界检查
if (x < width && y < height) {
int idx = y * width + x; // 转为一维索引
// 处理像素...
img[idx] = img[idx] * 2; // 简单示例：亮度加倍
}
}
// 调用方式
dim3 threadsPerBlock(16, 16); // 每个Block 16x16=256个线程
dim3 blocksPerGrid(
(width + threadsPerBlock.x - 1) / threadsPerBlock.x,
(height + threadsPerBlock.y - 1) / threadsPerBlock.y
);
imageProcess<<<blocksPerGrid, threadsPerBlock>>>(d_img, width, height);我的经验：Block的线程数通常选16x16或32x8这种组合。为什么？因为GPU的warp（线程束）大小是32，Block的线程数最好是32的倍数，这样能充分利用硬件资源。选256（16x16）是最稳妥的，几乎适用于所有GPU。
4.5 避坑指南：我踩过的那些坑
讲几个我实际项目中遇到的坑，希望能帮你省点时间：
- Block内线程数上限：每个Block最多1024个线程（不同架构略有差异）。我曾经想当然地设了2048，结果kernel直接启动失败，连个错误提示都没有——嗯，CUDA的错误检查一定要养成习惯。
- 三维索引别搞混：
threadIdx.z和blockIdx.z的默认值是0。如果你只用一维或二维，记得把没用的维度设成1，比如dim3 threads(256, 1, 1)。 - 边界检查不能省：总线程数可能超过实际数据量，
if (tid < N)这种检查一定要写。我见过有人偷懒不写，结果程序跑起来偶尔崩溃，查了三天才发现是越界访问。
重要提醒：threadIdx、blockIdx这些内置变量只能在kernel函数内部使用。在主机端（CPU代码里）访问它们会编译报错。这个错误新手经常犯，我自己刚开始学的时候也中过招。
4.6 总结：记住这三句话
好了，这一章的内容就这些。最后帮你总结三句口诀：
- Grid管全局，Block管局部，Thread干实事
- 全局ID = blockIdx * blockDim + threadIdx
- Block大小选256，性能稳妥不出错
下一章我们会聊共享内存——那个让Block内线程真正协作起来的神奇工具。到时候你会发现，理解了线程层次结构，共享内存用起来就顺理成章了。