当你运行一个 CUDA 内核时会发生什么？

来源：https://fergusfinn.com/blog/what-happens-when-you-run-a-gpu-kernel/

下面是一个简单的 CUDA 程序。它将两个向量相加。

__global__ void vadd(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}

int main() {
    int n = 1 << 20; // 一百万个浮点数 (1,048,576)
    size_t bytes = n * sizeof(float);
    float *a = (float*)malloc(bytes), *b = (float*)malloc(bytes), *c = (float*)malloc(bytes);
    for (int i = 0; i < n; i++) a[i] = b[i] = 1.0f;

    float *da, *db, *dc;
    cudaMalloc(&da, bytes);
    cudaMalloc(&db, bytes);
    cudaMalloc(&dc, bytes);
    cudaMemcpy(da, a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(db, b, bytes, cudaMemcpyHostToDevice);
    vadd<<<4096, 256>>>(da, db, dc, n); // 4096 * 256 = n 个线程，每个浮点数对应一个线程
    cudaMemcpy(c, dc, bytes, cudaMemcpyDeviceToHost);
    printf("c[0]=%f c[n-1]=%f\n", c[0], c[n-1]);
}

为 RTX 4090 编译并启动后，它确实正确计算出了 1+1=2，重复了一百万次（我没有全部检查）。

$ nvcc -arch=sm_89 -o vadd vadd.cu && ./vadd
c[0]=2.000000 c[n-1]=2.000000

告诉你这个过程涉及数千万条 CPU 指令、几个设备文件、九百多次 ioctl，以及一个内存映射的门铃寄存器。在这篇文章中，我们将跟随这一个内核，从代码一直追踪到 warp，再返回到结果。

顺便提一句，这篇文章是智能体引发的“可读性转变”的一个实例。凭借好奇心和（机器增强的）毅力，关于计算机你几乎没有什么不能发现的。关于可读性对 AI 帮助我们认知的影响，这里有一个有趣的讨论：https://resobscura.substack.com/p/ai-legibility-archives-future-of-research。

用 nvcc 编译我们的程序

§ (https://fergusfinn.com/blog/what-happens-when-you-run-a-gpu-kernel/#compiling-our-program-with-nvcc)

我们首先要做的事是如何把这个 CUDA 程序转换成设备能真正读取的东西。为此我们需要一个编译器。实际上，我们需要多个编译器。nvcc 是一个驱动程序，它运行几个其他编译器并把它们的输出合并。如果传入 --keep，它会将整个流水线保留在磁盘上供你阅读：

$ nvcc --keep -arch=sm_89 -o vadd vadd.cu && ls
...
vadd.ptx          # 设备代码的 PTX 格式（来自 cicc）
vadd.sm_89.cubin  # 设备代码的 SASS 格式（来自 ptxas）
vadd.fatbin       # cubin + PTX 打包在一起（来自 fatbinary）
vadd.cudafe1.stub.c  # 主机启动桩 + 内核注册
vadd.o            # 最终主机目标文件，其中嵌入了 fatbin
...

主机代码会交给你的主机编译器。设备代码（vadd）需要更多步骤：基于 LLVM 的编译器 cicc 先将其转换成 PTX，然后 ptxas 将 PTX 转换成 SASS。PTX 是一种虚拟指令集架构，拥有无限多个类型化寄存器，并且不关心硬件实际有多少寄存器。下面是我们 vadd 的 PTX 代码（做了删减）：

$ cat vadd.ptx
...
mad.lo.s32 %r1, %r3, %r4, %r5;  // 将寄存器 r1 设为 ctaid*ntid + tid
setp.ge.s32 %p1, %r1, %r2;      // 如果 i >= n 则设置谓词 p1
@%p1 bra $L__BB0_2;              // 如果越界，跳转到退出
cvta.to.global.u64 %rd4, %rd1;   // 将通用指针 %rd1 转换为全局地址，存入 %rd4
mul.wide.s32 %rd5, %r1, 4;       // r1 乘以 4，结果存入 %rd5
add.s64 %rd6, %rd4, %rd5;        // %rd4 + %rd5，结果存入 %rd6
ld.global.f32 %f2, [%rd6];       // 加载 a[i] 到 %f2
...
add.f32 %f3, %f2, %f1;           // %f1 + %f2，结果存入 %f3
st.global.f32 [%rd10], %f3;      // 将 c[i] 存入全局内存

虚拟寄存器看起来像 %rd1–%rd10、%f1–%f3。前缀表示类型：%r 是 32 位整数，%rd 是 64 位整数，%f 是 32 位浮点，%p 是 1 位谓词。PTX 比你想象的要“冗长”得多。例如，形成地址 %rd6 需要三条 PTX 指令。这是因为 PTX 是与设备无关的。为什么要三条？CUDA 指针默认是“通用”的，意味着它们可能指向全局、共享或局部内存。cvta.to.global 断言该指针位于全局窗口，这样后面就可以使用更便宜的 ld.global。mul.wide.s32 将索引 i 乘以 4（sizeof(float)）并一步将 32 位扩展到 64 位，得到字节偏移量。add.s64 将其加到基址指针上。

接下来，ptxas 将我们的 PTX——与设备无关——转换成针对你架构的 SASS，这是与设备相关的。它生成的 SASS 看起来不一样：

$ cuobjdump -sass vadd
/*0000*/ MOV R1, c[0x0][0x28] ;           // 设置栈指针（ABI 要求；此处未使用）
/*0010*/ S2R R6, SR_CTAID.X ;            // R6 = blockIdx.x
/*0020*/ S2R R3, SR_TID.X ;              // R3 = threadIdx.x
/*0030*/ IMAD R6, R6, c[0x0][0x0], R3 ;  // i = ctaid*ntid + tid
/*0040*/ ISETP.GE.AND P0, PT, R6, c[0x0][0x178], PT ; // P0 = (i >= n)
/*0050*/ @P0 EXIT ;                      // 如果越界则退出
/*0060*/ MOV R7, 0x4 ;                   // 加载立即数 4 (sizeof(float)) 作为乘数到 R7
/*0070*/ ULDC.64 UR4, c[0x0][0x118] ;    // 统一加载驱动程序提供的系统值
/*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ; // &b[i]
/*0090*/ IMAD.WIDE R2, R6, R7, c[0x0][0x160] ; // &a[i]
/*00a0*/ LDG.E R4, [R4.64] ;            // b[i]
/*00b0*/ LDG.E R3, [R2.64] ;            // a[i]
/*00c0*/ IMAD.WIDE R6, R6, R7, c[0x0][0x170] ; // &c[i]
/*00d0*/ FADD R9, R4, R3 ;              // a[i] + b[i]
/*00e0*/ STG.E [R6.64], R9 ;            // c[i] = ...
/*00f0*/ EXIT ;

S2R 指令在做什么？S2R 表示“特殊寄存器到寄存器”——它将硬件为每个线程维护的一个特殊寄存器（这里是 SR_CTAID.X（块的索引，blockIdx.x）和 SR_TID.X（线程在块内的索引，threadIdx.x））复制到普通寄存器，以便 IMAD 能对其进行算术运算。

十多个虚拟寄存器已经压缩成了七个真实寄存器。ncu 报告 launch__registers_per_thread = 16。反汇编中只出现了 R0 到 R9，但分配器为 ABI 和对齐预留了更多。

两条 mul.wide 加 add 的序列已经融合成了单一的 IMAD.WIDE。cvta 转换消失了，被合并到寻址模式中。c[0x0][...] 操作数位于常量 bank 0，这是一个由驱动程序管理的小区域。这些是内核的参数——指针 a、b、c 以及大小 n——以及启动几何参数。填充这个 bank 的是在启动时由驱动程序交给 GPU 的一个叫做 QMD 的结构，我们会在启动本身到达显卡时再讨论它。

为什么参数位于常量 bank 0 以及它们在哪儿？因为它们位于常量内存中，因为这是一种广播读取：网格中的每个线程都需要相同的指针，常量缓存能够在一次操作中服务于所有 32 条线程。布局是固定的——0x160、0x168、0x170 分别是指针 a、b、c，0x178 是 n，启动几何布局在它们旁边的 0x0（blockDim.x）。Bank 0 还保存 ABI 参数，例如 c[0x0][0x28]（栈基址），MOV R1, c[0x0][0x28] 在入口处加载它。当主机桩打包启动参数时，我们会再次看到相同的偏移量。

包含这些 SASS 的 cubin 文件是一个 ELF 文件——与 Linux 用于普通可执行文件和共享库相同的目标文件容器。cuobjdump -elf 会显示一个符号表、一个包含机器代码的 .text.vadd 段，以及特定于 CUDA 的段如 .nv.callgraph。

fatbinary 可执行文件将 cubin 与 PTX 打包成一个单独的“fatbin”，在结果上运行 cuobjdump 会显示嵌入在我们的二进制文件中的 fatbin 同时包含两者：

$ cuobjdump vadd
...
Fatbin elf code:
    arch = sm_89  # 我们刚才看到的 SASS
Fatbin ptx code:
    arch = sm_89
    compressed     # PTX 也一并携带

SASS 是实际在这块 4090 上运行的代码，但 PTX 作为向前兼容的备选方案随行。如果你后来把这个二进制文件拿到一个 cubin 不涵盖的 GPU 架构上，驱动程序可以在加载时即时将 PTX 编译成新的 SASS。

最后，那个 fatbin 嵌套在主机可执行文件中，readelf -S 发现它占用自己的段：

$ readelf -S vadd
...
[18] .nv_fatbin     PROGBITS ...
[19] __nv_module_id PROGBITS ...
[29] .nvFatBinSegment PROGBITS ...
...

nvcc 输出的 vadd 二进制文件是一个单一的可执行文件，包含主机代码、一个包含 Ada SASS 的完整 ELF 目标文件，以及一份 PTX 副本。由于 PTX 是冗长的纯文本，nvcc 默认会压缩它以保持二进制文件大小；只有当二进制文件在预编译的 SASS 不覆盖的架构上运行时，驱动程序才会解压缩并即时编译它。

主机如何触发 GPU

§ (https://fergusfinn.com/blog/what-happens-when-you-run-a-gpu-kernel/#how-the-host-triggers-the-gpu)

编译好的 GPU 机器代码现在静静地躺在 ./vadd 可执行文件的 .nv_fatbin 段中。当你在主机上运行该程序时，我们需要连接两个世界：主机 CPU 和位于 PCIe 总线另一端的 GPU。为了设置一个知道如何跨越这座桥的主机二进制文件，前端编译器（cudafe++）会在你的代码中插入一个隐藏的构造函数，它在 main 函数开始之前运行。它的工作是注册我们嵌入的 fatbinary 到 CUDA 运行时，并记录一个运行时稍后会用到的映射：将主机端的函数指针 vadd 与 fatbin 中编译好的设备内核的修饰名称关联起来。

当编译器遇到 vadd<<<4096, 256>>>(da, db, dc, n) 时，它会用生成的主机启动桩替换这个高级表达式。这个桩将我们的内核参数打包到主机内存的一个缓冲中。指针 da、db、dc 和整数 n 分别对齐到字节偏移 0、8、16 和 24。这些偏移量就是我们之前看到 SASS 机器代码从常量 bank 0 读取的常量 bank 偏移量 0x160、0x168、0x170 和 0x178。

// from vadd.cudafe1.stub.c
void __device_stub__Z4vaddPKfS0_Pfi(const float *__par0, const float *__par1, float *__par2, int __par3) {
    __cudaLaunchPrologue(4);
    __cudaSetupArgSimple(__par0, 0UL);  // 参数缓冲偏移 0
    __cudaSetupArgSimple(__par1, 8UL);  // 偏移 8
    __cudaSetupArgSimple(__par2, 16UL); // 偏移 16
    __cudaSetupArgSimple(__par3, 24UL); // 偏移 24
    __cudaLaunch((char*)(void(*)(const float*, const float*, float*, int))vadd);
}

参数打包后，桩调用 __cudaLaunch，传入主机端占位函数 vadd 的内存地址。因为此主机函数在 CPU 上只是一个空壳，其主机内存地址充当查找键。运行时用这个地址查询它的注册表，找到对应的设备端符号名，然后跨越边界进入闭源用户态驱动程序（libcuda.so.1）。用户态驱动程序是与 GPU 的内核驱动程序一起提供的，而不是 CUDA 工具包：strace 中看到的 libcuda.so.1 解析为 libcuda.so.590.48.01，即此机器上的驱动程序版本。运行时在程序第一次调用 GPU 时动态打开这个驱动程序，我们可以用 strace 捕获到：

$ strace -f -e trace=openat ./vadd
...
openat(..., "/lib/x86_64-linux-gnu/libcuda.so.1", O_RDONLY|O_CLOEXEC) = 3
...

当这第一次调用发生时，会创建一个“上下文”，包含驱动程序与设备通信所需的所有基础设施，包括 CPU 与 GPU 对话的通道。我们将在下一节进一步讨论这一点。

在这个阶段，编译好的机器代码仍未到达 GPU。从 CUDA 12.2 开始，模块加载默认是惰性的（通过 CUDA_MODULE_LOADING 控制；它在 CUDA 11.7 中作为可选功能引入，多年来默认是 EAGER；12.x 系列将默认值改为 LAZY，如果希望一次性支付加载成本，可以覆盖）。——驱动程序会推迟将内核的 SASS cubin 上传到显卡内存，直到该特定内核首次实际启动。

在 libcuda 之下是内核态驱动程序 nvidia.ko，libcuda 通过对设备文件发起 ioctl 调用来与之通信。当 cuLaunchKernel 最终需要将工作放在 GPU 上时，它就变成了与内核模块的一场对话。接下来是这场对话的机制。

将代码加载到 GPU

§ (https://fergusfinn.com/blog/what-happens-when-you-run-a-gpu-kernel/#getting-it-onto-the-gpu)

GPU 不像 CPU 那样进行函数调用。它没有可跳转的入口点，也没有可以从 CPU 压入参数的栈。GPU 位于 PCIe 总线的另一端，从主机内存中读取一串驱动程序命令。cuLaunchKernel 在此之后所做的每一件事都是为了将一个完整的启动命令放入那个命令流中，然后告诉 GPU 它已经这样做了。

首先需要做的是将 GPU 代码加载到设备上。第一次运行 vadd 时，驱动程序将内核的代码复制过去：它分配一个缓冲并将 SASS 复制进去。一旦代码在 GPU 上，CPU 需要让 GPU 读取它并开始执行。这通过一个跨越主机和设备内存的复杂舞蹈来完成。主机和 GPU 都可以映射对方内存空间的区域，但跨越 PCIe 总线的访问会付出代价。为了实现内核启动，两者都向位于两个空间的各种结构中写入内容。这些结构构成了通道——运行 GPU 操作的工作队列。在主机 RAM 中有两个重要的结构：pushbuffer 和 GPFIFO，它们共同代表 GPU 必须执行的工作列表。

pushbuffer 是一块内存区域，驱动程序将命令写入其中，称为methods。一个 method 是 GPU 本机命令编码中的一个寄存器地址和一个值——这一对定义了 GPU 应该执行什么动作。

GPFIFO 是一个指针环形缓冲，GPU 和 CPU 用它来协调 GPU 还需要读取什么以及已经读取了什么。GPFIFO 中的每个条目由两个 32 位字组成，描述 pushbuffer 的一段 (base, length)（这里 base 是指向主机内存的 GPU 虚拟地址）。GPU 不断遍历 GPFIFO 以寻找工作。

在驱动程序和 GPU 之间，需要维护两个游标：GP_GET（GPU 已经消费了多少）和 GP_PUT（驱动程序已经生产了多少）。两个游标都位于 USERD 中，这是一个每个通道的小型结构，这里位于设备内存中。

要启动一个内核，驱动程序会用相关方法填充一个 pushbuffer 段，将 GPFIFO 条目指向它，并推进 GP_PUT。一旦 GPU 消费了该条目，它就会推进 GP_GET。

各部分的位置示意图：

CPU                          GPU
│                            │
├─ host RAM                  │
│   ├─ pushbuffer            │
│   │   └─ methods + QMD     │
│   └─ GPFIFO ring           │
├─ writes                    ├─ DMA
├─ ioctl → nvidia.ko         ├─ USERD — GP_GET / GP_PUT
└─ doorbell (MMIO)           └─ HOST engine

PCIe: DMA, writes

我们的启动是由一系列 methods 触发的，首先是 SET_INLINE_QMD_ADDRESS_A/B（如何在源码中找到这个 method，鉴于 libcuda 是闭源的，请参见附录），接着是一连串的 LOAD_INLINE_QMD_DATA。这些 methods 用于将一个叫做“队列元数据”（QMD）的对象流入 pushbuffer。QMD 是计算网格的启动描述符。它包含网格和块的维度——我们的 4096 和 256，来自 .cu 代码——每个线程所需的寄存器和共享内存，以及两个地址：程序的起始地址（第一次启动时加载到 GPU 内存的 SASS）和常量基址。

QMD 的内部布局（对于 C6C0 类）大致如下（参考开源 GPU 内核模块）：

Offset | 大小 | 描述
0x00   | 4    | 信号量（启动完成时递减）
0x04   | 4    | 网格维度 X
0x08   | 4    | 网格维度 Y
0x0c   | 4    | 网格维度 Z
0x10   | 4    | 块维度 X
0x14   | 4    | 块维度 Y
0x18   | 4    | 块维度 Z
0x1c   | 4    | 每个线程的寄存器数
0x20   | 4    | 每个块的共享内存字节数（静态）
0x24   | 4    | 每个块的共享内存字节数（动态）
0x28   | 8    | 程序起始地址（GPU 虚拟地址）
0x30   | 8    | 常量基址（GPU 虚拟地址）
0x38   | 8    | 参数缓冲的 GPU 设备地址（或立即数内联？）
0x40   | 4    | 启动标志 / 类型

在我们的例子中，SET_INLINE_QMD_ADDRESS_A 和 B 合并设置 QMD 的基地址（指向 pushbuffer 中的某个位置），然后 LOAD_INLINE_QMD_DATA 一次一个 32 位字地流式写入整个结构。驱动程序没有单独分配 QMD 并传递其指针，而是将其写入 pushbuffer 本身，作为一个内联的数据负载。启动一个内核只需将这个 QMD 放入 pushbuffer 中即可。

门铃：将工作推给 GPU

一旦 QMD 在 pushbuffer 中，驱动程序的下一步是追加一个 method，告诉 GPU“处理这个命令”。这通过一个特殊的寄存器写入来完成，这个寄存器映射到 GPU 的 MMIO 空间。该寄存器称为门铃 (doorbell)。门铃位于一个预定义的物理地址，由 PCIe 配置空间中的 BAR 区域映射而来。驱动程序执行一条 4 字节的 MMIO 写入（例如，使用 movl 指令将通道编号写入该地址），GPU 上的主机引擎收到中断并开始在 GPFIFO 中提取好条目。

在实际的代码中，这是一个 write 或 mmap 操作。strace 可以显示相关的 mmap 调用，例如：

mmap(NULL, 4096, PROT_READ|PROT_WRITE, MAP_SHARED, 3 /* /dev/nvidia0 */, 0x1000000)

这个映射将 GPU 的 BAR 区域暴露给用户态，允许快速门铃写入而无需额外系统调用。

从 GPU 返回结果

在 GPU 处理完 pushbuffer 的最后一个 method（通常是 END 或 RETURN）后，它通过更新 GP_GET 游标并可能递减信号量来发出完成信号。主机上的驱动程序可以通过轮询 GP_GET 或者更高效地，通过 GPU 在完成时发出的一个中断来等待。在我们的简单示例中，cudaMemcpy(..., cudaMemcpyDeviceToHost) 是一个同步操作，它在复制数据之前确保所有先前启动的内核已经完成。设备到主机复制是阻塞的，直到结果可用。

内核完成后，结果已经位于 dc 指向的设备内存中。cudaMemcpy 启动一个 DMA 传输，将数据从设备内存复制到主机内存。DMA 引擎完成后再通知主机。printf 可以打印出结果。

性能考虑与开销

现在你看到，每次 <<<...>>> 启动都涉及大量的状态设置：填充 pushbuffer、写入 QMD、MMIO 门铃、等待完成。这就是为什么 CUDA 内核启动延迟在典型的 GPU 上大约为 3–10 微秒，并且为什么通常建议通过增大网格（每个内核做更多工作）或使用异步流来分摊这种开销。

此外，这篇分析解释了为什么所谓的“无操作”内核（立即返回的内核）可以作为测量启动开销的有效基准。

附录：如何查看启动闭源代码内部

鉴于 libcuda.so 是闭源的，你怎么知道实际发出的是哪些 methods？主要有两种方法：

1. 跟踪 ioctl 调用：使用 strace 并启用 -e trace=ioctl 和 -v 标志，你可以看到传递给 nvidia 驱动的参数。驱动不接受 raw register values，但你可以看到 NV_ESC_REGISTER_FD 等操作。
2. 使用 NV_VFIO 或 GPU 模拟器：通过虚拟化框架或 NVIDIA 的开源模块配合模拟，你可以设置断点并记录每次寄存器写入。
3. HW 调试器（如 NVDEC 的后门）：通常需要特殊权限，但对于研究目的可能可行。

在开源模块 nvgpu（src/common/sdk/nvidia/inc/class/clc6c0.h）中，你可以看到类的定义和相关的 method 常量。

深入阅读

• NVIDIA 开源 GPU 内核模块：https://github.com/NVIDIA/open-gpu-kernel-modules
• CUDA 编程指南：https://docs.nvidia.com/cuda/cuda-c-programming-guide/
• PTX ISA 参考：https://docs.nvidia.com/cuda/parallel-thread-execution/