3、线程层次结构：Grid、Block、Thread

好，咱们今天聊聊CUDA里最核心的一个概念——线程层次结构。说白了，就是GPU怎么组织它那成千上万个干活的小工（线程）。

我第一次接触CUDA时，最困惑的就是这个三层结构。CPU上我们写多线程，顶多就是线程池、任务队列。GPU倒好，上来就是Grid、Block、Thread三层。为什么要这么设计？

嗯，咱们一步步拆开看。

3.1 三个基本概念：Thread、Block、Grid

先给个直观的理解：

  • Thread（线程）：最基础的执行单元。每个线程执行同一个函数（叫kernel函数），但处理不同的数据。
  • Block（线程块）：一组线程的集合。同一个Block里的线程可以协作——它们能共享一块内存（Shared Memory），也能同步执行。
  • Grid（线程网格）：一组Block的集合。一个Grid对应一个kernel函数的启动。

你想想看，这就像一家公司：Thread是普通员工，Block是项目组，Grid是整个公司。项目组内部可以开会（共享内存）、同步进度（同步操作）。但不同项目组之间，基本各干各的，互不干扰。

关键点：同一个Block内的线程可以协作，不同Block的线程不能直接通信。这是硬件决定的，别想着绕过。

3.2 为什么需要三层结构？

我刚开始也觉得，直接搞个二维数组的线程不好吗？干嘛非要Grid包Block？

原因其实很实际：可扩展性。

GPU的硬件是分层的：

  • 一个SM（流多处理器）可以同时运行多个Block
  • 一个Block内的线程在同一个SM上执行
  • 不同Block可能被分配到不同的SM上

所以，Grid这个层级让你可以写出与硬件无关的代码。你只需要告诉GPU：我要处理100万个元素，分成1000个Block，每个Block里1024个线程。至于这1000个Block怎么分配到SM上，那是硬件调度器的事。

我在项目中遇到过一种情况：有人把Block数设得特别少，比如只有4个Block。结果GPU上几十个SM大部分都在闲着。性能惨不忍睹。后来改成几百个Block，立马跑满。

我的建议：Block数至少是SM数量的几倍，才能充分利用GPU。一般建议Block数在100-1000之间比较稳妥。

3.3 线程ID的计算

好，现在问题来了：每个线程怎么知道自己该处理哪部分数据？

CUDA提供了几个内置变量：

  • threadIdx：线程在Block内的索引（三维）
  • blockIdx：Block在Grid内的索引（三维）
  • blockDim：Block的维度大小（每个维度有多少线程）
  • gridDim：Grid的维度大小（每个维度有多少Block）

每个线程通过这些变量，就能算出自己的全局唯一ID。这个ID通常用来访问数组中的对应元素。

一维线程网格

最简单的情况。假设我们有一个长度为N的一维数组，想用一维Block和一维Grid来处理：

// 一维情况
int tid = blockIdx.x * blockDim.x + threadIdx.x;

这个公式很直观：先算出前面有多少个Block（blockIdx.x * blockDim.x），再加上当前Block内的偏移（threadIdx.x）。

举个例子：blockDim.x = 256，blockIdx.x = 2，threadIdx.x = 50。那么tid = 2 * 256 + 50 = 562。这个线程就处理数组的第562个元素。

二维线程网格

处理图像、矩阵这类二维数据时，用二维网格更自然。比如一张1024x768的图片：

// 二维情况
int col = blockIdx.x * blockDim.x + threadIdx.x;  // 列索引
int row = blockIdx.y * blockDim.y + threadIdx.y;  // 行索引

// 如果需要一维的全局ID（比如访问一维数组）
int tid = row * width + col;

我个人习惯在处理图像时用二维网格。这样代码可读性更好，col和row直接对应像素坐标，不用来回换算。

三维线程网格

三维网格用在体数据、3D卷积等场景。比如一个256x256x256的体数据：

// 三维情况
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;

// 一维全局ID
int tid = z * width * height + y * width + x;

注意：三维网格的Block总数不能太大。gridDim.x * gridDim.y * gridDim.z 不能超过65535（计算能力2.x及以下）或2^31-1（计算能力3.x及以上）。但实际中很少用到这么大的三维网格。

3.4 维度选择的实战经验

这里我分享一些实际项目中的经验：

数据类型 推荐维度 理由
一维数组（向量） 一维Grid + 一维Block 最简单，计算ID方便
二维数组（图像） 二维Grid + 二维Block 直接对应像素坐标，代码清晰
三维数组（体数据） 三维Grid + 三维Block 自然映射到体素位置
不规则数据 一维Grid + 一维Block 用一维ID做索引映射更灵活

我曾经在一个项目里处理稀疏矩阵，一开始用了二维网格，结果发现很多线程在空转。后来改成一维网格，每个线程处理多个元素，效率反而更高。所以维度选择不是死的，要看具体问题。

3.5 一个完整的例子

咱们写个简单的向量加法，把上面这些概念串起来：

// kernel函数
__global__ void vecAdd(float *A, float *B, float *C, int N) {
    // 计算全局线程ID
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 边界检查，防止越界
    if (tid < N) {
        C[tid] = A[tid] + B[tid];
    }
}

// 调用kernel
int N = 1000000;
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

注意那个边界检查 if (tid < N)。为什么需要？因为 blocksPerGrid * threadsPerBlock 可能大于N，多出来的线程什么都不做，直接返回。这是CUDA编程的常见模式。

小技巧：Block大小一般选32的倍数（比如128、256、512）。因为GPU的warp大小是32，这样能避免warp内线程浪费。我一般用256，不大不小，刚刚好。

3.6 避坑指南

最后说几个我踩过的坑：

  • Block内线程数上限：不同计算能力上限不同，一般是1024。别设成2048，会直接报错。
  • Grid维度上限：gridDim.x/y/z最大是2^31-1（新卡）或65535（老卡）。但实际用不到那么大。
  • 线程ID计算别搞混：我见过有人把blockIdx和threadIdx搞反，结果数据全乱套。记住：blockIdx是Block的索引，threadIdx是线程在Block内的索引。
  • 别忘了边界检查：数据量不是Block大小整数倍时，多出来的线程要处理。不加边界检查，轻则数据错误，重则段错误。

嗯，线程层次结构这块就这些。说白了就是三个层级、四个内置变量、一个ID计算公式。但就是这些基础的东西，决定了你的CUDA程序能不能跑得快、跑得稳。

下一章咱们聊聊内存层次结构——全局内存、共享内存、寄存器这些。到时候你会发现，线程怎么组织，和内存怎么用，是紧密相关的。

← 上一章 📖 返回目录 下一章 →