GPU 执行模型 1：CPU 与 GPU 前端

基础概念

Host and Device 概念

GPU 前端
- HIU（Host Interface Unit)：负责监听门铃，启动 DMA 从 PushBuffer 拉取指令放到 FIFO 中。
- FIFO 缓冲区：位于 GPU 硅片，按队列切分的片上缓存，存放 HIU 从 PushBuffer 拉来的指令。
- CP（Command Processor）：前端大脑。轮询所有活跃的 FIFO，取出指令分发到合适的执行引擎。
- Syncpoint：GPU 片上的纯硬件同步寄存器（整数计数器）。
GPU 后端
- GTE（GigaThread Engine）：处理 kernel launch 的引擎，把 kernel 计算任务分发到 SM。
- SM（Streaming Multiprocessor）：具体执行计算任务的单元。
- CE（Copy Engine）：负责内存 DMA 搬运。
Host 端内存
- context
  
  Context 非常像操作系统里的进程（Process），是资源容器。包括：
  - 显存页表
  - 硬件队列（见下）句柄
  - 与硬件队列（见下）对应的 PushBuffer
- PushBuffer
  
  PushBuffer 也叫 command buffer。
  
  其中存放：
  - GPU Operation Launch 指令
  - [Signal Syncpoint X]
    
    GPU 的 CP 执行某个 FIFO 的 signal syncpoint 指令后：
    - 底层执行引擎（如 GTE）会被注册上一个回调。
    - 当底层执行引擎中该 FIFO 的上一个操作完成后，底层执行引擎会使某一个 syncpoint 硬件寄存器的值置 1。
  - [Semaphore / Timestamp Write to Host Memory]
    
    GPU 的 CP 执行某个 FIFO 的 host 内存写指令后：
    - 底层执行引擎（如 GTE）会被注册上一个回调。
    - 当底层执行引擎中该 FIFO 的上一个操作完成后，底层执行引擎会通知 CP
    - CP 收到信号，使用 DMA 跨越 PCIe 总线，使 host 内存空间中的值发生变化
  - [Wait Syncpoint X]
    
    GPU 的 CP 执行某个 FIFO 的 wait syncpoint 指令时：
    - 先检查 syncpoint 硬件寄存器的值是否为 1
      - 如果不是 1
        
        该 FIFO 会被挂起，该 FIFO wait syncpoint 指令后面的指令不会被执行
        
        当 syncpoint 硬件寄存器的值变为 1 之后，该 FIFO 会被唤醒，CP 继续执行 wait syncpoint 指令后面的指令
      - 如果是 1，CP 继续执行该 FIFO
  - [Wait For Context Idle]
    
    GPU 的 CP 执行某个 FIFO 的 wait for context idle 指令时：
    - 先检查 context 的所有微观物理活动是否归 0
      - 如果不是
        
        所有 FIFO 被挂起
        
        CP 等待 context 的所有微观物理活动彻底归零
        
        唤醒所有 FIFO
        
        继续执行
      - 如果是，继续执行
  UMD（User Mode Driver）将指令存入 PushBuffer，PushBuffer 中的指令被 GPU 取走。
- PCIe 门铃寄存器
  
  实际上存在于 GPU 上，通过 MMIO 映射到 Host 地址空间中。
  
  KMD（Kernel Mode Driver）写 Host 地址空间中的门铃寄存器，通知 GPU。

硬件队列的概念

硬件队列的概念需要单独进行强调。

硬件队列

当你听到「这台 GPU 支持 32 个硬件队列」时，这里的「队列」不仅仅指 GPU 内部的一块缓存。它指的是一条端到端的、完整的通信上下文。

这个实体有很多叫法：hardware work queue、hardware channel、hardware connection。

一个完整的「硬件队列」包含三样东西：
- Host 端的 PushBuffer：UMD 会为每一个硬件队列，在 Host 内存中分配一块独立的 PushBuffer。
- GPU 端的门铃寄存器：每个硬件队列在 GPU 上有专属的 PCIe 门铃（Doorbell）寄存器，这类寄存器通过 PCIe 被映射（MMIO）到了 Host 的虚拟地址空间内。
- GPU 端的硬件状态：驱动分配给一个硬件队列的唯一 ID，以及它的运行状态记录，其中包括一个前端 FIFO 缓冲区。
  - 前端 FIFO 缓冲区
    
    这个缓冲区物理存在于 GPU 前端硅片上的一小块 SRAM 中。
    
    在现代 GPU 的前端控制器中，片上的 FIFO SRAM 空间是被动态切分的。
    
    当 CP 决定同时激活 4 个硬件队列时，这块 SRAM 就会在物理或逻辑上被切分成 4 个独立的 FIFO 槽。
    
    HIU 会并发地从 Host 端的 PushBuffer 1、2、3、4 往片上的 FIFO 1、2、3、4 里灌水。
    
    CP 则像一个打地鼠的玩家，飞快地在片上的 FIFO 1、2、3、4 之间轮询抽水，谁有数据就读谁。
