第3章 CUDA图创建方式

3、CUDA图创建方式：流捕获(Stream Capture)与手动构建(Manual API)

好，咱们今天来聊聊CUDA图的两种创建方式。说实话，我刚接触CUDA图那会儿，也觉得这东西有点玄乎。不就是把一堆kernel调用串起来嘛，能有多大区别？后来在项目里踩过坑，才真正体会到——选对创建方式，性能差距能到30%以上。

CUDA图本质上是一个计算任务的依赖关系图。你可以把它想象成一个流水线作业的蓝图。两种创建方式，说白了就是：一种是你手把手教CUDA怎么搭这个流水线（手动构建），另一种是让CUDA自己看你干活然后偷偷记下来（流捕获）。

3.1 流捕获（Stream Capture）——让CUDA偷师学艺

流捕获是我个人最常用的方式。为什么？因为它省事。你想想看，你本来就有现成的CUDA代码，各种kernel launch、cudaMemcpyAsync写得飞起。现在要改成CUDA图，难道要全部重写？不存在的。

流捕获的核心思想很简单：你正常写代码，CUDA在背后偷偷记录。就像你教徒弟做菜，你一边炒菜一边说“看好了啊”，徒弟在旁边记笔记。等记录完了，直接回放就行。

核心流程：

调用 cudaStreamBeginCapture(stream) 开始捕获
正常执行所有CUDA操作（kernel、内存拷贝等）
调用 cudaStreamEndCapture(stream, &graph) 结束捕获，得到CUDA图
用 cudaGraphInstantiate(&execGraph, graph, ...) 实例化
用 cudaGraphLaunch(execGraph, stream) 执行

来看个实际例子。这是我之前做图像处理项目时的一段代码，简化版：

cudaStream_t stream;
cudaStreamCreate(&stream);

// 开始捕获
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);

// 这些操作会被记录到图中
kernel_preprocess<<<grid, block, 0, stream>>>(d_input, d_temp);
cudaMemcpyAsync(d_temp2, d_temp, size, cudaMemcpyDeviceToDevice, stream);
kernel_filter<<<grid, block, 0, stream>>>(d_temp2, d_output);
kernel_postprocess<<<grid, block, 0, stream>>>(d_output, d_result);

// 结束捕获
cudaGraph_t graph;
cudaStreamEndCapture(stream, &graph);

// 实例化并执行
cudaGraphExec_t execGraph;
cudaGraphInstantiate(&execGraph, graph, NULL, NULL, 0);
cudaGraphLaunch(execGraph, stream);
cudaStreamSynchronize(stream);

嗯，这里要注意一点：流捕获模式下，你不能在捕获区域内调用任何会阻塞CPU的操作，比如 cudaDeviceSynchronize。我曾经犯过这个错，debug了一下午才发现是捕获期间调了同步函数，直接报错。

避坑指南：

我曾经在流捕获里用了 cudaMalloc，结果CUDA直接罢工。记住，捕获期间只能做异步操作。内存分配、同步操作都得放到捕获外面去。

3.2 手动构建（Manual API）——自己当建筑师

手动构建就完全是另一回事了。你得自己定义每个节点，自己搭边。听起来麻烦，但好处是——你拥有绝对的控制权。

我一般在两种情况下用手动构建：一是图结构需要动态变化，二是要对每个节点做精细控制（比如设置不同的共享内存大小）。

手动构建的核心API就几个：

API	作用
`cudaGraphCreate`	创建空图
`cudaGraphAddKernelNode`	添加kernel节点
`cudaGraphAddMemcpyNode`	添加内存拷贝节点
`cudaGraphAddEmptyNode`	添加空节点（用于依赖控制）
`cudaGraphAddDependencies`	添加节点间的依赖关系

来看个手动构建的例子。假设我们有三个kernel，A -> B -> C，但B和C之间还要等一个内存拷贝完成：

cudaGraph_t graph;
cudaGraphCreate(&graph, 0);

// 创建节点参数
cudaKernelNodeParams paramsA = {0};
paramsA.func = (void*)kernelA;
paramsA.gridDim = grid;
paramsA.blockDim = block;
paramsA.kernelParams = (void**)kernelArgsA;

// 添加节点
cudaGraphNode_t nodeA, nodeB, nodeC, nodeMemcpy;
cudaGraphAddKernelNode(&nodeA, graph, NULL, 0, ¶msA);
cudaGraphAddKernelNode(&nodeB, graph, NULL, 0, ¶msB);
cudaGraphAddMemcpyNode(&nodeMemcpy, graph, NULL, 0, &memcpyParams);
cudaGraphAddKernelNode(&nodeC, graph, NULL, 0, ¶msC);

// 建立依赖：A->B, A->Memcpy, Memcpy->C, B->C
cudaGraphNode_t dependencies[] = {nodeA};
cudaGraphAddDependencies(graph, &nodeB, dependencies, 1);

dependencies[0] = nodeA;
cudaGraphAddDependencies(graph, &nodeMemcpy, dependencies, 1);

cudaGraphNode_t depsForC[] = {nodeB, nodeMemcpy};
cudaGraphAddDependencies(graph, &nodeC, depsForC, 2);

你看，手动构建的代码量明显多了。但好处是——你能精确控制每个节点的依赖关系。流捕获做不到这么细，它只能按执行顺序来。

我的经验：

如果你刚开始用CUDA图，我建议先从流捕获入手。等你对图的行为摸透了，再考虑手动构建。别一上来就手撸依赖关系，容易把自己绕晕。

3.3 两种方式的对比与选择

说白了，选哪种方式取决于你的场景。我整理了个对比表，你一看就明白：

维度	流捕获	手动构建
代码改动量	小，几乎不改现有代码	大，需要重写
控制粒度	粗，按执行顺序	细，可任意指定依赖
动态图支持	差，每次改结构要重新捕获	好，可以动态增删节点
调试难度	低，代码逻辑不变	高，依赖关系容易出错
性能上限	高，CUDA自动优化	更高，可手动调优

我个人习惯是：80%的场景用流捕获，剩下20%用手动构建。比如做推理服务，图结构基本固定，流捕获就够用了。但做那种计算图会动态变化的科学计算，就得手动来。

3.4 混合使用——两全其美

其实你还可以把两种方式混着用。比如先用流捕获生成一个基础图，然后用手动API去修改它。CUDA是允许这么干的。

我记得有个项目，需要在一个固定流水线里动态插入一个校验节点。我的做法是：

先用流捕获生成主流水线图
用 cudaGraphGetNodes 获取所有节点
用 cudaGraphAddKernelNode 插入新节点
用 cudaGraphAddDependencies 调整依赖

这样既保留了流捕获的便捷性，又获得了手动构建的灵活性。说白了，工具是死的，人是活的。

总结一下：

流捕获：适合快速上手、图结构固定的场景
手动构建：适合需要精细控制、图结构动态变化的场景
混合使用：取两者之长，灵活应对复杂需求