PMPP第三章 Multidimensional grids & data 精简笔记
第二章主要使用一维 grid 处理一维数组。本章扩展到多维 grid/block，用于处理图像、矩阵、三维数组等多维数据。核心目标是掌握：如何组织线程，如何把线程坐标映射到数据坐标，如何把多维数据线性化为一维内存访问。
3.1 多维 Grid 组织
核心概念
CUDA 线程采用两级层次结构：
Grid → Block → Thread
一个 kernel 启动后会生成一个 grid，grid 中包含多个 block，每个 block 中包含多个 thread。所有 thread 执行同一个 kernel，但通过自己的索引确定处理哪一部分数据。
CUDA 中常用的内置变量：
blockIdx // 当前 block 在 grid 中的坐标
threadIdx // 当前 thread 在 block 中的坐标
blockDim // 每个 block 的尺寸
gridDim // grid 的尺寸
这些变量是 CUDA C++ 的内置变量，不能改名。
dim3 与执行配置
CUDA 的 grid 和 block 最多都是三维的：
dim3 dimGrid(x, y, z);
dim3 dimBlock(x, y, z);
kernel<<<dimGrid, dimBlock>>>(...);
其中：
第一个参数：grid 维度，即 block 数量
第二个参数：block 维度，即每个 block 的 thread 数量
如果只用一维或二维，不使用的维度设为 1。
一维配置可以简写：
vecAddKernel<<<ceil(n / 256.0), 256>>>(...);
等价于：
dim3 dimGrid(ceil(n / 256.0), 1, 1);
dim3 dimBlock(256, 1, 1);
全局线程坐标公式
一维：
int i = blockIdx.x * blockDim.x + threadIdx.x;
二维：
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
三维：
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
注意事项
一个 block 中的线程数量为：
blockDim.x * blockDim.y * blockDim.z
当前 CUDA 设备中，一个 block 通常最多支持 1024 个线程。因此：
dim3 block(32, 16, 2); // 1024，合法
dim3 block(32, 32, 2); // 2048，非法
grid 和 block 的维度不要求一致。例如 grid 可以是二维，block 可以是三维。维度如何设置，主要取决于数据结构和映射方式。
还要注意：dim3 的参数顺序是 (x, y, z)，但描述多维数组时经常习惯写成 (z, y, x) 或 (row, col)，这两种顺序容易混淆。
3.2 将线程映射到多维数据
核心思想
多维线程组织的目的，是让线程坐标自然对应数据坐标。例如：
二维图像：一个 thread 处理一个 pixel
二维矩阵：一个 thread 处理一个 matrix element
三维数组：一个 thread 处理一个 tensor element
常见二维映射公式：
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
其中 row 和 col 就是当前线程负责的数据位置。
边界检查
grid/block 的尺寸通常按 block 大小向上取整，因此线程数可能多于数据元素数。例如图像是 62 × 76，block 是 16 × 16，实际启动线程覆盖范围是 64 × 80。
所以 kernel 内部必须判断：
if (col < width && row < height) {
// 处理有效元素
}
这和一维数组中的：
if (i < n)
作用相同，都是为了防止多余线程访问越界数据。
多维数组线性化
GPU 内存是线性地址空间，多维数组最终都需要映射为一维地址。
C/CUDA 使用 row-major layout，即行主序存储：
index = row * width + col;
含义是：先存完第 0 行，再存第 1 行，以此类推。
三维数组线性化公式：
index = z * height * width + y * width + x;
含义是：先定位到第 z 个平面，再定位到该平面的第 y 行，最后定位到第 x 列。
RGB 转灰度示例
彩色图像中，一个像素有三个连续通道：
R, G, B
每个线程负责一个输出灰度像素：
int grayOffset = row * width + col;
int rgbOffset = grayOffset * CHANNELS;
读取 RGB：
unsigned char r = Pin[rgbOffset];
unsigned char g = Pin[rgbOffset + 1];
unsigned char b = Pin[rgbOffset + 2];
写出灰度值：
Pout[grayOffset] = 0.299f * r + 0.587f * g + 0.114f * b;
注意事项
row-major 和 column-major 不同：
| 存储方式 | 含义 | 常见场景 |
|---|---|---|
| row-major | 按行连续存储 | C/CUDA C++ |
| column-major | 按列连续存储 | FORTRAN、部分数学库 |
CUDA C++ 默认按 row-major 理解数组。调用某些来自 FORTRAN 生态的数学库时，需要注意矩阵存储顺序。
3.3 图像模糊：更复杂的 kernel
核心思想
前面的例子通常是：
一个线程 → 读取少量输入 → 计算一个输出
图像模糊更复杂：
一个线程 → 读取周围多个输入像素 → 计算一个输出像素
这类模式接近后面会讲的 convolution pattern，即一个输出元素依赖一片输入区域。
图像模糊计算
本节使用简化版 blur：对目标像素周围的 patch 求平均值。
例如 3 × 3 patch：
(row-1, col-1) (row-1, col) (row-1, col+1)
(row, col-1) (row, col) (row, col+1)
(row+1, col-1) (row+1, col) (row+1, col+1)
输出：
out[row][col] = 周围有效像素的平均值
BLUR_SIZE 表示模糊半径：
patch size = (2 * BLUR_SIZE + 1) × (2 * BLUR_SIZE + 1)
例如：
BLUR_SIZE = 1 → 3 × 3
BLUR_SIZE = 3 → 7 × 7
Kernel 逻辑
每个线程仍然负责一个输出像素：
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
核心计算：
int pixVal = 0;
int pixels = 0;
for (int blurRow = -BLUR_SIZE; blurRow < BLUR_SIZE + 1; ++blurRow) {
for (int blurCol = -BLUR_SIZE; blurCol < BLUR_SIZE + 1; ++blurCol) {
int curRow = row + blurRow;
int curCol = col + blurCol;
if (curRow >= 0 && curRow < h && curCol >= 0 && curCol < w) {
pixVal += in[curRow * w + curCol];
pixels++;
}
}
}
out[row * w + col] = (unsigned char)((float)pixVal / pixels);
注意事项
外层边界检查：
if (col < w && row < h)
只能保证输出像素合法。
但 blur 会访问周围 patch，因此还需要内部边界检查：
if (curRow >= 0 && curRow < h && curCol >= 0 && curCol < w)
对于 3 × 3 patch：
| 像素位置 | 实际参与平均的像素数 |
|---|---|
| 内部像素 | 9 |
| 边缘非角落 | 6 |
| 四个角落 | 4 |
所以不能固定除以 9，而应该除以实际有效像素数 pixels。
3.4 Matrix Multiplication
基本定义
矩阵乘法是 BLAS 中的重要操作，属于 Level-3 BLAS，即矩阵-矩阵运算。
若：
M: i × j
N: j × k
则：
P = M × N
P: i × k
输出矩阵 P 中每个元素是 M 的一行和 N 的一列的内积：
P[row][col] = sum(M[row][k] * N[k][col])
CUDA 线程映射
本节采用最直接的策略：
一个 thread 计算 P 的一个元素
一个 block 计算 P 的一个 tile
线程坐标：
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
这里 row 和 col 直接对应输出矩阵 P[row][col]。
Kernel 逻辑
__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;
for (int k = 0; k < width; ++k) {
Pvalue += M[row * width + k] * N[k * width + col];
}
P[row * width + col] = Pvalue;
}
}
其中：
M[row * width + k]
表示 M 的第 row 行、第 k 列。
N[k * width + col]
表示 N 的第 k 行、第 col 列。
P[row * width + col]
表示输出矩阵 P 的第 row 行、第 col 列。
注意事项
这个版本是基础实现，重点是理解线程到矩阵元素的映射，不是高性能矩阵乘法。
它的问题是：
每个线程都直接从 global memory 读取 M 和 N；
不同线程之间会重复读取大量数据；
没有使用 shared memory 复用数据。
后续优化会围绕 tiling、shared memory、访存合并等展开。
3.5 总结
本章核心内容可以压缩成三点：
第一，CUDA 的 grid 和 block 都可以是多维的，最多支持三维。多维组织便于把线程映射到图像、矩阵、张量等多维数据。
第二，线程通过 blockIdx、threadIdx、blockDim 计算全局坐标，再用这个坐标确定自己负责的数据元素。
第三，多维数据在内存中通常是线性存储的，因此访问时需要把多维坐标转换为一维 offset。CUDA C++ 中最常见的是 row-major layout：
offset = row * width + col;
一句话总结：
本章讲的是：如何用多维 CUDA grid/block 组织线程，并把线程坐标映射到多维数据，再转换成一维内存地址完成访问。