Hyper-Q

在 Hyper-Q 诞生之前，物理上只有 1 个硬件队列。这意味着，不管你在软件层写了多少个并行的 CUDA Stream（Stream 的概念见下），到了硬件前端，它们全部要像漏斗一样被硬塞进这唯一的队列里排队。只要属于 Stream A 的任务还在队列头部，Stream B 的任务就得死等，这就导致了严重的虚假序列化 (False Serialization)，无法实现真正的并发。

Kepler 架构中引入了 Hyper-Q (Hyper Queue) 技术。它在物理层面上，将硬件队列，从 1 个暴力扩充到了 32 个。有了这 32 个独立的队列，你在软件层创建的不同 Stream 就可以被映射到不同的硬件队列中。CP 可以同时监控这 32 条队列，从而让不同 Stream 的 Kernel 得以真正地并发进入 GigaThread Engine。

软件概念

stream

stream 是软件层面的抽象。CUDA runtime 和 UMD（User Mode Driver）将其暴露给程序员用来表达并发性（concurrency）和依赖关系。

event

如果说 Stream (流) 是装载 GPU 任务的传送带，那么 Event (事件) 就是你插在传送带上的“旗帜”。

相关 api
- 创建 event
  - cudaEventCreate(&event_ready)
  - cudaEventCreateWithFlags(&event_ready, cudaEventDisableTiming);
- 插入 event
  - cudaEventRecord(event_ready, stream)
- 计算时间差
  - cudaEventElapsedTime(&milliseconds, start, stop);
- 等待 event
  - CPU 等待 event：cudaEventSynchronize(event_ready)
  - 某个 stream 等待 event：cudaStreamWaitEvent(stream_B, event_ready, 0);

api 使用

在写 CUDA 代码时，Event 主要承担两大核心职能：

精准测速 (Profiling / Timing)

如果你想测量某个 Kernel 或者某几个 API 执行到底花了多少时间，千万不要用 CPU 的 timer (比如 std::chrono) 去算。因为 Kernel Launch 是异步的，CPU 测出来的只是“把指令塞进 PushBuffer 的时间”，而不是 GPU 真实计算的时间。

正确做法是用 Event 掐表：

cudaEvent_t start, stop;
cudaEventCreate(&start); cudaEventCreate(&stop);

cudaEventRecord(start, stream);    // 动作 1：在 stream 中插一个 start 信号弹
kernel<<<..., stream>>>();         // 动作 2：执行计算
cudaEventRecord(stop, stream);     // 动作 3：在 stream 中插一个 stop 信号弹

cudaEventSynchronize(stop);        // 动作 4：CPU 强制等待 stop 信号弹被触发
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop); // 计算时间差

跨流同步 (Cross-Stream Synchronization)

假设你有两个独立的流：Stream A 在做数据预处理，Stream B 在做模型推理。推理必须等预处理做完才能开始。

如果用 cudaDeviceSynchronize()，会让整个 GPU 停转，代价太大。用 Event 就可以实现纯 GPU 内部的轻量级依赖：
```
// Stream A 的流水线
kernel_preprocess<<<..., stream_A>>>();
cudaEventRecord(event_ready, stream_A); // 处理完后，插个旗子

// Stream B 的流水线
// 告诉 Stream B：在执行后续任务前，必须等 event_ready 这个旗子立起来
cudaStreamWaitEvent(stream_B, event_ready, 0);
kernel_inference<<<..., stream_B>>>();
```
注意：cudaStreamWaitEvent 是一个非阻塞 CPU 的调用，它只是向 Stream B 注入了一个等待指令。

cudaEventRecord 往 PushBuffer 中注入的指令

