索引与边界

写 kernel 的第一件事永远是:

int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid >= n) return;

第一行算自己处理哪个数据,第二行防越界。掌握这两行的所有变体,kernel 就基本不会写错入口。

1. 1D 索引模板

最常见的情况——数组、向量。

__global__ void process(float* data, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid >= n) return;
    data[tid] = ...;
}

启动时:

int block = 256;
int grid  = (n + block - 1) / block;     // 向上取整
process<<<grid, block>>>(data, n);

(n + block - 1) / block 是向上取整的标准写法。等价于 ceil(n / block),但只用整数运算。

2. tail guard 为什么必要

n = 1000,block = 256:

grid = (1000 + 255) / 256 = 4 个 block
启动线程总数 = 4 × 256 = 1024
处理元素数  = 1000
多出来的    = 24 个线程,tid = 1000..1023

这 24 个”多余”线程必须用 tid >= n 拦住——否则 data[1023] 越界,要么报错要么破坏相邻内存。

个别情况下,如果你确保 n 是 block 的整数倍(在 host 端 padding 过),可以不写 tail guard。但默认写,养成习惯。

3. 2D 索引(矩阵、图像)

__global__ void process_matrix(float* M, int rows, int cols) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    if (row >= rows || col >= cols) return;
    M[row * cols + col] = ...;       // 行主序
}

启动:

dim3 block(16, 16);
dim3 grid((cols + 15) / 16, (rows + 15) / 16);
process_matrix<<<grid, block>>>(M, rows, cols);

注意 col 用 x 维度,row 用 y 维度——这是 CUDA 约定。原因:threadIdx.x 变化最快(相邻 thread 差 1),希望它对应内存里最连续的维度(行主序矩阵的列方向),这样能合并访问。

4. 3D 索引(体数据、3D 卷积)

__global__ void process_volume(float* V, int D, int H, int W) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;   // W 方向
    int y = blockIdx.y * blockDim.y + threadIdx.y;   // H 方向
    int z = blockIdx.z * blockDim.z + threadIdx.z;   // D 方向
    if (x >= W || y >= H || z >= D) return;
    V[z * H * W + y * W + x] = ...;
}

启动:

dim3 block(8, 8, 8);
dim3 grid((W + 7) / 8, (H + 7) / 8, (D + 7) / 8);

注意:3D block 总线程数仍受 1024 上限约束(8*8*8 = 512,OK;16*16*16 = 4096 会编译失败)。

5. grid-stride loop:每个 thread 处理多个元素

如果你不想让 grid 跟着 n 缩放(比如想固定用某个 grid 大小),用 grid-stride loop:

__global__ void process_strided(float* data, int n) {
    int tid    = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = gridDim.x * blockDim.x;     // 总线程数
    
    for (int i = tid; i < n; i += stride) {
        data[i] = ...;                       // 自动跳到下一份工作
    }
}

每个 thread 处理 n / stride 个元素。好处:

	普通模板	grid-stride loop
数据量增大	grid 跟着变,可能爆 grid 上限	grid 固定,for 多跑几轮
启动开销	跟 n 成正比	跟选的 grid 大小成正比
适合	n 不太大	n 巨大或不确定

默认建议:n < 几百万 用普通模板,n 上亿或不确定用 grid-stride。

6. 常见索引错误

6.1 忘记 tail guard

// 错:n = 1000, blockDim = 256,起 4 个 block,最后 24 线程越界
int tid = blockIdx.x * blockDim.x + threadIdx.x;
data[tid] = ...;

报错:an illegal memory access was encountered,或者写坏邻近数据(更糟,默默错)。

6.2 row/col 写反

// 错:把 y 当 x 维度
int row = blockIdx.x * blockDim.x + threadIdx.x;     // ← 应该是 col 用 x
int col = blockIdx.y * blockDim.y + threadIdx.y;
M[row * cols + col] = ...;

代码能跑,但 warp 内 32 个 lane 跨 32 行访问(因为 threadIdx.x 变化最快,但被用作 row),完全打破合并访问。性能慢几十倍,功能正确,不报错——所以是最阴险的错。

6.3 grid 算错(忘了向上取整)

// 错:n = 1000, block = 256, grid = 1000/256 = 3,只处理 768 个元素,剩下 232 个不动
int grid = n / block;

// 对
int grid = (n + block - 1) / block;

7. 一句话

每个 kernel 的第一件事都是算 tid = blockIdx.x * blockDim.x + threadIdx.x 加 tail guard。2D 时 col 用 x、row 用 y(因为 threadIdx.x 变化最快,对应行主序的列方向 → 合并访问)。grid 大小用 (n + b - 1) / b 向上取整,不向上取整就丢数据。