附录 A · CUDA Graph 与 Stream

"If you're not using CUDA Graphs in 2026, you're leaving 20-30% latency on the table for free." ——LLM 推理工程的常见判断

A.1 Stream 与 Event：异步执行的基础

CUDA 的 host/device 异步模型基于两个原语：

Stream（流）：一个 FIFO 队列，里面装的是 GPU 任务（kernel launch、memcpy）。同一 stream 内的任务严格按顺序执行；不同 stream 之间默认并行（除非用 event 同步）。

Event（事件）：一个时间戳/同步点。可以记录到某个 stream 上，也可以让某个 stream 等待某个 event。

cudaStream_t s1, s2;
cudaStreamCreate(&s1);
cudaStreamCreate(&s2);

cudaEvent_t event;
cudaEventCreate(&event);

// stream 1: 跑 kernel A
kernel_a<<<..., 0, s1>>>(...);
cudaEventRecord(event, s1);  // 在 s1 上记录 event

// stream 2: 等 event, 然后跑 kernel B
cudaStreamWaitEvent(s2, event, 0);
kernel_b<<<..., 0, s2>>>(...);  // 等 kernel A 完成才会跑

这套机制让程序员可以构造任意复杂的 DAG（有向无环图）执行计划。

A.1.1 默认 Stream 的坑

CUDA 有一个"默认 stream"（也叫 NULL stream），所有不指定 stream 的调用都跑在它上。默认 stream 与所有用户 stream 互斥——意味着默认 stream 上的任务会阻塞其他 stream。

// 反例: 误用默认 stream
kernel_a<<<...>>>(...);                        // 默认 stream
cudaMemcpyAsync(..., user_stream);              // 想跟 kernel_a 并行?
// 实际上: cudaMemcpyAsync 必须等默认 stream 上的 kernel_a 完成

修复：所有调用都指定 stream，避免默认 stream。或者用 cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking) 创建"non-blocking"stream，它不会被默认 stream 阻塞。

A.2 为什么需要 CUDA Graph

LLM 推理的一次 forward 涉及几百次 kernel launch（attention、GEMM、LayerNorm 等）。每次 launch 在 host 端都有几μs 开销，包括：

• CPU 调度
• driver 级 dispatch
• GPU 接收 launch 指令
• 寄存器初始化

对小规模推理（batch_size=1, decoding 阶段），这些开销加起来可能占总延迟的 30%+。

单 token decoding 时间分解 (LLaMA-7B, H100):
  GPU 计算:           ~6 ms
  Kernel launch overhead: ~3 ms  ← 30%!
  其他:                 ~1 ms
  总:                  ~10 ms

CUDA Graph 解决这个问题：把一系列 kernel launch 录制成一个 Graph，每次只 launch 这个 Graph（一条命令），省去重复的 dispatch 开销。

效果：3ms 开销可以降到 0.5ms。整体延迟从 10ms 降到 ~7ms，30% 提升。

A.3 CUDA Graph 用法

A.3.1 录制方式

最常用的方式是 stream capture：

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

// 2. 跑一遍 forward (所有 kernel 都在这个 stream 上)
forward(input, output, stream);
// 内部:
//   layer_norm<<<..., 0, stream>>>(...);
//   gemm<<<..., 0, stream>>>(...);
//   attention<<<..., 0, stream>>>(...);
//   ...

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

// 4. 实例化 graph (生成可执行版本)
cudaGraphExec_t graphExec;
cudaGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0);

// 后续推理: 只 launch 这个 graph
for (int step = 0; step < N; ++step) {
    cudaGraphLaunch(graphExec, stream);
}

CUDA Graph 把整个 forward 当作一个原子任务提交。

A.3.2 Graph 的限制

Graph 不是万能的。它有几个关键限制：

1. 形状固定：录制时所有 kernel 的 launch 参数（grid_size、block_size）固定。如果 batch_size 或 seq_len 变化，graph 失效。
2. 指针固定：录制时的输入/输出指针被烧到 graph 里。每次 launch 必须用同样的指针（或用 cudaGraphExecKernelNodeSetParams 更新）。
3. 不能有 host 同步：graph 内的 kernel 之间不能有 cudaDeviceSynchronize。
4. 首次 launch 开销大：实例化和首次 launch 比单独 launch 慢，但后续 launch 极快。

A.3.3 Graph 与 LLM 推理

LLM 推理对 graph 来说是完美场景——decoding 阶段每个 token 的 forward 形状完全一样（都是 batch_size×1×hidden_size）。所以可以一次录制，反复 launch。

vLLM、TensorRT-LLM、SGLang 都默认启用 graph：

# vLLM 配置
llm = LLM(model="...", enforce_eager=False)  # enforce_eager=False 启用 graph

但 prefill 阶段每次 prompt 长度可能不同——这里 graph 不能用，必须 eager 模式。

实际推理中常用 graph pool 策略：预先录制几个不同 batch_size 的 graph，运行时按 input 选最合适的 graph。

A.3.4 更新 Graph 输入

如果不希望每次都重新录制，可以用 cudaGraphExecKernelNodeSetParams 在 graph 实例上修改某个 kernel 的参数：

// 取得 graph 中的某个 kernel node
cudaGraphNode_t kernel_node = ...;

// 准备新参数
cudaKernelNodeParams params = {...};
params.func = (void*)my_kernel;
params.gridDim = new_grid;
params.blockDim = new_block;
params.kernelParams = ...;

// 更新
cudaGraphExecKernelNodeSetParams(graphExec, kernel_node, &params);

这种 in-place 更新比重新实例化快得多。vLLM 用它支持动态 batch_size。

A.4 Stream 与多 GPU

多 GPU 场景下，stream 是基本协调单位：

cudaSetDevice(0);
cudaStream_t s_gpu0;
cudaStreamCreate(&s_gpu0);

cudaSetDevice(1);
cudaStream_t s_gpu1;
cudaStreamCreate(&s_gpu1);

// GPU 0 跑 attention
cudaSetDevice(0);
attention<<<..., s_gpu0>>>(...);

// GPU 1 跑 FFN
cudaSetDevice(1);
ffn<<<..., s_gpu1>>>(...);

// 跨 GPU 同步用 event
cudaEvent_t e0;
cudaEventCreate(&e0);
cudaSetDevice(0);
cudaEventRecord(e0, s_gpu0);
cudaSetDevice(1);
cudaStreamWaitEvent(s_gpu1, e0, 0);

NCCL 的 collective 操作（AllReduce 等）也都在 stream 上。

A.5 这个附录的小结

CUDA 异步执行模型是 LLM 推理优化的隐藏维度：

1. Stream + Event 让 host/device 异步并行。
2. CUDA Graph 把 launch 开销摊薄到几乎 0。
3. LLM 推理 decoding 阶段 是 graph 的完美场景，可以省 20-30% 延迟。
4. vLLM/TensorRT-LLM 默认启用 graph——如果你部署 LLM 推理服务，这是必须开的。