cudaEventRecord 会根据 event 创建方式的不同，往 PushBuffer 注入不同指令：
- 默认事件 cudaEventCreate
  
  如果是用默认方式 cudaEventCreate(&event_ready) 创建的事件，UMD 在面对 cudaEventRecord(event_ready, stream_A) 时，必须做一个保守的假设：「我不知道未来是 GPU 里的另一个 Stream 会等它，还是外面的 CPU 会等它，又或者 CPU 想计算它的耗时。所以我必须把两手准备都做好。”
  
  因此，UMD 会在 Stream A 的 PushBuffer 里塞入两个连续的硬件动作：
  - 动作 A (给 GPU 内部看的)：[Signal Syncpoint X]
  - 动作 B (给 CPU 留的后路)：[Semaphore / Timestamp Write to Host Memory]
- 轻量级事件 cudaEventCreateWithFlags(&event, cudaEventDisableTiming)
  
  如果你非常确定你创建的这个 event 只用于 cudaStreamWaitEvent 进行流之间的同步，而绝对不会被 CPU 用来做 cudaEventSynchronize，也绝对不需要用 cudaEventElapsedTime 算时间，你应该用下面这种方式创建事件：cudaEventCreateWithFlags(&event_ready, cudaEventDisableTiming);。这样就可以告诉驱动，这是一个纯粹的、无计时的、轻量级同步事件。
  
  加上 cudaEventDisableTiming 这个 Flag 后，UMD 瞬间懂了你的意图。当你再调用 cudaEventRecord(event_ready, stream_A) 时，UMD 会直接砍掉那个昂贵的动作 B。
  
  此时，PushBuffer 里只会被注入孤零零的一条指令：
  - [Signal Syncpoint X]
  至此，没有任何 PCIe 流量产生，也没有多余的 Host 内存写入。

整体执行流程

第零阶段：Context

当你的程序启动并调用 CUDA 时，驱动已经在 GPU 上为你创建了一个 Context（上下文）。

这个 Context 是资源的沙箱，

它不仅掌管你的显存页表，
还预先从 GPU 前端申请了一批硬件队列和门铃寄存器，
并在主机的锁页内存中为每一个队列分配了一个专属的 PushBuffer。

第一阶段：Host 端

在这个阶段，CPU 上的用户态驱动（UMD）像是一个编译器和调度员，负责把高层的并发意图转化为底层的物理队列流。

1 Kernel Launch
- 你在代码中写下 kernel<<<grid, block, smem, stream_A>>>(args)。
- NVCC 编译器将其拆解，通过 CUDA Runtime API，最终穿透到底层的 Driver API：cuLaunchKernel。
  
  此时，Kernel 的参数、网格维度配置以及关联的软件 Stream 句柄，全部交给了用户态驱动（UMD）。
2 UMD 将 Stream 映射到硬件队列 1
- UMD 检查你传入的 stream_A。对于 UMD 来说，Stream 只是一个用来表达依赖关系的软件标签。
- UMD 维护着一个映射表，它会将这个软件 Stream 映射到当前 Context 分配的某一个物理硬件队列（例如队列 1）上。
- 注意：如果你的软件 Stream 数量超过了硬件队列上限，UMD 会让多个 Stream 共享同一个队列（退化为串行）。但只要队列够用，不同的 Stream 会被映射到不同的队列。
3 UMD 将信息写入硬件队列 1 对应的 PushBuffer
- UMD 知道队列 1 对应着 Host 内存中一块特定的 PushBuffer 1。
- UMD 将执行该 Kernel 所需的硬件配置寄存器状态（如 SM 资源需求），连同 Kernel 的启动指令，一起打包写入 PushBuffer 1。
- 可能的内存写指令
  
  在打包完 Kernel Launch 指令后，UMD 可能会悄悄在 PushBuffer 1 的末尾追加一条“内存写指令”（Semaphore Write / Release Fence），用于在未来通知 CPU 任务已完成。
  - PushBuffer 批量提交
    
    你 launch 一次 kernel，UMD 不会立刻去敲门铃惊动内核态（KMD）和 GPU，它会把指令缓存在 PushBuffer 里。
    
    当 PushBuffer 缓存积攒到一定大小，或者 UMD 决定这是一个合适的提交边界时（Kickoff）。UMD 会在这个批次的末尾自动追加一个进度围栏（Progress Fence / Semaphore Write）。
    
    这样做的目的是让驱动底层随时掌握 GPU 执行到了哪个批次，方便在后台回收资源，并在 Nsight Systems 中对齐时间轴。
  - 调用 cudaEventRecord 时
    
    cudaEventRecord 根据 event 创建方式的不同，可能会往 PushBuffer 注入内存写指令。具体见上。
  - 调用 cudaDeviceSynchronize 或 cudaStreamSynchronize 时
    
    具体见另一篇。
