你在 nvprof 里看到一个 kernel 跑了 10 微秒。直觉上这已经很快了——但你的同事说，按照纸面带宽算，这个向量加法应该能在 8 微秒内完成。差的 2 微秒去哪了？你开始怀疑是不是 launch overhead，还是 occupancy 不够，还是 DRAM 没跑满。
6 月底，一篇名为「What happens when you run a CUDA kernel」的技术长文登上了 Hacker News 首页。作者 Fergus Finn 用一台 RTX 4090、一个只有 20 行代码的向量加法 kernel，从头到尾追踪了一次 CUDA 启动的完整路径——从 nvcc 的编译管线到 GPU 硬件上 warp 调度器的 stall 计数器。这场追踪最终发现：为了在屏幕上打印 c[0]=2.000000，主机执行了数千万条 CPU 指令、打开了两个设备文件、发出了 948 次 ioctl、写了一次内存映射的门铃寄存器。我们跟随这篇文章的路径，把每个环节拆开来看。
编译：四层编译器一条龙
一个 .cu 文件在变成 GPU 能执行的机器码之前，要经过三道编译器。nvcc 先调用 cicc（基于 LLVM）把 device code 编译成 PTX——一种设备无关的虚拟 ISA，拥有无限个类型化寄存器。然后 ptxas 将 PTX 转换为当前架构的 SASS 汇编，在这个阶段虚实寄存器被映射到硬件寄存器（vadd 从 10 个虚拟寄存器坍缩为 7 个物理寄存器），多次地址运算被融合为单条 IMAD.WIDE。最后 fatbinary 将 SASS 与 PTX 一同打包进 ELF 的 .nv_fatbin 区段，由链接器嵌入最终的可执行文件。PTX 作为后备随行——如果你把同一个二进制拿到 cubin 不兼容的新架构 GPU 上跑，驱动会在加载时 JIT 编译 PTX。
值得留意的是，kernel 代码的 GPU 上传并不是在程序启动时发生的。自 CUDA 12.2 起，模块加载默认懒执行：驱动只在第一次实际 launch 该 kernel 时，才把 cubin 里的 SASS 拷贝到显存。在此之前，编译好的机器码就静静躺在宿主二进制里。
主机侧：从 chevron 语法到门铃
当编译器遇到 vadd<<<4096, 256>>>(da, db, dc, n) 这个 chevron 语法时，它将其替换为一个自动生成的 launch stub。这个 stub 把四个参数按固定偏移（0、8、16、24 字节）打包进参数缓冲区，然后调用 __cudaLaunch。
__cudaLaunch 以宿主端 dummy 函数指针为 key，在注册表中查到对应的 mangled device 符号名，随后进入闭源的用户态驱动 libcuda.so.1。从这一刻起，事情进入了「跟内核模块对话」的领域——驱动通过 /dev/nvidiactl 发出 948 次 ioctl（大部分是一次性初始化），将一次 launch 转化为 GPU 能理解的命令流。
命令流写入的是宿主内存中的两块核心结构：pushbuffer 和 GPFIFO。pushbuffer 承载着 GPU 原生命令编码的 method（寄存器地址 + 值对），GPFIFO 是一个环形指针缓冲，每个条目指向 pushbuffer 的一个跨度。驱动填充 pushbuffer、更新 GP_PUT 指针后，向 GPU 的一个 MMIO 寄存器——也就是门铃——写入 channel 的 work-submit token。GPU 的 host engine 被唤醒，通过 DMA 拉取 method 流，从中解析出一个 QMD（Queue Meta Data）——包含了 grid/block 维度、寄存器/共享内存需求、SASS 入口偏移和常量 bank 地址的启动描述符。
cuLaunchKernel 在门铃响起的那一刻就返回了。整个 launch 是异步的——CPU 继续往前跑，GPU 还没开始干活。
GPU 侧：分发、调度、等待
host engine 将 QMD 交给 GPU 上唯一的 compute work distributor（有时仍称 GigaThread Engine）。面对一个 4096 blocks × 256 threads 的启动配置和 128 个 SM，它的任务是把工作均匀铺开。
每个 SM 能容纳多少线程块由三项硬件上限共同决定：最大活跃线程数（AD102 为 1536）、寄存器文件大小（65536 个 32 位寄存器）和共享内存总量（100 KB）。vadd 每个 block 256 线程、每线程 16 寄存器，寄存器侧能塞 16 个 block，但线程数上限只允许 6 个。最终每个 SM 常驻 6 个 block、48 个 warp。这 48 个 warp 被均分给 SM 的 4 个 warp 调度器——每个调度器管理 12 个候选 warp，每周期从中选一个发射一条指令。
一个 warp 能否被发射，取决于硬件上一个极简的调度机制。SASS 指令以 128 位编码，其中高 21 位是 ptxas 写入的控制负载，包含三类信号：
- 静态 stall 计数：对于延迟可预测的算术指令，编译器直接编码等待周期数。FADD 带着
stall=5，意味着发射后 warp 被冻结 5 个周期，等 R9 写完再继续。 - yield 提示位：单比特，告诉调度器该 warp 是否应该让出优先级。
- 依赖屏障：6 个硬件 scoreboard 屏障，编号 0-5。
LDG.E（全局内存加载）是不可预测延迟操作——两条LDG.E各自写入 B2，而FADD携带waits-on B2。B2 不清，调度器就跳过这个 warp。
这种「编译器预测一切可预测的，硬件仅覆盖编译时无法预判的部分」的设计，使得 GPU 能以极小的芯片面积代价实现大规模线程级并行。GPU 不用乱序执行，它用海量 warp 的快速切换来隐藏延迟。这也是为什么 occupancy 对性能至关重要——只有足够多的活跃 warp，调度器才有足够的「其他选择」来填满那些等待的时钟周期。
访存：合并、缓存、还有你看到的带宽
当 warp 最终发出 LDG.E 时，32 个线程各自计算地址。因为它们访问连续的 float 数组，每线程 4 字节，整个 warp 请求的是一个连续的 128 字节块。SM 的 load/store 单元将 32 个 4 字节请求合并为 4 个 32 字节扇区请求——恰好填满缓存行。如果地址不连续，同样的数据量可能产生 32 次独立访存。
合并后的请求先查 L1，未命中则通过 crossbar 路由到 L2。RTX 4090 有 72 MB 的分布式 L2。在 vadd 的 ncu 剖析中，DRAM 利用率达到 79.65% 峰值——原因是这个 kernel 的算术强度极低：每传输 12 字节数据（两次 load 一次 store），只执行一次浮点加法。性能天花板就是内存总线。
另外，剖析显示写入的 4 MB 输出 c 从未真的到达 DRAM。它们滞留在 L2 中，直到后续的 cudaMemcpy Device→Host 才被直接读出并跨 PCIe 送回宿主，省了一趟显存往返。
驱动是 GPU 的「脏活」承揽者
这篇文章在 HN 社区引发的一个焦点讨论，恰恰是关于路径上最不透明的那一段——驱动。一位评论者直言，在大规模运行 CUDA 时，处理 NVIDIA 驱动和库的 bug 占据了「令人厌恶的巨大比例的工程师时间」。而曾在 Qualcomm 编写 OpenCL 驱动的 david-gpu 则反驳说，他在任期间收到的所有驱动 bug 报告，最终都定位到应用代码的 bug。
这两种说法可能都对。驱动承担了大量硬件 workaround（用 david-gpu 的话说，这是「行业的脏秘密」），其复杂度本身就构成了 bug 的温床。另一个评论则指出，GPU 驱动对错误 kernel 代码的鲁棒性远不如 CPU 操作系统——一个死循环的 kernel 可以卡死显示输出，而现代驱动的看门狗机制本质上只是「症状治疗」。对 GPU 编程者来说，理解驱动在做什么既关乎性能调优，也关乎理解故障边界。
回到那个 2 微秒
有了全路径的视角，那差的 2 微秒就不那么神秘了：launch overhead 确实有，但一次冷启动之外它很小；occupancy 是满的（48 warps/SM）；问题出在算术强度——这个 kernel 的生命就是等内存。唯一的优化方向是融合算子，让数据在寄存器里多待一会儿。
10 微秒跑完一百万个加法，其实已经很快了。但知道为什么是 10 微秒而不是 8 微秒，才是这篇文章的价值所在。
本文的素材来自公开信息和社区讨论。如果你对这个话题有更深入的一手经验，欢迎指出文中的不足。