CUDA Graph 完全拆解：从逐次 Kernel Launch 到可重放推理执行图

显式 Graph API：自己创建节点和边#

显式方式先用 cudaGraphCreate 创建空图，再为 kernel、memcpy 等操作构造节点参数，用依赖节点数组表达边。其优势是节点句柄、拓扑和参数都在应用手中，适合需要精确更新少数节点、拼接子图或生成非线性拓扑的底层库。代价也直接：原本一段清晰的 stream 代码，要改写成繁琐的图构造代码；一旦库调用内部发射多少个 kernel 并不透明，手工复刻几乎不可维护。

较新的 CUDA Programming Guide 推荐统一的多态接口 cudaGraphAddNode()，节点类型由 cudaGraphNodeParams.type 决定；早期为不同节点分别提供的 cudaGraphAddKernelNode()、cudaGraphAddMemcpyNode() 等接口将逐步退居兼容位置。核心逻辑可以抽象为：

1
cudaGraph_t graph;
2
cudaGraphNode_t load, compute, store;
3
cudaGraphCreate(&graph, 0);
4

5
// 省略各节点 cudaGraphNodeParams 的填写。
6
cudaGraphAddNode(&load, graph, nullptr, nullptr, 0, &load_params);
7
cudaGraphAddNode(&compute, graph, &load, nullptr, 1, &compute_params);
8
cudaGraphAddNode(&store, graph, &compute, nullptr, 1, &store_params);

这段图表达 load → compute → store，而不是要求 CPU 在每次运行时依次提交三项工作。若两个节点只共享前驱但互不依赖，它们可以在前驱完成后并行。图 launch 所在 stream 只负责图与外部异步工作的顺序，不会把图内部所有节点强制压成单 stream 串行。

Stream Capture：把现有异步代码录成图#

Stream Capture 更适合框架和已有 CUDA 程序。cudaStreamBeginCapture 到 cudaStreamEndCapture 之间提交到 stream 的工作不会立刻执行，而是追加到正在构造的图中：

1
cudaGraph_t graph;
2
cudaGraphExec_t exec;
3

4
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
5
embedding<<<grid1, block1, 0, stream>>>(...);
6
rmsnorm<<<grid2, block2, 0, stream>>>(...);
7
attention_library_call(stream, ...);
8
cudaStreamEndCapture(stream, &graph);
9

10
cudaGraphInstantiate(&exec, graph, 0);
11
for (int step = 0; step < decode_steps; ++step) {
12
    cudaGraphLaunch(exec, stream);
13
}

attention_library_call 内部即使调用了多个可捕获的 CUDA 库操作，也能被一起录入，因此改造面远小于显式建图。但「录制」这个词容易制造两个误解。第一，捕获期不是执行一遍再记住结果，而是记录未来要执行的 CUDA 操作；第二，图不会记录 CPU 分支背后的所有可能路径，只会得到捕获时实际提交的那条路径。

跨 stream 也能捕获，但必须把 fork–join 完整表达出来。原始 stream 记录 event，其他 stream 等待该 event 后进入同一 capture graph；结束前，分支 stream 再记录 event，并由原始 stream 等待，把所有分支汇合回来。cudaStreamEndCapture 必须在最初调用 begin 的 origin stream 上执行。漏掉 join 会使整个捕获失败。

捕获期间不能对被捕获的 stream 或 event 做同步/状态查询，因为此时它们代表的是尚未排队执行的图节点。某些分配、隐式同步、CPU 读取 GPU 结果以及不支持捕获的第三方库调用，也会破坏捕获。出现非法操作后，capture graph 会进入 invalidated 状态；程序仍应调用 cudaStreamEndCapture 退出捕获，但拿到的是错误和空图。这里最危险的做法是吞掉返回码，直到第一次 replay 才寻找不存在的执行图。

混合建图：捕获现有库，再手工改图#

两种方法可以组合。程序先用 cudaStreamBeginCaptureToGraph() 把一段现有库调用接入指定图，再通过显式 API 添加条件节点、额外依赖或可单独更新的节点。这个路径适合中间层：既不想重写 cuBLAS、cuDNN 或框架内部的提交逻辑，又需要拿到某些节点句柄做细粒度控制。

实际选型可压缩为三条：完整掌握底层节点且拓扑需要精确编辑，用显式 API；已有一段成熟的异步 stream 代码，用 capture；大部分流程来自库调用、少数节点需要定制，用混合建图。不要因为显式 API 看起来「更底层」就默认它更快。两条路径最后都会得到 cudaGraph_t，性能差别主要来自最终拓扑、实例化结果和重放方式，而不是图最初如何构造。

