CUDA程序性能瓶颈有时并不来自GPU算力，而是来自CPU与GPU之间的命令提交管道。当一个数组被拆分成几十个小kernel依次执行时，每一次kernel调用都需要CPU端走完运行时检查、参数打包、驱动转发、队列排序等一整套流程。这些开销在GPU计算时间较长时可以忽略，但当单个kernel只有几微秒甚至更短时，CPU提交命令的成本就会反过来主导整个程序的执行时间。CUDA Graph正是针对这一问题设计的：它把多次kernel启动捕获成一张静态执行图，之后无论图中有多少个节点，CPU端只需要发起一次launch动作。
为什么kernel launch开销会成为性能瓶颈？
传统CUDA程序使用的是流式执行模型。代码中每写一次kernel<<<grid, block, 0, stream>>>，主机端就要调用一次CUDA运行时API。这个调用会穿越多层软件栈：运行时层会做参数校验并转换成本地描述符，驱动层会把这些描述符写入命令缓冲区，然后命令缓冲区会被提交给GPU硬件队列。对于一个大kernel而言，这个过程可能只需要几微秒，完全可以被GPU执行时间掩盖。但如果kernel本身只有1到5微秒的执行时间，那么启动路径上的开销就可能占全程耗时的50%甚至80%以上。
问题更严重的地方在于，这种开销并不会随着GPU性能提升而自动消失。GPU的核心数量增加后，单个轻量kernel的执行时间反而可能进一步缩短，使得CPU提交命令的时间占比变得更高。很多涉及逐层处理、逐元素变换、小型规约或者细粒度算子融合的程序，都会在连续启动大量小kernel时表现出明显的CPU-bound特征。此时即使GPU占用率很低，整体帧率或吞吐量也上不去。
要判断自己的程序是否受kernel launch开销影响，一个简单的方法是使用CUDA profiler观察CPU端的API调用耗时与GPU端执行耗时的比例。如果CPU时间线上一长串小块的提交动作连成密集区域，而GPU侧每个kernel都很短，那基本上就是启动开销在作怪。CUDA Graph正是为这类重复、固定、小粒度kernel序列提供了一个批量化提交方案。
CUDA Graph捕获的内部机制
CUDA Graph的核心思路是把原本分散的、逐个提交的操作记录成一张有向无环图。图中的节点代表kernel、内存拷贝、事件等操作，边代表它们之间的依赖关系。捕获过程发生在某个CUDA stream上：调用cudaStreamBeginCapture之后，该stream上随后发出的所有操作都不会立即执行，而是被拦截下来写入图中。当调用cudaStreamEndCapture时，捕获结束并返回一个cudaGraph_t句柄。
捕获模式有两种：cudaStreamCaptureModeThreadLocal和cudaStreamCaptureModeRelaxed。前一种只允许捕获当前线程在该stream上发出的操作，后一种则允许多个线程向同一个stream提交操作，并在捕获中进行合并。大多数应用使用默认的线程局部模式已经足够。捕获期间可以正常调用kernel启动、cudaMemcpyAsync、事件记录等异步API，但同步操作例如cudaStreamSynchronize或cudaDeviceSynchronize是不被允许的，它们会导致捕获失败。
捕获完成后得到的是一个未实例化的图结构。图中的节点数量、kernel参数、内存依赖关系都在这个阶段确定下来。接下来需要调用cudaGraphInstantiate将图实例化成可执行对象cudaGraphExec_t。实例化的过程会做额外的优化，例如合并相邻的轻量节点、检查资源分配、生成显卡可以直接消费的命令包。实例化之后的图可以反复launch，每次launch只消耗一次CPU提交成本。
从原理上看，CUDA Graph相当于把多次kernel launch的CPU侧工作一次性预处理，然后在GPU端形成一个紧凑的执行计划。当图中包含100个小kernel时，传统方式需要CPU发送100次命令，而CUDA Graph只需要一次cudaGraphLaunch调用。这对于CPU线程负担很重的应用来说，节省效果非常明显。
代码实战：捕获、实例化与launch
下面用一个简单示例串起CUDA Graph的基本流程。代码中先创建stream和内存，然后在该stream上开始捕获，执行两个向量加法kernel，接着结束捕获并实例化，最后循环launch同一个图1000次。
#include <cuda_runtime.h>
#include <stdio.h>
__global__ void addKernel(float *a, float *b, float *c, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
c[i] = a[i] + b[i];
}
}
int main() {
const int N = 1024;
float *a, *b, *c;
cudaMallocManaged(&a, N * sizeof(float));
cudaMallocManaged(&b, N * sizeof(float));
cudaMallocManaged(&c, N * sizeof(float));
for (int i = 0; i < N; ++i) {
a[i] = 1.0f;
b[i] = 2.0f;
}
cudaStream_t stream;
cudaStreamCreate(&stream);
// 开始捕获
cudaStreamBeginCapture(stream, cudaStreamCaptureModeThreadLocal);
addKernel<<<N / 256, 256, 0, stream>>>(a, b, c, N);
addKernel<<<N / 256, 256, 0, stream>>>(c, a, b, N);
cudaGraph_t graph;
cudaStreamEndCapture(stream, &graph);
// 实例化图
cudaGraphExec_t instance;
cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);
// 重复launch 1000次
for (int i = 0; i < 1000; ++i) {
cudaGraphLaunch(instance, stream);
}
cudaStreamSynchronize(stream);
cudaGraphExecDestroy(instance);
cudaGraphDestroy(graph);
cudaStreamDestroy(stream);
cudaFree(a);
cudaFree(b);
cudaFree(c);
return 0;
}
这段代码展示了CUDA Graph最核心的四个步骤：cudaStreamBeginCapture开始捕获，cudaStreamEndCapture结束并得到cudaGraph_t，cudaGraphInstantiate生成可执行实例，最后用cudaGraphLaunch提交整个图。其中cudaGraphInstantiate的参数可以传入cudaGraphInstantiateFlagDeviceLaunch等标志，以便在某些平台上允许从GPU device端发起图launch，这属于进阶用法。
需要注意的是，捕获期间对内存的访问模式必须保持稳定。图实例化时记录的指针地址在后续launch中仍然有效，如果这些指针指向的内存被释放或重新分配，launch时会引发未定义行为。对于参数固定的重复计算场景，这个问题不大；但如果每轮计算需要更新数据，可以使用cudaGraphExecKernelNodeSetParams更新kernel节点的参数，然后再重新实例化部分子图。
适用场景与避坑指南
CUDA Graph并不是万能的，它最适合那些由大量短小kernel组成的、依赖关系固定且需要反复执行的序列。典型的例子包括深度学习推理中的算子融合、图像处理中的多级滤波、科学计算中的迭代求解以及通信集合操作后的连续计算。如果整个序列只执行一次，图上预处理的成本可能超过直接流式launch的成本，收益微乎其微。如果序列本身只有一两个大kernel，启动开销占比很小，也没有必要引入图结构。
使用CUDA Graph时有一个常见的误区：在捕获过程中使用同步API。很多开发者习惯在创建buffer后调用cudaDeviceSynchronize确认内存可用，但这种同步调用在stream捕获期间是不允许的。遇到这种情况，应该把同步操作移到捕获开始之前或者结束之后。另一个坑是关于多stream依赖：默认捕获只针对单个stream，如果图中需要跨多个stream执行，必须使用事件节点来显式描述跨stream依赖，否则图会丢失部分并行性或者产生错误顺序。
内存管理也要格外小心。使用统一内存时，捕获前需要确保数据已经预热到GPU侧，否则第一次launch可能触发页错误和迁移，导致图执行时间异常增长。此外，图的销毁顺序必须是先销毁cudaGraphExec_t实例，再销毁cudaGraph_t图句柄，最后销毁stream，顺序错乱可能造成资源泄漏或崩溃。
如果图中的kernel参数需要动态调整，不建议每次重新捕获整张图，因为捕获和实例化的代价相对较高。更合理的做法是使用cudaGraphExecKernelNodeSetParams更新指定kernel节点的参数，然后调用cudaGraphExecUpdate让实例与修改后的图保持一致。这样可以保留大部分图的优化结果，同时支持有限的动态性。
最后，CUDA Graph还能与CUDA事件、stream优先级组合使用。在复杂的异步流水线中，将整张图launch到一个高优先级stream上，可以让GPU调度器更高效地分配执行单元。实际项目里可以先对传统流式版本和图版本做基准对比，通常在小kernel密集、CPU线程压力大的情况下，CUDA Graph能带来2到5倍的提交性能提升。
