decode 每个 token 只算一点，但整模型仍有大量 kernel 要由 CPU 逐个提交；batch 一小，GPU 就在等 CPU。CUDA Graph 把一段固定 GPU 工作录成图，之后一次 cudaGraphLaunch 提交整张图。代价是形状必须静态、输入地址必须固定，runtime batch 要 padding 到捕获档位。
decode 的瓶颈常常不在 GPU 算力，而在提交次数。每生成一个 token，模型仍要跑完所有层；每层有多个小 kernel。batch 小时每个 kernel 的 GPU 时间很短，CPU 的 Python、框架分发和 CUDA launch 反而成了关键路径。PyTorch 文档把 graph replay 的收益写成「greatly reduced CPU overhead」：它跳过 Python、C++ 和 CUDA driver 的参数设置与逐 kernel 分发。
CUDA Graph 的主要收益不是「更快的 kernel」，而是「更少的提交」。stream capture 只记录、不执行；instantiate 生成可执行图；replay 时同一张图用同一批地址反复提交。vLLM V1 把这件事拆成两套 routine：FULL 捕获整个 forward（包括 attention，前提是 backend 支持），PIECEWISE 在 attention 前切开、attention 走 eager、其余进 graph。
在默认的 piecewise 编译（-O2）下，默认不是 FULL，而是 FULL_AND_PIECEWISE；关闭编译时 resolve 可落到 NONE。uniform decode 走 FULL，prefill 与 mixed batch 走 PIECEWISE；-O2 默认就是它。-O1 只开 PIECEWISE，--enforce-eager 全关。档位越多，padding 越小，但捕获数量、启动时间和显存占用越高。
max_num_seqs=256（或 ≥512，如 1024）时，基础 token 捕获网格的档位数（cap=512）；若 max_num_seqs 在 (256,512) 且不在网格上（如 264），源码会追加一档，总数为 52读完这节你应该知道：为什么 batch 小、序列短时，CUDA Graph 能省掉的不是计算，而是「把计算交给 GPU」的过程。
自回归 decode 每一步通常只生成一个 token（投机解码时多几个）。对模型来说，这一步仍要做完所有层：RMSNorm、QKV 投影、RoPE、attention、输出投影、MLP，张量并行时还有 all-reduce。每个操作都是一个或几个 CUDA kernel。batch 很小时，每个 kernel 的 GPU 执行时间很短，但 CPU 仍要逐个完成 Python 层的张量元数据、框架分发、CUDA driver 的参数校验与入队。
K 是这一步的 GPU 操作数，B 是 batch 大小。B 越小，tkernel 越短，Tgpu 下降；K 和 tlaunch 不随 B 缩小，于是 max 的取值从 GPU 转向 CPU。
这就是「launch-bound」的含义：不是 GPU 不会算，而是 GPU 经常在等下一批 kernel 被提交。连续批处理把多个请求拼成一个 batch，能提高 B、拉长每个 kernel 的 GPU 时间；但在小模型、低并发、短输出或 P/D 分离的 decode 实例上，B 仍然可能很小，提交开销占比就高。反过来，如果 B 很大、每个 kernel 都能跑满 GPU，那么 launch 开销被摊薄，CUDA Graph 的收益也会变小。
还有一个容易混淆的边界：CUDA Graph 只记录 GPU 工作。PyTorch 文档写得很直接：「CPU work is not captured」；graph replay 跳过的是 Python、C++ 与 CUDA driver 的分发，不是调度器、分词、采样这些 CPU 逻辑。vLLM 的 scheduler 仍然每步选 batch、更新 KV cache 元数据；graph 只负责把选定 batch 的前向计算一次性提交。因此，graph 不能替代调度优化，也不能消除采样和 detokenization 的 CPU 时间。
下一节先看没有 graph 时两条常见路径：纯 eager，以及只做 torch.compile。
读完这节你应该知道：为什么「编译过」不等于「图化」，以及 vLLM 为什么把编译与 CUDA Graph 做成两个可独立开关的维度。
不做捕获时，模型 forward 就是普通的 PyTorch 执行：Python 逐层调用，框架逐个提交 kernel。启动最快，调试最直接；vLLM 的 --enforce-eager 就是这条路径，官方文档说它跳过编译与 CUDA graph 捕获，用稳态 decode 性能换启动时间（docs/configuration/optimization.md:25）。
问题也在这里：kernel 数不随 batch 缩小，batch 一小，CPU 提交就成为关键路径。eager 不会自动把多个 kernel 合并成一次提交。
Dynamo 抓图、Inductor 融合和代码生成能减少 Python 开销、减少部分算子，但编译产物仍然要在每个 step 逐个启动 kernel。vLLM 的 piecewise compilation 还会在 attention 等位置切图：attention 留在 eager，其余片段编译（docs/design/torch_compile.md:250-254）。
切图的目的之一，是让这些片段可以被 CUDA Graph 捕获；编译本身并不等于图化。vLLM 文档明确说，full cudagraph 可以脱离编译单独使用，piecewise cudagraph 才要求 piecewise compilation（vllm/config/compilation.py:722-726）。
把两者混为一谈，会得出错误的调优结论。编译解决的是「生成什么样的 GPU 代码、减少多少 Python 开销」；CUDA Graph 解决的是「这些 GPU 代码怎样被提交」。前者可以没有后者，后者也可以没有前者（FULL 模式）。vLLM 的 -O1 是 PIECEWISE，-O2 默认是 FULL_AND_PIECEWISE（docs/configuration/optimization.md:13-14；vllm/config/vllm.py:337），这两个档位同时改变了编译与图化策略，但不能把收益全部记在编译上。
那 CUDA Graph 具体怎样把 K 次提交变成一次？下一节拆 capture、instantiate、replay 三步，以及「静态 buffer」这个硬约束。
读完这节你应该知道：CUDA Graph 的三段生命周期各自做什么、哪些东西一旦捕获就不能再变，以及 vLLM 如何用 descriptor 把「变」挡在 graph 外面。
NVIDIA 的 CUDA Programming Guide 把图的生命周期分成定义、实例化、执行三段（§4.2.2）。stream capture 用 cudaStreamBeginCapture() 打开捕获模式，之后发到该 stream 上的工作「is not enqueued for execution」，而是被记成节点；cudaStreamEndCapture() 返回图。实例化用 cudaGraphInstantiate() 生成 cudaGraphExec_t，把校验和准备工作提前；执行用 cudaGraphLaunch(graphExec, stream)，同一张可执行图可以反复 launch，不必每次重新实例化（§4.2.2.2–§4.2.2.3）。
PyTorch 的封装对应同样的语义：torch.cuda.graph 进入捕获，捕获期间「CUDA work issued to a capturing stream doesn't actually run on the GPU」；torch.cuda.CUDAGraph.replay() 提交整张图，一次 replay 对应一次 cudaGraphLaunch。PyTorch 文档还强调，replay 时「the same memory addresses are used」，所以输入必须写进捕获时用的那块张量，输出也从同一块地址读。
cudaGraphExec_t（拓扑固定，参数可有限更新）cudaGraphLaunch（一次提交整张图）
依据：NVIDIA CUDA Programming Guide §4.2.2；PyTorch CUDA semantics 的 CUDA Graphs 一节。
「参数可有限更新」有明确边界。NVIDIA 文档说，拓扑或节点类型变了必须重新实例化；只有参数变化时可以用 cudaGraphExecUpdate() 或逐节点 setter 就地更新，但要求图拓扑完全相同，且 kernel 的 context、memcpy 的设备和 cudaMemcpyKind、外部 semaphore 计数等不能变（§4.2.3）。vLLM 的选择更简单：不动态改图，而是为每个需要的形状各捕获一张图，运行时按 descriptor 选图。
图里的 kernel 参数在捕获时就固定了，包括输入、输出和中间 buffer 的地址。PyTorch 文档的典型用法是：捕获时持有输入/输出张量，replay 前用 copy_ 把新数据写进捕获输入，replay 后从捕获输出读结果。这样每次 replay 仍然读写同一组虚拟地址，图不需要改参数。
vLLM 在 FULL 模式里按同样的方式组织：捕获前用 dummy 输入准备 attention metadata，捕获时把模型输出复制到持久的 hidden_states buffer；replay 时只换输入 buffer 的内容，不换地址（vllm/v1/worker/gpu/cudagraph_utils.py:699-732、:816-846、:604-617）。对 PIECEWISE，vLLM 还提供 cudagraph_copy_inputs：调用方不能保证输入 buffer 固定时，编译器把输入复制到内部管理的 buffer（vllm/config/compilation.py:803-810）。
这也解释了为什么 graph 对「动态形状」如此敏感：形状变了，buffer 大小和 kernel 的 grid/block 配置就变了，旧图的参数不再成立。vLLM 的应对不是让一张图适配所有形状，而是把 runtime batch padding 到最近的捕获档位，然后用 BatchDescriptor 去查表。descriptor 只带少量字段：num_tokens、num_reqs、uniform、has_lora、num_active_loras（vllm/forward_context.py:30-50）。
在 vLLM V1 里，CUDAGraphWrapper 包住一个可调用对象，负责捕获/重放；CudagraphDispatcher 维护 FULL 与 PIECEWISE 两套合法 key，并在每步根据 batch 选 runtime mode。文档把优先级写得很清楚：FULL > PIECEWISE > None；key 不存在就退回 eager（docs/design/cuda_graphs.md:98-118）。代码里的 dispatch() 也是这个顺序：先查 FULL key，再查放宽了 num_reqs 与 uniform 的 PIECEWISE key，最后返回 NONE（vllm/v1/cudagraph_dispatcher.py:233-322）。
FULL 与 PIECEWISE 的差别不只是「捕获得多不多」，而是 attention 在不在图里。PIECEWISE 用 splitting_ops 把 attention 等不适合捕获的 op 切成断点，断点之间才进 graph；FULL 则把整个 forward 放进一张图，因此要求 attention backend 自己支持 CUDA Graph。这个约束是下一节与第 07 节的主线。
读完这节你应该知道：默认捕获网格怎样生成、一个 runtime batch 会被 pad 到哪一档、每多一档要付出什么，以及为什么没有一张可直接引用的“每档固定 MiB”表。
vLLM 的默认捕获上限不是直接取 max_num_seqs，而是先算一个平台上限：非 SM100 为 512，SM100（如 B200）为 1024；再取 min(max_num_seqs × decode_query_len × 2, 平台上限)，最后再与 max_num_batched_tokens 取小（vllm/config/vllm.py:2619-2625、:2669）。decode_query_len 通常为 1；投机解码时为 1 + num_speculative_tokens。
网格本身是分段的：[1, 2, 4]，然后 8 到 248 每 8 一个档，256 以上每 16 一个档，直到 cap（vllm/config/vllm.py:2696-2708）。cap 较小时高段自然为空。按这个基础网格复算，非 SM100 的典型配置（max_num_seqs=256 或 max_num_seqs ≥ 512，max_num_batched_tokens ≥ 512）会得到 cap=512、共 51 档；若 max_num_seqs 在 (256,512) 且不在网格上（如 264），源码会把它追加进 uniform_decode_sizes（vllm/config/vllm.py:2666-2667、:2718-2720），总数为 52 档。SM100 的 cap=1024 默认 max_num_seqs=1024 时为 83 档；off-grid 时同理追加一档。
| cap（tokens） | 默认档位数 | 档位构成 | 典型来源 |
|---|---|---|---|
| 32 | 7 | [1, 2, 4] + 8..32 步 8（4 个） | max_num_seqs=16 之类的小配置 |
| 64 | 11 | 3 + 8..64 步 8（8 个） | max_num_seqs=32 |
| 128 | 19 | 3 + 8..128 步 8（16 个） | max_num_seqs=64 |
| 256 | 35 | 3 + 31 + 1 | max_num_seqs=128 |
| 512 | 51 | 3 + 31 + 17（基础网格；max_num_seqs=264 时再 +1 = 52） | 非 SM100、max_num_seqs=256（或 ≥512，如 1024） |
| 1024 | 83 | 3 + 31 + 49（基础网格；off-grid 时追加一档） | SM100（如 B200），默认 max_num_seqs=1024 |
注意 cap 还受 max_num_seqs 与 max_num_batched_tokens 约束：如果 max_num_seqs=128，cap 只有 256，不是 512。vLLM 的 EngineArgs 在真实服务里会把其余档（含 A100 40G）的 max_num_seqs 默认成 256，≥70GB 且非 A100/非 GB10 的档默认成 1024（vllm/engine/arg_utils.py:3014-3038）；因此非 SM100 的典型服务配置（默认 256 或 1024）通常落在 cap=512、基础网格 51 档；若自定义成 264 这类 off-grid 值，捕获数为 52。
dispatcher 会预先算一张 _bs_to_padded_graph_size 表：如果 runtime batch 恰好等于某个捕获档，就用该档；否则用下一个更大的档；超过最大档则没有匹配，走 eager（vllm/v1/cudagraph_dispatcher.py:71-90）。这就是静态形状约束的实际表现：图不重新编译，batch 被 pad 到已捕获的形状。
S 是捕获档位集合，b 是 runtime batch 的 token 数。padding 只改变 padding 位置的内容，不改变真实 token 的结果；代价是多余的计算和 attention 槽位。
默认网格在小档位很密、在大档位变粗，所以 padding 浪费不是均匀的。下面的交互按 4.1 的默认生成式算档位数，并按 4.2 的映射算 padding；它复现的是「非 interactivity、decode_query_len=1、无 LoRA、无投机」的简化路径。
cap = min(max_num_seqs × 2, 512) = 512；网格 51 档。batch 500 → padded 512，浪费 12 tokens（2.34%）。；超过最大捕获档，dispatcher 返回 NONE，走 eager。
简化模拟：只复现非 interactivity 的默认 token 网格与「向上取档」映射；未覆盖 Blackwell 的 1024 上限、投机解码的 query_len > 1 取整、LoRA 分档、用户自定义 cudagraph_capture_sizes、max_num_batched_tokens 截断，以及 v2 runner 的 request-keyed 解码图。
口径：cap = min(max_num_seqs × 2, 512)，且假设 max_num_batched_tokens ≥ cap；档位集合 = [1, 2, 4] + 8..min(cap,248) 步 8 + 256..cap 步 16；若 max_num_seqs ≤ cap 且不在集合中，再补上 max_num_seqs（vllm/config/vllm.py:2666-2667、:2718-2720）；若 max_num_tokens ≤ cap 且不在集合中，还会补上 max_num_tokens（:2709-2714）。padding 浪费比例 = (padded − batch) / padded × 100%。这是页面按源码生成式做的示例，不是仓库公布的基准。
组 A：max_num_seqs=64，batch=5。cap = min(128, 512) = 128。档位 = [1, 2, 4] + 8..128 步 8（16 个，8, 16, …, 128），共 19 档；max_num_seqs=64 已在集合中。batch=5 向上取到 8，浪费 8 − 5 = 3 tokens，3 / 8 = 37.50%。
组 B：max_num_seqs=256，batch=500。cap = min(512, 512) = 512。档位 = 3 + 31（8..248）+ 17（256..512 步 16）= 51 档；256 已在集合中。batch=500 向上取到 512，浪费 512 − 500 = 12 tokens，12 / 512 = 2.34%。
JS 在这两组参数下分别输出 19 档 / 37.50% 与 51 档 / 2.34%，与上面的手算一致。滑块初始值（max_num_seqs=256、batch=500）对应组 B。组 A 的 max_num_seqs=64 与 batch=5 可直接拖到对应位置复核。
未覆盖：max_num_seqs 很大时 cap 会被 512 截断；SM100 的 cap=1024 不在本模拟中；投机解码会把尺寸向上取整到 query_len 的倍数（vllm/config/compilation.py:1690-1718）；LoRA 会按 active adapter 数复制档位（vllm/config/compilation.py:811-818）。
vLLM 的 graph 共享 memory pool：捕获时先捕 PIECEWISE，再捕 FULL；FULL 图不比最宽 PIECEWISE 图宽时，可以复用已经分配的 buffer。同一模式内也按 num_tokens 从大到小捕获，让窄图复用宽图先占的 pool（vllm/v1/worker/gpu/cudagraph_utils.py:480-482、vllm/v1/cudagraph_dispatcher.py:324-345）。因此「第 n 档新增多少 MiB」不是常数，取决于该档的静态输入/输出、attention metadata、workspace，以及它与前面图的重叠程度。
仓库文档与源码中未找到按档位列出的 MiB 表（未获取）。它在启动时用 start_free_gpu_memory − end_free_gpu_memory 实测所有捕获的总增量，打印 Graph capturing finished ... took X GiB（vllm/v1/worker/gpu_model_runner.py:6567-6611）；代码注释给出的经验范围是「This usually takes 5~20 seconds」（:6605）。在真正分配 KV cache 之前，profile_cudagraph_memory 只捕获最大的两张 FULL 图做采样，用第二张的增量作为每图成本、并以 1 MiB 为下限外推总成本（vllm/v1/worker/gpu/cudagraph_utils.py:986-988、:1101-1109）。
这个估计默认会从可用 KV cache 预算里扣除：VLLM_MEMORY_PROFILER_ESTIMATE_CUDAGRAPHS 默认是 1，available_kv_cache_memory_bytes = requested_memory − non_kv_cache_memory − cudagraph_memory_estimate（vllm/v1/worker/gpu_worker.py:803-839；vllm/envs.py:2339-2341）。所以 graph 的显存代价不是「额外送」，而是从 KV cache 配额里扣；档位越多、图越宽，可用于 KV cache 的块越少。V1 用户指南也提示，V1 的 CUDA graph 捕获比 V0 更占显存（docs/usage/v1_guide.md:43）。
| 组成 | 随档位/宽度怎么变 | 说明 |
|---|---|---|
| 静态输入/输出 buffer | 随 num_tokens 增长 | 输入 ids、positions、hidden states、slot mapping 等；replay 原地覆盖 |
| 中间激活 / workspace | 随 num_tokens 增长，但 pool 可复用 | graph pool 内的分配；窄图可复用宽图先占的块 |
| attention metadata | 随 num_reqs、max_query_len 增长 | block table、seq lens、slot mapping 等；FULL 还要让 backend 生成可捕获的 metadata |
| graph executable / driver 结构 | 每图一份，近似固定 | 相对 buffer 通常较小；外推时以 1 MiB 为下限 |
| KV cache 机会成本 | 档位越多，扣得越多 | 默认从 KV cache 预算扣除 estimate |
所以第 04 节的结论是：档位不是免费的。基础网格 51 档覆盖 1–512，最坏 padding 出现在刚超过某个档位的位置；小档位密、大档位粗，调度器无法通过「多加几档」无限逼近零浪费，因为每加一档都要多一次 capture、多一份 pool 占用和启动时间。
读完这节你应该知道：capture、instantiate、replay 在时间轴上分别发生在哪里，以及一次 decode step 从 scheduler 到 graph replay 的完整路径。
vLLM 在 worker 初始化时调用 capture_model()。捕获前先做一次 eager warmup；_warmup_and_capture 按 cudagraph_num_of_warmups 跑若干次 CUDAGraphMode.NONE 的 dummy run，再用非 NONE 模式捕获（vllm/v1/worker/gpu_model_runner.py:6613-6660）。捕获顺序是先 PIECEWISE、再 FULL；同一模式内按 num_tokens 从大到小，先占大 pool（vllm/v1/worker/gpu/cudagraph_utils.py:480-482、vllm/v1/cudagraph_dispatcher.py:324-345）。
捕获完成后，vLLM 记录 start_free_gpu_memory − end_free_gpu_memory 作为 graph 总占用，并打印耗时与 GiB（vllm/v1/worker/gpu_model_runner.py:6567-6611）。这一步发生在 KV cache 初始化之后、服务开始之前（vllm/v1/worker/gpu_worker.py:1016-1164）。
一次 decode step 的路径可以压成四步：
1 + num_speculative_tokens。num_tokens pad 到最近的捕获档，带上 num_reqs、uniform、LoRA 计数，形成 BatchDescriptor。num_reqs、uniform 后查 PIECEWISE；都没有就返回 NONE，走 eager（vllm/v1/cudagraph_dispatcher.py:233-322）。run_fullgraph()，内部先 get_offloader().sync_prev_onload()，再 self.graphs[desc].replay()（vllm/v1/worker/gpu/cudagraph_utils.py:604-617）；PIECEWISE 走编译后的片段 + wrapper；NONE 走普通 forward。这条路径解释了为什么 graph 不改变模型的数学结果：dispatcher 只决定「用哪张已捕获的图」，图里的 kernel 序列与 eager 执行的是同一套计算；padding 位置的内容被标记为无效，attention 与后续层不会把它们当成真实 token。真正的风险在于 metadata 与 buffer 地址是否仍然一致，而不在于图本身会不会「算错」。
读完这节你应该知道：这组数字在什么条件下成立、为什么平均 TPOT 变好而 TTFT 和 P99 TPOT 可能变差，以及为什么默认模式不是 FULL。
vLLM 设计文档没有把性能数字写进正文，只在「About the Performance」链接到 PR #20059 的评论。下面的表 2 转录该评论 Serving Benchmark Result 中的关键指标（未转录 successful requests、benchmark duration、input/generated tokens、median/P99 TTFT、median TPOT、median/P99 ITL）；作者注明这是 PR 的早期版本，CLI 后来改过，数据「just kept for reference」。条件必须在同一句里交代：A100 40G、torch 2.6、CUDA 12.4、Qwen2.5-7B-Instruct-GPTQ-Int4、ShareGPT 100 prompts、request-rate 20、gpu-memory-utilization 0.9。
| 指标 | FA2 · piecewise | FA2 · FULL+piecewise | FlashInfer · piecewise | FlashInfer · FULL+piecewise |
|---|---|---|---|---|
| Output tok/s | 1898.67 | 1992.27 | 1907.38 | 1946.35 |
| Total tok/s | 3937.88 | 4132.01 | 3955.65 | 4036.48 |
| Request throughput（req/s） | 8.77 | 9.20 | 8.81 | 8.99 |
| Mean TTFT（ms） | 76.37 | 78.69 | 73.61 | 76.03 |
| Mean TPOT（ms） | 17.08 | 16.57 | 16.85 | 16.34 |
| Mean ITL（ms） | 13.45 | 12.83 | 13.34 | 13.11 |
| P99 TPOT（ms） | 67.68 | 78.21 | 65.75 | 58.86 |
按表 2 复算，FA2 的 output tok/s 相对提升 4.9%（1992.27 vs 1898.67），mean TPOT 相对下降 3.0%（17.08 → 16.57 ms）；FlashInfer 的 output tok/s 相对提升 2.0%（1946.35 vs 1907.38），mean TPOT 相对下降 3.0%（16.85 → 16.34 ms）。作者在评论里把这两个吞吐变化概括为「5%」与「2%」，把 TPOT 变化概括为「2.9%」与「3.1%」；按表复算分别是 4.9% / 2.0% 与 3.0% / 3.0%，作者未说明取整口径；推断（依据：2.9% 可由截断到一位复现，3.1% 可由改用 16.34 作分母复现，但作者未注明）。
反直觉的地方在首 token 与尾延迟。FA2 的 mean TTFT 从 76.37 ms 升到 78.69 ms（相对 +3.0%），P99 TPOT 从 67.68 ms 升到 78.21 ms；FlashInfer 的 mean TTFT 从 73.61 ms 升到 76.03 ms（相对 +3.3%），但它的 P99 TPOT 从 65.75 ms 降到 58.86 ms。也就是说，同一组「FULL + piecewise」配置，平均 decode 间隔变好，并不保证首 token 和尾延迟一起变好；FA2 与 FlashInfer 的尾延迟方向甚至相反。
作者对 FA2 的 TTFT 退化给了一个解释：full graph 在 mixed prefill-decode batch 上会让 FA2 的 attention 启动更多 block，而 eager attention 的 block 数少得多；他把原因归到 full graph 对 mixed batch 的启动配置，而不是 kernel 本身变慢。这个解释与「默认用 FULL_AND_PIECEWISE、只让 uniform decode 走 FULL」的设计是一致的：mixed batch 留在 PIECEWISE，避开 full graph 在 prefill/mixed 上的额外开销。
这组数据也说明收益的尺度：在这台 A100、这个 7B GPTQ 模型、这个 ShareGPT 负载下，吞吐提升是个位数百分比，不是数量级。CUDA Graph 省的是提交开销；当 GPU 已经是瓶颈、或模型/backend 的 kernel 本身很重时，省下来的 CPU 时间占比自然下降。
读完这节你应该知道：哪些情况下 graph 会退回 eager、哪些 backend 只能支持部分模式，以及 torch.compile 与 CUDA Graph 各自管什么。
图捕获的是固定形状和固定地址。runtime batch 只要超过最大捕获档，dispatcher 就找不到 key，直接返回 NONE 走 eager（vllm/v1/cudagraph_dispatcher.py:233-322）。这意味着「最大档必须覆盖常见 batch」是一个显式的容量选择：档位太小，频繁退回 eager；档位太大，显存与启动成本上升。PyTorch 文档还把 data-dependent control flow 列为 graph 的限制（torch.cond() 是例外），并禁止在捕获期间做 CPU 同步（如 .item()）和在 graph 私有 pool 之外分配。
投机解码会把这件事变得更复杂：uniform decode 的 query 长度是 1 + num_speculative_tokens，捕获尺寸要向上取整到 query_len 的倍数（vllm/config/compilation.py:1690-1718）。如果长度动态变化，vLLM 还有 varlen decode graph 的路径，但它要求 backend 能重放可变长度的 decode batch；不满足时仍然回到固定档位或 eager。
vLLM 用 AttentionCGSupport 描述 backend 的捕获能力，等级从高到低是 ALWAYS > UNIFORM_BATCH > UNIFORM_SINGLE_TOKEN_DECODE > NEVER（vllm/v1/attention/backend.py:563-578）。FULL 的 mixed prefill-decode 需要 ALWAYS；uniform decode 至少需要 UNIFORM_SINGLE_TOKEN_DECODE（query_len=1）或 UNIFORM_BATCH（query_len>1）。混合模型（如 Mamba mixer）取所有 backend 的最低等级。
| Attention backend | cudagraph_support | 备注 |
|---|---|---|
| FlashAttention v2 | UNIFORM_BATCH | 文档注释：实际为 ALWAYS，但为性能回退到 FULL_AND_PIECEWISE |
| FlashAttention v3 | ALWAYS | prefill/mixed 与 decode 共用 routine |
| Triton Attention | ALWAYS | 文档建议 FULL_AND_PIECEWISE，因为 prefill 与 decode 的 kernel 不同 |
| AITER FlashAttention | UNIFORM_BATCH | |
| FlashInfer | UNIFORM_SINGLE_TOKEN_DECODE | Blackwell 上使用 TRTLLM attention 时变为 UNIFORM_BATCH |
| FlashMLA | UNIFORM_BATCH | |
| FlashInferMLA | UNIFORM_BATCH | |
| FlashInferMLASparse | UNIFORM_BATCH | |
| AITER MLA | UNIFORM_SINGLE_TOKEN_DECODE | |
| CUTLASS MLA | UNIFORM_SINGLE_TOKEN_DECODE | |
| Mamba attention | UNIFORM_SINGLE_TOKEN_DECODE | |
| 未列出的 backend | NEVER | 文档原文：Unlisted backends are all declared as NEVER |
模式不是硬选，vLLM 会按 backend 能力降级。下面的交互复现 resolve_cudagraph_mode_and_sizes 的主路径：给定请求模式和 backend 等级，看最终模式。它假设 VLLM_COMPILE、attention 在 splitting_ops 中、uniform_decode_query_len=1、piecewise capture 可用；不覆盖混合 backend 取最小值、Mamba block 检查、v2 runner、LoRA 与用户自定义尺寸。
解析结果：ALWAYS + FULL → FULL
backend 支持 mixed full，保持请求模式。
简化模拟：按 resolve_cudagraph_mode_and_sizes 的前三个主分支复现降级；未覆盖 hybrid backend 取最小值、Mamba 每序列 block 检查、FULL 在 NEVER 下的最终二次检查、v2 runner 的 request-keyed 尺寸、LoRA 分档与 FULL 之外的 mode 组合细节。
这张表解释了一个常见误解：FULL 不是「总是最好」。如果 backend 只支持单 token decode，mixed batch 的 full graph 不被允许，vLLM 会把它降成 FULL_AND_PIECEWISE（attention 在 splitting_ops 时）或 FULL_DECODE_ONLY；如果 backend 是 NEVER，请求 FULL 会直接报错，而 FULL_DECODE_ONLY / FULL_AND_PIECEWISE 在 VLLM_COMPILE 且 attention 在 splitting_ops 中（或启用 use_inductor_graph_partition）时退到 PIECEWISE，否则退到 NONE（vllm/config/compilation.py:1576-1601、:1637-1649）。
多卡时，每张卡都要捕获自己的图，且各 rank 的 batch 形状必须对齐。数据并行会把 num_tokens pad 到所有 DP rank 的最大值，再通过 num_tokens_across_dp 传给 forward context（vllm/v1/worker/dp_utils.py:190-223、vllm/v1/worker/gpu_model_runner.py:3840-3856；tensor 创建见 vllm/v1/worker/gpu/cudagraph_utils.py:744-748）；张量并行与流水线并行则要求各 rank 的捕获顺序、通信 buffer 和中间张量地址保持一致。通信本身也要能被图捕获；vLLM 在 DP>1 时对 DeepEP 高吞吐后端直接关闭 CUDA Graph，因为它的 kernel 与 graph 不兼容（vllm/config/compilation.py:1397-1413）。
不可捕获的算子通常落在四类：需要 CPU 同步的算子（.item()、动态控制流）、动态分配或数据依赖形状的算子、被 backend 声明为 NEVER 的 attention、以及被显式排除在 splitting_ops 之外的 op。vLLM 的 splitting_ops 默认包含一组 attention / linear-attention / kv-cache 相关 op（vllm/config/compilation.py:921-943），并且在 Dynamo 分割路径下还会把 unified_kv_cache_update 等加入断点（:1367-1368）。
编译与图化的交互边界也在这里：fuse_attn_quant 需要看整张图，无法与 piecewise compilation 共存，vLLM 会把 splitting_ops 置空并切到 FULL（vllm/config/compilation.py:1415-1429）；use_inductor_graph_partition 是 torch>=2.9 的实验路径，在 Inductor codegen 阶段切分，而不是在 Dynamo 之后切 FX 图（vllm/config/compilation.py:820-838、:1159-1166）。cudagraph_specialize_lora 默认开启，会为「无 LoRA」与不同 active LoRA 数分别捕获图，增加图数量和启动时间；关掉它可以省图，但每次都要跑 LoRA 算子（vllm/config/compilation.py:811-818）。
最后，--enforce-eager 仍然是最终退路：它同时关掉编译和 CUDA graph，启动最快，但稳态 decode 失去提交开销的优化（docs/configuration/optimization.md:25）。如果你的 backend 是 NEVER、batch 形状高度动态、或者模型里有大量数据依赖控制流，eager 或 PIECEWISE 往往比强行 FULL 更稳。