一、线程层次结构：Grid、Block、Thread

各位同学，咱们今天聊点实在的。CUDA 编程里最基础也最容易搞混的东西——线程层次结构。说白了，就是 GPU 怎么组织它那成千上万个干活的小工（线程）。

我刚开始学 CUDA 的时候，看着 Grid、Block、Thread 这三个词，脑子里全是问号。后来在项目里调一个图像处理的核函数，死活跑不对，才发现是线程索引算错了。嗯，从那以后，我对这块的理解就深刻多了。

你可以这么理解：

Thread（线程）：最底层的执行单元。每个线程执行相同的指令，但处理不同的数据。
Block（线程块）：一组线程的集合。同一个 Block 里的线程可以共享显存（Shared Memory），也能通过同步指令（__syncthreads()）互相等待。
Grid（网格）：所有 Block 的集合。一个核函数启动时，就对应一个 Grid。

我习惯用「工厂流水线」来类比：

核心要点：Block 内的线程可以协作，Block 之间的线程是独立的。这个特性决定了你设计核函数时的数据划分策略。

说实话，我见过太多人在索引计算上栽跟头。包括我自己。有一次做矩阵乘法，算出来的结果全是乱的，debug 了一下午，最后发现是线程 ID 算错了。

咱们先看一维的情况：

// 一维网格，一维块
int tid = blockIdx.x * blockDim.x + threadIdx.x;

这个公式很简单：blockIdx.x 是当前 Block 在 Grid 中的编号，blockDim.x 是每个 Block 有多少线程，threadIdx.x 是线程在 Block 内的编号。

二维的情况稍微复杂点：

// 二维网格，二维块
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int tid = y * width + x;  // 转成一维索引

三维的就更不用说了，但道理是一样的。你想想看，GPU 本质上还是线性地执行线程，多维索引只是方便我们映射到多维数据（比如图像、体素）。

我的小技巧：写核函数之前，先在纸上画个网格，标出每个线程的坐标。别嫌麻烦，这能省下你两小时的 debug 时间。

这个问题没有标准答案。我个人习惯是：

我曾经在一个项目里，处理 4K 分辨率的图像。一开始偷懒用一维网格，结果索引计算绕来绕去，性能还不好。改成二维网格后，代码清爽了，性能也提升了 15%。

注意：Block 的维度（blockDim.x, blockDim.y, blockDim.z）乘积不能超过 1024（对于现代 GPU）。这是硬件限制，别想着一个 Block 塞 10000 个线程。

下面这张图，是我自己总结的线程层次与索引计算的关系。你看一眼，应该能明白整个逻辑。

讲几个我踩过的坑，你们记一下：

Block 大小不是越大越好。我曾经以为一个 Block 塞 512 个线程肯定比 128 个快。结果呢？寄存器不够用，性能反而下降。合适的 Block 大小需要根据你的核函数资源占用情况来定。
别忘了检查 grid 边界。数据量不一定是 Block 大小的整数倍。我习惯在核函数里加个 if 判断：if (tid < N)，防止越界。
二维索引的宽高别搞反。图像处理时，x 对应列（width），y 对应行（height）。我见过有人把 x 和 y 写反了，结果图像旋转了 90 度。

一个小建议：刚开始写 CUDA 时，先用一维网格把逻辑跑通。性能优化是后面的事。先保证正确，再谈效率。

线程层次结构是 CUDA 编程的基石。Grid 定义全局范围，Block 定义协作范围，Thread 是实际干活的人。索引计算就是找到每个线程对应的数据位置。一维、二维、三维的选择，取决于你的数据形状。

嗯，这一章就到这里。记住我说的：画图、检查边界、别搞反坐标。这些习惯养成了，后面写复杂的核函数会顺手很多。

想要更多资料，联系可可