4. CUDA图捕获机制：使用Stream Capture API捕获操作序列、图实例化与更新、图内存管理

核心原理：从流到图的“录制-回放”范式

CUDA图（CUDA Graph）的本质是将GPU上的一系列kernel启动、内存拷贝、事件同步等操作，从传统的“逐次提交到流”模式，转换为“一次性构建依赖图，然后整体提交执行”的模式。Stream Capture API（流捕获API）正是实现这一转换的桥梁。它允许你在一个特殊的捕获流（capture stream）上，像平时写CUDA代码一样启动kernel和cudaMemcpy，但此时这些操作并不会真正执行，而是被CUDA运行时“录制”为图中的节点，节点间的依赖关系则根据流上的顺序和事件同步自动推导。

捕获完成后，调用cudaGraphInstantiate()将捕获到的图实例化为可执行的图执行句柄（cudaGraphExec_t）。实例化过程会完成kernel参数的固定、内存分配的优化、依赖关系的编译等准备工作。之后，每次需要执行这组操作时，只需调用cudaGraphLaunch()，即可将整个图提交到指定流上执行，从而大幅降低CPU启动开销，并让GPU驱动有机会对图内操作进行全局调度优化。

方法流程：Stream Capture API的完整使用步骤

使用Stream Capture API捕获CUDA图，需要遵循一套严格的流程。首先，必须选择一个非默认流（即不是cudaStreamLegacy或cudaStreamPerThread）作为捕获流，并调用cudaStreamBeginCapture()进入捕获模式。从这一刻起，所有在该流上提交的操作（包括kernel启动、cudaMemcpyAsync、事件记录等）都不会真正执行，而是被记录为图的候选节点。

// 步骤1：创建捕获流
cudaStream_t captureStream;
cudaStreamCreateWithFlags(&captureStream, cudaStreamNonBlocking);

// 步骤2：开始捕获
cudaStreamBeginCapture(captureStream, cudaStreamCaptureModeGlobal);

// 步骤3：在捕获流上提交操作（不会真正执行）
myKernelA<<<grid, block, 0, captureStream>>>(d_data, N);
cudaMemcpyAsync(d_temp, d_data, bytes, cudaMemcpyDeviceToDevice, captureStream);
myKernelB<<<grid, block, 0, captureStream>>>(d_temp, N);

// 步骤4：结束捕获，得到图对象
cudaGraph_t graph;
cudaStreamEndCapture(captureStream, &graph);

// 步骤5：实例化图
cudaGraphExec_t graphExec;
cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0);

// 步骤6：启动图执行
cudaGraphLaunch(graphExec, stream);

// 步骤7：清理资源
cudaGraphExecDestroy(graphExec);
cudaGraphDestroy(graph);
cudaStreamDestroy(captureStream);

捕获模式有两种：cudaStreamCaptureModeGlobal和cudaStreamCaptureModeThreadLocal。全局模式要求所有线程的流操作都被捕获，适合单线程控制流；线程局部模式只捕获当前线程的流操作，适合多线程场景。在捕获期间，如果其他线程向捕获流提交了操作，会导致捕获失败，因此需要特别注意线程同步。

图实例化与更新：静态结构与动态参数

图实例化（cudaGraphInstantiate）是图从“逻辑描述”变为“可执行实体”的关键步骤。实例化过程中，CUDA驱动会分析图中所有节点的依赖关系，为每个节点分配固定的GPU资源（如kernel参数缓冲区、内存地址等），并生成最优的执行计划。实例化后的cudaGraphExec_t对象可以被多次启动，但默认情况下，其内部参数是固定的。

然而，实际应用中往往需要动态调整参数（如每次推理的输入数据指针、batch size等）。CUDA提供了两种更新机制：整体更新和节点更新。整体更新是重新捕获并重新实例化整个图，开销较大；节点更新则通过cudaGraphExecKernelNodeSetParams()或cudaGraphExecMemcpyNodeSetParams()等API，只修改特定节点的参数，然后再次启动图。这种方式避免了重新实例化的开销，非常适合动态形状场景。