- 可能的 syncpoint 指令
  - 两个 kernel 前后调用
    
    当你在同一个 stream 中调用两个 kernel 时：
```
kernel_A<<<...>>>();
kernel_B<<<...>>>();
```
    CPU 端的 UMD 在向这个 Stream 对应的 PushBuffer 写入指令时，绝对不会只是简单地把 A 和 B 的 Launch 指令挨着放。为了强制硬件串行，UMD 会在它们之间强行插入硬件级的依赖屏障。
    
    PushBuffer 里的真实指令流大概长这样：
```
[Launch Kernel A]
[Signal Syncpoint X]
[Wait on Syncpoint X]
[Launch Kernel B]
```
  - cudaEventRecord cudaStreamWaitEvent
    
    cudaEventRecord 一定注入 signal 指令，cudaStreamWaitEvent 注入 wait 指令，具体见上和另一篇。
4 KMD 通过 PCIe 写入硬件队列 1 的 GPU 寄存器（敲门）
- 指令打包完毕后，UMD 通过系统调用陷入内核，将控制权交给内核态驱动（KMD）。
- KMD 写 host 端地址空间内，队列 1 的门铃寄存器。底层的硬件总线控制器会自动把这个写操作转换成 PCIe 事务，发给 GPU。队列 1 在 GPU 上的门铃寄存器最后被写入。
- CPU 敲完门铃后立刻返回（异步执行），去处理下一行代码。

第二阶段：GPU 硬件前端

在这个阶段，指令跨越 PCIe 总线抵达 GPU 硅片。

1 HIU 监听到队列 1 的门铃，主动通过 DMA 把队列 1 对应的 PushBuffer 内容拉到队列 1 的前端 FIFO 中
- GPU 前端的 HIU (主机接口单元) 侦听到队列 1 的门铃响了。
- HIU 知道队列 1 的 PushBuffer 在 Host 内存的具体物理地址。它启动内部的 DMA 引擎，主动跨过 PCIe，将 PushBuffer 1 里的二进制指令流拉取到 GPU 上队列 1 对应的前端 FIFO 缓冲区中。
2 CP 从所有 FIFO 中取指令，并进行识别与分发
- Command Processor (命令处理器) 是前端的大脑。得益于 Hyper-Q 技术，它可以并发地从当前 context 下所有活跃的硬件队列的 FIFO 缓冲区中读取指令。方式为轮询（现代 CP 不仅仅是无脑轮询，它内部带有硬件调度器，会根据不同 Context 或 Stream 的优先级，比如你用 cudaStreamCreateWithPriority 创建的高优先级流，来决定多给哪个 FIFO 时间片）或者内部仲裁机制并行感知。
- 指令解码与类型识别
  
  CP 在读取到指令后，会解析指令的头部操作码 (Opcode)。它能够精准识别出当前这条指令属于哪种 GPU 操作：
  - 这是一个 Compute 操作？（比如我们写的 Kernel Launch）
  - 这是一个 Memory Transfer/Copy 操作？（比如 cudaMemcpyAsync）
  - 这是一个 Graphics 操作？（比如 OpenGL/Vulkan 的 Draw Call）
  - 还是一个 Synchronization 操作？（比如 Event Record / Wait）
- 分发：识别出指令类型后，CP 就变成了交通警察。它不会把所有东西都丢给同一个地方，而是根据指令类型，将任务的元数据分发给 GPU 内部不同的专属硬件引擎。
3 专属硬件引擎执行任务
- 分支 A：计算任务 (Compute) 交给 GigaThread Engine (GTE)
  
  如果 CP 发现这是一条 Kernel Launch 指令，它会提取出网格维度、寄存器需求等元数据，推送给全局硬件调度器 GigaThread Engine。
- 分支 B：内存拷贝 (Transfer) 交给 Copy Engine (CE / DMA 引擎)
  
  如果 CP 发现这是一条 cudaMemcpyAsync 指令（例如从 Host 拷贝数据到 Device），它绝对不会把它发给 GTE，因为 SM 是用来算数的，做内存搬运太浪费了。
  
  CP 会将这条拷贝指令交给 GPU 上专门负责搬砖的硬件模块——Copy Engine (CE)。现代 GPU 通常有多个独立的 CE（例如一个专门负责 PCIe 上行，一个负责下行，还有负责 NVLink 传输的）。
  
  CE 会在后台默默地通过 DMA 直接搬运内存，完全不需要 SM 的参与。
