@kalyan_kpl: 当你运行一个CUDA内核时会发生什么 CUDA内核是一种专门编写的代码，用于在…

当你运行一个 CUDA Kernel 时会发生什么
CUDA kernel 是一种专门编写的代码，用于在 NVIDIA GPU 设备上执行并行计算。自定义 CUDA kernel 能让你的模型获得显著的性能优势。这篇出色的文章详细解释了运行一个 CUDA kernel 时到底发生了什么。
文章 - https://fergusfinn.com/blog/what-happens-when-you-run-a-gpu-kernel/

vadd.ptx # 设备代码，PTX 格式（来自 cicc）
vadd.sm_89.cubin # 设备代码，SASS 格式（来自 ptxas）
vadd.fatbin # 打包后的 cubin + PTX（来自 fatbinary）
vadd.cudafe1.stub.c # 主机启动桩 + kernel 注册
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` 将索引 `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`。反汇编只命名到 `R9`，但分配器额外预留了几个用于 ABI 和对齐）。两个 `mul.wide` 加 `add` 序列融合为单个 `IMAD.WIDE`。`cvta` 转换消失了，已被吸收到寻址模式中。`c[0x0][...]` 操作数是**常量 bank 0**，位于一个小的、驱动程序管理的区域。这些是 kernel 的参数——指针 `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 (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 # 压缩后的 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  
编译后的 GPU 机器代码现在静置在 `./vadd` 可执行文件的 `.nv_fatbin` 段中。当你在主机上运行该程序时，我们必须连接两个世界：主机 CPU 和 PCIe 总线另一侧的 GPU。为了构建一个知道如何跨越这座桥的主机二进制文件，前端编译器（`cudafe++`）在你的代码中插入了一个隐藏的构造函数，在 `main` 函数启动之前运行。其任务是向 CUDA 运行时注册我们嵌入的 fatbinary，并记录一个映射，运行时稍后将使用该映射：将主机端的函数指针 `vadd` 与 fatbin 中编译后的设备 kernel 的混淆名称关联起来。  

当编译器遇到 `vadd<<<4096, 256>>>(da, db, dc, n)` 时，它会将这个高级表达式替换为生成的主机启动桩。该桩将我们的 kernel 参数打包到主机内存中的一个缓冲区中。指针 `da`、`db`、`dc` 和整数 `n` 分别以字节偏移量 `0`、`8`、`16` 和 `24` 对齐（这些偏移量就是我们之前看到的 SASS 机器代码从常量 bank 0 读取的常量 bank 偏移量 `0x160`、`0x168`、`0x170` 和 `0x178`）。  
```c  
// 来自 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，即本机上的驱动程序版本），以启动该 kernel。

运行时在程序中首次 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（如果需要，可以覆盖以提前支付加载成本））——驱动程序会将 kernel 的 SASS cubin 上传到显存的操作推迟到该特定 kernel 首次启动时。

在 libcuda 之下是内核模式驱动程序 nvidia.ko，libcuda 通过对设备文件调用 ioctl 来访问它。当 cuLaunchKernel 最终需要将工作放到 GPU 上时，它就成了与内核模块的对话。下面就是这次对话的机制。

将其加载到 GPU 上

GPU 不像 CPU 那样接受函数调用。没有可以跳转的入口点，也没有 CPU 可以将参数推入的栈。GPU 位于 PCIe 总线另一侧，并读取主机内存中的驱动程序命令流。cuLaunchKernel 在此之后所做的所有事情都是为了将一条完整的启动命令放入该流中，然后通知 GPU 已完成。

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

pushbuffer 是一块内存，驱动程序将命令写入其中，这些命令称为方法。一个方法包含 GPU 本地命令编码中的一个寄存器地址和一个值——该对定义了 GPU 应执行的操作。GPFIFO 是一个指针的环形缓冲区，用于 GPU 和 CPU 协调 GPU 还需要读取什么以及已经读取了什么。GPFIFO 中的每个条目由两个 32 位字组成，描述了 pushbuffer 的一个跨度（在这种情况下，基址是一个指向主机内存的 GPU 虚拟地址 (base, length)）。GPU 持续遍历 GPFIFO 以寻找工作。在驱动程序和 GPU 之间，需要维护两个游标：GP_GET（GPU 已消耗了多少）和 GP_PUT（驱动程序已生产了多少）。两个游标都位于 USERD 中，这是一个每个通道的小结构，此处位于设备内存中。要启动一个 kernel，驱动程序用相关方法填充一个 pushbuffer 跨度，将 GPFIFO 条目指向它，并推进 GP_PUT。一旦 GPU 消耗了该条目，它就推进 GP_GET。

各个部分的位置：
CPU · host RAM — pushbuffer（方法 + QMD），GPFIFO 环形缓冲区
GPU — USERD（GP_GET / GP_PUT），门铃（MMIO）
HOST engine — DMA
CPU 通过 PCIe 写入（write()）或通过 DMA 写入；GPU 通过 PCIe 读取（DMA 或直接）

我们的启动由一系列方法触发，首先是 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)。这些方法用于…