4. 线程的诞生：从Grid到Block再到Thread的层级结构
好，咱们今天聊聊CUDA里最核心的一个概念——线程是怎么来的。
你写一个kernel函数，调用它的时候，指定了<<
4.1 三层结构：Grid、Block、Thread
说白了，CUDA把线程分成了三个层级：
- Thread（线程）：最底层的执行单元。每个线程执行相同的kernel代码，但处理不同的数据。
- Block（线程块）：一组线程的集合。同一个block里的线程可以共享数据（通过shared memory），也能同步（通过__syncthreads()）。
- Grid（网格）：所有block的集合。一个grid对应一次kernel调用。
你可以把Grid想象成一个工厂，Block是车间，Thread就是车间里的工人。每个工人干一样的活，但手里的零件不同。
关键点：线程是并行执行的，但block内部的线程可以协作，block之间不行（除非用全局内存+原子操作，但那很慢）。
4.2 为什么要有Block？
你可能会问：为什么不直接搞一堆线程，非要中间加个Block？
嗯，这里有两个原因。
第一，硬件限制。 GPU里的SM（流多处理器）一次能调度的线程数是有限的。比如一个SM最多支持2048个线程。如果你要跑100万个线程，SM得分批处理。Block就是分批的单位。
第二，数据共享。 我在项目中遇到过，需要让相邻的线程交换数据。如果没有Block，每个线程只能读写自己的寄存器，没法共享。有了Block，就能用shared memory做快速交换。我曾经做过一个矩阵转置的优化，用shared memory把性能提升了5倍多。
个人经验：Block的大小一般选128或256。太小了浪费SM资源，太大了shared memory不够用。我习惯先试256，不行再调。
4.3 线程的索引：threadIdx, blockIdx, blockDim, gridDim
每个线程都有自己唯一的ID。CUDA提供了几个内置变量：
| 变量 | 含义 | 例子 |
|---|---|---|
| threadIdx | 线程在block内的索引 | threadIdx.x = 0, 1, 2, ... |
| blockIdx | block在grid内的索引 | blockIdx.x = 0, 1, 2, ... |
| blockDim | block的维度（大小） | blockDim.x = 256 |
| gridDim | grid的维度（大小） | gridDim.x = 100 |
要计算全局唯一的线程ID，公式是：
int tid = blockIdx.x * blockDim.x + threadIdx.x;这个公式我闭着眼睛都能写出来。你想想看，每个block有blockDim.x个线程，第blockIdx.x个block的起始位置就是blockIdx.x * blockDim.x，再加上线程在block内的偏移threadIdx.x，就是全局ID了。
4.4 一维、二维、三维：你选哪个？
CUDA支持一维、二维、三维的线程组织。什么意思？
- 一维：处理数组、向量。最简单，最常用。
- 二维：处理图像、矩阵。用(row, col)索引更直观。
- 三维：处理体数据、3D网格。用得少，但需要时很爽。
我个人习惯，能用一维就用一维。除非数据本身是二维的（比如图像），用二维会让代码更清晰。
举个例子，处理一张1024x1024的图像：
dim3 blockDim(16, 16); // 每个block 16x16 = 256个线程
dim3 gridDim(64, 64); // 总共64x64 = 4096个block
// kernel里这样算坐标
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
int idx = row * width + col; // 转成一维索引你看，用二维组织，代码里直接操作行和列，不用自己算偏移。但底层执行时，硬件还是把它当成一维的线程来调度。说白了，维度只是方便你写代码。
4.5 线程的诞生过程：一张图说清楚
下面这张图展示了从Grid到Block再到Thread的完整层级。我画了个SVG，你一看就明白。
这张图里，Grid里有3个Block，每个Block里有256个线程。实际项目中，Block的数量可能成百上千，线程数更是百万级。但结构是一样的。
4.6 避坑指南：Block大小怎么选？
我曾经踩过一个坑。刚开始写CUDA时，觉得线程越多越好，把Block设成1024。结果程序跑得比CPU还慢。
为什么？因为一个SM里的寄存器、shared memory是有限的。Block太大，每个线程分到的资源就少，反而限制了并行度。
我现在的经验是：
- Block大小：128~512之间。256是个好起点。
- Block数量：至少是SM数量的几倍。比如你的GPU有20个SM，那至少搞80个Block，让SM有事干。
- 总线程数：覆盖你的数据量就行。别为了凑数多开线程。
注意：Block的线程数不能超过1024（MaxThreadsPerBlock）。这是硬件限制，你硬写超过会编译报错。
4.7 小结
好了，咱们总结一下：
- Grid是最大的单位，对应一次kernel调用。
- Block是中间层，负责组织线程，支持共享和同步。
- Thread是最小单位，真正干活的就是它。
- 用threadIdx、blockIdx、blockDim、gridDim这四个变量，就能算出每个线程的全局ID。
- 维度（1D/2D/3D）只是方便你写代码，底层都一样。
下一节咱们聊聊线程怎么被调度到SM上执行。嗯，那才是真正有意思的部分。