CUDA 线程层级:grid → block → thread → warp

你已经熟练 C/C++。好消息:CUDA C 的"语言"部分小得惊人 —— 真正要换的是心智模型。 普通 C 函数你从"一个执行流"的视角写;CUDA kernel 你从单个线程的视角写一段代码, 然后硬件把它复制成成千上万份同时跑。这一课就专攻这个视角切换。

1. 从 C 函数到 kernel:__global__ 与 <<< >>>

把一个普通函数变成在 GPU 上跑的 kernel,只需两件事:

• 在函数前加 __global__ —— 它就成了 kernel:device(GPU)执行、host(CPU)可调用、返回类型必须 void。¹
• 调用时用三尖括号执行配置:kernel<<<numBlocks, threadsPerBlock>>>(args)。第一个参数是 grid 里的 block 数,第二个是每个 block 的线程数。¹

普通调用 f(x) 跑一次;f<<<2,8>>>(x) 启动 2×8 = 16 个线程,每个都把同一段 kernel 代码跑一遍。

2. 三层层级:grid 装 block,block 装 thread

CUDA 把并行计算组织成三层。每次 kernel 启动创建恰好一个 grid;grid 由多个 block 组成(同一 grid 内所有 block 大小相同);每个 block 由若干 thread 组成。

3. 每个线程怎么知道"我是谁":四个内建变量

kernel 里有四个内建变量,每个线程都能读到,用来定位自己:

四个变量都有 .x / .y / .z 三个分量(grid/block 可 1D/2D/3D,用 dim3 指定²)。本课只用 .x。

4. 核心公式:全局索引

这是 CUDA 里最惯用的一行代码 —— 务必背下来:

直觉:blockIdx.x * blockDim.x 是"我这个 block 的起点偏移"(第几个 block × 每 block 多少线程),再加 threadIdx.x(我在 block 内的位置)。¹ 代入上图高亮线程:block 1、blockDim.x=8、threadIdx.x=3 → 1 × 8 + 3 = 11。

5. warp:32 个线程一组的硬件执行单位

block 是逻辑分组;硬件上,block 里的线程每 32 个为一组叫 warp。³ warp 是 SIMT(单指令多线程)的调度/执行单位:一个 warp 的线程同时执行同一条指令,lane 编号 0–31。 所以 block 线程数最好取 32 的倍数(如 256),否则最后一个 warp 有空 lane 被浪费。

这一点你优化 Open-Qwen2VL 时会反复用到:同一 warp 内连续、对齐的全局内存访问能被合并成一次事务(memory coalescing)—— 第 6 课展开。

动手:在 A100 上跑出来 🚀

下面这段程序我已经在你的 A100 上编译运行过(sm_80 / CUDA 11.4),真实输出贴在后面。轮到你复现:

1. 把代码存成 hello_threads.cu,scp 到远程机:
   # 在本机 $ scp hello_threads.cu user@<workshop-ip>:/tmp/
2. SSH 上去,用 nvcc 编译(-arch=sm_80 针对 A100 的 compute capability 8.0):
   $ ssh user@<workshop-ip> $ cd /tmp && /usr/local/cuda/bin/nvcc -arch=sm_80 hello_threads.cu -o hello_threads $ ./hello_threads

hello_threads.cu —— 每个线程报出身份并自算全局索引:

#include <cstdio>

// Kernel: 每个线程报出自己的 block/thread 坐标，
// 并用惯用公式 blockIdx.x * blockDim.x + threadIdx.x 算全局索引。
__global__ void report_identity(int n)
{
    int globalIndex = blockIdx.x * blockDim.x + threadIdx.x;
    if (globalIndex < n) {                       // 边界保护
        int lane   = threadIdx.x % warpSize;     // warpSize == 32 (A100)
        int warpId = globalIndex / warpSize;
        printf("block %d | thread-in-block %d | blockDim %d -> globalIndex %2d (warp %d, lane %2d)\n",
               blockIdx.x, threadIdx.x, blockDim.x, globalIndex, warpId, lane);
    }
}

int main()
{
    const int numBlocks       = 2;
    const int threadsPerBlock = 8;
    const int n = numBlocks * threadsPerBlock;   // 16 个工作项

    report_identity<<<numBlocks, threadsPerBlock>>>(n);   // 启动 grid

    // kernel 启动是异步的：必须同步，否则程序退出时
    // device 端 printf 可能还没刷回 host。
    cudaError_t err = cudaDeviceSynchronize();
    if (err != cudaSuccess) {
        printf("CUDA error: %s\n", cudaGetErrorString(err));
        return 1;
    }
    return 0;
}

在 A100 上的真实输出(我跑出来的):

检验:凭记忆作答(别翻上文)

回忆比重读更能形成长期记忆。先盖住上面,凭脑子答完再看解析。

常见误区(核验时专门标注)

📖 首选精读(primary source)

引用与延伸

1. NVIDIA, An Even Easier Introduction to CUDA — __global__、执行配置、全局索引公式。
2. NVIDIA, CUDA Refresher: The CUDA Programming Model — 三层层级、1024 上限、3D 内建变量、block 独立执行。
3. Simon Boehm, How to Optimize a CUDA Matmul Kernel — warp=32 与 SIMT、Volta+ 独立线程调度、coalescing。