线程层次结构：Grid、Block、Thread
好，咱们直接进入正题。CUDA 的线程层次结构，说白了就是 GPU 怎么组织它那成千上万个干活的小工（线程）。你如果搞不懂这个，后面写核函数基本就是瞎蒙。
我个人习惯把 GPU 想象成一个大型工厂：Grid 是整个工厂，Block 是车间，Thread 就是车间里具体的工人。工人（Thread）在车间（Block）里干活，多个车间组成一个工厂（Grid）。就这么简单。
核心要点：CUDA 采用三级层次结构——Grid → Block → Thread。每个层次都有对应的索引，用来唯一标识每个线程。
1. 三个层次的概念
Thread（线程）：最基础的执行单元。每个线程执行相同的核函数代码，但处理不同的数据。嗯，这就是 SIMT（单指令多线程）的精髓。
Block（线程块）：一组线程的集合。同一个 Block 里的线程可以共享一块内存（Shared Memory），也能通过同步操作（__syncthreads()）来协作。我在项目中遇到过，很多人以为 Block 越大越好，其实不是——每个 Block 能用的资源（寄存器、共享内存）是有限的。
Grid（网格）：所有 Block 的集合。一个核函数启动时，就对应一个 Grid。Grid 可以是一维、二维或三维的，看你处理的数据结构。
我的经验：处理图像时用二维 Grid 最自然，因为图像本身就是二维的。处理一维数组就用一维 Grid。别为了炫技硬上三维，除非你真的需要。
2. 线程索引计算
这是新手最容易晕的地方。每个线程怎么知道自己该处理哪块数据？靠的就是内置变量：threadIdx、blockIdx、blockDim、gridDim。
咱们直接看代码，比说一万句都管用：
// 一维 Grid，一维 Block
int tid = blockIdx.x * blockDim.x + threadIdx.x;
// 二维 Grid，二维 Block（处理图像时常用）
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int idx = y * width + x; // 转成一维索引
// 三维 Grid，三维 Block（处理体数据）
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;公式其实就一个：全局线程ID = blockIdx * blockDim + threadIdx。你想想看，每个 Block 里有 blockDim 个线程，当前是第 blockIdx 个 Block，那这个 Block 的起始位置就是 blockIdx * blockDim，再加上线程在 Block 内的偏移 threadIdx，就得到了全局唯一的线程 ID。
避坑指南：我曾经犯过一个低级错误——在二维 Grid 里直接用 threadIdx.x 当全局索引。结果呢？不同 Block 里的线程索引重复了，数据被反复覆盖，跑出来的结果一塌糊涂。记住：threadIdx 只在 Block 内唯一，跨 Block 会重复。
3. 线程块大小选择
这是性能调优的关键一步。Block 大小选错了，性能可能差好几倍。我一般遵循以下几个原则：
| 因素 | 说明 | 建议值 |
|---|---|---|
| Warp 大小 | GPU 执行的基本单位是 Warp（32线程）。Block 大小最好是 32 的倍数 | 128、256、512 |
| 寄存器压力 | 每个线程用的寄存器越多，能同时活跃的线程越少 | 根据 Occupancy Calculator 算 |
| 共享内存 | 每个 Block 能用的共享内存有限，太大就限制活跃 Block 数 | 通常 16KB-48KB |
| SM 资源 | 每个 SM 能同时容纳的 Block 和线程数有上限 | Max Blocks/SM: 16-32 |
我个人习惯从 256 开始试。为什么？256 是 32 的倍数，而且大多数 GPU 上能获得不错的 Occupancy（占用率）。如果核函数逻辑简单、寄存器用得少，我会试试 512。如果逻辑复杂、寄存器压力大，128 可能更合适。
一个小技巧：用 cudaOccupancyMaxPotentialBlockSize 这个 API 帮你算。它会根据你的核函数自动推荐最优的 Block 大小。我刚开始做优化时全靠手动试，后来发现这个 API 真香。
4. 知识体系总览
下面这张图把线程层次结构的核心逻辑串起来了，你一看就明白：
这张图里，Grid 包含 3 个 Block，每个 Block 包含 4 个线程（2×2）。Block (1,0) 里的 Thread (1,1)，它的全局 x 索引就是 1×2 + 1 = 3，全局 y 索引是 0×2 + 1 = 1。嗯，就是这么算的。
5. 实际项目中的选择策略
说了这么多理论，咱们来点实际的。我一般按这个流程来选 Block 大小：
- 先看数据量：总线程数要能覆盖所有数据。比如处理 100 万个元素，总线程数至少 100 万。
- 确定 Block 大小：从 256 开始，保证是 32 的倍数。
- 计算 Grid 大小：Grid 大小 = (数据量 + Block 大小 - 1) / Block 大小。向上取整。
- 检查资源限制：用 Occupancy API 验证，确保每个 SM 有足够的活跃 Warp。
- 微调：如果性能不理想，试试 128、192、256、512，看哪个最快。
记住：Block 大小不是越大越好。太大的 Block 会导致每个 SM 能同时处理的 Block 数减少，反而降低 Occupancy。我见过有人把 Block 设成 1024，结果性能还不如 256 的——因为寄存器不够用，线程被频繁换进换出。
好了，线程层次结构这块就讲这么多。你只要记住 Grid、Block、Thread 的关系，会算线程索引，知道怎么选 Block 大小，后面写核函数就顺了。
最后说一句：别怕试错。我刚开始调 Block 大小时，每个核函数至少试 5 种配置。跑一遍 profiler，看 Occupancy 和带宽利用率，哪个高就用哪个。时间长了，你一眼就能看出该用多大。
想要更多资料，联系可可