4. CUDA线程组织与索引计算

CUDA 的并行执行模型把一次 kernel 启动展开成三层结构：grid（网格）→ block（线程块）→ thread（线程）。理解这三层以及它们与数据下标之间的映射关系，是写出正确且高效 kernel 的第一步。很多初学者写出的 kernel 结果错误或越界访问，根源往往不是算法本身，而是线程索引算错、边界没检查、或者多维数据映射到线程的方式与内存布局不匹配。

一、grid 与 block 的维度设计

启动 kernel 时使用 <<<gridDim, blockDim>>> 语法，其中 gridDim 和 blockDim 都可以是 1D、2D 或 3D。它们既可以用整数表示（等价于 1D），也可以用 dim3 结构体表示多维。例如处理一张 1920×1080 的灰度图像，可以这样组织：

dim3 block(16, 16);                       // 每个 block 16×16 = 256 线程
dim3 grid((1920 + 15) / 16, (1080 + 15) / 16); // 120 × 68 个 block
kernel<<<grid, block>>>(d_img, width, height);

这里 block 取 16×16 是常见选择：256 个线程正好是 warp（32 线程）的整数倍，且每个 block 覆盖一个 16×16 的二维数据块，便于共享内存分块和缓存局部性。grid 维度用向上取整公式 (N + block - 1) / block 计算，保证即使 N 不是 block 的整数倍也能覆盖全部数据——代价是最后一个 block 中会有部分线程越界，必须靠边界检查剔除。

block 大小的选择需要权衡：太小则 block 数量多、调度开销上升；太大则占用寄存器/共享内存多、可能降低 occupancy。经验上 128～512 线程/block 是常用区间，且尽量取 32 的整数倍以避免 warp 内线程浪费。grid 维度上限通常很大（x 方向 2³¹−1），一般不会成为瓶颈。

二、全局线程索引的计算

kernel 内部通过内置变量获取线程位置：threadIdx（block 内坐标）、blockIdx（grid 内坐标）、blockDim（block 尺寸）、gridDim（grid 尺寸）。一维情况下全局索引非常直接：

int i = blockIdx.x * blockDim.x + threadIdx.x;

二维情况下，需要把二维坐标展平成一维线性下标。这里有一个关键约定：行优先（row-major），即 x 方向是连续维度，y 方向是步长维度。对于宽度为 width 的二维数据，全局线性索引为：

int x = blockIdx.x * blockDim.x + threadIdx.x;  // 列
int y = blockIdx.y * blockDim.y + threadIdx.y;  // 行
int idx = y * width + x;                        // 展平后的线性下标

这个 y * width + x 的公式必须与主机端数据的存储顺序一致。如果主机端用 img[y][x] 或一维数组按行优先存放，那么 kernel 中就必须用同样的公式；若写成 x * height + y，访问的将是转置后的位置，结果完全错乱。三维情况同理，在 z 方向再乘一层：idx = (z * height + y) * width + x。

三、边界检查

由于 grid 维度向上取整，实际线程总数往往大于数据元素总数。例如 1920 能被 16 整除，但 1080 不能被 16 整除（1080/16 = 67.5），grid.y 取 68，最后一个 block 的 y 方向有 8 行线程是多余的。如果不加检查，这些线程会读写越界地址，导致段错误或静默破坏其他数据。标准做法是在计算完索引后立即判断：

if (x < width && y < height) {
    int idx = y * width + x;
    out[idx] = in[idx] * 2.0f;
}

边界检查应尽量放在最前面，让越界线程尽早退出，避免它们参与后续的共享内存写入或同步操作——否则 __syncthreads() 可能因部分线程提前返回而死锁。如果 kernel 中使用了共享内存分块，通常的做法是：越界线程仍然参与共享内存的加载（用 0 填充），但在写回全局内存前才做边界判断。

四、多维数据映射到线程的完整示例

下面给出一个把 2D 图像做 3×3 均值滤波的 kernel 骨架，展示维度设计、索引计算、边界检查与多维映射如何协同：

__global__ void blur3x3(const float* in, float* out, int w, int h) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x >= w || y >= h) return;          // 边界检查
    float sum = 0.0f;
    for (int dy = -1; dy <= 1; ++dy)
        for (int dx = -1; dx <= 1; ++dx) {
            int nx = x + dx, ny = y + dy;
            if (nx >= 0 && nx < w && ny >= 0 && ny < h)
                sum += in[ny * w + nx];     // 行优先线性索引
        }
    out[y * w + x] = sum / 9.0f;
}

注意内层循环里对邻域也做了边界判断，这是图像处理中常见的“halo”处理。若改用共享内存分块，可以把 16×16 的 block 扩展为 18×18 的共享 tile，先协作加载（越界补 0），再在共享内存上做滤波，能显著减少全局内存访问次数。

把上面这张图记牢，就能在写任何 kernel 时快速定位索引问题：先确认 grid/block 维度是否覆盖全部数据，再确认展平公式是否与主机端布局一致，最后确认边界检查是否放在正确位置。这三步做完，绝大多数“结果不对”和“非法访问”的问题都能提前避免。