CUDA Graph 是 NVIDIA CUDA Runtime 提供的一种机制，用来把一系列 GPU 操作（kernel launch / memcpy / stream sync / event / …）一次性“捕获”成一个有向无环图（DAG），之后可以像一个整体一样“实例化并重复执行”。

  • 拓扑预定义：将计算流程编码为显式 DAG，一次性提交整个计算图
  • 零成本启动，GPU 直接执行预编译的指令序列，消除 99% 的启动开销
  • 但是强依赖于静态计算图

核心价值在于：减少调度开销，避免每次都走 CPU -> GPU 的 syscall/调度路径（kernel launch、memcpy 等），特别在循环执行同一序列时收益明显；提高可重复性，多个 kernel/拷贝之间的依赖由图管理，避免手工同步；潜在优化，CUDA 运行时可以对图做优化、融合甚至跨 kernel scheduling。

bubble: 空泡，空闲时间

CUDA Graph

创建一个 CUDA Graph

1
2
cudaGraph_t graph;
cudaGraphCreate(&graph, 0);

接下来，CUDA Graph 需要获取计算图。我们可以选择手动添加 DAG

1
2
3
4
5
6
7
8
9
// 添加计算内核
cudaGraphNode_t kernelNode;
cudaKernelNodeParams kernelParams = { /* ... */ };
cudaGraphAddKernelNode(...);

// 添加内存拷贝内核
cudaGraphNode_t memcpyNode;
cudaMemcpy3DParams memcpyParams = { /* ... */ };
cudaGraphAddMemcpyNode(...);

我们也可以选择让 CUDA 自行根据运行的 kernel 进行捕获：

1
2
3
4
5
6
7
8
cudaStream_t stream;
cudaStreamCreate(&stream);
cudaGraphBeginCapture(stream, cudaStreamCaptureModeGlobal); // 开始捕获，接下来进行正常的 cuda kernel launch

myKernel<<..., stream>>(...);
cudaMemcpyAsync(..., stream);

cudaGraphEndCapture(stream, &graph); // 结束捕获，生成计算图

然后实例化可执行计算图

1
2
cudaGraphExec_t graphExec;
cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0);

最后执行

1
2
cudaGraphLaunch(graphExec, stream); // 提交执行
cudaStreamSynchronize(stream); // 可选：同步等待完成

一般来说，CUDA Graph 相关的流程都是用 flag 标志是否实例化了 CUDA Graph，否的话就首先进行实例化，是的话就直接执行实例化的 CUDA Graph．

CUDA Graph