更新方式	API	开销	适用场景
整体重新实例化	cudaGraphDestroy + cudaGraphInstantiate	高（毫秒级）	图结构变化频繁
节点参数更新	cudaGraphExecKernelNodeSetParams	低（微秒级）	仅参数变化，结构不变
节点替换	cudaGraphExecNodeReplace	中	替换单个节点为另一个子图

图内存管理：捕获期间的分配策略与注意事项

CUDA图的内存管理是捕获过程中最容易出错的部分。在捕获期间，如果代码中调用了cudaMalloc、cudaFree等同步内存分配函数，会导致捕获失败，因为这些操作无法被记录为图中的节点。正确的做法是使用图感知的内存分配器（如cudaMallocAsync）或预先分配好所有内存，在捕获期间只使用已有的内存指针。

// 错误示例：捕获期间调用cudaMalloc
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
float *d_ptr;
cudaMalloc(&d_ptr, size);  // 错误！会导致捕获失败
kernel<<<grid, block, 0, stream>>>(d_ptr);
cudaStreamEndCapture(stream, &graph);

// 正确示例：使用cudaMallocAsync（图感知）
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
float *d_ptr;
cudaMallocAsync(&d_ptr, size, stream);  // 正确！会记录为图节点
kernel<<<grid, block, 0, stream>>>(d_ptr);
cudaStreamEndCapture(stream, &graph);

另一个关键点是内存池的使用。CUDA 11.2+支持cudaDeviceSetMemPool()为流绑定内存池，捕获期间的内存分配会从池中获取，并在图执行时复用。这避免了每次图启动时的内存分配开销。对于动态形状，可以创建多个不同大小的内存池，根据实际输入尺寸选择对应的池，从而减少内存浪费。

案例：动态batch的图捕获与更新

假设我们有一个推理模型，输入batch size在1~8之间动态变化。使用CUDA图加速时，可以按以下策略实现：

1. 预先为最大batch size（8）创建图，捕获所有kernel和内存操作。
2. 实例化图，得到graphExec。
3. 每次推理时，根据实际batch size，调用cudaGraphExecKernelNodeSetParams()更新每个kernel的grid维度（block数）和输入指针偏移。
4. 调用cudaGraphLaunch()执行。

// 更新kernel节点参数的示例
cudaGraphNode_t node;
cudaGraphExecKernelNodeSetParams(graphExec, node, &newParams);
// newParams中设置新的gridDim、blockDim和kernel参数指针

// 启动更新后的图
cudaGraphLaunch(graphExec, stream);

这种方式的优势在于，图的结构（节点间的依赖关系）完全不变，只是每个节点的执行参数（如grid大小）发生变化。实例化后的图执行计划仍然有效，因此更新开销极小。实测表明，对于包含10个kernel的图，节点参数更新的开销约为2~3微秒，而重新实例化整个图则需要50~100微秒。

误区与复盘：捕获失败的常见原因

• 误区1：在捕获流上使用同步API。 cudaMemcpy、cudaMalloc、cudaDeviceSynchronize等同步操作会破坏捕获，必须全部替换为异步版本。
• 误区2：忽略捕获模式的选择。 多线程环境下使用cudaStreamCaptureModeGlobal可能导致其他线程的流操作被意外捕获，应使用ThreadLocal模式。
• 误区3：图更新时修改了节点类型。 cudaGraphExecKernelNodeSetParams只能更新kernel节点的参数，不能改变节点类型（如将kernel节点改为memcpy节点）。
• 误区4：内存生命周期管理不当。 图执行时引用的内存必须在图启动期间保持有效。如果使用cudaMallocAsync，需要确保内存池的生命周期覆盖图的执行周期。
• 误区5：忽略错误检查。 捕获失败时，CUDA可能不会立即报错，而是在cudaStreamEndCapture时返回错误。务必检查所有API的返回值。