第4章 CUDA程序结构：从grid到thread

4.1 Kernel函数——这才是GPU的“入口”

很多初学者会把CUDA程序理解成「写一个函数，放到GPU上跑」。这个说法没错，但少了最关键的一步——你得明确告诉GPU：这个函数怎么切分任务？

kernel函数就是用 __global__ 修饰的C函数。我最早写的时候，总忘加这个修饰符，编译半天报错「找不到入口」。后来养成了习惯：写任何一个要在GPU上并行的函数，第一行先写 __global__。

一个典型的kernel长这样：

__global__ void harmonyCorrectionKernel(float* input, float* output, int width, int height) {
    // 每个线程独立处理一个像素
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int idy = threadIdx.y + blockIdx.y * blockDim.y;
    if (idx >= width || idy >= height) return;
    // 执行校正逻辑
    output[idy * width + idx] = applyCorrection(input[idy * width + idx]);
}

注意调用时的尖括号语法：

dim3 gridDim(16, 16);
dim3 blockDim(32, 32);
harmonyCorrectionKernel<<>>(d_input, d_output, width, height);

尖括号里的两个参数，就是本章的核心——grid和block。

4.2 Grid、Block、Thread——物理还是逻辑？

我见过新手把thread当成GPU硬件线程。其实不是。GPU硬件有SM、SP/Warp这些物理单位。而grid、block、thread是**编程模型**里的逻辑概念。说白了，是为了让你方便地组织并行任务。

为什么要有这三层？

想想看：一张图像可能有几百万像素。如果让所有线程直接平铺，你根本没法管理。分层之后：

  • Thread：最小的执行单元，对应一个数据元素。比如一个像素。
  • Block：一组线程，它们可以共享shared memory，也能同步。我常把block对应到图像的一个tile（图块）。
  • Grid：所有block的集合，覆盖整个数据域。比如整张图像。

说个经验：block内的线程数不要超过1024（计算能力3.0+的限制）。但更实际的是，block内线程数最好是32的倍数——因为warp大小是32。我曾经在一个项目里设了31线程/block，结果SM利用率掉了将近一半。这坑，你们别踩。

4.3 内置变量——threadIdx, blockIdx, blockDim, gridDim

这些是CUDA自动提供的变量，类型都是 dim3。不需要你声明，直接拿来用。

变量含义取值范围
threadIdx线程在block内的索引[0, blockDim - 1]
blockIdxblock在grid内的索引[0, gridDim - 1]
blockDimblock内每维的线程数由<<<>>>第二个参数设定
gridDimgrid内每维的block数由<<<>>>第一个参数设定
提示：这些变量在kernel内是只读的。别尝试修改它们——编译不会报错，但结果不可预测。

4.4 线程索引计算——把逻辑坐标映射到数据

有了这些内置变量，怎么算出每个线程负责的数据位置？

我个人习惯用一维映射处理一维数组（比如音频数据），用二维映射处理图像。二维情况最典型：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int linearIdx = y * width + x;  // 全局一维索引

为什么这样算？

把grid想象成一块大棋盘，每个block是棋盘上的格子，每个thread是格子里的点。你从grid左上角出发，先找到你在哪个格子（blockIdx），再找到格子里的位置（threadIdx），最后加上格子的偏移（blockIdx * blockDim.x）。

注意边界检查：实际数据大小可能不是block大小的整数倍。比如图像宽度为1013像素，你设blockDim.x=32，那么最后一个block会多出几个线程。不加边界检查，这些线程会访问越界内存，轻则读脏数据，重则cudaErrorUnknown。

边界检查的样板代码：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x >= width || y >= height) return;

嗯，一个return就搞定了。

4.5 一张图看懂层级关系

下面这张图是我自己经常给团队新人讲的：线程、block、grid构成三层树形结构，每个节点都是逻辑概念，但决定了你的并行粒度。

CUDA线程层次结构 Grid (<<>>) 逻辑上覆盖整个数据域 Block (0,0) blockIdx.x=0 blockIdx.y=0 Block (1,0) blockIdx.x=1 blockIdx.y=0 Thread (0,0) threadIdx.x=0 threadIdx.y=0 Thread (1,0) threadIdx.x=1 threadIdx.y=0 ... ... 关键对应关系： 1个Grid对应整张图像 → 1个Block对应一个tile → 1个Thread对应一个像素 索引公式: globalID = blockIdx * blockDim + threadIdx

4.6 多维配置——dim3的精妙

上面例子用的都是二维grid和二维block。实际上你可以用一维、二维、三维。什么时候用三维？

比如处理3D体数据（医学CT、地震波），或者卷积核需要跨通道时。我个人项目中很少用到三维block，但偶尔为了匹配数据维度会设定三维grid（z方向对应时间帧）。

创建dim3的方式：

dim3 block(8, 8, 1);   // block = (8,8,1) 等效于 block = dim3(8,8)
dim3 grid(128, 128);   // grid = (128,128,1)

没指定Z维时，默认是1。我记得初学那会儿，以为dim3会自动填充0，结果用了0维……程序直接崩溃。默认是1，不是0。

4.7 实践小结——一个完整的Harmony校正调用

把前面所有串起来，看一个更完整的例子。假设我要对一张1920×1080的图像做色差校正：

const int W = 1920;
const int H = 1080;
dim3 blockDim(32, 32);
dim3 gridDim((W + blockDim.x - 1) / blockDim.x,
             (H + blockDim.y - 1) / blockDim.y);

harmonyCorrectionKernel<<>>(d_in, d_out, W, H);

这里用 (W + blockDim.x - 1) / blockDim.x 向上取整，保证覆盖所有像素。多余线程在kernel内部被边界检查过滤掉。

一句话总结：grid管全局，block管辖区，thread管元素。索引公式是你的地图，边界检查是你的保险。
← 上一章 📖 返回目录 下一章 →