CUDA Graph 在 LLM 推理里的实战：从 launch overhead 到 ragged batch

前情说明：技术深度 D4。前几篇讲了 attention、KV cache、MoE。都不涉及一个隐性但关键的性能杀手——CUDA kernel launch overhead。这一篇专门讲 CUDA Graph——它是怎么解决 launch overhead 的、为什么"variable-length"的 LLM 推理特别难用它、以及各家推理引擎的实战。素材来自 vLLM Optimization 文档、SGLang DSpark blog（07-06）里那段 ragged verify + full CUDA graphs 描述、NVIDIA CUDA Graphs 官方文档。

起手：一个隐性的性能杀手

你打开 nsight profiler 观察 vLLM 的 decode，会看到很奇怪的东西：

时间轴（一次 decode，Llama 70B, batch=1）:
GPU 计算:  ▓▓▓▓ ▓▓▓▓ ▓▓▓▓ ▓▓▓▓ ...
GPU 空闲:      ██   ██   ██   ██   ...
              5μs  5μs  5μs  5μs

每个 kernel 之间都有 5-10μs 的空隙——GPU 在等 CPU 发下一个 kernel 的指令。

一次 Llama 70B decode 涉及80 层 × 每层 ~10 个 kernel = 800 个 kernel。launch overhead 累积 4-8 ms。

这是什么概念？ 一次 decode 本身也就 10-20 ms。launch overhead 占 20-40%——接近一半时间 GPU 在等 CPU。

这就是 kernel launch overhead——CUDA 生态最古老、最普遍、最容易被忽视的性能杀手。

第一层：为什么 kernel launch 这么贵

cudaLaunchKernel 这个 API 表面只是"告诉 GPU 跑一个 kernel"。但底层要做很多事：

参数打包：把 kernel 参数序列化成 GPU 能理解的格式
命令入队：写入 GPU 的命令队列
驱动开销：CUDA 驱动做 validation + 调度
CPU-GPU 同步：轻量的 handshake

在现代硬件上（H100 + Xeon Sapphire Rapids），这一整套下来大约 3-10μs。看似很少，但 LLM 推理一秒钟要 launch 几万次 kernel——累积起来非常可观。

特别是 decode——每一层几十个小 kernel，都是 memory bound（很快跑完），launch 时间可能比 kernel 本身还长。

GEMM 大 kernel 为什么不受影响？

Prefill 时的一个大 GEMM 可能跑 100μs，launch 的 5μs 只占 5%——可以忽略。

问题出在 decode + 各种融合后的小 kernel——每个 kernel 本身几微秒到几十微秒，launch overhead 可以占 30-50%。

这就是为什么 CUDA Graph 对 decode 至关重要——它就是给这种"小 kernel 密集调用"场景设计的。

第二层：CUDA Graph 的核心机制

CUDA Graph（2018 年 CUDA 10 引入）的核心思想：一次录制，多次回放。

三阶段流程

阶段 1：Capture（一次性）

cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);

// 正常写 CUDA 代码
kernel_a&#x3C;&#x3C;&#x3C;...>>>();
kernel_b&#x3C;&#x3C;&#x3C;...>>>();
kernel_c&#x3C;&#x3C;&#x3C;...>>>();
// ... 100 个 kernel ...