参数变了，不一定要重新捕获#

图的「静态」不等于所有数值永远相同。CUDA 把变化分成三档：

• 数据内容变化：地址和形状不变，只把新 token、position、slot mapping 等写入静态缓冲区，直接 replay；
• 节点参数变化：拓扑相同，但 kernel 参数或内存地址变化，用 individual node update 或 whole graph update；
• 拓扑/节点类型变化：通常必须重新实例化，必要时重新捕获。

少量已知节点变化时，可用 cudaGraphExecKernelNodeSetParams()、cudaGraphExecMemcpyNodeSetParams() 等接口直接更新 executable graph，跳过无关节点的拓扑比较。大量节点变化或节点来自 opaque library capture 时，可重新得到一张拓扑相同的 cudaGraph_t，再调用 cudaGraphExecUpdate() 更新已有实例。如果更新失败，应用应回退到重新实例化，而不是把失败当成「CUDA 会自动调整」。

更新能力有明确边界。例如，kernel 的函数拥有者不能任意跨 CUDA context 改变；memcpy 的源/目的内存类型和传输类型不能随意替换；conditional node 的主体结构和句柄对应关系也受约束。更新解决的是「同一工作流的参数变化」，不是把一张图原地变成任何另一张图。

这一层区分对推理服务尤其关键。把 batch size 从 8 改成 13，可能同时改变 grid、attention 元数据长度、KV slot 数和某些后端选择，未必是廉价参数更新能覆盖的变化。生产框架更常见的做法不是每轮 update 整张图，而是为一组代表性 shape 建立图缓存，以 shape/batch 描述符查找可重放实例；缓存未命中时选择 padding、捕获新图或回退 eager。

指针稳定：replay 不会重新做 Python 参数绑定#

以 PyTorch 的 raw graph API 为例，捕获时 kernel 参数里记录了输入、输出和临时张量的设备地址。下一次 replay() 不会重新遍历 Python 对象并把一个新 tensor 的地址塞进去。因此正确模式是保留捕获时的长生命周期 static_input，每轮把真实输入复制进这块地址，再重放：

1
static_input.copy_(real_input)
2
graph.replay()
3
result = static_output.clone()  # 仅在结果需要跨下一次 replay 保存时复制

把 static_input = real_input 重新绑定成另一个 Python 对象没有作用，图里的 kernel 仍会读取旧地址。类似地，static_output 通常也复用同一地址；若调用方把它长期保存，而下一轮 replay 又覆盖这块内存，就会出现结果悄悄变化的别名 bug。

图内临时分配也必须有稳定语义。PyTorch caching allocator 会在 capture 期间从 graph-private memory pool 满足分配，使 replay 可以复用相同虚拟地址。多个图可以共享私有 pool，但只有在它们的重放顺序与捕获顺序一致、且不会并发时才安全。用一份显存换低 CPU 开销，是 CUDA Graph 在推理框架里的真实成本之一；shape bucket 越多，图实例、静态输入和私有内存池占用越值得监控。

原生 CUDA 的 graph memory node 提供了更底层的机制：allocation node 创建的虚拟地址在图生命周期内固定，物理内存可在生命周期不重叠的 allocation 之间复用；分配与释放按 GPU 执行顺序而不是 CPU API 调用顺序生效。固定虚拟地址正是让其他节点无需每次更新指针的基础，但销毁图不等价于自动释放所有仍存活的 graph allocation，生命周期管理不能省略。

形状稳定：固定的是执行变体，而非业务请求#

在线服务的 batch size、prompt 长度和每个请求的解码状态不断变化，看起来与 CUDA Graph 天生冲突。常用解法是 bucket 化。设实际 token 数为 $b$，已捕获尺寸集合为 $S=\{s_1,s_2,\ldots,s_m\}$，运行时选不小于 $b$ 的最小 bucket：

真实数据写入前 $b$ 个槽位，其余 $s(b)-b$ 个槽位用 padding、mask 或空请求占位。这样 graph 看到的是固定 shape，调度器看到的仍是动态流量。代价是 padding 工作量：

bucket 很密，$\rho_{\mathrm{pad}}$ 小，但捕获时间和图缓存显存上升；bucket 很稀，图少了，却让 GPU 计算大量无效槽位。最佳集合不是固定的 $1,2,4,8,\ldots$，而应由线上 batch 分布、不同尺寸的 kernel 效率和显存预算共同决定。命中频繁的尺寸可以单独捕获，长尾尺寸则向上取整或回退 eager。