- 并发执行的终极体现 (Overlap)
  
  正是因为 CP 做了这种“分类移交”，才使得不同类型的操作可以在硬件底层完美重叠 (Overlap)。
  
  假设 Stream A 在执行 Kernel，Stream B 在执行 Memcpy：
  
  CP 会把 Stream A 的指令发给 GTE（占用 SM 算力），把 Stream B 的指令发给 CE（占用 PCIe 和显存带宽）。两者在物理电路上互不干扰，实现了 Compute 和 Data Transfer 的完美隐藏与并行。
4 硬件引擎执行完毕以后，CP 执行内存写指令或者 syncpoint 指令
- 当 GTE 报告某个 Kernel 的所有 Block 都在 SM 上执行完毕，或者 CE 执行完内存操作，后。
- 前端的 CP 会继续执行之前 UMD 打包在 PushBuffer 尾部的“内存写指令”或者 syncpoint 指令。

NSYS 讨论

Kernel Launch 和 Kernel Execution 对应关系

CPU 上执行的每一个 launch API 调用都和一个 GPU kernel execution 对应。

两者通过 correlationId 关联，可以在 nsys 中观察到这个关联。这个 ID 实际上就是在第一阶段第 3 步，UMD 往 PushBuffer 里写指令时打上的一个序列号标签。

另外值得注意的是，nsys 里 nvtx range 不一致的问题。即：一个 process 的 CUDA HW 和 Thread 对同一个 NVTX range 显示的时间范围不一样。

Thread 的 range 显示的是 CPU 上执行 range_push 和 range_pop 的时刻。由于异步执行，如果 range_pop 之前 CPU 没有和 GPU 同步（即等待 GPU），那么 range_pop 时刻之后，range_pop 之前的 launch 的 kernel 可能还在 GPU 上跑。

CUDA HW 的 range 显示的是 thread 的 range 中 launch 的 kernel 在 GPU 上执行的时刻。CUDA HW range 的开始时刻，是 CPU 在 thread range 里 launch 的第一个 GPU 操作（kernel 或 memcpy），在 GPU 上开始执行的时刻。CUDA HW range 的结束时刻，是 CPU 在 thread range 里 launch 的最后一个 GPU 操作，在 GPU 上结束执行的时刻。

Thread range 和 CUDA HW range 通过 correlationId 进行关联。

需要注意的是，两个 kernel 在 GPU 上执行时，中间会有空隙。如果空隙比较大，是因为 PushBuffer 经常空，GPU 拿不到新任务，只能空闲。CUDA HW 的 range 开始时刻减去结束时刻的时间值，包括空隙。

PushBuffer 满

Nsys 在 CUDA Overhead 那一行报的 command buffer full，实际上就是第一阶段第 3 步 UMD 写 PushBuffer 的时候，PushBuffer 满了。

这时候，CPU 会阻塞在第 3 步进行等待，等 GPU HIU 去取 PushBuffer 的内容。PushBuffer 腾出空间以后，CPU 上的 UMD 才能写入新的内容，然后返回。

PushBuffer 空

以下几种情况可能导致 PushBuffer 经常是空的：

CPU 计算太重：两次 launch 之间，CPU 正在做极其耗时的准备工作（比如处理复杂的逻辑、查表等），导致写 PushBuffer 的速度变慢。
Kernel 计算太轻：假设 CPU 把一个任务写入 PushBuffer 需要 10 us，GPU 执行一个简单 kernel 只需要 1 us，那么剩下 9 us，GPU 都在等待。如果两次 launch 不是 back-to-back，中间 CPU 还要进行其他操作，那更是雪上加霜。
CPU 被抢占：操作系统的调度器把你的 CPU 线程挂起了，没人在写 PushBuffer。
触发了隐式同步：如果你在 launch 之间不小心插了一个会导致管线排空的 API（比如非异步的内存分配），即使 CPU 没闲着，GPU 也会被迫停下来等，从而产生巨大的 Gap。

导致的结果就是，kernel 执行之间会出现 gap，GPU 利用率不高。

CUDA Graph

CUDA Graph 的物理本质：把成百上千次“写 PushBuffer + 敲门铃”的 CPU 动作，压缩成“一次写 + 一次敲门铃”，然后把剩下的分发工作全部甩锅给 GPU 硬件自己去干。

以前的流模式 (Stream Model) 有多笨

