4、流捕获入门：cudaStreamBeginCapture、cudaStreamEndCapture、捕获模式详解

好，咱们今天聊聊流捕获（Stream Capture）。

这东西说白了，就是CUDA提供的一个“录屏”功能。你想想看，平时我们写CUDA代码，都是手动把一个个kernel launch、cudaMemcpy调用排好队，扔进流里。但有时候，我们希望把一段已有的、非CUDA Graph的代码，自动“翻译”成一个CUDA Graph。流捕获就是干这个的。

我个人习惯，在需要做动态图优化或者反复执行相同计算模式时，一定会用流捕获。它比手写Graph API要省心得多，尤其是当你的计算流程比较复杂的时候。

4.1 核心API：开始与结束

流捕获的入口和出口，就是两个函数：cudaStreamBeginCapture 和 cudaStreamEndCapture。

cudaStreamBeginCapture

这个函数告诉CUDA驱动：“嘿，从现在开始，往这个流里丢的所有操作，你都给我记下来，别真执行。” 它会返回一个 cudaError_t，一般我们检查一下就行。

cudaStream_t stream;
cudaStreamCreate(&stream);
cudaError_t err = cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
// 检查err...

cudaStreamEndCapture

这个函数告诉CUDA：“好了，录完了，把刚才录下来的东西打包成一个Graph给我。” 它会返回一个 cudaGraph_t 对象。

cudaGraph_t graph;
err = cudaStreamEndCapture(stream, &graph);
// 检查err...
// 现在graph里就包含了从Begin到End之间所有的操作

嗯，这里要注意：Begin和End必须在同一个流上调用。而且，在Begin之后、End之前，你不能对这个流做任何“非捕获”的操作，比如 cudaStreamSynchronize 或者 cudaStreamQuery。我曾经犯过这个错，在捕获中间加了个同步，结果CUDA直接报错，说流状态不对。

4.2 捕获模式详解

你可能会问：“录屏的时候，能不能录到其他流里的操作？” 这就是捕获模式要解决的问题了。

cudaStreamBeginCapture 的第二个参数，就是捕获模式。目前有三种：

模式	枚举值	行为描述
全局模式	cudaStreamCaptureModeGlobal	捕获所有流上的操作。只要有一个流在捕获，其他流上的操作也会被捕获进去。
线程局部模式	cudaStreamCaptureModeThreadLocal	只捕获当前线程发起的、且属于被捕获流的操作。其他线程的操作不影响。
中继模式	cudaStreamCaptureModeRelaxed	最宽松的模式。允许在捕获期间，其他流上有未被捕获的操作。但要注意，这可能导致Graph依赖关系不完整。

全局模式：我最常用。它最严格，也最安全。你想想看，如果你的计算图里涉及多个流（比如一个流做计算，一个流做数据传输），用全局模式能保证所有依赖关系都被正确捕获。我在项目中遇到过，用线程局部模式时，因为另一个线程偷偷往别的流里丢了个操作，结果Graph执行时数据还没到位，直接崩了。后来换成全局模式，问题就解决了。

线程局部模式：适合单线程、单流的简单场景。如果你确定你的计算只在一个线程里、一个流上，用这个模式可以减少一些不必要的检查开销。

中继模式：这个模式我建议新手慎用。它允许你在捕获期间，其他流上可以自由操作。但代价是，这些自由操作不会被记录到Graph里。如果你后续的Graph执行依赖了这些操作的结果，那就出大问题了。说白了，这是给高级用户准备的，他们能精确控制哪些操作进图，哪些不进。

4.3 一个完整的捕获示例

光说不练假把式。咱们看个完整的例子，把向量加法用流捕获做成Graph。

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void vecAdd(float *a, float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int n = 1024;
    size_t size = n * sizeof(float);
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);

    // 创建流
    cudaStream_t stream;
    cudaStreamCreate(&stream);

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

    // 这些操作会被捕获
    vecAdd<<<1, 256, 0, stream>>>(d_a, d_b, d_c, n);
    // 可以加更多kernel或者cudaMemcpyAsync

    // 结束捕获，得到Graph
    cudaGraph_t graph;
    cudaStreamEndCapture(stream, &graph);

    // 实例化Graph（编译成可执行形式）
    cudaGraphExec_t graphExec;
    cudaGraphInstantiate(&graphExec, graph, NULL, NULL, 0);

    // 执行Graph
    cudaGraphLaunch(graphExec, stream);
    cudaStreamSynchronize(stream);

    // 清理
    cudaGraphExecDestroy(graphExec);
    cudaGraphDestroy(graph);
    cudaStreamDestroy(stream);
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);

    return 0;
}

你看，代码结构很清晰。Begin、丢操作、End、Instantiate、Launch。就这么几步。

4.4 避坑指南

我曾经在一个项目里，为了图方便，在捕获中间调用了 cudaMalloc 来分配临时缓冲区。结果CUDA直接返回 cudaErrorStreamCaptureImplicit 错误。后来我改成在捕获前一次性分配好所有内存，问题就解决了。

另外，捕获期间不要切换CUDA上下文。比如不要调用 cudaSetDevice。这会导致捕获状态混乱。

4.5 什么时候用流捕获？

我个人觉得，流捕获最适合以下场景：

• 你的计算流程是固定的，但每次执行时参数（比如数据指针）会变。
• 你有一个复杂的、多kernel的计算链，手动构建Graph太麻烦。
• 你想把现有的、基于流的代码快速迁移到Graph上，获得性能提升。

说白了，流捕获就是CUDA给咱们提供的一条“捷径”。它让你不用手写Graph API那些繁琐的节点和边，而是通过“录制”的方式自动生成。对于大多数场景，这已经足够了。

好，关于流捕获的入门，咱们就聊到这里。下一节，我会讲讲如何对捕获到的Graph进行修改和优化，让它更贴合你的实际需求。