控制流稳定：CPU 不能在 replay 中间重新做决定#

如果模型先运行一个 GPU kernel，再把结果同步回 CPU，最后由 Python if 决定执行专家 A 还是专家 B，那么整段无法按传统方式捕获为一张固定图。数据相关控制流的选择发生在 replay 内部，而 CPU 在这次 cudaGraphLaunch 中没有重新参与逐节点调度。

有四类处理办法：

1. 把分支改写成 predication 或 mask，固定执行拓扑；
2. 只捕获分支两侧稳定的子区间，让动态边界保持 eager；
3. 为有限分支分别缓存图，由 CPU 在图外选择；
4. 在合适的 CUDA 版本和节点约束下使用 conditional graph node，让设备端条件驱动 IF、WHILE 或 SWITCH 子图。

第四种能力扩展了图能表示的动态性，但不是「任意 Python 控制流自动可捕获」。conditional body 允许的节点类型、设备范围、内存操作和更新方式仍有限制，高层框架也需要显式映射。工程上更稳妥的原则是先划出静态岛，再决定是否值得把动态边界继续下沉。

Prefill 和 decode 为什么要区别对待#

Prefill 的 token 数跨度大，attention 计算量高，常由少数持续时间较长的 kernel 主导。此时 launch overhead 占比小，强行 padding 到大 bucket 可能得不偿失。Decode 每个请求通常只新增一个 token，batch 内 query length 更整齐；单步会重复数十层短 kernel，又要执行几百到几千步，是图重放最理想的区域。

P/D 分离部署进一步强化了这一区别：decode 实例可以只准备 uniform-decode 图，避免为复杂 prefill/mixed batch 保存大量图和私有池；prefill 实例则保留 eager 或只图化稳定子段。图优化的边界由工作负载决定，不必追求「全模型、全阶段、全流量」这一种形式。

vLLM 的官方设计把这种取舍显式化为多种运行模式：NONE 完全 eager；PIECEWISE 只捕获兼容区段，不兼容的 attention 等操作留在图外；FULL 捕获整个 forward；FULL_DECODE_ONLY 只为 uniform decode 使用完整图；FULL_AND_PIECEWISE 则按批次在完整图与分段图之间切换。运行时用包含 token 数、请求数、是否 uniform、是否含 LoRA 等信息的 batch descriptor 选择模式和图实例。这个设计的重点不是枚举值本身，而是承认 CUDA Graph 是调度器的一种执行资源，必须和 batch 形态协同。

图源：PyTorch 官方 CUDA Graph 文章。绿色 CPU 工作已经占满一侧，蓝色 GPU kernel 之间仍有大段空白，这是典型的 CPU-bound submission。不能只看 GPU kernel 自身 duration；真正影响端到端延迟的是从上一段有用设备工作结束，到下一段开始之间的空洞。

图源同上。图化区间通过一次重放提交后，kernel 在设备时间线上明显紧凑。这个现象不代表每个 kernel 的算法变快，而是 Driver 不再等 CPU 为相同序列逐一准备 launch。PyTorch 官方的 Llama 3 FP8 推理案例也给出过更贴近生成模型的微观数据：其模拟 Llama 3 70B attention 层中，相邻 GEMM 的间隙为 165 微秒，而 matmul 本身约 12 微秒；图重放后间隙降到 6.65 微秒。该实验报告单层相对朴素实现约 6.4 倍加速，但这是刻意暴露 launch overhead 的特定 attention 层，不应外推成完整 LLM 的固定加速比。

与连续批处理、算子融合、persistent kernel 的关系#

连续批处理在每个迭代边界重新组合活跃请求，提高 GPU 工作量和 KV 利用率；CUDA Graph 减少组成一个迭代的 kernel 提交成本。二者会发生张力：调度越灵活，batch descriptor 的组合越多，图缓存越复杂。关于迭代级调度的基本机制可参考 连续批处理完全拆解 和 vLLM 调度器源码拆解。

算子融合减少节点数、内存流量和 launch 次数；CUDA Graph 降低剩余节点的批量提交成本。persistent kernel 或 megakernel 则更进一步，把调度器和多个操作常驻到 GPU 内部。三者大致处于不同层级：融合改变 kernel 边界，Graph 改变提交方式，persistent kernel 改变执行组织。Graph 的改造成本通常低于重写 megakernel，也保留了库 kernel 的模块性；但它仍有每次 graph launch、bucket 分派和图内节点调度成本，不等于 GPU 内部常驻 runtime。

