这篇“CUDA Graph”主题的博客，标题是“CUDA Graph 省下的，真的是 launch 时间吗？”，需要写一篇不少于5000字的中文技术分享。要像资深从业者口吻，在社区里发的那种。不重复、不套路。要覆盖：CUDA Graph原理、launch开销到底指什么、graph捕获与重放、与常规kernel launch对比、什么时候用值得、什么时候没用、常见误解、性能对比、注意事项等。还要结合热搜词？热搜词基本不相干（accelerate launch、quartus等），我们不要使用那些，但可以适度提到launch的普遍含义。注意安全，无违规。
结构规划：
开头：直接以问题引入，说很多人拿CUDA Graph说省launch时间，但实测后发现收益不明显，为什么？我自己的理解。
主体章节：
先把"launch时间"拆清楚：CPU发kernel的调用路径、命令提交到GPU、launch到底包含哪些开销；与执行时间对比。解释CUDA Graph按图重放时减掉的其实是CPU侧的调度与驱动调用开销，而不是GPU执行时间。
CUDA Graph工作机制：capture/instantiate/launch，stream capture原理、依赖关系的记录，graph节点合并和依赖优化；graph launch一次的CPU耗时远低于上千个kernel逐个launch。
为什么实测收益不如预期：小kernel越密集收益越大；单一大kernel几乎没收益；memory copy、event同步也会被捕获；driver侧已经是异步批处理（前几个launch后CPU返回），但graph减少了每launch的固定CPU成本和kernel间的间隙。还要分析典型基准：比如多个小kernel、有依赖链、有间隙，graph才能把kernel间隙压缩掉。
什么场景真正值得用CUDA Graph：大模型训练/推理的step里大量小kernel（如transformer layer很多elementwise、attention算子），推荐用graph；短时间重复执行且kernel数量多；计算图本身就是静态的。不适用场景：动态shape、分支、每次变化大的图，或者重捕获开销大于收益（capture本身有开销，graph instantiate的编译优化，内存拷贝开销）。
实操要点与避坑：capture时一些api限制（不能同步等待、不能使用malloc、必须用stream capture机制），捕获到错误怎么办；如何测量launch时间（用nsys、event timing），建议做microbenchmark；graph的更新方式（graphExecutable update，不要每次重新instantiate）；记忆体分配与pool；如何判断是否需要。
常见误解与FAQ：launch overhead是0吗？CPU时间与GPU时间重叠；graph并不能减少GPU kernel本身时间；graph快是因为合并了kernel？不准确；还要注意多stream caution端。最终以个人经验收尾：我自己在生产里用它收益最大的地方是静态训练step的图，动态场景别强行套。
需要满足：H2编号，H3编号，禁mermaid，禁emoji，禁总结性套话。以满足5000字以上。开篇不要“随着”。要直接有吸引力：“一个老生常谈的问题”。
内容要具体，像真实的人写的。示例代码或bash命令可以，如循环launch和graph launch伪代码。可以加入表格比较。需要多段每段至少150字，总字数要够。我大约需要写到6000-7000字。我们不需要真的引用外部资料，但要可信。
开始撰写。注意避免AI痕迹。不用“本文将”。标题不需要，直接从 ## 1. 开始？规范：不需要主标题，直接从二级标题开始。但用户给了标题，我们可以不重复，直接从正文段落开始再进入第一个H2？我需要先有开头段（≥200字），然后主体用 H2。开头不需要标题，直接从文字开始。按要求“不需要主标题，直接从二级标题开始”，但从开头段开始也可以？之前类似任务：开头几个自然段，然后## 1.。这样OK。不要前置标题，直接文字。开始。
要注意每个H2标题编号：## 1. xx。H3格式 ### 1.1。
开始正文。 先说结论：CUDA Graph 真正省下的，不只是 launch 那一下 CPU 调用，更多是 kernel 与 kernel 之间那段“看不见的空白期”。但这不代表所有场景都能拿到收益。我自己第一次用 CUDA Graph 做 benchmark 的时候，单跑一个大 kernel，跑完一看耗时几乎没变，一度怀疑 API 用错了。后来把图拆成几十个小 kernel 的链，同样的计算量，性能差距才真正拉开。这篇文章就从 launch 开销到底是什么开始讲，把 CUDA Graph 的收益来源、适用边界和常见坑一次说清楚。
1. launch 时间到底包含了什么
很多人说“kernel launch 开销”的时候，脑海里想的是 CPU 调用 cudaLaunchKernel 到 GPU 开始执行之间的那段时间。这个理解没错，但太粗糙了。实际的 launch 开销由好几个部分组成，如果不拆开看，就很难解释为什么 CUDA Graph 有时候收益巨大，有时候毫无帮助。
1.1 CPU 侧的固定成本
在 CPU 侧，每次调用 cudaLaunchKernel 都要经过 CUDA runtime 库，完成参数校验、命令打包、写入命令缓冲区等操作。这一套流程里有不少固定成本：函数调用栈、参数结构体拷贝、stream 内部锁、安全检查。单个 kernel 的 launch 在普通机器上大概要花 3 到 10 微秒，具体看驱动版本、操作系统和调用方式。如果你在循环里连续 launch 100 个 kernel，光是 CPU 侧发出这些命令可能就要几百微秒。
这还没算更糟的情况。如果每次 launch 之间还要插入 cudaStreamSynchronize、cudaDeviceSynchronize 或者 cudaMemcpyAsync 的隐式同步，CPU 会被迫等待 GPU 追上来，那 launch 开销就会成倍放大。很多性能不达标的代码，问题并不是 kernel 本身慢，而是 CPU 在 launch 循环里被各种同步点卡住，GPU 一直在等命令。
1.2 GPU 侧的调度与执行延迟
GPU 收到命令后，也不是瞬间就能开始执行。命令先要到 GPU 前端队列，由工作调度器解析，分配硬件资源，然后才真正启动 kernel。这个“前端解析”的延迟通常在 1 微秒左右。如果命令流里有多个 kernel，它们之间的启动间隙往往比 kernel 本身的执行时间还长。
特别是那种每个 kernel 只跑几微秒的小算子场景。你 launch 了 1000 个 elementwise 操作，每个 GPU 上只跑 2 微秒，但每个之间的调度间隙可能有 5 微秒。结果就是 70% 的 GPU 时间都空转在调度上。这种情况下，CUDA Graph 的价值根本不是优化某个 kernel 的性能，而是把“kernel 间空隙”压缩到几乎为零。
1.3 launch 与执行的流水线重叠
还有一个常见误解：launch 开销和 kernel 执行时间可以重叠，所以无所谓。确实，CPU 可以连续发命令，GPU 在后端异步执行，两者在时间轴上可以流水线化。但这种重叠是有限度的。CPU 如果不小心在某个 launch 之后做了 cudaStreamQuery、cudaEventQuery 或者依赖 host 数据的判断，流水线立刻断掉。
即使没有显式同步，CPU 侧 launch 命令的生成速度也可能跟不上 GPU 的执行速度。假设一个 kernel 在 GPU 上跑 5 微秒，CPU launch 需要 8 微秒，那 CPU 就成了瓶颈，GPU 每执行完一个 kernel 就要等新的命令。这时候 CUDA Graph 能做的，就是把 CPU 发命令的时间从“每个 kernel 8 微秒”降成“整张图一次 20 微秒”，GPU 侧就能连续跑完整个计算链。
打个比方，常规 launch 就像你站在流水线旁边，每加工完一个零件，你才手动放上去下一个。CUDA Graph 相当于把一整批加工顺序提前写成了程序，机器可以一口气跑完。这个类比虽然粗糙，但能帮你抓住关键：graph 优化的是“编排”，不是“单个零件的加工速度”。
2. CUDA Graph 是怎么把 launch 省下来的
CUDA Graph 的核心机制是捕获（capture）和重放（replay）。它不是通过某种魔法加快了 cudaLaunchKernel 的调用速度，而是把原本在 host 侧逐一发出的命令流，在捕获阶段就固化成一张图，之后每次执行只调用一次图启动接口。
2.1 stream capture 捕获过程
用起来不复杂。先创建一个 capture 模式的 stream，调用 cudaStreamBeginCapture，然后像平常一样在这个 stream 上 launch kernel、拷贝内存、记录 event。这一系列操作不会真的在 GPU 上执行，而是被 CUDA 驱动记录下来，变成图中的节点和依赖边。最后调用 cudaStreamEndCapture，得到一张 cudaGraph_t。
这里有第一个坑：捕获期间，你不可以在 host 侧做任何依赖于 kernel 执行结果的同步操作。比如 cudaMemcpy 到 host 内存、cudaStreamSynchronize，甚至某些带默认同步语义的 API，都会导致捕获失败。因为捕获阶段 GPU 不执行，任何需要读回数据的操作都没有意义。我第一次捕获时就因为在 capture stream 里加了一次 cudaMemcpy 到 pinned memory，直接报错，后来改成 cudaMemcpyAsync 到 device memory 才顺利。
捕获得到的图本质上是一个有向无环图（DAG）。节点可以是一个 kernel、一次 memcpy、一个 event 记录，节点之间有依赖边，表示执行顺序。这个依赖关系不仅来自你调用 API 的顺序，还来自 kernel 之间实际访问的内存依赖。CUDA 会尽量自动分析，但如果你用同一个 stream 顺序调用，那依赖关系就是串行的。
3.2 实例化与图执行
得到 cudaGraph_t 之后，还需要调用 cudaGraphInstantiate 生成一个可执行的图（cudaGraphExec_t）。这一步会做很多编译和优化工作，包括把相邻 kernel 的启动参数打包、分析依赖关系、尝试合并某些节点。实例化可能比较耗时，通常需要几十到几百微秒，对于复杂图可能更久。所以生产环境里应该只在图结构发生变化时重新实例化，不要每次 inference 都做。
执行图时，调用 cudaGraphLaunch(exec, stream)，整张图被提交给 GPU。CPU 侧这一个函数调用只花几十微秒，GPU 侧则按照图内部的依赖关系自动调度所有节点。由于 CPU 只需要发出一次命令，所有 kernel 的执行顺序由 GPU 前端自己编排，避免了 host 逐条 launch 带来的空隙。
对比一下：
| 执行方式 | CPU 命令次数 | CPU launch 开销 | kernel 间间隙 |
|---|---|---|---|
| 常规循环 launch | N 次 | N × 固定开销 | 每次调度都有间隙 |
| CUDA Graph 一次执行 | 1 次 | 单次图启动开销 | 依赖内节点间隙极小 |
注意，即使使用 CUDA Graph，各个 kernel 在 GPU 上仍然需要逐个调度执行，不可能真正并行。但如果两个节点没有依赖关系，CUDA Graph 的调度器可能让它们在支持并行的 stream 上重叠执行，这是额外收益，不过前提是硬件资源足够、kernel 之间没有内存冲突。
3. 什么场景收益最大，什么场景白折腾
判断你的程序适不适合 CUDA Graph，不要看“是不是深度学习模型”，而要看 kernel 粒度、kernel 数量和依赖结构。
3.1 小 kernel 密集、依赖链深的场景
收益最大的场景是这样的：一次 forward 或者一次迭代里，有成百上千个计算量很小的算子，比如 layer norm 里的均值方差、残差相加、激活函数、softmax 的中间步骤。这些算子单独执行时间可能只有 1 到 3 微秒，CPU 逐个 launch 的开销占比极高。
这种场景下，CUDA Graph 带来的端到端加速可能达到 20% 甚至更高。注意，这 20% 不是 GPU 算得更快了，而是 GPU 空转的时间少了。指令吞吐没有提升，但有效利用率提升了。我跑过一个 transformer block 的对比，常规 launch 版本每个 block 大约 280 微秒，graph 版本大约 230 微秒，省下的 50 微秒大部分就是 launch 和调度间隙。
3.2 大 kernel 占比高的场景
反过来，如果你的程序里每个 kernel 都跑 50 微秒以上，一共只有两三个 kernel，那 CUDA Graph 基本没什么用。launch 开销相对于执行时间微不足道，即使省下几十微秒，在总耗时里占比也不到 1%。这种场景强行上 graph，反而要承担捕获和实例化的额外开销，得不偿失。
有个很常见的误区：以为 CUDA Graph 通过某种方式把多个 kernel 融合成一个 kernel。不是的。CUDA Graph 不会去改写你的 kernel 代码，它只是改变了命令提交和调度方式。如果 kernel 本身写得非常低效，graph 也救不了你。
3.3 动态 shape 和分支逻辑：别硬上
CUDA Graph 要求图结构在执行时保持稳定。如果你的模型有动态 shape、batch size 经常变化、或者控制流里有 if 语句导致执行路径不同，那捕获出来的图可能失效。每个不同的 shape 或路径都要重新捕获和实例化，而这本身是有开销的。
我遇到过最典型的情况：一个推理服务，输入长度不定，导致 attention mask 形状变化，整个计算图中一部分 kernel 启动参数也跟着变。一开始想用 graph 加速，但每个请求都要重新 instantiate，图实例化的时间比省下的 launch 时间还长，最后放弃了。正确做法是固定 batch 和序列长度，或者用 CUDA Graph 的更新机制在已有图上只修改参数。
CUDA Graph 提供 cudaGraphExecKernelNodeSetParams 这样的接口，可以只更新某个节点的 kernel 参数，而不用重新实例化整张图。但要注意，更新节点参数也有开销，而且只适用于参数变化不改变节点数量和依赖关系的情况。如果图中节点数都变了，那还是重新实例化吧。
3.4 与 CPU 端计算交错的场景
如果你的程序是 CPU 和 GPU 交替工作：CPU 算一段，launch kernel 算一段，然后 CPU 又需要读取 GPU 结果继续运算，这种模式用 CUDA Graph 收益很小。因为图只是把 GPU 侧命令打包了，CPU 端的计算和同步等待并没有消失。图的优势在于一次性提交多条 GPU 命令，但 CPU 在提交后如果立刻同步等待图执行完，那和逐个 launch 后同步等待，总耗时可能没区别。
这种情况下应该优先考虑把 CPU 计算也流水线化：提前准备下一轮输入、使用双缓冲、用 cudaEvent 异步通知，而不是把宝都押在 graph 上。
4. 实操过程：从捕获到性能对比的全流程
下面这一套流程是我自己验证 CUDA Graph 收益时用的标准做法，你可以直接照着抄。步骤不复杂，但每个细节都会影响结论。
4.1 准备一个公平的 benchmark
不要直接拿整个应用跑，先做一个微基准。选一组 kernel 序列，比如 100 个 elementwise 操作加 20 个 reduction，模拟真实计算链。分别实现两个版本：常规循环 launch 和 CUDA Graph launch。两者执行的计算完全相同，唯一的区别是启动方式。
由于 CUDA 是异步执行，计时一定要用 cudaEvent，放在 CPU 循环外面。常规版本的代码结构大致是：
cudaEventRecord(start, stream); for (int i = 0; i < N; ++i) { kernel_i<<<grid, block, 0, stream>>>(args_i); } cudaEventRecord(stop, stream); cudaEventSynchronize(stop); float ms; cudaEventElapsedTime(&ms, start, stop);graph 版本则是：
// 捕获阶段，只需一次 cudaStreamBeginCapture(stream, cudaStreamCaptureModeThreadLocal); for (int i = 0; i < N; ++i) { kernel_i<<<grid, block, 0, stream>>>(args_i); } cudaStreamEndCapture(stream, &graph); cudaGraphInstantiate(&graphExec, graph, 0); // 计时并执行 cudaEventRecord(start, stream); cudaGraphLaunch(graphExec, stream); cudaEventRecord(stop, stream); cudaEventSynchronize(stop); cudaEventElapsedTime(&ms, start, stop);注意常规循环版本里，如果每个 kernel_i 都有不同的参数，你还需要在循环里做一些 CPU 端的判断、索引计算，这其实也包含在 launch 开销里。graph 版本在捕获时就把参数固定下来了，执行时不再有这些 CPU 成本。如果你的真实程序里 launch 之间还有 host 计算或日志打印，那 graph 的收益还会更大。
4.2 用 Nsight Systems 看时间线
只用 cudaEvent 统计总耗时还不够。我强烈建议用 nsys 生成时间线，直观地看 kernel 之间的空隙到底有多大。命令很简单：
nsys profile -o my_app --trace=cuda ./my_app打开生成的 .nsys-rep，你会看到 GPU stream 上的时间轴。常规 launch 版本里，每个 kernel 之间往往有一段空白，短的几百纳秒，长的可能几微秒。graph 版本里，这些空白明显变短，整个 stream 看起来连续得多。
有一次我排查一个模型，发现常规版本里 kernel 之间空白平均有 8 微秒，但单个 kernel 平均才 4 微秒。也就是说 GPU 一半时间在等命令。用 graph 重写后，端到端时间缩短了 40%。这些空白不是靠猜的，而是在时间线上肉眼可见的。
4.3 捕获时的内存与 stream 注意事项
捕获阶段特别容易踩内存相关的坑。在 capture stream 上使用 cudaMalloc 是不被允许的，至少很多驱动版本会有问题。建议使用 CUDA 的流序内存池（stream-ordered memory allocator），也就是 cudaMallocAsync。这样分配内存本身就参与 capture，执行图时可以恢复到固定状态。
另外，如果你需要在 graph 执行时动态改变某个 kernel 的输出 buffer 地址，必须用 cudaGraphExecKernelNodeSetParams 更新节点参数，并在更新前用 cudaGraphKernelNodeGetParams 拿到当前参数。很多网上抄来的代码没有这一步，直接把新指针传给原 kernel 的启动参数，结果图执行时仍然用的旧地址，导致数据错乱。
这个问题我 debug 过一整天。图捕获时 kernel 参数的指针如果是临时变量的地址，捕获后临时变量销毁，graph 执行时就是悬空指针。正确做法是让参数指向长期存活的内存，或者捕获结束后用更新接口设置最终地址。
4.4 如何判断 graph 是否真的有效
判断有没有效，不要只看总时间变化。建议拆三个指标：CPU 侧 launch 命令所花时间、GPU 侧 kernel 总执行时间、GPU 侧端到端时间。
如果 GPU 侧 kernel 总执行时间（把每个 kernel 的执行时长相加）没变，但端到端时间变短了，说明省下的就是 launch 和调度间隙。如果连 kernel 总执行时间都变了，要么是 scheduling 发生了重叠，要么是优化器做了节点合并，要么是你的 benchmark 有缺陷。
还有一个容易被忽略的因素：小数据量测试时，CPU 初始化和 PCIe 传输可能主导耗时，graph 的收益会被淹没。所以 benchmark 要建立在 GPU 计算为主的前提下，数据量也不能太小。
5. 常见问题与排查心得
5.1 捕获报错 cudaErrorStreamCaptureImplicit
最常见的是在 capture 作用域里调用了 cudaMemcpy（sync 版本），或者 cudaStreamSynchronize。捕获期间不能被阻塞，因为图没有实际执行。把同步版本全部改成 Async 版本，并把 host 数据的拷贝移出 capture 区域，通常能解决。
还有一种情况：使用了默认 stream 的隐式同步。CUDA 里默认 stream 与其他 stream 同步的规则很隐蔽，一旦在捕获过程中触发了隐式同步，就会报错。稳妥的办法是不要用默认 stream，所有操作都在同一个非默认 stream 上执行，捕获模式选择 cudaStreamCaptureModeThreadLocal。
5.2 性能不升反降
如果你发现 graph 版本比常规版本还慢，先检查是不是每次执行前都重新实例化了图。cudaGraphInstantiate 可能涉及设备端代码编译，开销不小。应该把 graphExec 缓存起来复用。
另一个原因是图里的节点过多，且每个节点之间都有依赖边，导致 GPU 调度器需要处理更多的元数据。虽然单次 launch 的 CPU 成本降低了，但 GPU 前端解析图结构的成本可能反而升高。这种情况下，考虑减少节点数量：可以把某些不关键的 kernel 合并进相邻 kernel 里，或者把多个 elementwise 操作融合成一次 kernel launch。
5.3 graph 里的 kernel 没有按预期顺序执行
虽然图捕获遵循 stream 顺序，但如果没有显式添加依赖，某些节点可能被调度器优化成乱序执行。如果你依赖 node A 的结果去计算 node B，但两者使用的是不同 stream 或者没有事件依赖，可能出错。解决方法是确保所有有依赖的节点都在同一个 capture stream 里按顺序调用，或者使用 cudaGraphAddDependencies 手动添加边。
我自己就吃过这个亏：捕获时借用了一个临时 stream 去执行某个库内部的 kernel，结果图执行时那个 kernel 跑到了错误的位置。后来把所有库调用都统一到 capture stream 上，问题消失。
5.4 动态输入尺寸怎么办
如果输入尺寸偶尔变化，但不频繁，可以把 CUDA Graph 做成一个 LRU 缓存，按尺寸或 shape 存多份 graphExec。每次请求查缓存，命中就执行，没命中重新捕获。注意限制缓存数量，否则显存里堆一堆实例化图，占用可能超过你预期。
如果输入尺寸每个请求都变，那真的不建议用 CUDA Graph。去优化 kernel 本身，或者用动态并行，都比强行套 graph 实在。
5.5 CUDA Graph 与深度学习框架配合
在 PyTorch 里，torch.cuda.CUDAGraph 封装了这套机制。它的正确用法是 capture 一个 forward + backward 全流程，之后反复执行。但需要注意给模型分配静态显存缓冲区，不能有任何动态 new/delete。PyTorch 的捕获对动态 shape 很敏感，建议用静态 shape 跑一次 warmup，再开始捕获。
框架生成的 graph 不一定是最优的，因为框架层面可能插入了很多额外节点，比如 dtype 转换、梯度清零、内存拷贝。我遇到过一次，PyTorch 捕获出来的图里居然有一个 CPU 同步点，直接让 graph 收益归零。排查方式还是用 nsys 看时间线，找到那个明显不该存在的空隙。
6. CUDA Graph 之外：真正值得投入的优化方向
聊完了 CUDA Graph 的边界，我想说句实在话：很多项目不敢上 CUDA Graph，卡在动态 shape 和自动微分灵活性上。如果你的场景实在不适合 graph，也有一条曲线救国的路线——减少 launch 次数的镜像优化思路。
6.1 kernel 融合，从根上减少 launch 数量
CUDA Graph 解决的问题是“launch 次数多”，但更彻底的方案是“launch 次数少”。把多个 elementwise 操作写进同一个 kernel，每个线程处理完整的数据链，读写都留在寄存器或共享内存里。这样做既能省 launch 开销，还能减少中间数据的全局内存读写，收益往往比 graph 更大。
所以我在实践里的决策顺序通常是：先看有没有 kernel 融合空间；融合不了的，再考虑 CUDA Graph。如果时间允许，两个都做也常见：先融合一部分算子，降低节点数，然后用 graph 把剩下的节点编排起来。
6.2 用多 stream 并行要谨慎
CUDA Graph 支持多 stream 捕获，图内可以包含多个 stream 上的节点。但多 stream 并行带来的收益取决于硬件资源是否充足，以及节点之间是否存在真实可并行性。如果只是把一串串行 kernel 分散到多个 stream，反而可能引入额外的同步开销。
正确做法是利用 CUDA 事件来跨 stream 建立依赖，让没有依赖的节点真正并行。graph 在多 stream 场景下的调度比手动管理省心不少，但复杂度也高。建议先从单个 stream 开始用 graph，跑通了再扩展到多 stream。
6.3 测量先行，优化后置
没有数的优化都是耍流氓。你问“CUDA Graph 省下的真的是 launch 时间吗”，答案可以通过实验自己验证。写一个小的 benchmark，测量单次 kernel launch 的平均耗时、kernel 间平均间隙、graph launch 的 CPU 耗时。有了这些数，你再判断要不要引入 graph。
以我自己的经验，单次 kernel launch 的 CPU 开销在不同驱动版本上有明显差异。旧驱动可能 8 微秒，新驱动优化到 3 微秒。这就导致同一个 graph 优化，在旧驱动上收益 30%，在新驱动上只有 10%。但 kernel 间调度间隙的缩减是实打实的，基本不受驱动版本影响。所以如果你问“最稳定的收益来自哪里”，我的回答是：来自压缩 kernel 间的气泡，而不是单纯减少 CPU launch 次数。
6.4 共享内存和 bank conflict 别指望 graph 帮你
这个一定要提醒。CUDA Graph 不会优化单个 kernel 的实现。如果你的 kernel 因为共享内存 bank conflict、全局内存访问不合并等原因跑得慢，那是 kernel 本身的问题。graph 只能让慢的 kernel 更紧密地排列，但每个 kernel 的墙钟时间不会缩短。
我见过有人用 graph 包了一个性能很烂的 kernel，发现端到端时间只下降了几个百分点，于是得出结论“CUDA Graph 没用”。实际上 graph 已经尽全力把 launch 开销降到底了，剩下的瓶颈在 kernel 的效率。这时候应该回头优化 kernel，而不是质疑 graph。
7. 最后的实操体会
个人观点：CUDA Graph 是我见过最容易在 benchmark 里给人幻觉的技术之一。因为捕获阶段把一次计算链的启动过程“冻结”了，看起来高端且聪明，但如果你不理解 launch 时间的组成，很容易用错场景、量错指标、得出结论。
我自己的生产经验是：在固定 shape 的静态训练 step 上，CUDA Graph 收益稳定；在动态推理服务里，收益波动大，甚至经常负优化。所以在你的项目里，如果 kernel 平均执行时间小于 5 微秒且数量超过 500，请务必花一个下午试试 CUDA Graph。如果每个 kernel 都跑几十微秒，那我不太建议折腾，不如把时间花在 kernel 融合和内存访问优化上。
最后分享一个实用的小技巧：在已经能正常执行的 graph 上，用 cudaGraphExecKernelNodeSetParams 更新输入 buffer 指针时，先调用 cudaGraphExecKernelNodeSetParams 会触发一次参数校验，耗时可能不可忽略。如果你只需要更换输入地址，一个更快的办法是直接用 cudaGraphExecMemcpyNodeSetParams 更新 memcpy 节点，或者把输入指针通过 kernel 参数间接传入——比如指向一个 device 全局变量的地址。这样更新参数时改动最小，执行时也不会因为参数校验而增加额外延迟。
踩过几次坑之后，我给团队的规矩是：凡是要用 CUDA Graph 的模块，必须先在 nsys 里展示出“kernel 间有可见空隙”的截图。没有空隙，就不要上 graph；有空隙，graph 才能帮你把它们钉死。
