写 CUDA 的人，遇到的第一个真正意义上的门槛，往往不是语法，而是空间感。
CPU 上写循环写惯了，你的脑子里是一条线：for (int i = 0; i < n; i++)，从头到尾，一件事接一件事。而 CUDA 要求你换一副眼镜：你要把一整个数组，一次性地摊开在几千、几万个执行单元面前，然后回答一个问题——
第 i 号线程，负责哪一块数据？
这个问题回答不清楚，后面所有的优化都无从谈起。而《Programming Massively Parallel Processors》（PMPP）第五版的第三章，整整一章就是在回答它。
1. 两级结构：网格（grid）与线程块（block）
CUDA 的线程组织是一个两级结构：
- 一个 网格（grid） 由若干个 线程块（block） 组成；
- 一个 线程块 又由若干个 线程（thread） 组成；
- 同一个网格里的所有线程，执行的是同一个 kernel 函数——它们靠坐标区分彼此，而不是靠代码路径。
每个线程身上带了两组坐标：
| 内建变量 | 含义 |
|---|---|
threadIdx | 线程在自己块内的坐标 |
blockIdx | 这个块在网格中的坐标 |
blockDim | 每个块有多少线程 |
gridDim | 网格中有多少个块 |
它们在 kernel 内部是只读的，由你启动时的执行配置参数决定。
注意这个结构的含义：块是调度的单位，线程是执行的单位。后面第二章（架构篇）我们会看到，GPU 正是以"块"为粒度把它们分派给不同的 SM 去跑。这也是为什么"跨块同步"在 CUDA 里是一条铁律上的禁忌——这一点留到下一篇细讲。
2. dim3：一个"默认值为 1"的甜头，和它的陷阱
CUDA 用一个叫 dim3 的类型来描述网格和块的形状。它本质上是三个整数 x, y, z。
dim3 dimBlock(16, 16); // 每个块 16×16 = 256 个线程
dim3 dimGrid(ceil(W/16.0), // x 方向需要多少个块
ceil(H/16.0)); // y 方向需要多少个块
matmulKernel<<<dimGrid, dimBlock>>>(M, N, P, W);
<<<...>>> 里的两个参数，就是执行配置参数：第一个是网格有多少个块，第二个是每个块有多少个线程。
CUDA 在这里埋了一个很舒服的简化：dim3 没写的分量默认是 1。所以一维数组的处理可以写成极简的样子：
vecAddKernel<<<ceil(n / 256.0), 256>>>(A, B, C, n);
这行代码等价于"网格是 ceil(n/256) 个块，每个块 256 个线程，y 和 z 方向都是 1"。
甜头背后是陷阱。因为单个数值会被塞进 x，如果你的本意是二维形状却只传了一个数，编译器不会报错——你会得到一个"一维网格"，然后 kernel 里的 blockIdx.y 永远是 0。这种错误不会崩溃，只会算错。
另外两条边界值得记一下：gridDim.x 的取值范围是 1 到 $2^{31}-1$，而单个块的线程数有硬件上限（目前通常是 1024）——不是你想开多大就开多大。
3. 线程到数据的映射：两行代码的全部秘密
二维数据（图像、矩阵）最自然的映射是这样两行：
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
它的逻辑朴素得不能再朴素：先算出我所在的块"起点"在哪，再加上我在块内的偏移。
这说明了一件重要的工程习惯：写 kernel 时，先选块的形状，再倒推网格的形状。块的形状决定了线程在块内如何排列，直接影响到访存效率（threadIdx.x 是变化最快的维度，相邻线程应当访问相邻内存——这是第六章 coalescing 的主题，先埋个伏笔）。
4. 行主序：线性化是你自己的事
这里有一个 C/C++ 的硬事实：动态分配的多维数组，在内存里其实是一维的。所谓 M[row][col]，编译器帮你做的运算不过是：
这就是行主序（row-major）。而在 CUDA 里，一旦你用 cudaMalloc 申请一块显存，编译器就再也不知道"宽度"是多少了——它只看到一个扁平的指针。
所以，线性化这件事，责任在程序员身上。
我们拿书上的算例走一遍。假设图像宽 62、高 76，块取 16×16：
$$ \text{grid} = \left(\lceil 62/16 \rceil,\ \lceil 76/16 \rceil\right) = (4, 5) $$某个线程的 blockIdx = (1, 2)、threadIdx = (3, 5)，那么：
它在内存中的扁平偏移就是：
$$ 37 \times 62 + 19 = 2313 $$请务必亲手推导一次这个算术。它的意义不在于数字本身，而在于：如果这里想错了，你的程序会在方方正正、能整除的测试用例上完美运行，然后在别的地方悄无声息地出错。
5. 由浅入深的三个例子
PMPP 这一章用三个例子把上面的机制串了起来，复杂度递增：
| 例子 | 线程负责什么 | 数据映射 |
|---|---|---|
| 彩色图转灰度 | 一个线程算一个像素 | 二维网格，每个线程读一组 RGB |
| 图像模糊 | 一个线程算一个输出像素 | 二维网格，读周边一个邻域 |
| 矩阵乘法 | 一个线程算一个输出元素 | 二维网格，沿共享维度循环 |
前两个是"点对点"和"点对邻域"，第三个才是真正意义上的计算模式。它就是全书后续反复出现的那个基线 kernel：
#define BLOCK_SIZE 16
__global__ void matrixMulKernel(float* M, float* N, float* P, int Width) {
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row < Width && col < Width) {
float Pvalue = 0.0f; // 在寄存器里累加
for (int k = 0; k < Width; ++k)
Pvalue += M[row * Width + k] * N[k * Width + col];
P[row * Width + col] = Pvalue;
}
}
代码本身毫无花招，但它藏着一个巨大的性能问题，就在那个 for 循环里。
要算出一个输出元素，需要 $N$ 次乘加；而每一次乘加，都要从全局内存里取两个数——M 的一个元素和 N 的一个元素。于是每个输出元素要读 $2N$ 个数据。整个矩阵有 $N^2$ 个输出元素，所以这一份代码在全局内存上开出了：
次读取。这个数字就是第五章的靶子。PMPP 把它作为"基线"摆在这里，并不是让你去用，而是让后面所有优化都有一个可以对照的起点：先用最直白的写法保证正确，再去测量、再去优化。
6. 边界检查：别小看那个 if
上面代码里有一行 if (row < Width && col < Width)，初学者常常觉得多余。它为什么必须在？
因为网格尺寸是向上取整算出来的。宽 62、块 16，需要 4 个块覆盖 64 个位置——多出来那 2 个线程是"凑数"的。如果放任它们去读写，轻则算错，重则非法访存。
所以规则很明确：
- 计算网格用
ceil，永远不要假设网格能整除数据规模； - kernel 内部必须对每一条数据访问做边界保护；
- 如果是后面要讲的分块（tiled）版本，越界的元素通常要填 0而不是跳过，否则会污染累加结果。
7. 小结
这一章没有讲任何性能技巧，它在建立坐标直觉。记住四件事：
- CUDA 的线程是两级结构：网格 → 块 → 线程，所有线程执行同一个 kernel。
dim3未写的分量默认为 1，这带来简洁，也带来"一维/二维写错却不报错"的隐患。- 线程到数据的映射，核心就是
blockIdx * blockDim + threadIdx；先定块形状，再倒推网格形状。 - 行主序线性化
row * Width + col必须由程序员亲手推导，边界检查不是可选项。
以及一个已经被埋下的悬念：那个看起来人畜无害的矩阵乘法基线，会在第五章被拆开重写。它的 2N^3 次全局读取，正是"内存墙"送给每一位 CUDA 程序员的见面礼。
下一篇，我们向下钻一层：GPU 到底是个什么样的车间？你启动的那些线程块，究竟被谁、以什么顺序、放到了哪里去跑。
参考链接
Programming Massively Parallel Processors: A Hands-on Approach (5th Edition) — Wen-mei W. Hwu, David B. Kirk, Izzat El Hajj. 第 3 章 Multidimensional grids and data.