为什么必须 warmup#

第一次运行模型可能触发 cuBLAS/cuDNN 句柄初始化、JIT 编译、autotune、显存池扩展或其他一次性动作。这些工作有的不能被捕获，有的即使能捕获也不应成为每次 replay 的一部分。warmup 的任务是把「只该做一次」的状态建立完，再记录稳定热路径。

warmup 放在 side stream 上，不只是风格偏好。PyTorch 官方 CUDA semantics 要求在用于捕获的上下文之外完成这些准备，并用 stream dependency 保证它结束后才开始 capture。少了前后的 wait_stream，可能出现默认 stream 尚未完成模型初始化，捕获 stream 就开始读写同一资源的竞态。

输入复制为何可以留在图外#

上例的 static_input.copy_(real_input) 没有被捕获。这样做保留了请求接入层的灵活性：真实输入对象可以每次不同，只要形状、dtype 和 device 匹配，复制后图内仍读取固定地址。复制本身也是异步 CUDA 操作；它和 graph.replay() 提交到同一当前 stream 时，stream 顺序保证先复制再计算，不需要中间 torch.cuda.synchronize()。

如果 H2D 输入路径也高度稳定，可以把 memcpy 作为图节点捕获；但 host 地址、pinned memory 生命周期和请求并发会变得更难管理。推理服务通常先图化 device-side model runner，把网络接入、tokenization、CPU scheduler 和不规则 H2D 路径留在图外。

多个 shape 怎么管理#

生产代码应把单个 graph 扩展为缓存项，每项至少包含：bucket key、CUDAGraph、静态输入、静态输出、额外元数据缓冲区和估算显存。伪代码如下：

1
class GraphEntry:
2
    def __init__(self, graph, inputs, outputs, bucket_size):
3
        self.graph = graph
4
        self.inputs = inputs
5
        self.outputs = outputs
6
        self.bucket_size = bucket_size
7

8
def run(batch):
9
    key = choose_bucket(batch)
10
    entry = graph_cache.get(key)
11
    if entry is None:
12
        return run_eager(batch)  # 也可低频捕获，但不要阻塞关键请求
13

14
    pack_into_static_buffers(batch, entry.inputs)
15
    entry.graph.replay()
16
    return unpack_valid_rows(entry.outputs, batch.size)

缓存淘汰不能只按图对象数量。一个图实例背后可能关联 graph-private pool 和多块静态 tensor，真正约束是总显存、捕获耗时与命中率。若在线临时捕获，首个命中某 bucket 的请求会承担明显延迟尖峰；更可控的方式是在模型加载后根据流量画像预热热点 bucket，长尾保持 eager。

torch.compile(mode="reduce-overhead") 会在适合的区段利用 CUDA Graph 来降低 Python 和 launch overhead，适合不想手工管理 raw graph 的场景；但它仍受 mutation、动态 shape、CPU-GPU 同步和 graph break 影响。自动化封装减少了 API 工作量，不会取消底层约束。遇到性能不符预期，仍要查看实际捕获区间、重编译次数和 eager fallback。

先看时间线，再跑端到端基准#

推荐按以下顺序测量：

1. 用 Nsight Systems 检查 CPU API 与 GPU kernel 时间线，确认 kernel 间存在可见空洞；
2. 把 capture、instantiate、first launch 和 steady replay 分开计时；
3. 分别报告 batch、输入长度、输出长度、并发、dtype、量化、GPU、CPU、CUDA 与框架版本；
4. 对在线系统同时比较 TTFT、TPOT、P50/P95/P99 和吞吐，不只比较平均值；
5. 检查图缓存命中率、padding 比例和额外显存，避免把成本藏在总 tokens/s 外。

图捕获会改变 profiler 的观察方式，profiler 本身也可能给微秒级事件增加扰动。稳态延迟应使用足够多次重复和 CUDA event，端到端延迟则用 CPU wall clock 并在计时边界正确同步。下面是区分首次执行与稳态 replay 的简化测量：

1
start = torch.cuda.Event(enable_timing=True)
2
end = torch.cuda.Event(enable_timing=True)
3

4
# 不把 warmup、capture 和第一次可能的 upload 混入稳态数据。
5
for _ in range(10):
6
    graph.replay()
7
torch.cuda.synchronize()
8

9
start.record()
10
for _ in range(1000):
11
    graph.replay()
12
end.record()
13
end.synchronize()
14

