在深入阅读本文之前，了解以下基本概念将大有裨益： 1、**Host 与 Device（主机与设备）：** 在 CUDA 编程中，“Host” 指代 CPU 及其主存（RAM），而 “Device” 指代 GPU 及其显存（VRAM）。 2、**Kernel（内核函数）：** 程序员编写的、在 GPU 阵列上并行执行的函数（通过 `__global__` 修饰符定义）。 3、**PTX 与 SASS：** * **PTX** (Parallel Thread Execution) 是一种类似于汇编的中间语言，它与具体的 GPU 硬件架构无关（虚拟 ISA）。 * **SASS** (Streaming Assembler) 是特定于某一代 GPU 硬件（如 Ampere, Ada Lovelace）的最终底层机器码。 4、**CUDA 线程层级 (Thread Hierarchy)：** * **Thread (线程)：** 最小的执行单元。 * **Warp (线程束)：** 硬件调度的基本单位，通常由 32 个线程组成。它们以锁步 (lockstep) 的方式执行同一条指令。 * **Block (线程块)：** 由多个线程（或线程束）组成，同一个 Block 内的线程可以共享内存并同步。 * **Grid (网格)：** 由多个 Block 组成，代表整个 Kernel 的启动规模。 5、**SM (Streaming Multiprocessor, 流多处理器)：** GPU 内部的核心计算单元。一个现代 GPU 包含几十到上百个 SM。每个 SM 负责接收并调度多个 Block 与 Warp。 6、**Ioctl & MMIO：** CPU 与 GPU 驱动通信的底层方式。`ioctl` 用于向内核模块发送控制命令；MMIO (内存映射I/O) 允许 CPU 直接像写内存一样向 GPU 的寄存器（如“Doorbell / 门铃”）写入数据来触发任务。 --- # What happens when you run a CUDA kernel # 当你运行一个 CUDA 内核时会发生什么 Here’s a simple CUDA program. It adds two vectors. 这是一个简单的 CUDA 程序。它用于将两个向量相加。 ```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; // a million floats (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 threads, one per float / n个线程，每个浮点数一个 cudaMemcpy(c, dc, bytes, cudaMemcpyDeviceToHost); printf("c[0]=%f c[n-1]=%f\n", c[0], c[n-1]); } ``` Compiled for an RTX 4090, and launched, it does correctly work out that $1+1=2$, a million times[^1]. 将它编译为适用于 RTX 4090 的程序并启动，它确实正确地算出了一百万次 $1+1=2$ [^1]。 ```txt $ nvcc -arch=sm_89 -o vadd vadd.cu && ./vadd c[0]=2.000000 c[n-1]=2.000000 ``` Telling you that involved tens of millions of CPU instructions, a couple of device files, nine hundred ioctls, and one memory-mapped doorbell register. In this post, we’ll follow this one kernel from the code down to the warps, and back up to the answer[^2]. 为了给你呈现这个结果，计算机执行了上千万条 CPU 指令，访问了几个设备文件，调用了 900 次 ioctl 系统调用，并使用了一个内存映射的“门铃 (doorbell)”寄存器。在这篇文章中，我们将跟随这个内核，从代码层面一直下探到线程束（warps），再一路追踪回到最终答案的生成[^2]。 ## Compiling our program with `nvcc` ## 使用 `nvcc` 编译我们的程序 We ought to start with how to turn this CUDA program into something that the device can actually read. To do that we need a compiler. Really, we need many compilers. 我们首先需要弄清楚，如何将这个 CUDA 程序转换成设备（GPU）真正能读懂的内容。为此我们需要一个编译器。实际上，我们需要多个编译器。 `nvcc` is a driver program that runs several other compilers and combines their output. If you pass `--keep` it leaves the whole pipeline on disk for you to read: `nvcc` 是一个驱动程序，它会运行其他几个编译器并整合它们的输出。如果你在编译时加上 `--keep` 参数，它会将整个编译流水线的中间产物保留在磁盘上供你查阅： ```txt $ nvcc --keep -arch=sm_89 -o vadd vadd.cu && ls ... vadd.ptx # device code as PTX (from cicc) / PTX 格式的设备代码（出自 cicc） vadd.sm_89.cubin # device code as SASS (from ptxas) / SASS 格式的设备代码（出自 ptxas） vadd.fatbin # cubin + PTX, bundled (from fatbinary) / 捆绑在一起的 cubin + PTX（出自 fatbinary） vadd.cudafe1.stub.c # host launch stub + kernel registration / 主机启动桩代码 + 内核注册代码 vadd.o # final host object, fatbin embedded / 最终的主机对象文件，已嵌入 fatbin ... ``` The host code goes to your host compiler. The device code (`vadd`) takes more steps: `cicc`, an LLVM-based compiler, turns it into PTX, and then `ptxas` turns the PTX into SASS. 主机代码（CPU 侧）交由你的主机编译器处理。而设备代码（GPU 侧，即 `vadd`）则需要经过更多步骤：基于 LLVM 的编译器 `cicc` 首先将其转换为 PTX，然后 `ptxas` 将 PTX 转换为 SASS。 PTX is a virtual ISA. It has infinitely many typed registers, and no notion of how many of them the hardware actually has. Here is the (elided) body of `vadd` in PTX: PTX 是一种虚拟的指令集架构 (ISA)。它具有无限多的强类型寄存器，且完全不关心实际硬件究竟有多少个寄存器。以下是 PTX 格式的 `vadd` 主体代码（部分省略）： ```ptx $ cat vadd.ptx ... mad.lo.s32 %r1, %r3, %r4, %r5; // set register r1 to ctaid*ntid + tid / 将 r1 设为全局线程索引 i setp.ge.s32 %p1, %r1, %r2; // set predicate p1 if i >= n / 如果 i >= n 则设置谓词 p1 @%p1 bra $L__BB0_2; // if out of bounds, skip to exit / 越界则跳转至结束 cvta.to.global.u64 %rd4, %rd1; // convert generic pointer %rd1 to a global address, store in %rd4 / 转换通用指针 mul.wide.s32 %rd5, %r1, 4; // multiply r1 by 4, store the result in %rd5 / i 乘以 4 算出字节偏移 add.s64 %rd6, %rd4, %rd5; // add %rd4, %rd5, result in %rd6 / 加上基地址 ld.global.f32 %f2, [%rd6]; // load a[i] into %f2 / 读取 a[i] 到 %f2 ... add.f32 %f3, %f2, %f1; // add %f1 and %f2, result in %f3 / 浮点加法 st.global.f32 [%rd10], %f3; // store c[i] = ... in global memory / 将结果写回 c[i] ``` The virtual registers look like `%rd1`–`%rd10`, `%f1`–`%f3`[^3]. 其中的虚拟寄存器看起来像 `%rd1`–`%rd10`、`%f1`–`%f3` [^3]。 PTX is more ‘longhand’ than you might expect. For example, forming one address in `%rd6` takes three PTX instructions. This happens because PTX is device agnostic. PTX 比你想象的要更“繁琐”。例如，在 `%rd6` 中生成一个地址就需要三条 PTX 指令。这是因为 PTX 是与具体设备无关的。 > **Why three? / 为什么需要三条指令？** > CUDA pointers are "generic" by default, meaning they could name global, shared, or local memory. `cvta.to.global` asserts the pointer lives in the global window, so a cheaper `ld.global` can be used later. `mul.wide.s32` then turns the index `i` into a byte offset by multiplying by 4 (`sizeof(float)`) and widening 32→64 bits in one step. `add.s64` adds that to the base pointer. > > CUDA 指针默认是“通用”的，意味着它们可以指向全局、共享或本地内存。`cvta.to.global` 断言该指针位于全局内存空间，这样后续就可以使用更低开销的 `ld.global` 指令。接着 `mul.wide.s32` 将索引 `i` 乘以 4 (`sizeof(float)`) 转换为字节偏移量，并一步完成 32 位到 64 位的展宽。最后 `add.s64` 将其加到基指针上。 Next, `ptxas` transforms our PTX, which is device agnostic, into the SASS for your architecture, which isn’t. The SASS it emits looks different: 接下来，`ptxas` 将与设备无关的 PTX 转换为适用于你特定架构的 SASS（这与设备息息相关）。它生成的 SASS 看起来截然不同： ```sass $ cuobjdump -sass vadd /*0000*/ MOV R1, c[0x0][0x28] ; // set up the stack pointer (ABI; unused here) / 设置栈指针 /*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 ; // if so, exit / 若是，则退出 /*0060*/ MOV R7, 0x4 ; // load literal 4 (sizeof(float)) into R7 as multiplier / 载入常数4 /*0070*/ ULDC.64 UR4, c[0x0][0x118] ; // uniform load of a driver-provided system value / 加载驱动系统值 /*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 ; ``` > **What the S2R lines are doing / S2R 指令在做什么** > `S2R` is "special register to register": it copies a _special_ register the hardware maintains per thread — here `SR_CTAID.X` (the block's index, `blockIdx.x`) and `SR_TID.X` (the lane's index within the block, `threadIdx.x`) — into an ordinary register so `IMAD` can do arithmetic on it. > > `S2R` 意为“特殊寄存器到通用寄存器”：它将硬件为每个线程维护的_特殊_寄存器——此处是 `SR_CTAID.X` (`blockIdx.x`) 和 `SR_TID.X` (`threadIdx.x`) ——复制到普通寄存器中，以便 `IMAD` 能够对其进行算术运算。 Ten-odd virtual registers have collapsed onto seven real ones[^4]. The two `mul.wide` plus `add` sequences have fused into a single `IMAD.WIDE`. The `cvta` conversions are gone, absorbed into the addressing. 十几条虚拟寄存器被压缩映射成了 7 个真正的物理寄存器[^4]。两次 `mul.wide` 加 `add` 的序列被融合为了单条 `IMAD.WIDE` 指令。`cvta` 转换也不见了，直接被吸收进寻址操作中了。 The `c[0x0][…]` operands are **constant bank 0**, in a small, driver-managed region. These are the kernel’s arguments — the pointers `a`, `b`, `c` and the size `n` — along with the launch geometry. Filling the bank is the job of a structure called the QMD that the driver hands the GPU at launch, which we’ll come to once the launch itself reaches the card. `c[0x0][…]` 操作数代表 **常量存储区 0 (constant bank 0)**，这是由驱动程序管理的一小块内存区域。这里面存放了内核的参数——指针 `a`, `b`, `c` 和大小 `n` ——以及启动配置（launch geometry）。在内核启动时，驱动程序会交给 GPU 一个称为 QMD 的数据结构，向该存储区填入数据的任务正是由它负责，稍后当启动命令到达显卡时我们将详细探讨它。 > **Why the arguments sit in constant bank 0, and where / 为什么参数存放在常量存储区 0 以及具体位置** > They’re in constant memory because this is a _broadcast_ read: every thread in the grid needs the identical pointers, and the constant cache is able to serve all 32 lanes in one shot. The layout is fixed — `0x160`, `0x168`, `0x170` are the pointers `a`, `b`, `c`, and `0x178` is `n`, with the launch geometry alongside them at `0x0` (`blockDim.x`). Bank 0 also holds ABI parameters such as `c[0x0][0x28]`, the stack base that `MOV R1, c[0x0][0x28]` loads at entry. We’ll see these same offsets again when the host stub packs the arguments for launch. > > 它们之所以在常量内存中，是因为这是一种*广播*读取：网格 (grid) 中的每个线程都需要相同的指针，而常量缓存能够一次性满足同一线程束内所有 32 个线程的读取需求。其布局是固定的——`0x160`、`0x168`、`0x170` 是指针 `a`、`b`、`c`，`0x178` 是变量 `n`，启动配置则与它们同在，比如 `0x0` 存放了 `blockDim.x`。存储区 0 同样存放 ABI 参数，比如 `c[0x0][0x28]` 表示栈基址，这是由于入口处有 `MOV R1, c[0x0][0x28]`。稍后我们在看到主机端桩代码 (host stub) 打包启动参数时，会再次见到这些偏移量。 The ‘cubin’ file holding this SASS is an ELF file — the same object-file container Linux uses for ordinary executables and shared libraries[^5]. The `fatbinary` executable bundles the cubin together with the PTX into a single ‘fatbin’, and `cuobjdump` on the result reveals that the fatbin embedded in our binary contains *both*: 存放这些 SASS 代码的 'cubin' 文件是一个 ELF 文件——即 Linux 用于普通可执行文件和共享库的同一类对象文件容器[^5]。可执行程序 `fatbinary` 会将该 cubin 和 PTX 打包进单个 'fatbin' 里。对生成的文件运行 `cuobjdump` 就能发现，嵌入在我们二进制文件中的 fatbin *同时*包含了两者： ```txt $ cuobjdump vadd ... Fatbin elf code: arch = sm_89 # the SASS we just read / 我们刚才阅读的 SASS 代码 Fatbin ptx code: arch = sm_89 compressed # the PTX, shipped too / 一同附带的 PTX 代码 ``` The SASS is what actually runs on this 4090, but the PTX rides along as a forward-compatibility fallback. If you then take this binary to a GPU whose architecture the cubin doesn’t cover, the driver can JIT the PTX into fresh SASS at load time. 这些 SASS 就是在此 RTX 4090 上真正运行的机器码，而 PTX 作为向前兼容的后备选项随车同行。如果你随后将这个二进制文件放到一个架构不包含在此 cubin 中的 GPU 上运行，驱动程序可以在加载时通过 JIT（即时编译）将 PTX 编译成崭新的 SASS 码。 Finally, that fatbin is nested in the host executable, where `readelf -S` finds it occupying its own sections: 最后，这个 fatbin 被嵌套在主机可执行文件（你的主程序）内部，通过 `readelf -S` 可以发现它占据了专属的 section： ```txt $ readelf -S vadd ... [18] .nv_fatbin PROGBITS ... [19] __nv_module_id PROGBITS ... [29] .nvFatBinSegment PROGBITS ... ... ``` The `vadd` binary that nvcc spits out is a single executable containing host code, a complete ELF object containing the Ada SASS, and a copy of the PTX. Because PTX is verbose plain text, `nvcc` compresses it by default to keep the binary size small; the driver will only decompress and JIT-compile it if the binary is run on an architecture that the pre-compiled SASS doesn’t cover. `nvcc` 输出的 `vadd` 二进制文件是一个单一的可执行文件，包含了主机端 CPU 代码、一个完整的包含 Ada 架构 SASS 代码的 ELF 对象，以及一份 PTX 的副本。由于 PTX 是冗长的纯文本，`nvcc` 默认会对其进行压缩以保持较小的二进制体积；只有当该二进制文件运行在没有预编译 SASS 覆盖的架构上时，驱动程序才会对其进行解压和 JIT 编译。 ## How the host triggers the GPU ## 主机 (Host) 是如何触发 GPU 的 The compiled GPU machine code is now sitting inert inside the `.nv_fatbin` section of our `./vadd` executable. When you launch the program on the host, we have to bridge two worlds: the host CPU, and the GPU sitting across the PCIe bus. 现在，编译好的 GPU 机器码正静静地待在我们的 `./vadd` 可执行文件的 `.nv_fatbin` 区段中。当你在主机上启动程序时，我们必须在两个世界之间搭建桥梁：一个是主机的 CPU，另一个是横跨在 PCIe 总线另一端的 GPU。 To set up a host binary that knows how to cross the bridge, the frontend compiler (`cudafe++`) inserts a hidden constructor into your code, running before the `main` function starts. Its job is to register our embedded fatbinary with the CUDA runtime and record a mapping that the runtime will later use: associating the host-side function pointer `vadd` with the compiled device kernel’s mangled name in the fatbin. 为了构建一个知道如何跨越这座桥梁的主机二进制文件，前端编译器 (`cudafe++`) 在你的代码中插入了一个隐藏的构造函数，它会在 `main` 函数启动之前运行。它的任务是将我们嵌入的 fatbinary 注册到 CUDA 运行时环境 (Runtime) 中，并记录一个运行时后续会用到的映射关系：将主机侧的函数指针 `vadd` 与 fatbin 中已编译好的、经过名称修饰（mangled）的设备内核名称关联起来。 When the compiler encounters `vadd<<<4096, 256>>>(da, db, dc, n)`, it replaces that high-level expression with a generated host launch stub. This stub packs our kernel arguments into a buffer in host memory. The pointers `da`, `db`, `dc` and the integer `n` are aligned at byte offsets `0`, `8`, `16`, and `24` [^6]: 当编译器遇到 `vadd<<<4096, 256>>>(da, db, dc, n)` 这种启动语法时，它会将这种高级表达式替换为生成的“主机启动桩代码 (host launch stub)”。这个桩代码负责将我们的内核参数打包进主机内存中的一个缓冲区内。指针 `da`, `db`, `dc` 以及整数 `n` 分别在字节偏移量 `0`, `8`, `16`, `24` 处对齐 [^6]： ```c // 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); // arg buffer offset 0 / 参数缓冲区偏移 0 __cudaSetupArgSimple(__par1, 8UL); // offset 8 / 偏移 8 __cudaSetupArgSimple(__par2, 16UL); // offset 16 / 偏移 16 __cudaSetupArgSimple(__par3, 24UL); // offset 24 / 偏移 24 __cudaLaunch((char*)(void*)(const float*, const float*, float*, int))vadd); } ``` Once the arguments are packed, the stub calls `__cudaLaunch`, passing it the memory address of the host-side dummy `vadd` function. Because this host function is just an empty shell on the CPU, its host memory address serves as a lookup key. The runtime queries its registration table with this address to find the corresponding device-side symbol name, and then crosses the boundary into the closed-source user-mode driver (`libcuda.so.1`)[^7] to initiate the launch of that kernel. 一旦参数打包完成，桩代码就会调用 `__cudaLaunch`，并将主机侧伪函数 `vadd` 的内存地址传递给它。因为这个主机函数在 CPU 上只不过是一个空壳，所以它的主机内存地址仅仅被用作一个查找键值 (lookup key)。运行时通过这个地址查询其注册表，以找到对应的设备侧符号名称，然后跨越边界，进入闭源的用户态驱动程序 (`libcuda.so.1`)[^7] 以启动该内核。 The runtime opens this driver dynamically on the first GPU call in our program, which we can catch using `strace`: 在我们的程序首次进行 GPU 调用时，运行时会动态打开这个驱动程序，我们可以通过 `strace` 捕捉到这一行为： ```txt $ strace -f -e trace=openat ./vadd ... openat(..., "/lib/x86_64-linux-gnu/libcuda.so.1", O_RDONLY|O_CLOEXEC) = 3 ... ``` When this first call is performed, a ‘context’ is created, containing all the infrastructure the driver needs to talk to the device, including the *channel* through which the CPU speaks to the GPU. We’ll talk more about that in the next section. 在执行这首次调用时，程序会创建一个“上下文 (context)”，其中包含驱动程序与设备通信所需的所有基础设施，这其中包括 CPU 向 GPU 发送指令的“通道 (channel)”。我们将在下一节中详细介绍这一点。 At this stage, the compiled machine code still hasn’t reached the GPU. Since CUDA 12.2, module loading is lazy by default[^8]—the driver defers uploading a kernel’s SASS cubin to the card’s memory until the very first time that specific kernel is actually launched. 在这个阶段，编译好的机器码依然还没有被传输到 GPU 上。从 CUDA 12.2 开始，模块的加载默认是惰性的 (lazy loading) [^8]——驱动程序会推迟将内核的 SASS cubin 上传到显卡内存的时间，直到该内核被实际启动的第一刻。 Underneath `libcuda` sits the kernel-mode driver, `nvidia.ko`, which `libcuda` reaches by invoking `ioctl` on device files. When `cuLaunchKernel` finally needs to put work on the GPU, it becomes a conversation with that kernel module. What follows is the mechanics of that conversation. 在 `libcuda` 之下是内核态驱动程序 `nvidia.ko`，`libcuda` 通过对设备文件触发 `ioctl` 来与其通信。当 `cuLaunchKernel` 最终需要向 GPU 派发工作时，这就演变成了一场与内核模块的对话。接下来的内容便是这场对话背后的底层机制。 ## Getting it onto the GPU ## 将任务推送到 GPU 上 A GPU does not take function calls like a CPU does. There is no entry point to jump to, and no stack to push arguments onto from the CPU. The GPU sits across a PCIe bus and reads a stream of driver commands out of host memory. Everything `cuLaunchKernel` does past this point is in service of getting one fully formed launch command into that stream, and then telling the GPU it has done so. GPU 接受指令的方式与 CPU 处理函数调用的方式大不相同。GPU 没有可供跳转的入口点，也没有用来压入 CPU 传来参数的执行栈。GPU 横跨 PCIe 总线，它的工作方式是从主机内存中读取驱动程序生成的命令流。在此之后，`cuLaunchKernel` 所做的一切都是为了将一条完整成型的“启动命令”塞进这个命令流中，然后通知 GPU：“我已经放好指令了。” The first thing that needs to be done is loading the GPU code onto the device. The first time you run `vadd`, the driver copies across the kernel’s code: it allocates a buffer and copies the SASS in. 要做的第一件事是将 GPU 代码加载到设备上。你第一次运行 `vadd` 时，驱动程序会将内核的代码复制过去：它在显存里分配了一个缓冲区，并将 SASS 机器码复制进去。 Once the code is on the GPU, the CPU needs to get the GPU to read it and start executing it. It does so via a complex dance, across host and device memory. Both the host and the GPU can map regions of each other’s memory spaces, but accesses across the PCIe bus pay a penalty. To achieve a kernel launch, both write to various structures, living across both spaces. These structures comprise the *channel* — the work queue that runs the GPU’s operations. 代码到了 GPU 上后，CPU 还需要让 GPU 去读取并开始执行它。这需要跨越主机和设备内存进行一场极其复杂的“双人舞”。主机和 GPU 都可以映射对方的部分内存空间，但是跨 PCIe 总线的访问存在性能开销。为了实现内核的启动，双方都在读写散布在两个空间内的各种数据结构。这些数据结构共同构成了 *通道 (channel)* ——即驱动 GPU 操作的工作队列。 There are two important such structures living in host RAM: the **pushbuffer**, and the **GPFIFO**, representing between them the list of work the GPU has to perform. 在主机 (Host) RAM 中存在两个非常重要的结构：**推送缓冲区 (pushbuffer)** 和 **GPFIFO**。它们二者共同表征了 GPU 必须执行的工作列表。 The **pushbuffer** is a region of memory into which the driver writes commands to the GPU, called *methods*. A method is a register address and a value in the GPU’s native command encoding — the pair defines what action the GPU should perform. **推送缓冲区 (pushbuffer)** 是一块内存区域，驱动程序向此处写入针对 GPU 的命令，这些命令被称为 *methods (方法)*。一个 method 其实就是一个寄存器地址加上 GPU 原生指令编码的键值对——这一组合定义了 GPU 应该采取的动作。 The **GPFIFO** is a ring buffer of pointers, used by the GPU & CPU to coordinate what the GPU still needs to read, and what it’s read already. Each entry in the GPFIFO is made up of two 32-bit words, describing a span of the pushbuffer[^9] `(base, length)`. **GPFIFO** 是一个存储指针的环形缓冲区 (ring buffer)，GPU 和 CPU 通过它来协调“GPU 还需要读取什么”以及“GPU 已经读取了什么”。GPFIFO 中的每一项由两个 32 位的字组成，用来描述 pushbuffer 中的一段内存区间[^9] `(基址, 长度)`。 The GPU continually walks the GPFIFO to find work. Between the driver and the GPU, two cursors need to be maintained: `GP_GET` (how far the GPU has consumed), and `GP_PUT` (how far the driver has produced). Both cursors live in USERD, a small per-channel structure that here sits in device memory. To launch a kernel, the driver fills a pushbuffer span with the relevant methods, points a GPFIFO entry at it, and advances `GP_PUT`. Once the GPU consumes the entry, it advances `GP_GET`. GPU 会持续遍历 GPFIFO 来寻找工作。在驱动程序和 GPU 之间，必须维护两个游标：`GP_GET`（GPU 已经消耗/读取了多少）以及 `GP_PUT`（驱动程序生产/写入了多少）。这两个游标都驻留在 USERD 内，这是一个每个 channel 专属的小型数据结构，在此处它位于设备的显存中。为了启动一个内核，驱动程序用各种 methods 填满一段 pushbuffer 区间，将一个 GPFIFO 条目指向它，并向前推进 `GP_PUT` 游标。一旦 GPU 消费了该条目，它就会推动更新 `GP_GET` 游标。 *(图解注：在物理跨度上，主机的 CPU 写入位于 Host RAM 上的 pushbuffer 与 GPFIFO 环，然后通过 PCIe 总线跨界写入 GPU 侧的 GP_PUT 游标与 Doorbell 门铃寄存器，最后 GPU 端的主机引擎 (HOST engine) 通过 DMA 读取主机端上的 methods 指令)* Our launch is triggered by a burst of methods, first `SET_INLINE_QMD_ADDRESS_A/B`[^10] followed by a run of `LOAD_INLINE_QMD_DATA`. These methods serve to stream an object called the “Queue Meta Data” (**QMD**) into the pushbuffer. 我们的启动过程是由爆发写入的一组 methods 触发的：首先是 `SET_INLINE_QMD_ADDRESS_A/B` [^10]，随后跟进一系列 `LOAD_INLINE_QMD_DATA`。这些 methods 的目的是将一个名为“队列元数据”(**QMD**, Queue Meta Data) 的对象流式传输到 pushbuffer 中。 The QMD is the launch descriptor for a compute grid. It holds the grid and block dimensions — our 4096 and 256, from the `.cu` code — the registers per thread and shared memory it needs, and two addresses: the program’s start (the SASS the first launch loaded into GPU memory) and the constant bank holding the kernel’s arguments. That bank is where the arguments the host stub packed land: the driver copies them in and records the bank’s address in the QMD. The QMD tells the GPU where the SASS is, how to turn that SASS into a parallel program, and where to signal its completion of that program. QMD 是整个计算网格 (compute grid) 的启动描述符。它保存着网格和线程块的维度——即我们从 `.cu` 代码传来的 4096 和 256——以及它所需的每个线程寄存器数量和共享内存。更重要的是，它包含两个关键地址：程序的起始地址（即首次启动时加载到显存里的 SASS）以及保存内核参数的常量存储区 (constant bank)。常量存储区正是主机桩代码打包好参数的落脚点：驱动程序将参数拷贝进去，并在 QMD 中记录下该存储区的地址。简而言之，QMD 告诉 GPU：SASS 机器码在哪里、如何将这些 SASS 代码调度成并行程序，以及程序完成时该去哪里发送完工信号。 Everything is now in place for the GPU to start running. The problem is that the GPU’s **host engine**[^11] hasn’t acted: it doesn’t watch the cursor on modern cards[^12], so the change to `GP_PUT` just sits there until something tells the engine to look. 万事俱备，GPU 随时可以开跑了。但问题在于，GPU 的 **主机引擎 (host engine)** [^11] 目前还没有动作：在现代显卡上，它并不会主动去轮询监视游标[^12]，因此对 `GP_PUT` 的修改就那样静止着，直到有东西明确通知主机引擎去看一眼。 It is told to look through the **doorbell**. The GPU maps a small window of its registers into the process, and one of them is the doorbell; the driver writes the channel’s *work-submit token* to it. The token tells it which channel has new work. 让它“看一眼”的通知手段就是 **门铃 (doorbell)**。GPU 会将其一小块寄存器窗口内存映射进进程中，其中一个寄存器就是门铃；驱动程序向门铃寄存器写入该通道的 *任务提交令牌 (work-submit token)*。这个令牌明确告诉 GPU 哪个通道有了新活儿。 When its doorbell gets rung, the host engine reads the updated `GP_PUT`, follows the new GPFIFO entry to the pushbuffer span, and pulls the methods out of it by DMA. When it reaches the compute method carrying our QMD, it hands that descriptor to the “compute work distributor”, about which more shortly. 当“门铃被按响”时，主机引擎就会读取更新后的 `GP_PUT`，顺着新的 GPFIFO 条目找到对应的 pushbuffer 内存段，然后通过 DMA（直接内存访问）将这些 methods 提取出来。当它遇到携带我们 QMD 信息的计算 method 时，就会把这个描述符移交给“计算工作分配器 (compute work distributor)”，我们马上就会讨论到它。 From the CPU’s side the launch is done: `cuLaunchKernel` returned the moment the doorbell was rung. The call was asynchronous, so control returns to the program and the CPU runs on while the GPU works; we pick the host side back up once the kernel has run. 从 CPU 的视角来看，内核启动 (launch) 到此已经完毕：在门铃响起的瞬间，`cuLaunchKernel` 函数就已返回。这个调用是异步的，因此程序的控制权交还给了 CPU，CPU 可以继续往下跑，而 GPU 同时开始干活；我们会在内核跑完后，再回到主机端的故事线。 It’s time for the GPU to start doing its job. 现在，轮到 GPU 登场发力了。 ## Instruction by instruction ## 一条条指令的执行 The host engine hands the QMD to the **compute work distributor**[^13]. There is one of these on the whole GPU. There is one linear list of SASS instructions in VRAM, and the compute work distributor + the QMD is the first step in telling the hardware how to make that linear list of thread instructions into a massively parallel program across all the **Streaming Multiprocessors** (SMs). 主机引擎把 QMD 移交给了 **计算工作分配器 (compute work distributor)** [^13]。整个 GPU 内只有一个这样的全局调度分配器。显存 (VRAM) 中保存着一条线性的 SASS 指令序列，而“工作分配器 + QMD”的组合正是告诉硬件：如何将这线性的单线程指令列表，跨越所有 **流多处理器 (SMs)**，转化为真正的大规模并行程序的第一步。 In our journey down the stack, our compute work distributor now has a QMD describing 4096 blocks of 256 threads. The card we are targeting is a GeForce RTX 4090 chip with **128 SMs**[^14]. The distributor’s task is to keep all 128 saturated with work. 在我们的调用栈下探之旅中，计算工作分配器拿到了一份 QMD，里面描述着 4096 个线程块，每个块有 256 个线程。我们当前指向的显卡是一块搭载了 **128 个 SM** 的 GeForce RTX 4090 [^14]。分配器的核心使命就是：用任务把这 128 个 SM 全部喂饱。 The compiled machine code sits as a single linear sequence in global memory. Each SM contains its own local Instruction Cache (I-cache), and every active warp on the GPU maintains its own private Program Counter (PC)[^15]. Schedulers on the SM then fetch instructions from that linear sequence independently, allowing different warps to execute the same SASS code at different speeds, or down different branch paths. 编译好的机器码作为一条线性序列静卧在全局内存中。每个 SM 都包含自己的本地指令缓存 (I-cache)，GPU 上每个处于活跃状态的线程束 (warp) 都维护着自己专属的程序计数器 (Program Counter, PC)[^15]。紧接着，SM 上的各个调度器 (schedulers) 会独立地从那条线性序列中抓取指令。这种独立调度允许不同的线程束以不同的步调来执行同样的 SASS 代码，或者顺着完全不同的条件分支路径走下去。 *(图解注：VRAM 中的线性指令流通过缓存供各 SM 独立提取，每个 SM 并行管理多达 48 个活跃线程束，多数线程束可能由于正处于内存读取等待而挂起，极少数处于发射执行状态)* The hardware constraints of our SMs set the number of blocks that can run at the same time[^16]: SM 硬件层面的制约条件决定了在同一时刻最多能有几个线程块并发运行[^16]： ```txt +------------------------------------------------------------+ | AD102 SM Resource Caps | +------------------------------------------------------------+ | Max Active Threads/SM | 1,536 threads (48 warps) | | Register File/SM | 65,536 32-bit registers (256 KB) | | Shared Memory/SM | 100 KB | +------------------------------------------------------------+ ``` Our launch configuration specifies blocks of **256 threads (8 warps)**, and `ptxas` reserved **16 registers per thread**. 我们的启动配置明确每个块有 **256 个线程（相当于 8 个线程束）**，而 `ptxas` 编译器为每个线程保留了 **16 个寄存器**。 1. **Register capacity (寄存器容量限制)**: Each block needs $256 \times 16 = 4,096$ registers. On registers alone, an SM could fit $65,536 / 4,096 = 16$ resident blocks. 每个块需要 $256 \times 16 = 4,096$ 个寄存器。单从寄存器容量来看，一个 SM 可以容纳 $65,536 / 4,096 = 16$ 个常驻块 (resident blocks)。 2. **Thread capacity (线程数限制)**: The hardware caps each SM at 1,536 active threads. Divided by our block size, this yields $1,536 / 256 = 6$ resident blocks. 硬件限制每个 SM 的活跃线程上限为 1,536 个。除以我们的块大小，得出 $1,536 / 256 = 6$ 个常驻块。 Because thread capacity is the tighter bottleneck, each SM holds at most **6 blocks (48 warps) at once**. 因为线程数限制是更加严苛的瓶颈，故每个 SM 同一时间最多只能容纳 **6 个块（即 48 个线程束）**。 The distributor assigns these 6 resident blocks to an SM. Each SM is divided into **four processing blocks (sub-partitions)**. Each sub-partition is a self-contained execution pipeline. 分配器将这 6 个常驻块分配给一个 SM。每个 SM 在内部被划分为 **四个处理子分区 (sub-partitions)**。每个子分区都是一条自包含的完整执行流水线。 The SM distributes our 48 resident warps evenly across these four sub-partitions, so when the SM is full each warp scheduler has **12 active warps** ($48 / 4$) to manage. Every cycle, a warp scheduler evaluates its 12 candidates, selects one *eligible* warp, and dispatches its next instruction across the 32 physical lanes of its execution slice. SM 会将这 48 个常驻线程束均匀分布到这四个子分区里。因此当 SM 满载时，每个线程束调度器 (warp scheduler) 手下会管理 **12 个活跃的线程束** ($48 / 4$)。在每一个时钟周期，线程束调度器会评估手下的这 12 个“候选人”，挑选出一个*符合资格 (eligible)*的线程束，然后把它的下一条指令发射给所属切片内的 32 条物理执行车道（计算核心）。 ### What does it mean for a warp to be eligible? ### 一个线程束“符合资格 (eligible)”意味着什么？ A GPU decides when an instruction is ready to run differently from a CPU. A modern out-of-order CPU discovers dependencies dynamically at runtime, with reorder buffers and rename logic spending silicon on extracting parallelism from a single thread. A GPU doesn’t need that: it hides latency by keeping many warps resident and switching between them when they stall. With parallelism the order of the day, too much heavyweight dependency machinery is the wrong use of silicon. So the hardware leans on the compiler to schedule everything whose timing it can predict, falling back to lightweight hardware scoreboards for whatever it can’t. 对于决定一条指令何时就绪可运行，GPU 的逻辑与 CPU 大不相同。现代乱序执行 CPU 会在运行时动态发掘指令依赖关系，它耗费大量的晶体管资源在乱序缓冲区 (reorder buffers) 和寄存器重命名 (rename logic) 上，试图从单个线程中榨取指令级并行性。GPU 不需要这一套：它的杀手锏是通过保持大量常驻线程束，在某个线程束阻塞 (stall) 时迅速切换到另一个来掩盖延迟。在并行化为主旋律的背景下，过多的重量级依赖关系检测机制简直是在浪费硅片面积。因此，GPU 硬件极其依赖编译器去预先安排好所有时序可预测的调度，而在遇到不可预测的情况时，则后备退化到使用极轻量级的硬件记分板 (scoreboards)。 Every 128-bit SASS instruction carries a packed control-code payload written by `ptxas`[^17]. These scheduling control bits dictate hardware timing directly and contain three key directives: `ptxas` 编译器在生成每条 128 位的 SASS 指令时，都会在其中打包附带一段控制码 (control-code payload) [^17]。这些调度控制位直接指挥硬件的时序动作，主要包含三大关键指令： 1. **A static stall count (静态阻塞周期数)**: For fixed-latency instructions—like standard integer or floating-point maths—the compiler knows exactly when the ALUs will write back. It encodes a precise cycle count telling the scheduler exactly how long to park this warp before issuing its very next instruction. 对于那些执行延迟固定的指令（例如标准的整型或浮点运算），编译器清清楚楚地知道 ALU 会在什么时候完成计算并将结果写回。它会在指令里编码一个精确的时钟周期数，告诉调度器：在这个线程束执行下一条指令前，需要让它暂停等待多少个周期。 2. **A yield hint (让出优先级提示)**: A single bit telling the scheduler whether this warp should yield its scheduling priority. If the compiler knows this warp is about to hit a bottleneck, it sets this hint to let the scheduler prioritize other active warps on the next clock cycle. 仅仅用 1 个比特位，用来提示调度器这个线程束是否应该主动让出其调度优先级。如果编译器预判该线程束马上就要撞上执行瓶颈，它会设置这个提示位，好让调度器在下一个时钟周期优先考虑其他活跃的线程束。 3. **Dependency-barrier indices (依赖屏障索引)**: For variable-latency operations whose duration cannot be predicted at compile time—most notably global memory loads (`LDG`) and special functions (`MUFU`)—the hardware provides **six physical scoreboard barriers (numbered 0 to 5)** per warp. 对于那些在编译阶段完全无法预测所需周期的变延迟操作——最典型的就是全局内存读取 (`LDG`) 和特殊函数执行 (`MUFU`) ——硬件为每个线程束提供了 **6 个物理的记分板屏障 (scoreboard barriers，编号从 0 到 5)**。 > **Why you won't see these bits in the disassembly / 为什么在反汇编代码中看不到这些控制位** > When you disassemble a binary using NVIDIA’s standard `nvdisasm` tool, the raw control codes are hidden by default; the tool strips them away to show you standard, clean SASS mnemonics. However, they are stored directly alongside the instructions. If you inspect the raw binary using `cuobjdump -sass` and look closely at the hexadecimal instruction comments (e.g., `/* 0x... */`), you will see the packed, raw hex words that house these control bits. > What we know about their exact layout comes from the microbenchmarking community’s reverse-engineering efforts. Although the bit fields have shifted and evolved between Maxwell, Volta, Ampere, and Ada Lovelace, the core architectural concept remains identical: compile-time static scheduling metadata is packed directly into the instruction stream to keep the SM hardware as simple and power-efficient as possible. > > 当你使用 NVIDIA 官方的 `nvdisasm` 工具进行反汇编时，这些原始的控制码默认是被隐藏的；工具会剥离它们，向你展示干净纯粹的 SASS 助记符。但是，它们实际上是跟在指令旁边直接存储的。如果使用 `cuobjdump -sass` 检查原始二进制文件，并仔细观察那些十六进制的指令注释 (如 `/* 0x... */`)，你就能看到包裹着这些控制位的原始 16 进制打包词 (hex words)。 > 我们对它们具体布局的了解，得益于微基准测试 (microbenchmarking) 社区不懈的逆向工程努力。尽管从 Maxwell, Volta, Ampere 一路进化到 Ada Lovelace，比特位字段发生了诸位偏移与演化，但其核心架构理念亘古不变：**将编译期的静态调度元数据直接打包进指令流中，从而尽可能地保持 SM 硬件的极致精简与高能效。** Running `cuobjdump -sass` on our `vadd`, each instruction comes with its raw 128-bit encoding as two 64-bit words, and the *second* word of each pair carries the control payload: 在我们的 `vadd` 程序上运行 `cuobjdump -sass`，每条指令的原始 128 位编码以两个 64 位字的形态呈现，而每对编码的*第二个字*承载的正是上述控制码： ```sass $ cuobjdump -sass vadd # control payload / 控制码部分 /*00a0*/ LDG.E R4, [R4.64] /* 0x000ea8000c1e1900 */ /*00b0*/ LDG.E R3, [R2.64] /* 0x000ea2000c1e1900 */ /*00c0*/ IMAD.WIDE R6, R6, R7, c[0x0][0x170] /* 0x000fe200078e0207 */ /*00d0*/ FADD R9, R4, R3 /* 0x004fca0000000000 */ /*00e0*/ STG.E [R6.64], R9 /* 0x000fe2000c101904 */ ``` Pulling out the control payloads[^18] — their bit layout is in the appendix — you can see the schedule that `ptxas` wrote, with each directive in action: 将这些控制载荷抽丝剥茧（其位布局信息参见附录）[^18]，你就可以清楚地看见 `ptxas` 为你写下的执行时刻表，每项指令都在发挥着作用： | instruction (指令) | stall (阻塞周期) | yield (让出权限) | sets (设置屏障) | waits-on (等待屏障) | | --- | :---: | :---: | :---: | :---: | | `LDG.E` | 4 | yes | `B2` | — | | `LDG.E` | 1 | yes | `B2` | — | | `IMAD.WIDE` | 1 | yes | — | — | | `FADD` | 5 | no | — | `B2` | | `STG.E` | 1 | yes | — | — | The two loads leverage directive 3, each “set”-ing the **same** scoreboard barrier, `B2`. The `FADD`, the first instruction that needs the loaded `R4` and `R3`, carries a wait on `B2`: until both loads have returned and the barrier clears, the warp is **ineligible**, and the scheduler skips it for one of the other eleven warps in the sub-partition. 表格里前两行的 Load (加载) 指令运用了我们上文讲的规则三，它们分别“设置 (set)”了**同一个**记分板屏障 `B2`。紧随其后的 `FADD` 是第一条急需加载回来的 `R4` 和 `R3` 寄存器内容的指令，因此它携带着一个“等待 `B2` (wait on B2)”的标记：在两笔内存读取结果彻底返回、清除该屏障之前，这个线程束一直被视为 **不符合资格 (ineligible)**，调度器会果断跳过它，转而去调度该子分区里其余 11 个线程束之一。 The `FADD`→`STG` hand-off is directive 1. A floating-point add has a fixed latency, so there is no barrier: `FADD` just carries `stall=5`, which parks the warp for the few cycles it takes for `R9` to land before `STG` reads it. 而从 `FADD` 到 `STG` 的交接依靠的是规则一。一次浮点数加法拥有完全固定的执行延迟，根本用不着依赖屏障：`FADD` 指令只需带上 `stall=5` 的阻塞标记即可，这会把该线程束原地挂起几个周期，待运算结果稳稳落地落进 `R9` 后，随后的 `STG` (写存) 自然就能顺利读取它。 The yield bit, directive 2, toggles on & off across the sequence, as the compiler nudges scheduling priority around the operations that are about to wait. 至于让权标志 (yield bit，即规则二)，它在这段指令序列中被时而开启、时而关闭。每当编译器察觉接下来的操作马上要陷入等待时，它就会悄然拨动优先级开关。 Each cycle the scheduler reads the warp’s six-bit barrier state and a small stall counter, and makes the eligibility decision for each warp. This is how a GPU hides latency with close to zero hardware-scheduling overhead. 每一轮时钟周期，调度器仅仅去读取该线程束区区 6 个 bit 的屏障状态标志位，外加一个小小的倒数计时器 (stall counter)，接着就能火速决定该线程束是否符合资格。这就是 GPU 为何能以近乎于零的硬件调度开销，实现绝妙的延迟掩盖 (latency hiding)。 ### Loading the data ### 数据的加载 When a warp scheduler does find an eligible warp and issues the `LDG.E` loads, we can follow the hardware requests down the memory hierarchy. Each of the 32 threads in the warp computes an address. Because our threads access consecutive elements of `float` arrays (each 4 bytes), the warp requests a contiguous block of 128 bytes ($32 \times 4$ bytes). 当某个线程束调度器终于找到一个“符合资格”的线程束并将其 `LDG.E` 加载指令发射出去时，我们便可以顺藤摸瓜，跟随这批硬件请求一路俯冲进层次化的内存体系中。这个线程束内的 32 个线程，每个都独立计算出一个地址。由于我们的线程严格按顺序访问连续的 `float` 数组元素（每个 `float` 占 4 字节），这整个线程束发出的正是一块连续的 128 字节 ($32 \times 4$) 的内存请求区块。 The SM’s load/store unit detects this consecutive access pattern and performs **request coalescing**. It merges the 32 per-thread 4-byte requests into four 32-byte sector requests. Fetches are in units of 32 bytes, so this is perfect — if the reads were not consecutive & coalesced like this we’d end up loading more data than we needed. SM 内的加载/存储单元 (load/store unit) 敏锐地侦测到了这种连续的访问模式，随即执行了 **请求合并 (request coalescing)**。它巧妙地把 32 个分别要 4 字节的散碎请求，拼凑成了 4 个 32 字节的扇区级大请求。底层显存的抓取恰好是以 32 字节为单位的，一切完美契合——设想如果读取不连续、没被合并的话，我们势必会被迫读入大量并不需要的多余数据。 The coalesced requests first check the SM’s local L1 Data Cache. If they miss, they are routed through a high-bandwidth crossbar interconnect that links all 128 SMs to the distributed slices of the 72 MB L2 Cache. If the requests miss in the L2 cache as well, they descend further to the memory controllers and travel across the memory bus to the physical GDDR6X VRAM chips[^19]. The `STG.E` store that writes `c[i]` at the end of the loop follows the exact same path in reverse[^20]. 这些合并后的访存请求首先会向 SM 本地的 L1 数据缓存发起探查。若发生缓存未命中 (cache miss)，请求会被转送至极高带宽的交叉路互连网 (crossbar interconnect)，这条主干道连接着所有 128 个 SM 和分布式切片的 72MB L2 缓存系统。倘若倒霉到 L2 缓存也未命中，请求就会进一步下沉到内存控制器，跨过显存总线，径直冲入物理级的 GDDR6X 显存颗粒[^19]。在循环尾声负责把结果写回 `c[i]` 的 `STG.E` 存储指令，遵循的完全是一条原路返回的反向路径[^20]。 If we run our compiled kernel under the NVIDIA Nsight Compute profiler (`ncu`), we can get some telling metrics: 如果我们用 NVIDIA 官方的性能分析利器 Nsight Compute (`ncu`) 来跑跑我们编译好的内核，就能获得一些极其直白的数据： ```txt $ ncu --metrics \ launch__grid_size,launch__block_size,launch__registers_per_thread,\ launch__waves_per_multiprocessor,sm__warps_active.avg.pct_of_peak,\ smsp__issue_active.avg.pct_of_peak,dram__throughput.avg.pct_of_peak,\ gpu__time_duration.sum \ ./vadd ... ---------------------------------------------------------- Metric Name Unit Value ---------------------------------------------------------- launch__grid_size 4,096 launch__block_size 256 launch__registers_per_thread 16 launch__waves_per_multiprocessor 5.33 sm__warps_active.avg.pct_of_peak % 82.77 smsp__issue_active.avg.pct_of_peak % 5.17 dram__throughput.avg.pct_of_peak % 79.65 gpu__time_duration.sum us 10.78 ---------------------------------------------------------- ``` 82.77% of warps were active over the run. The warps were issuing instructions 5.17% of the time. The DRAM was running at 79.65% of its maximum utilization. 在运行期间，高达 82.77% 的线程束处于活跃状态。然而，这些线程束真正发射执行指令的时间仅占 5.17%。同时，DRAM 显存跑到其峰值带宽利用率的 79.65%。 The kernel has an extremely low **arithmetic intensity**: it performs exactly one floating-point addition (`FADD`) and a tiny amount of pointer arithmetic for every 12 bytes of data it transfers (two 4-byte loads and one 4-byte store). 这说明这个内核有着极低的 **算术强度 (arithmetic intensity)**：它每搬运 12 个字节的数据（两次 4 字节的加载加上一次 4 字节的存储），仅仅执行了一次凄凉的浮点加法运算 (`FADD`) 外加一小搓指针算术运算。 So the `10.78` $\mu$s just comes down to how fast the DRAM bus can feed the kernel its inputs, here about four-fifths of peak[^21]. 因此，区区 `10.78` $\mu$s 的执行耗时，彻头彻尾地归结于 DRAM 总线投喂内核输入数据的速度，在这里大约达到了理论峰值带宽的五分之四[^21]。 ## Back to the CPU ## 回到 CPU 世界 The result is now sitting in the GPU’s L2 cache. The CPU is what runs our terminal, so it needs to get the result in order to show it to us. We return to its view of events. 历经周折，计算结果此刻正躺在 GPU 的 L2 缓存里。跑着我们终端程序的可是 CPU，所以它必须得把结果要回去才能显示在屏幕上。让我们把视角切回 CPU 这边。 The launch returned control to the CPU the moment the doorbell rang. So the GPU needs to tell the CPU it’s done. When the last of our 4096 blocks retires, the GPU does so by posting a completion semaphore the QMD carried (the fence fields at words 23–24). 别忘了，早在一开始“门铃”按响的那一瞬间，内核启动调用就已经把控制权交还给了 CPU。因此，现在 GPU 得想办法通知 CPU 活干完了。当 4096 个线程块中的最后一个光荣退休 (retire) 时，GPU 按照 QMD 当初携带来的“完工信号量 (completion semaphore)”地址（即围栏/栅栏字段），发布了一个任务竣工的信号。 The device-to-host `cudaMemcpy(c, dc, …)`[^22] copy sits behind the kernel on the default stream, so the GPU’s copy engine (which performs the transfer) is gated on the semaphore. Once the value appears the GPU performs the DMA. Because `c` is still sitting dirty in the 72 MB L2 — the `STG.E` stores never had to spill it to DRAM — the engine’s reads are served straight from L2, and the data crosses PCIe without a DRAM round trip. 在默认的执行流 (default stream) 中，设备到主机的 `cudaMemcpy(c, dc, …)` [^22] 拷贝指令就排在内核执行任务的后头，所以负责传输的 GPU 拷贝引擎 (copy engine) 正卡在这个信号量上死等。信号量的值一旦更新，GPU 拷贝引擎立马拉响 DMA 传输。得益于 `c` 数组的脏数据全数静卧在 72MB 大小的 L2 缓存里——此前的 `STG.E` 写存压根没被挤回 DRAM ——引擎的读取将直接由 L2 缓存响应，数据顺着 PCIe 飞奔而去，彻彻底底省下了跑一趟 DRAM 内存回路的开销。 Once the copy finishes, it posts its *own* semaphore, which the host was waiting on in `cudaMemcpy`. `cudaMemcpy` completes on the host, `c` is ordinary host memory again, and `printf` loads `c[0]` and `c[n-1]` out of RAM, formats them into a string, and hands them to a `write` syscall on stdout. 一旦这份拷贝也大功告成，拷贝引擎同样会发布一个*它自己*的完工信号量，而主机那头的 `cudaMemcpy` 正在眼巴巴地等着它。此时主机的 `cudaMemcpy` 宣告结束，数组 `c` 又回归为平淡无奇的主机内存，接着 `printf` 从 RAM 里把 `c[0]` 和 `c[n-1]` 捞出来，格式化成字符串，最后一把塞给往标准输出的 `write` 系统调用。 ## The whole path ## 全程回顾 The kernel source went through `cicc` to PTX and through `ptxas` to SASS, which `fatbinary` packed with a fallback copy of the PTX into a cubin-bearing fatbin that the linker welded into an ordinary Linux executable. A constructor registered that fatbin before `main`, mapping a host stub to a mangled device name. The first launch lazily uploaded the cubin to the GPU. `cuLaunchKernel` built a QMD from the launch configuration, wrote it into a pushbuffer as GPU methods, advanced `GP_PUT`, and rang a doorbell with a single MMIO store, at which point the GPU’s host engine fetched the work and handed the QMD to the compute work distributor. The distributor spread 4096 blocks across 128 SMs at full occupancy, four warp schedulers per SM issued 128-bit instructions whose stall counts the compiler had written, and a coalesced memory path pulled the inputs through DRAM at four-fifths of peak bandwidth to compute, in each of a million lanes, a single sum. A completion semaphore and a copy engine then carried that result back across the bus to where `printf` was waiting, and we learnt that: 综上所述，一段微不足道的内核源码走完了这波澜壮阔的一生：它历经 `cicc` 蜕变为 PTX，再借 `ptxas` 之手化作 SASS，随后 `fatbinary` 将其与一份 PTX 备胎打包成包裹着 cubin 的 fatbin，最后链接器粗暴地把它焊进了一个普普通通的 Linux 可执行文件里。一个藏在幕后的构造函数赶在 `main` 函数之前给这块 fatbin 上了户口，将主机端的桩代码映射到一个名字被扭曲修饰过的设备名上。首次运行时才慵懒地把 cubin 上传给 GPU。随后 `cuLaunchKernel` 根据启动配置捏出了一份 QMD 投递单，将其以 GPU methods 的形式写入 pushbuffer，顺手把 `GP_PUT` 游标往前一推，紧接着用区区一个 MMIO 写操作摁响了门铃。只听叮咚一声，GPU 的主机引擎应声取货，转手便把 QMD 丢给了计算工作分配器。分配器雷厉风行，将 4096 个块满员塞进 128 个 SM，每个 SM 内部的四个线程束调度器精准发射着由编译器算好阻塞周期的 128 位指令；与此同时，一条合并访存的高速通道正以巅峰带宽五分之四的速度把 DRAM 中的输入数据玩命吸取上来……就这样，在一连串庞杂机制的保驾护航下，一百万条车道同时开算了一次可怜的加法。完工信号量和拷贝引擎带着战利品跨过总线回到了苦苦等候的 `printf` 手里，于是乎，屏幕打出了那行字，让我们恍然大悟： ```txt c[0]=2.000000 c[n-1]=2.000000 ``` ## Appendix: how to look inside the launch ## 附录：如何窥探内核启动的内部机制 Claude & I used a lot of different tricks to see the different parts of the kernel launch happen here. Some of it comes from painstakingly reading the open kernel modules. Claude 和我使用了许多不同的黑客小把戏来窥探内核启动过程中各部分的实际动向。其中部分信息来源于对开源内核模块代码夜以继日的痛苦啃读。 A few claims in this post can’t be read off the open source, because `libcuda` is closed source. To figure them out, there are a few useful diagnostic hooks. 本文中的某些断言无法直接从开源代码里读出来，因为 `libcuda` 是无情闭源的。想要搞明白它们，有几个非常好用的诊断挂钩 (diagnostic hooks) 可以帮大忙。 ### An interposition hook ### 拦截挂钩 (Interposition hook) Driver method writes never go through a syscall (the driver writes them straight into a write-combined buffer it has already mapped), so to find them you need to read the memory. We used an `LD_PRELOAD` shim that wraps `mmap`, records every region the driver maps from the `/dev/nvidia*` file, and exposes a function a test program calls just after the launch returns to dump them: 向 GPU 写入 method 指令的行为绝不会通过系统调用 (syscall) 来完成（驱动程序会直接把它们拍进它早就映射好的内存写合并缓冲区内），所以为了找到它们，你必须直接读取内存。我们动用了 `LD_PRELOAD` 这个大杀器，写了一个垫片 (shim) 包裹住了系统的 `mmap` 函数，借此记录下驱动程序从 `/dev/nvidia*` 文件映射出来的每一个内存区域，并暴露了一个导出函数，在测试程序的内核启动刚返回时就被调用，把内存数据转储 (dump) 下来： ```c #define _GNU_SOURCE #include #include #include #include #include #include // Dynamic linker function pointers / 动态链接器函数指针 static void* (*orig_mmap)(void*, size_t, int, int, int, off_t) = NULL; // Store captured channel mappings / 存储捕获的通道映射信息 struct Map { void* addr; size_t length; off_t offset; char path[256]; } maps[128]; static int map_count = 0; void* mmap(void* addr, size_t length, int prot, int flags, int fd, off_t offset) { if (!orig_mmap) { orig_mmap = dlsym(RTLD_NEXT, "mmap"); } void* ret = orig_mmap(addr, length, prot, flags, fd, offset); if (ret != MAP_FAILED && fd != -1 && map_count < 128) { char proclink[256]; char path[256]; sprintf(proclink, "/proc/self/fd/%d", fd); ssize_t len = readlink(proclink, path, sizeof(path) - 1); if (len != -1) { path[len] = '\0'; // We care about NVIDIA device files / 我们只关心 NVIDIA 的设备文件 if (strstr(path, "/dev/nvidia")) { maps[map_count].addr = ret; maps[map_count].length = length; maps[map_count].offset = offset; strcpy(maps[map_count].path, path); map_count++; } } } return ret; } // Expose a function to dump memory ranges holding the pushbuffer / 暴露用于转储包含推送缓冲区的内存段的方法 void dump_pushbuffer() { printf("\n=== [Shim] Dump of Mapped Pushbuffers ===\n"); for (int i = 0; i < map_count; i++) { // User-space channels/pushbuffers are mapped at large sizes / 用户态通道/推送缓冲区的映射一般都有一定规模 if (maps[i].length >= 0x1000) { unsigned int* ptr = (unsigned int*)maps[i].addr; printf("Mapping %d: %s, at %p (%zu bytes), offset 0x%lx\n", i, maps[i].path, maps[i].addr, maps[i].length, (long)maps[i].offset); // Walk the words looking for a method-header burst / 遍历内容，寻找一波 method 头部数据的爆发 for (size_t j = 0; j < maps[i].length / 4; j++) { unsigned int word = ptr[j]; unsigned int opcode = (word >> 29) & 0x7; // 1 = INC unsigned int count = (word >> 16) & 0x1FFF; // payload words / 负载字数 unsigned int method = (word & 0xFFF) << 2; // register offset / 寄存器偏移 // 0x318 is SET_INLINE_QMD_ADDRESS_A, the start of the inline burst / 0x318 是内联爆发写入的起始标志 if (opcode == 1 && method == 0x318) { printf(" [+] Method burst at word %zu: header = 0x%08X\n", j, word); printf(" INC, count %d, offset 0x%04X\n", count, method); for (unsigned int k = 1; k <= count && (j + k) < (maps[i].length / 4); k++) { printf(" word %02u: 0x%08X\n", k, ptr[j + k]); } } } } } } ``` Compile it into a shared library: 把它编译成共享库： ```bash $ gcc -shared -fPIC -o shim.so shim.c -ldl ``` and then call `dump_pushbuffer()` from the test program just after the kernel launch, and run it with the shim preloaded so this `mmap` runs in place of libc’s: 随后我们在测试程序的内核启动函数之后调用 `dump_pushbuffer()`，并且在运行时使用这个 `LD_PRELOAD` 垫片来让它替代系统 libc 原生的 `mmap`： ```bash $ LD_PRELOAD=./shim.so ./vadd ``` The driver maps a write-combined buffer for the channel; the dump walks it and prints the launch’s method burst. Which we then need to decode. 如此一来，驱动为这个通道映射的内存写合并缓冲区就被截获了；我们的转储程序将其一通遍历，打印出了启动时的那波 method 指令暴雨。接下来的任务就是解码它们。 ### Decoding the pushbuffer command stream ### 解码推送缓冲区的命令流 A pushbuffer method is a header word followed by data words. The header packs four fields (defined as `NVC46F_DMA_INCR_*` macros in `clc46f.h`): 一条 pushbuffer method 命令本质上就是一个头部字 (header word) 跟在几个数据字 (data words) 前面。头部字中塞满了四个字段（定义在 `clc46f.h` 的 `NVC46F_DMA_INCR_*` 宏里面）： * **bits 31:29** — opcode (操作码): `0x1` is an increasing-method write (`INC_METHOD`/`INCR_OPCODE_VALUE`), `0x3` is a non-increasing-method write (`NON_INC_METHOD`), and `0x4` is an immediate-data write (`IMMD_DATA_METHOD`). / `0x1` 表示递增 method 写入，`0x3` 表示非递增写入，`0x4` 为立即数写入。 * **bits 28:16** — count (数量): the number of payload words (`NVC46F_DMA_INCR_COUNT`). / 代表负载中包含多少个字。 * **bits 15:13** — subchannel index (子通道索引): routes the commands to a specialized backend engine context (`NVC46F_DMA_INCR_SUBCHANNEL`). / 将指令路由至特定的后端引擎上下文中。 * **bits 11:0** — the method’s register offset (寄存器偏移), divided by four (`NVC46F_DMA_INCR_ADDRESS`), as the shim shifts it back. / method 目标的寄存器偏移，以 4 字节为单位。 There are two launch paths that seem relevant here. The methods are defined per compute class in `src/common/sdk/nvidia/inc/class/` — `clc3c0.h` (Volta), `clc5c0.h` (Turing), `clc6c0.h`/`clc7c0.h` (Ampere), `clc9c0.h` (Ada), `clcbc0.h` (Hopper), `clcdc0.h` (Blackwell). The Ada header (`clc9c0.h`) is a 29-line stub that only defines the class number `0xC9C0` and inherits the Ampere method set, so the definitions we actually read live in the Ampere headers: 在这里看起来有两条启动路径相关。每一个计算等级 (compute class) 都有它们各自的 methods 定义—— `clc3c0.h` (Volta)，`clc5c0.h` (Turing)，`clc6c0.h`/`clc7c0.h` (Ampere)，`clc9c0.h` (Ada) 等等。因为 Ada 架构的文件仅仅是一个只声明了自身类编号并且继承于 Ampere 架构方法集的 29 行打酱油的文件，所以我们真正阅读查阅的是 Ampere 的定义文件： * `0x0318` — `SET_INLINE_QMD_ADDRESS_A` (defined as `NVC6C0_SET_INLINE_QMD_ADDRESS_A` in the Ampere header, inherited unchanged by Ada), which opens an inline-QMD burst streamed straight into the pushbuffer via `LOAD_INLINE_QMD_DATA(i)` (offset `0x0320 + i * 4`). / 此指令在 pushbuffer 中开启了一段流式的内联 QMD 写入爆发。 * `0x02b4` — `SEND_PCAS_A`, the out-of-line path, which carries only a pointer to a QMD that lives elsewhere in VRAM. / 这是一条离线路径指令，它只附带一个指向存放在显存其他地方的 QMD 的指针。 From the dump, we can figure out which one ends up in the pushbuffer. The dump shows the inline path: one increasing-method burst, count 66, opening at `SET_INLINE_QMD_ADDRESS_A`. The 66 words are the two address words (`SET_INLINE_QMD_ADDRESS_A`/`_B`, `0x0318`/`0x031c`) followed by 64 `LOAD_INLINE_QMD_DATA` words (`0x0320` onward) — a 256-byte QMD carried inline. Within it, word 12 is `0x1000` and word 18 is `0x100`: the 4096 and 256 of `vadd<<<4096, 256>>>`. 从转储数据中，我们能找出哪一条指令真正进了 pushbuffer。数据显示是内联路径 (inline path)：一个连续的增量 method 爆发写入，字数为 66，在 `SET_INLINE_QMD_ADDRESS_A` 处拉开序幕。这 66 个字其实由两个地址字组成，后面紧紧跟随 64 个 `LOAD_INLINE_QMD_DATA` 字 —— 这正好是一个妥妥的 256 字节的内联 QMD 结构体。在其内部，第 12 个字是 `0x1000`，第 18 个字是 `0x100`：即 `vadd<<<4096, 256>>>` 里的 4096 与 256。 ### Reading device memory & QMD Layout ### 读取设备内存与 QMD 内存布局 The Queue Meta Data (QMD) structure is represented as a multi-word layout, with fields defined as multi-word (MW) bits spanning 32-bit boundaries inside `src/common/sdk/nvidia/inc/class/cla0c0qmd.h`. The QMD stores several address-like fields, but they aren’t all the same kind of value: QMD 结构表现为一个多字 (multi-word) 布局，在其定义头文件里用夸张的跨界比特位 (MW) 描述。QMD 里存放着好几个类似地址的字段，但它们并不都是同一种值： * `PROGRAM_OFFSET` — MW(287:256) (Word 8) is a **32-bit** entry-point offset relative to the channel’s code base, not a 64-bit pointer. / 是一个相当于通道代码基址的 32 位入口偏移，并非完整的 64 位指针。 * `CONSTANT_BUFFER_ADDR_LOWER(i)` / `ADDR_UPPER(i)` — MW(959+i*64:928+i*64) (e.g. Constant Bank 0 holding our arguments sitting in Words 29–30). / 比如持有我们参数的常量存储区 0 就呆在字数索引的 29-30 处。 * `RELEASE0_ADDRESS_LOWER/UPPER` — MW(767:736) (Words 23–24), used for fences/semaphores. / 用于完工信号的围栏栅栏标志位。 * `CIRCULAR_QUEUE_ADDR_LOWER/UPPER` — MW(319:288) (Words 9-10). / 循环队列地址。 These point into device memory the CPU can’t read directly: a plain load faults, and both `cudaMemcpy` and `cuMemcpyDtoH` reject the address. 这些都指向 CPU 无法直接染指的设备内存：一般的普通读取会直接发生段错误崩溃，而正规的拷贝 API 也冷漠地拒绝接受这些地址。 So we need to read it with the GPU. A small kernel copies 512 bytes from a raw pointer into a buffer the host can fetch: 于是我们只能出动 GPU 自己来读了。我们写了一个小巧的内核，专门负责从一个裸指针那生抠出 512 个字节，抄进一个主机端能够触及的缓冲里： ```cuda __global__ void peek(const unsigned char* src, unsigned char* dst) { for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < 512; i += blockDim.x * gridDim.x) { dst[i] = src[i]; } } ``` Pointed at each of the QMD fields, exactly one returns all 512 bytes of the SASS. If we run a memory-scanning shim to look for valid GPU virtual addresses inside the QMD, we see a match at Word 48: 拿着这玩意儿对准 QMD 里的各大字段挨个扫描，结果只有一处准确无误地吐出了那 512 个字节的 SASS 真身。当我们在 QMD 内跑了个寻址探雷器时，发现在第 48 处字偏移上对上了号： ```txt qmd[48] -> 0x74167b272300 512 / 512 bytes match ``` Why is SASS matched at Word 48 (`qmd[48]`) when the driver’s program field is `PROGRAM_OFFSET` at Word 8? 为什么 SASS 真身能在第 48 个字处对上号，而在驱动程序明面标着程序入口的地方（第 8 个字处的 `PROGRAM_OFFSET`）却不行？ Word 8 holds only a 32-bit offset set by the driver, whereas words 48/49 are the hardware-owned `HW_ONLY_INNER_GET` (`MW(1566:1536)`) and `HW_ONLY_INNER_PUT` (`MW(1598:1568)`) fields. In the post-launch dump those words hold a full 64-bit GPU virtual address, and dereferencing the word-48 value returns the kernel SASS. The simplest reading is that the scheduler resolves the program offset into these scheduler-owned fields at launch. 那是因为第 8 个字只装了一个由驱动程序丢在那儿的 32 位干瘪偏移量。反观第 48 和 49 个字，那是硬件大人独享专属的领地（标注为 `HW_ONLY`）。在启动转储数据里，那两个词囊括了完完整整的 64 位 GPU 虚拟地址，将它解引用后，鲜活的内核 SASS 代码便呼之欲出。最合理的推断就是：在内核启动之际，调度器大笔一挥，将程序偏移量彻底解析成实际物理地址，并填进了它自己的私人辖区。 ### Decoding the driver’s ioctls ### 窥探驱动程序的 ioctls The command stream needs to be read from the memory, but `libcuda` sets up its memory and GPU objects the ordinary way: by running `ioctl` on the driver’s device files. `strace` on the one-kernel program records 948 of them[^23], almost all on two file descriptors — `/dev/nvidiactl` and `/dev/nvidia-uvm`: 尽管命令流需要靠读取内存来获取，但 `libcuda` 在初始化其内存与 GPU 对象时还是老老实实走了传统老路：通过对驱动设备文件砸下无数个 `ioctl` 调用。在这仅此一核的测试程序上，`strace` 足足录下了 948 次狂轰滥炸[^23]，且几乎全砸在两个文件描述符头上 —— `/dev/nvidiactl` 和 `/dev/nvidia-uvm`： ```txt $ strace -f -e trace=ioctl ./vadd ... ioctl(8, _IOC(_IOC_READ|_IOC_WRITE, 0x46, 0x2a, 0x900), ...) # /dev/nvidiactl ioctl(8, _IOC(_IOC_READ|_IOC_WRITE, 0x46, 0x2b, 0x30), ...) # /dev/nvidiactl ioctl(9, ...) # /dev/nvidia-uvm ... ``` The magic byte `0x46` is `'F'`, the NVIDIA resource manager’s ioctl magic[^24]. The command numbers decode against the open kernel modules’ `nv_escape.h`: `0x2A` is `NV_ESC_RM_CONTROL` and `0x2B` is `NV_ESC_RM_ALLOC`. 幻数 (magic byte) `0x46` 等于 ASCII 码的 `'F'`，正是 NVIDIA 资源管理器的 ioctl 专属印记[^24]。拿着开源内核模块中的头文件 `nv_escape.h` 一对暗号，便水落石出：`0x2A` 对应着 `NV_ESC_RM_CONTROL`，而 `0x2B` 则是 `NV_ESC_RM_ALLOC` 分配调用。 ### Decoding the SASS control words ### SASS 控制码揭秘 The stall counts, barriers and yield bits from the eligibility section come from a 21-bit control field `ptxas` packs into the top of each instruction’s second 64-bit word, which `cuobjdump -sass` prints next to the mnemonic: 我们此前在高能的“资格判定”章节里聊到的阻塞周期数、内存屏障与让权标志，全数来自 `ptxas` 打包进指令后半段那 64 个位顶部的一个 21 位的控制位域。`cuobjdump -sass` 在打印汇编助记符时，会在旁边以十六进制默默将其打出： ```txt 20 17 16 11 10 8 7 5 4 3 0 ┌────────┬───────────┬──────┬──────┬─┬──────┐ │ reuse │ wait mask │ read │write │Y│stall │ │ (4) │ (6) │ barr │ barr │ │ (4) │ └────────┴───────────┴──────┴──────┴─┴──────┘ ``` The two 3-bit indices name the scoreboard barriers the instruction sets, the 6-bit mask is the barriers it waits on, `Y` is the yield bit, and `stall` is the static cycle count. The layout is undocumented and reconstructed from microbenchmarking[^25]. 其中两个 3 位的索引名对应了该条指令“设置”的记分板读/写屏障，6 位的掩码则是它所“等待”的屏障组合，`Y` 是让权标志位，而 `stall` 便是那精确的静态时钟等待周期。这份神秘的位布局图从未被官方文档收录，纯靠外界硬核老哥们的微基准测试反向推理而来[^25]。 ### NVCC host registration callbacks ### NVCC 主机注册回调函数 If you want to see the exact code the compiler generates to register your GPU code at startup, compiling with `nvcc --keep` lets you inspect `vadd.cudafe1.stub.c`. 如果你想欣赏编译器在程序启动时为你生成的用于注册 GPU 核心代码的真容，加上 `nvcc --keep` 编译，然后去尽情赏鉴 `vadd.cudafe1.stub.c` 即可。 The start-of-process registration is handled by an automatically generated constructor: 程序刚起步那会的注册事宜，全权由一段自动生成的构造函数代劳： ```c // from vadd.cudafe1.stub.c static void __sti____cudaRegisterAll(void) __attribute__((__constructor__)); static void __nv_cudaEntityRegisterCallback(void **__T4) { __cudaRegisterEntry(__T4, (void(*)(const float*, const float*, float*, int))vadd, _Z4vaddPKfS0_Pfi, -1); } static void __sti____cudaRegisterAll(void) { __cudaRegisterBinary(__nv_cudaEntityRegisterCallback); } ``` The `__attribute__((__constructor__))` directive tells the linker to execute `__sti____cudaRegisterAll` before `main` starts. It registers our device binary with the CUDA runtime and schedules the callback. When executed, `__cudaRegisterEntry` maps the host function pointer `vadd` to the mangled device entry point `_Z4vaddPKfS0_Pfi`, building the hash table that `cudaLaunchKernel` queries at launch time. `__attribute__((__constructor__))` 编译制导指令勒令链接器：在 `main` 踏出第一步之前，无论如何给我先把 `__sti____cudaRegisterAll` 跑完。它向 CUDA 运行时庄重地挂号注册了我们的设备二进制文件，并排布了回调函数。当执行起来时，`__cudaRegisterEntry` 粗暴地把主机端函数指针 `vadd` 直接绑死在了被修饰过名字的设备端入口点 `_Z4vaddPKfS0_Pfi` 上。这正是 `cudaLaunchKernel` 之后在启动时查阅的那张神秘哈希表。 --- ### 注释 (Footnotes) [^1]: I didn’t check all of them. (作者注：我也没无聊到去检查那一百万个计算结果是不是全对了。) [^2]: An aside, this post is an instance of the ‘legibility transition’ that agents have engendered. (顺带一提，这篇文章本身就是一个关于现代极客们通过“抽丝剥茧式的解析化手段”探究底层黑盒的缩影。) [^3]: The prefix is the type: `%r` is a 32-bit integer, `%rd` a 64-bit one, `%f` a 32-bit float, `%p` a one-bit predicate. (前缀代表数据类型：`%r` 是 32 位整型，`%rd` 是 64 位，`%f` 是 32 位浮点，`%p` 是单比特的布尔谓语。) [^4]: `ncu` reports `launch__registers_per_thread = 16`. The disassembly only names up to `R9`, but the allocator reserves a few more for the ABI and alignment. (`ncu` 报告每个线程占用了 16 个寄存器。反汇编代码最高只点名用到了 `R9`，但分配器因为 ABI 和对齐原因多预留了几个。) [^5]: `cuobjdump -elf` shows a symbol table, a `.text.vadd` section holding the machine code, plus CUDA-specific sections like `.nv.callgraph`. (`cuobjdump -elf` 会展示符号表，一段存放机器码的 `.text.vadd`，以及 `.nv.callgraph` 这种 CUDA 专享节区。) [^6]: These offsets are the constant bank offsets `0x160`, `0x168`, `0x170`, and `0x178` that we saw our SASS machine code reading from constant bank 0 earlier. (这些偏移量跟我们早些时候在 SASS 机器码中看到从常量存储区 0 扒拉数据时的 `0x160` 到 `0x178` 严丝合缝。) [^7]: The usermode bit of the driver comes with the GPU’s kernel driver, not with the CUDA toolkit: the `libcuda.so.1` from the `strace` resolves to `libcuda.so.590.48.01`, the driver release on this machine. (驱动程序的用户态部分是随着 GPU 的内核驱动程序捆绑安装的，而不是 CUDA toolkit 的一部分。) [^8]: Controlled by `CUDA_MODULE_LOADING`. It shipped opt-in in CUDA 11.7 and defaulted to `EAGER` for years; the 12.x series flipped the default to `LAZY` (which can be overridden if you want loading costs paid up front). (受 `CUDA_MODULE_LOADING` 环境变量控制。11.7 引入时默认是 `EAGER` 提前加载，12.x 系列翻转了默认值为 `LAZY` 惰性加载。) [^9]: In this case, base is a GPU virtual address pointing to host memory. (在此案例下，基址 (base) 实际上是一个直戳戳指回主机内存的 GPU 虚拟地址。) [^10]: How I know it’s this method, given that `libcuda` is closed source: see the appendix. (既然 `libcuda` 闭源，我是怎么知道是这个 method 呢：详见附录揭秘。) [^11]: The part of the GPU’s control logic that interfaces with the host. (GPU 控制逻辑结构中专门负责与主机 (CPU) 接口交际的那部分大脑。) [^12]: They used to: older GPUs snooped USERD, so writing `GP_PUT` was enough. Turing and later don’t, so the driver rings the doorbell instead. (以前是主动监控的：老派 GPU 会主动嗅探 USERD，因此单单写入 `GP_PUT` 就足够了。Turing 架构之后就不干了，所以驱动程序必须摁响门铃叫醒它。) [^13]: Sometimes still called the GigaThread Engine. (有时候依然被人们敬称为 GigaThread Engine - 十亿级线程引擎。) [^14]: NVIDIA’s AD102-300-A1 SKU disables 16 of the physical 144 SMs on the full die to maximize manufacturing yield, as detailed in the NVIDIA Ada GPU Architecture whitepaper. (白皮书里提到：老黄在 AD102 满血版核心上强行阉割掉了 16 个 SM 来确保良品率。) [^15]: Since Volta, the model goes finer still — each *thread* carries its own program counter and call stack (Independent Thread Scheduling), letting threads in a warp diverge and reconverge freely. (自从 Volta 架构开始，硬件模型精细得丧心病狂——每个单独的*线程*都背着自己的程序计数器和调用栈，让同一个线程束内的线程放飞自我地分化甚至重聚合。) [^16]: `cudaGetDeviceProperties` tells you this information. (通过 `cudaGetDeviceProperties` API 可以探听虚实。) [^17]: The clearest public reconstructions are the Citadel microbenchmarking papers and these `maxas` control-code notes for Maxwell. (最清晰的民间结构逆向来自于学术界的微基准论文以及 `maxas` 的公开文档笔记。) [^18]: The clearest public reconstructions are the Citadel microbenchmarking papers and these `maxas` control-code notes for Maxwell. (同上。) [^19]: The RTX 4090 uses GDDR6X memory rather than the High-Bandwidth Memory (HBM) found in datacenter-class GPUs like the A100 or H100. (因为它是 RTX 4090，所以它吞下的是 GDDR6X，而不是 A100/H100 高端机房卡上那尊贵的 HBM 显存。) [^20]: In principle anyway, we’ll see later that `c[i]` never hits VRAM. (纯理论上说它是原路返回的，只不过稍后我们会发现 `c[i]` 的写存压根就没被憋到 VRAM 这一层。) [^21]: Only the two inputs cross the bus, not the full 12 MB. `ncu` shows 8.4 MB read from DRAM and essentially nothing written: the 4 MB output `c` fits in the 72 MB L2 and isn’t flushed to DRAM until the later device-to-host copy reads it back. (只有输入数据跨过了总线。输出的 4MB 数组 `c` 舒舒服服地躺在 72MB 大小的 L2 缓存里，还没来得及刷入底层 DRAM，就被后续的回传程序端走了。) [^22]: A pinned-memory `cudaMemcpyAsync` would skip the wait and let the host run ahead. (如果你有钱任性地用了锁页内存搭配 `cudaMemcpyAsync` 的话，主机老大哥就根本不会停下来等，而是直接一溜烟跑了。) [^23]: Almost all are one-time setup; a steady launch loop makes far fewer. (它们绝大部分全是在做一次性的初始化开荒；如果你搞个稳定的无限循环启动，调用量就会急剧收缩。) [^24]: The ‘magic’ byte is a value every NVIDIA ioctl carries as a sanity check; see the Linux kernel documentation. (这是为了完整性自检所携带的神奇印记，详情可见 Linux 内核文档。) [^25]: The clearest public reconstructions are the Citadel microbenchmarking papers and these `maxas` control-code notes for Maxwell. (前文已提，致敬民间逆向大神。)