在没有 Graph 的时候，假设你有 100 个互相有依赖的小 Kernel（比如 A 算完算 B 和 C，B 和 C 算完算 D）。

哪怕你用了 Stream 和 Event 实现了完美的并发设计，在底层依然是：

CPU (UMD) 把 Kernel A 的状态和元数据打包写进 PushBuffer，敲 1 次门铃。
CPU 打包 Event 记录，打包 Kernel B、C 的 launch，注入 Wait Semaphore，敲 N 次门铃。
即使 GPU 端的 CP 很聪明能并发，但它吃的每一口饭，都得是 CPU 现喂的。只要 CPU 稍微打个盹（被操作系统挂起），或者 CPU 喂饭的动作（API 开销，单次 5~15 微秒）比 GPU 吃饭的速度慢，GPU 就得饿着（出现上面说的 Gap）。

CUDA Graph 怎么运作

CUDA Graph 将执行分成了两个截然不同的阶段：捕获/实例化 (Capture & Instantiate) 和发射 (Launch)。

1 预构建与编译 (Instantiate)：一切都在 CPU 端提前“杀青”

当你定义好一个包含 100 个 Kernel 的 Graph 后。驱动 (UMD) 会在 Host 端进行一次极其深度的预编译。

UMD 会把这 100 个 Kernel 需要的所有硬件状态（需要多少寄存器、Grid 维度、资源分配、依赖关系带来的 Wait/Release Semaphore）全部计算好，并静态打包成一个庞大且高度优化的命令列表 (Command List)。

这个包含整个大局的结构被直接放进了内存里。此时，GPU 还没动。
2 极速发射 (Graph Launch)：只 Launch 一次，敲一次门铃

当你在代码里调用 cudaGraphLaunch(graphExec, stream) 时。

CPU 只做一件事：往那个 Stream 的 PushBuffer 里写一条极其简短的触发指令（包含指向那个庞大 Command List 的指针），然后敲 1 次门铃。

cudaGraphLaunch 会在极短的时间内（通常 < 2 微秒）返回。
3 GPU 前端硬件接管 (Hardware Autonomous Scheduling)

GPU 端的 Host Interface Unit (HIU) 拉取到这条触发指令。

Command Processor (CP) 以及更高级的硬件调度器（在近代的 NVIDIA 架构中被称为 Grid Management Unit - GMU 或 Hardware Host Scheduler）直接接管了这整个庞大的 Command List。

硬件自己会看图纸：CP 内部的硬件电路直接读取 Graph 的依赖关系。它自己知道 A 执行完后，不需要等 CPU 下令，直接把 B 和 C 的块甩给 GigaThread Engine (GTE) 进行 CKE。

在整个 Graph 跑完之前，GPU 完全实现了闭环自治，再也不需要和 CPU 走 PCIe 进行频繁的握手。

CUDA Graph 解决的问题

1 彻底消灭 API Bound (Launch 开销受限)

这是它最核心的使命。对于 Transformer 这种有着极其海量的小算子（比如做 Tensor Parallelism 时各种切碎的 GEMM、Softmax、LayerNorm）的模型，每次运算的时间可能只有几微秒。

用了 Graph，不管图里有 10 个还是 1000 个 Kernel，CPU 端的 Launch 成本都被坍缩成了 1 次（约 2 微秒）。它消灭了 Nsys 上那些因为 CPU 喂饭太慢导致的 Gap。在 MLSys 的推理优化（比如 KV Cache 优化场景）中，这是提升小 Batch Size 吞吐量的终极杀器。
2 将 CPU 抖动 (Jitter) 从 GPU 隔离

操作系统是非常不可控的。在传统的流模型中，如果 CPU 在连续 Launch 几个 Kernel 的间隙，突然被操作系统切走去处理网卡中断了，那 PushBuffer 就断供了，GPU 就会被连累着空转。

而 CUDA Graph 是一次性下发。哪怕 CPU 刚敲完这次 Graph 的门铃就立刻被系统挂起了，GPU 照样能照着“图纸”全速满载把活干完。
3 驱动级的全局拓扑优化 (Whole-Graph Optimization)

在传统的流模式中，UMD 是“近视眼”，它一次只能看到你当前 Launch 的这一个 Kernel，只能做局部打包。

而在 Graph 实例化阶段，UMD 是“上帝视角”。它知道接下来整个管线要做什么，所以它可以在底层对硬件资源的分配进行极致压缩，比如重用某些隐式的寄存器配置，剔除冗余的同步屏障等。