15
print("steady replay ms:", start.elapsed_time(end) / 1000)

CUDA event 给出 GPU 时间域里的区间，不包含请求在 Python 队列、CPU 调度器和网络栈中的等待。若优化目标是单请求端到端延迟，还需在更外层用 wall clock 测量；若目标是隔离 graph 对设备提交的影响，则应同时保留 event 和 wall clock，两者差距本身就有诊断价值。

计算摊销阈值#

设逐项 eager 每轮为 $T_e$，稳定 graph replay 每轮为 $T_g$，捕获、实例化与预热总成本为 $C$。至少运行多少轮才能回本，可由下式估算：

若 $T_e\le T_g$，分母非正，这段工作没有图化收益；若单轮只省几微秒而捕获要数百毫秒，低频 bucket 可能永远无法回本。反过来，自回归 decode 的热点 bucket 每秒重放大量次数，即使单轮收益不大，累积 CPU 释放和尾延迟稳定性也可能很可观。

NVIDIA 在 CUDA 12.6、RTX 3060 与特定主机平台上的官方微基准显示：Ampere 上含 10 个以上节点的直线 kernel 图，重复 launch 的 CPU 开销约为 $2.5$ 微秒再加每节点约 $1$ 纳秒；同一测试在 CUDA 11.8 上近似为 $2$ 微秒加每节点 $200$ 纳秒。这个结果说明新版 Runtime/Driver 已把直线图的重复提交做得接近常数时间，但仅适用于文中测试的拓扑、架构与软件版本。多 root、非 kernel 节点、旧架构和复杂图都有不同曲线，部署决策仍须在目标机器上实测。

捕获直接报错或返回空图#

先检查 cudaStreamEndCapture 的返回码，并定位捕获区间内的同步、legacy default stream、动态分配和不支持 capture 的库调用。捕获中的 stream 并没有真正执行工作，所以对它调用 cudaStreamSynchronize 或查询状态在语义上就不成立。跨 stream 时再检查所有分支是否通过 event 汇回 origin stream。

排查时应从最小区间开始：先捕获一个已知安全的 kernel，确认基础链路；再逐段扩大范围，直到找到破坏 capture 的调用。一次包住整个 model forward 后只看到「capture invalidated」，通常很难直接定位根因。

replay 能跑，但结果错误#

优先审计地址和生命周期：是否把新 tensor 重新绑定给变量，却没有复制进静态输入；捕获时使用的临时 tensor 是否已被释放；输出是否在下一次 replay 后仍被外部持有；多个请求或多个 graph instance 是否并发写入同一静态 buffer；padding 区是否被正确初始化和 mask。

第二类错误来自捕获时的 CPU 分支。若 if tensor.item() > 0 在捕获前把 GPU 值同步到 CPU，图里只留下当时选择的分支；后续数据改变不会让 Python 重新选择。解决方式是把分支留在图外、分别缓存分支图，或用真正受支持的设备端条件表达。

replay 正确，却没有变快#

先确认瓶颈是否真在 launch。若 GPU 时间线已经被长 GEMM 填满，图能省的只是被完全掩盖的 CPU 工作。然后检查：捕获范围是否太小；每轮是否仍在图前后同步；是否为输出 clone 付出比 launch gap 更大的代价；是否因为 padding 执行了过多无效 token；图缓存是否频繁 miss；是否把捕获或第一次 upload 混进每轮计时。

还要避免把「CPU API 更短」误认为「请求一定更快」。CPU 被释放后可以去做调度、网络和下一批准备，这可能提升系统吞吐或降低抖动；但如果 GPU 本来就是满载计算，单请求设备延迟未必显著下降。优化目标必须在 benchmark 中明确。

显存突然增加#

检查捕获尺寸数量、每个图的静态输入输出、graph-private pool、workspace 和并发实例数。多个图共享内存池只在重放顺序固定且不并发时安全；为了省显存错误共享，代价可能是数据破坏。对于部署系统，更合理的是设置图缓存预算，记录每个 bucket 的命中率与增量显存，并淘汰长期不命中的实例。

多线程或多 stream 出现串行化#

同一个 executable graph 不能与自身并发，这是常被忽略的串行化来源。即使把它 launch 到不同 stream，CUDA 也会让后一次执行排在前一次之后。并发服务要么让调度器合并请求后用一个实例运行更大 batch，要么维护多个独立 graph exec 和不冲突的缓冲区。图 launch 所在 stream 只规定外部顺序；图内部并行由依赖边决定，不会因为换一个 stream 就重新解释拓扑。