一个最小 CUDA 例子：vadd<<<4096, 256>>>(da, db, dc, n)，对 1,048,576 个 float 做加法。

看上去，这只是 CPU 调了 GPU 上的一个函数。麻烦也在这里：它最不像函数调用的地方，恰恰是 CUDA 最容易被低估的地方。

这行代码背后，至少经过三段接力：编译器生成设备代码和主机端 stub，CUDA runtime 和 driver 查表、打包参数、提交命令，GPU 再从队列里取活，调度 warp 执行。

我更在意的判断是：kernel launch 不是“一次调用”，而是一套跨 CPU/GPU 的执行协议。理解这一点，很多小 kernel 慢、首次运行慢、低延迟推理成本高的问题，才有正确的排查方向。

nvcc 做的不是一次编译，而是拆成几种产物

原文用 RTX 4090 和 -arch=sm_89 编译这个向量加法。用 nvcc --keep 留下中间文件后，会看到一个很关键的事实：nvcc 更像调度器，不是单一编译器。

它把 CUDA 源码拆成主机端和设备端两条线。主机端给 CPU 编译，设备端再走 PTX、SASS、fatbin 这一套。

环节主要产物作用该怎么理解
ciccPTX生成设备侧虚拟 ISA便于兼容，不是最终硬件机器码
ptxasSASS / cubin面向 sm_89 生成机器码RTX 4090 实际执行的是这层
fatbinaryfatbin打包 cubin 和 PTX兼顾当前架构性能和后续兼容
host compiler主机对象文件嵌入 fatbin，生成 launch stub让 CPU 侧能发起 kernel launch

PTX 常被叫成“GPU 汇编”，但这个说法容易误导。更准确地说，PTX 是 NVIDIA 的虚拟 ISA。它隐藏了很多真实硬件细节，比如具体寄存器数量、指令调度和部分寻址细节。

SASS 才是面向具体架构的机器码。对 sm_89 来说，也就是 RTX 4090 真正执行的那一层。

fatbin 里同时放 cubin 和 PTX，也不是为了好看。cubin 服务当前目标架构，PTX 留给驱动在缺少匹配 cubin 时做 JIT。好处是兼容，代价是首次加载或首次运行可能多出一段等待。

对 CUDA 开发者来说，这里有一个很实用的动作：遇到“第一次慢、后面快”，不要只盯 kernel 代码。要把模块加载、PTX JIT、上下文初始化一起纳入排查。

<<< >>> 后面，是 stub、参数缓冲和驱动查表

vadd<<<4096,256>>> 在源码里很短。编译之后，它会变成主机端 launch stub。

这个 stub 不负责做加法。它负责把 da、db、dc、n 这些参数打包，交给 CUDA runtime。runtime 再根据注册信息，找到 fatbin 里对应的设备端 kernel。

这里最容易想错的是：主机端的 vadd 不是 CPU 真的去执行一个同名函数。它更像一个索引入口。runtime 和 driver 通过这个入口，定位设备代码、准备参数、提交任务。

随后流程会进入用户态驱动 libcuda.so.1，再和内核态 nvidia.ko 通过设备文件、ioctl 等机制交互。

原文提到，一次运行会牵涉大量 CPU 指令、多个设备文件、数百次 ioctl，以及一个内存映射的 doorbell 寄存器。这个数字要谨慎读。它会受驱动版本、CUDA 版本、上下文状态、模块加载策略影响，不能直接拿来当性能结论。

能下的结论是另一件事：CUDA launch 不是一次普通系统调用，也不是 CPU 把参数压栈后让 GPU 跳过去执行。

这对两类人最具体。

一类是写高性能 CUDA 的工程师。小 kernel 很多、launch 很频繁时，应该测 launch 开销，而不是只调 occupancy、访存合并和寄存器数量。该融合 kernel 就融合，该用 CUDA Graphs 就用 Graphs，计算量太小还硬拆，往往得不偿失。

另一类是做推理服务和框架的人。小 batch、低延迟场景里，PyTorch、TensorRT、Triton 这类上层工具能不能减少 launch 次数，会直接影响尾延迟和机器利用率。团队要调的未必是某条 SASS 指令，可能是图捕获、算子融合、常驻 kernel 和调度粒度。

GPU 不是被调用，而是从队列里取任务

CPU 最后并不是“命令 GPU 立刻执行函数”。更接近事实的说法是：驱动把任务写进队列，CPU 敲一下门，GPU 自己来取。

这条路径里有几个关键词：pushbuffer、GPFIFO、QMD、doorbell。

pushbuffer 里放的是给 GPU 的命令。GPFIFO 指向这些命令。QMD 描述一次 kernel 启动需要的信息，包括入口、grid/block 形状、参数位置等。doorbell 则用来通知 GPU：队列里有新任务。

SASS 里从 c[0x0][...] 读取参数的动作，也不是凭空来的。那些参数位置，是 launch 过程中由 runtime 和 driver 布置好的。

把它和 AMD ROCm 的 HSA queue，或者 Linux 图形栈里的命令缓冲放在一起看，会更好理解。现代加速器大多不是“被函数调用”，而是“被投递任务”。CPU 负责准备任务，GPU 负责消费队列。

限制也要说清楚。NVIDIA CUDA 的用户态和内核态驱动长期比较封闭，开源内核模块、公开资料和反汇编工具只能拼出一部分机制。pushbuffer、GPFIFO、QMD 这些结构能帮助理解路径，但不该被当成稳定、公开、可依赖的 ABI。

接下来真正该盯的，也不是某次 ioctl 到底有多少个。更该看三个变量。

观察变量为什么重要对开发者的动作
lazy module loadingCUDA 12.2 之后默认延后部分模块加载成本首次请求、首次 kernel 运行要单独测
CUDA Graphs / 图捕获能摊薄重复 launch 的固定成本推理服务和小 kernel 管线应优先评估
kernel fusion / persistent kernel减少提交次数或让任务常驻适合小 batch、低延迟、碎片化算子场景

这条链路不会让每个 kernel 自动变快。但它能减少误判。

如果瓶颈在 launch 路径上，继续抠单个 kernel 的几条指令，收益会很有限。工欲善其事，先要知道刀钝在哪里。