3. 线程层次结构
说到GPU编程，线程层次结构是绕不开的核心概念。我刚开始接触CUDA时，总觉得Grid、Block、Thread这三层关系有点绕。后来在项目中踩过几次坑，才真正理解了它们的划分逻辑。今天咱们就把它彻底讲透。
3.1 三层结构：从大到小的组织方式
GPU的线程组织方式，说白了就是三级管理架构：
- Grid：最顶层，一个kernel启动对应一个Grid
- Block：Grid里的工作组，每个Block内的线程可以协作
- Thread：最底层的执行单元，每个线程执行相同的指令
你想想看，这就像公司组织架构：Grid是总公司，Block是部门，Thread是员工。每个部门内部可以开会协作，但跨部门沟通就比较麻烦。
关键点：同一个Block内的线程可以共享数据（通过共享内存），也可以同步（通过__syncthreads()）。不同Block的线程之间，原则上不能直接通信。
3.2 划分逻辑：怎么分才合理？
我个人习惯把划分逻辑总结成三个问题：
- 数据规模多大？——决定Grid大小
- 需要多少协作？——决定Block大小
- 硬件限制在哪？——决定最终配置
举个例子。假设你要处理一张1024x1024的图像，每个像素需要一个线程。那么总线程数就是1,048,576个。怎么分？
我建议这样思考：
- 每个Block放256个线程（16x16），这是比较常见的配置
- 那么Block数量就是1,048,576 / 256 = 4,096个
- Grid就包含这4,096个Block
代码实现是这样的：
dim3 blockDim(16, 16); // 每个Block 16x16 = 256个线程
dim3 gridDim(64, 64); // Grid 64x64 = 4096个Block
kernel<<<gridDim, blockDim>>>(d_data, width, height);3.3 映射关系：线程如何对应到数据？
嗯，这里要注意。线程的索引计算是很多新手容易搞混的地方。我见过有人把blockIdx和threadIdx搞反，结果数据全乱了。
标准做法是用内置变量：
blockIdx.x：当前Block在Grid中的x索引blockIdx.y：当前Block在Grid中的y索引threadIdx.x：当前线程在Block中的x索引threadIdx.y：当前线程在Block中的y索引
全局线程ID的计算公式：
int global_x = blockIdx.x * blockDim.x + threadIdx.x;
int global_y = blockIdx.y * blockDim.y + threadIdx.y;
int global_id = global_y * width + global_x;我曾经在项目中遇到过一个bug：图像处理时边缘像素总是对不上。排查了半天，发现是忘记考虑Block边界对齐的问题。后来我养成了一个习惯——每次写kernel前，先在纸上画一遍线程到数据的映射图。
3.4 硬件层面的映射：Warp与SM
讲完了逻辑划分，咱们看看硬件怎么执行。这里有个重要概念——Warp。
Warp是GPU调度的最小单位，通常32个线程一组。同一个Warp里的线程执行相同的指令（SIMT模型）。
| 概念 | 逻辑层面 | 硬件层面 |
|---|---|---|
| Grid | kernel启动的全体线程 | 分配到多个SM |
| Block | 可协作的线程组 | 分配到单个SM |
| Warp | 无对应逻辑概念 | 32线程一组，SM调度单位 |
| Thread | 最小执行单元 | SP/CUDA Core执行 |
为什么Block大小最好是32的倍数？因为Warp是32线程。如果Block大小不是32的倍数，最后一个Warp会有空闲线程，造成资源浪费。我一般选128或256，既保证是32的倍数，又不会太大导致寄存器压力。
小技巧：Block大小建议在128-512之间。太小了SM利用率低，太大了寄存器不够用。具体值可以通过profiling工具调优。
3.5 避坑指南：我踩过的那些坑
做GPU开发这些年，我在线程层次上栽过不少跟头。分享几个典型的：
- Block数量太少：曾经写过一个kernel，只分了8个Block。结果GPU上几十个SM，只有8个在工作，利用率惨不忍睹。后来改成至少几十个Block，性能翻了好几倍。
- 线程ID计算错误：处理三维数据时，忘记考虑z维度的索引。排查了整整两天，最后发现是global_id算错了。
- 共享内存越界：Block内线程访问共享内存时，索引超出了BlockDim的范围。这种bug很难复现，因为有时候数据刚好没被覆盖。
警告：千万不要假设Block数量等于SM数量。现代GPU的SM数量从几十到上百不等，你的代码应该能适应不同硬件。用occupancy API来查询最佳配置，而不是硬编码。
3.6 核心逻辑图：线程层次结构总览
下面这张图展示了从Grid到Thread的完整映射关系，以及它们如何对应到硬件资源：
这张图把逻辑层次和硬件映射放在一起对比。你看，Grid里的Block会被分配到不同的SM上，每个SM再以Warp为单位调度线程。理解了这个映射关系，你就能明白为什么Block大小和数量会影响性能。
3.7 总结：记住这三句话
线程层次结构其实不复杂，记住三点就够了：
- Grid划分Block：根据数据规模决定，保证有足够多的Block填满SM
- Block划分Thread：根据协作需求决定，大小是32的倍数，128-512最佳
- Thread映射数据：用blockIdx和threadIdx计算全局ID，注意边界检查
嗯，这些就是我多年实战中总结的经验。下次写kernel时，先想清楚这三层怎么划分，能省去很多调试时间。