运行CUDA内核时会发生什么？

# 当你运行一个 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=21\+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`）需要更多步骤：`cicc`，一个基于LLVM (https://en.wikipedia.org/wiki/LLVM) 的编译器，将其转换为PTX (https://developer.nvidia.com/blog/understanding-ptx-the-assembly-language-of-cuda-gpu-computing/)，然后 `ptxas` 将 PTX 转换为SASS (https://modal.com/gpu-glossary/device-software/streaming-assembler)。PTX 是一种虚拟ISA (https://en.wikipedia.org/wiki/Instruction_set_architecture)。它拥有无限多个带类型的寄存器，并且不考虑硬件实际有多少个寄存器。以下是 `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` 通过乘以 4（`sizeof(float)`）并将结果从 32 位扩展到 64 位，一步完成将索引 `i` 转换为字节偏移量。`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`。反汇编只命名到 `R9`，但分配器为 ABI 和对齐保留了更多一些。两个 `mul.wide` 加 `add` 的序列已经融合成单个 `IMAD.WIDE`。`cvta` 转换消失了，被吸收到寻址中。`c[0x0][...]` 操作数是**常量库 0**，位于一个由驱动程序管理的小区域内。这些是内核的参数——指针 `a`、`b`、`c` 和大小 `n`——以及启动几何信息。填充这个库是名为 QMD 的结构的任务，该结构由驱动程序在启动时交给 GPU，我们将在启动本身到达显卡时讨论。为什么参数位于常量库 0，以及它们的位置它们位于常量内存中，因为这是一个*广播*读取：网格中的每个线程都需要相同的指针，而常量缓存能够一次服务所有 32 个通道。布局是固定的——`0x160`、`0x168`、`0x170` 是指针 `a`、`b`、`c`，`0x178` 是 `n`，启动几何信息在它们旁边的 `0x0` 处（`blockDim.x`）。库 0 也保存 ABI 参数，例如 `c[0x0][0x28]`，即 `MOV R1, c[0x0][0x28]` 在入口处加载的栈基址。当主机桩代码为启动打包参数时，我们将再次看到这些相同的偏移量。包含此 SASS 的 'cubin' 文件是一个ELF (https://en.wikipedia.org/wiki/Executable_and_Linkable_Format) 文件——与 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 架构上，驱动程序可以在加载时 JIT 将 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 不覆盖的架构上运行时，驱动程序才会解压缩并 JIT 编译它。 ## 主机如何触发 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 机器代码从常量库 0 读取的常量库偏移量 `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** 是一片内存区域，驱动程序向其中写入给 GPU 的命令，称为 *方法*。一个方法是 GPU 本地命令编码中的一个寄存器地址和一个值 —— 这对定义了 GPU 应该执行什么动作。**GPFIFO** 是一个指针的环形缓冲区，由 GPU 和 CPU 用来协调 GPU 还需要读取什么，以及已经读取了什么。GPFIFO 中的每个条目由两个 32 位字组成，描述 pushbuffer 的一个跨度 (在这种情况下，基址是一个指向主机内存的 GPU 虚拟地址) `(base, length)` 。GPU 不断遍历 GPFIFO 以找到工作。在驱动程序和 GPU 之间，需要维护两个游标：`GP_GET`（GPU 已消费了多少）和 `GP_PUT`（驱动程序已产生了多少）。两个游标都位于 USERD 中，这是一个每通道的小型结构，这里位于设备内存中。要启动一个内核，驱动程序用相关的方法填充一个 pushbuffer 跨度，将一个 GPFIFO 条目指向它，并推进 `GP_PUT`。一旦 GPU 消费了该条目，它会推进 `GP_GET`。不同部分的位置。| CPU | GPU ||------|------|| PCIe | 主机 RAM | pushbuffer — 方法 + QMD | GPFIFO 环形缓冲区 || USERD — GP_GET / GP_PUT | 门铃 (MMIO) || HOST 引擎 || DMA | 写入 | 读取 || CPU · 主机 RAM | pushbuffer — 方法 + QMD | GPFIFO 环形缓冲区 || PCIe | 写入 | DMA | GPU | USERD — GP_GET / GP_PUT | 门铃 (MMIO) || HOST 引擎 | 我们的启动由一系列方法触发，首先是 `SET_INLINE_QMD_ADDRESS_A/B` (https://github.com/NVIDIA/open-gpu-kernel-modules/blob/590.48.01/src/common/sdk/nvidia/inc/class/clc6c0.h#L403-L409) (我怎么知道是这个方法，因为 `libcuda` 是闭源的：见附录 (https://fergusfinn.com/blog/what-happens-when-you-run-a-gpu-kernel/#appendix-how-to-look-inside-the-launch)) ，随后是一系列 `LOAD_INLINE_QMD_DATA` (https://github.com/NVIDIA/open-gpu-kernel-modules/blob/590.48.01/src/common/sdk/nvidia/inc/class/clc6c0.h#L409-L410)。这些方法的作用是将一个称为「队列元数据」(QMD) 的对象流式传输到 pushbuffer 中。QMD 是计算网格的启动描述符。它包含网格和块的维度——我们 `..cu` 代码中的 4096 和 256——每个线程的寄存器和所需的共享内存，以及两个地址：程序的起始地址（第一次启动加载到 GPU 内存中的 SASS）和常量库 0 的基地址。