cudaGraph_t graph;
cudaStreamEndCapture(stream, &#x26;graph);

CUDA 会记录所有 kernel launch（但不真正执行），生成一个 cudaGraph_t 对象——包含 100 个 kernel 的依赖关系和参数。

阶段 2：Instantiate（一次性）

cudaGraphExec_t exec;
cudaGraphInstantiate(&#x26;exec, graph, ...);

Instantiate 是 CUDA Graph 的核心魔法——它把 graph 编译成 GPU 能"一次性执行"的形式。所有 launch overhead 都在这一步付掉。

阶段 3：Launch（多次）

cudaGraphLaunch(exec, stream);  // 这一次 API 调用 = 执行整个 100 个 kernel

只有一次 API 调用——launch overhead 从 N × 5μs 降到 1 × 5μs。

收益到底有多大？

vLLM 官方数据：

Decode 场景（Llama 70B, batch=1）：CUDA Graph on 提速 20-30%
Small batch decode：最高 50%+ 提速
Prefill 场景：几乎无收益（大 kernel 主导）

对 decode 越关键——这也是为什么 vLLM 默认开启 CUDA Graph（--enforce-eager 才关闭）。

第三层：LLM 推理里 CUDA Graph 的四大挑战

上面听起来完美——但 LLM 推理里用 CUDA Graph 是痛苦得多的事情。四个挑战：

挑战 1：图必须是"静态"的

CUDA Graph 要求 kernel 参数、shape、依赖关系全部固定——一次 capture 决定，之后每次 replay 都一样。

LLM 推理是动态的：

Batch size 每一步都可能变（新 request 加入、旧 request 完成）
Sequence length 变（decode 一次+1）
甚至 request 数量 变（continuous batching）

朴素做法崩溃——每种 shape 都 capture 一个 graph，几百种组合 = 几百个 graph，memory 爆炸。

挑战 2：控制流不能捕获

If/while 循环不能在 CUDA Graph 里表达——比如"如果 speculative decoding 全部接受就多跑一步"这种逻辑必须在 CPU 侧决定。

这就限制了什么能进 graph：只有纯前向、无条件分支的部分。

挑战 3：Prefill / Decode / Mixed 混合模式

vLLM V1 支持一个 batch 里同时有 prefill 和 decode 请求。这两种模式的 kernel 行为差异巨大：

Prefill: query 长度 = seq len（几百-几千）
Decode: query 长度 = 1

一个 mixed batch 的 shape 组合无数——不可能每个都 capture 一个 graph。

挑战 4：Variable-length verify（Speculative Decoding）

Speculative decoding 场景下，每个 request 的 verify 长度可能不同（DSpark 的 confidence-driven 每 request 不一样）。批内 shape 参差不齐——CUDA Graph 完全懵。

第四层：解决方案矩阵

各家推理引擎的解决方案都在同一个思路上做变奏：Bucket + Padding + Piecewise。

方案 A：Bucket + Round Up（vLLM 早期）

思想：预定义几个 "size tier"（比如 batch=1/2/4/8/16/32/64/128），每次找最贴合的那个用。

# Capture 阶段
for tier in [1, 2, 4, 8, 16, 32, 64, 128]:
    graph[tier] = capture_decode_step(batch_size=tier)

# Runtime 阶段
actual_batch = 5
tier = round_up_to_bucket(actual_batch)  # 5 → 8
launch_graph(graph[tier], padded_batch=8)  # padding 3 个假 request

优点：简单缺点：padding 浪费——batch=5 padding 到 8 就是浪费 37.5% 算力

方案 B：Piecewise Graph（vLLM V1）

思想：只把最热的 kernel 段捕获成 graph——比如 attention + FFN 这种稳定的核心，剩下的 CPU 侧动态处理。

一次 decode:
    [CPU: 决定 batch] → [Graph: attention + FFN × 80 layers] → [CPU: sample + update]
        ↑                    ↑                                     ↑
     动态                   静态部分（capture 成 graph）           动态

vLLM V1 的 FULL_AND_PIECEWISE 模式同时准备了：

Full graphs：几个常用 batch size 完整 capture
Piecewise graphs：只捕获中间稳定段，两端动态

Runtime 根据情况选：常见 batch 用 full（最快），罕见 batch 走 piecewise（灵活）。

方案 C：Front-Pack + Round Total（SGLang DSpark 2026-07）

这是最优雅的一个——SGLang 在 DSpark blog 里介绍。

问题：DSpark 的每 request verify 长度不同——比如 batch 里一个 request 要 verify 6 个 token、另一个只要 2 个、第三个要 4 个。没有 "batch size" 这个统一维度。

天真做法：把每个 request 都 pad 到最大 verify length（比如 6）——又变成 batched paded 那种浪费。

SGLang 的做法：Front-pack + total-round：

原始 batch (verify lengths [6, 2, 4]):
    Request 0: [tok0, tok1, tok2, tok3, tok4, tok5]   ← 6 tokens
    Request 1: [tok0, tok1]                             ← 2 tokens
    Request 2: [tok0, tok1, tok2, tok3]                 ← 4 tokens

Front-pack（拼接成一维）:
    [tok0_0, tok0_1, ..., tok0_5, tok1_0, tok1_1, tok2_0, ..., tok2_3]
     └─── Request 0 (6) ──────┘└── Req 1 (2) ┘└──── Req 2 (4) ─────┘

Total tokens: 6 + 2 + 4 = 12
Round up to nearest captured tier: 16

关键：Graph 只按 total token 数（12 → tier 16）分类，不管每个 request 具体多长。

如何做 attention？ 用 cu_seqlens（cumulative sequence lengths）传给 attention kernel：

cu_seqlens = [0, 6, 8, 12] → attention kernel 知道每个 sequence 的边界
复用 varlen attention kernel（FlashAttention 天然支持）

收益：

Padding 只到最近 tier（12 → 16，浪费 25%）
Graph 数量控制在 ~10 个 tier 内
多种 verify 长度组合都能命中

DSpark blog 说：这个方案让 DeepSeek V4-Pro 上 batch=1 decode 跑到 383.7 tok/s（accept length ~5，TP=8，B300）——不用 CUDA Graph 大约要减 20-30%。

方案 D：多 Tier + DP 协调（DP attention 场景）

DP attention（数据并行 attention，如 DeepSeek MLA）多张卡各自处理不同 batch，每张卡的 tier 可能不同。

问题：如果 GPU 0 用 tier 8、GPU 1 用 tier 16，同步会出问题——DP 需要 all-gather，两边 shape 不一致会 hang。

SGLang 的解：所有 rank 用同一个 tier（选最大的）。这样有小浪费，但保证同步正确。

第五层：多 Stream + CUDA Graph 的组合

前面 DeepSeek V4 那篇讲了 vLLM 用 两个 CUDA stream（default + indexer）overlap 不同分支。

问题：CUDA Graph 能不能捕获多 stream的执行？

答案：能，但比较麻烦。CUDA 提供了 cudaGraphAddEventRecordNode 和 cudaGraphAddEventWaitNode 来在 graph 里跨 stream 同步。

vLLM 对 DeepSeek V4 的实现里就用了这个——default stream 捕获主 attention 计算，indexer stream 捕获 indexer 计算，两个 stream 通过 event 同步。整个多 stream 结构一起 replay。

这是 CUDA Graph 的高阶用法——大多数教程只教单 stream，实战里多 stream 才是常态。

第六层：CUDA Graph 的"隐性坑"

不是说 CUDA Graph 一开就完美——实际用起来有几个反直觉的坑。

坑 1：Instantiate 很慢

第一次 cudaGraphInstantiate 可能要几百毫秒到几秒——CUDA 要把整个图编译成 GPU 能高效执行的形式。

推理侧影响：冷启动很慢。vLLM 启动时会 pre-warm 所有 tier 的 graph——启动时间 30秒-2分钟主要花在这里。

坑 2：内存无法动态分配

Graph 里的所有 tensor 必须预分配——runtime 不能 malloc。这就是为什么 vLLM 的 KV cache 是启动时一次性分配一个大 pool——中途申请新 tensor 会破坏 graph。

坑 3：CPU-GPU 数据依赖会破坏 graph

如果某个 kernel 的参数依赖 CPU 计算结果（比如"根据 hidden state 决定用哪种 sampling"），这个部分不能进 graph——只能走 piecewise 或者 event-based 同步。

坑 4：Update Node 是双刃剑

CUDA 12 引入了 cudaGraphExecKernelNodeSetParams——可以在 graph replay 之间修改 kernel 参数，不用重新 capture。

理论上极好——runtime 更灵活。实际上很脆弱——参数改错会导致 graph 静默出错（不 crash，出错误结果）。vLLM 用得很保守。

坑 5：Debug 极难

CUDA Graph 里的 kernel 错误几乎不可能单独定位——你只能看到"整个 graph 跑挂了"。调试基本靠 disable graph（--enforce-eager）复现。

这就是为什么 vLLM 有 --enforce-eager flag——生产上不用（性能差 20-30%），但 debug 时必开。

一张表对比几家的 CUDA Graph 策略

引擎	策略	特点
TensorRT-LLM	全静态 graph，capture 所有 tier	峰值性能最好，但 memory 占用大、启动慢
vLLM V1 FULL_AND_PIECEWISE	Full + Piecewise 混合	灵活性 vs 性能平衡
SGLang v0.4+	Zero-overhead + ragged pack	支持 DSpark variable-length verify
TGI	只 capture 主 attention	最保守，最少 corner case

没有绝对的最优——看你的 workload 特征：

长期稳定 workload（同一 model、同一 batch pattern）→ 全静态 graph 最强
多样 workload → piecewise 更适合
Speculative + agent 混合 → 需要 ragged pack

三个通用启示

启示 1：Launch overhead 是"隐性成本"，但可以显性化

在 GPU 编程里，人们经常问"我的 kernel 快不快"——但从来不问"我 launch 一个 kernel 要多久"。

其实 launch 时间也是可测量的——nsight profiler 里就有 kernel launch 那一条 timeline。把它做成可视化，你的 40% 时间浪费就现形了。

性能优化的第一步永远是 profile——看不见的东西优化不了。

启示 2：静态 + 动态的平衡艺术

CUDA Graph 完全静态。真实业务完全动态。推理引擎的核心工程课题就是——如何把动态 workload 塞进静态 graph。

各家的答案都在 "bucket / padding / piecewise / front-pack" 这几个技巧的组合里做变奏。这是"约束下的优化"的绝佳案例——不是消除约束，是绕开它。

启示 3：抽象层次决定优化空间

CUDA Graph 是"底层原语"——你可以往上叠任何自定义调度逻辑。SGLang 的 front-pack 就是在 CUDA Graph 之上加了一层"packed varlen batch"的抽象。

好的推理引擎 = 好的抽象层次设计——底层露出足够多的原语（graph、stream、event），上层用这些原语组装出针对具体 workload 的最优解。

vLLM 的 Model Runner V2、Connector Interface、Modular Kernel 都是这个哲学——"我提供工具，你按需组合"。这才是长期主义。

附：素材来源

主线

vLLM Docs · Optimization and Tuning · Chunked Prefill
SGLang Blog · DSpark in SGLang（2026-07-06，ragged verify + CUDA graph 部分）
NVIDIA · CUDA Graphs API

代码

vLLM V1 compilation config - FULL_AND_PIECEWISE 模式
SGLang zero-overhead scheduler - spec-v2 runtime

姊妹篇

《DeepSeek V4 稀疏注意力完全解剖》 - 多 stream + CUDA Graph 实战
《MoE 推理完整机制》 - MegaMoE persistent kernel
《Speculative Decoding 全家桶》 - DSpark ragged verify

整理于 2026-07-31 深夜。原文技术细节归属 vLLM/SGLang 团队与 NVIDIA。中文表达可能有解读偏差，请以原文为准。

CUDA Graph 在 LLM 推理里的实战：从 launch overhead 到 ragged batch

CUDA Graph 在 LLM 推理里的实战：从 launch overhead 到 ragged batch

起手：一个隐性的性能杀手

第一层：为什么 kernel launch 这么贵

GEMM 大 kernel 为什么不受影响？

第二层：CUDA Graph 的核心机制

三阶段流程

收益到底有多大？

第三层：LLM 推理里 CUDA Graph 的四大挑战

挑战 1：图必须是"静态"的

挑战 2：控制流不能捕获

挑战 3：Prefill / Decode / Mixed 混合模式

挑战 4：Variable-length verify（Speculative Decoding）

第四层：解决方案矩阵

方案 A：Bucket + Round Up（vLLM 早期）

方案 B：Piecewise Graph（vLLM V1）

方案 C：Front-Pack + Round Total（SGLang DSpark 2026-07）

方案 D：多 Tier + DP 协调（DP attention 场景）

第五层：多 Stream + CUDA Graph 的组合

第六层：CUDA Graph 的"隐性坑"

坑 1：Instantiate 很慢

坑 2：内存无法动态分配

坑 3：CPU-GPU 数据依赖会破坏 graph

坑 4：Update Node 是双刃剑

坑 5：Debug 极难

一张表对比几家的 CUDA Graph 策略

三个通用启示

启示 1：Launch overhead 是"隐性成本"，但可以显性化

启示 2：静态 + 动态的平衡艺术

启示 3：抽象层次决定优化空间

附：素材来源

评论