1. CUDA Graph 是什么？

CUDA Graph 是 NVIDIA CUDA 中的一个重要功能，用于优化 GPU 任务的调度和执行效率。它的核心思想是将一系列 GPU 操作（如内核启动、内存拷贝等）预先记录成一个静态的“执行图”，然后一次性提交整个图到 GPU 执行，从而减少 CPU 与 GPU 之间的交互开销（降低 API 调用和同步的开销），提升性能（尤其适合重复执行的固定模式任务）。

构成CUDA图的节点是任意异步CUDA操作，可被 graphs 的的操作（graphs 节点）：

• Kernel Launch：CUDA kernel running on GPU（GPU CUDA 操作）
• CPU Function Call：Callback function on CPU（CPU 操作）
• Memcopy / Memset：GPU data management（GPU 数据传输）
• Memory Alloc / Free：Inline memory allocation（内存分配）
• Sub-Graph：Graphs are hierarchical（子图 Graphs）



2. CUDA Graph 发展历程

2017 年：CUDA 9.0 引入基础 CUDA Graphs

• 核心特性 ：首次提出 “计算图” 概念，允许将一系列 CUDA 操作（核函数、内存复制、同步等）定义为一个静态图结构
• 功能局限：
  • 仅支持完全静态的图（结构和参数在构建后不可修改）
  • 图的执行依赖主机端启动，且不支持动态分支或条件执行
• 价值：通过 “一次构建、多次启动” 减少核函数启动的 CPU 开销（避免重复的参数传递和调度指令），适用于固定流程的重复计算（如深度学习推理）



2018 年：CUDA 10.0 增强图的可修改性

• 核心改进：允许对已创建的图进行有限修改，如更新核函数参数、替换部分节点（cudaGraphUpdateNodeParams 等 API）
• 局限：
  • 修改后需重新实例化（cudaGraphInstantiate），开销较大
  • 仍不支持动态分支，图的整体结构需固定
• 价值：一定程度上提升了静态图的灵活性，适合参数需动态调整但流程固定的场景（如迭代计算中更新输入数据）



2020 年：CUDA 11.0 引入 Graph Segment Table

• 核心突破：支持将大图拆分为多个独立 “图段（Segment）”，通过表格管理段间跳转关系，实现设备端动态分支
• 关键能力：
  • 运行时可根据条件在设备端直接切换图段，无需主机端介入
  • 图段修改后无需重新构建整个图，仅需更新对应段
• 价值：首次解决了静态图的动态控制流问题，兼顾低启动开销与分支灵活性，适用于需要动态调整路径的场景



2021 年：CUDA 11.4 引入图编译器（Graph Compiler）

• 核心优化：在图实例化阶段加入全局优化模块，自动对图结构进行深度优化
• 关键功能：
  • 全局指令重排（合并内存操作、平衡计算与内存密集型任务）
  • 自动分片超大规模图，实现设备端并行调度
  • 优化 Graph Segment Table 的段间切换效率
• 价值：大幅提升复杂图的执行效率，尤其对包含数百个节点的大图，性能提升可达 10%-30%



2022 年：CUDA 12.0 增强动态图与硬件协同

• 核心改进：
  • 扩展 Graph Segment Table 支持多输入 / 多输出分支，适配更复杂的动态逻辑
  • 与新 GPU 架构（如 Hopper）的异步拷贝、线程块调度特性深度集成
  • 优化图的序列化 / 反序列化，支持跨进程图复用
• 价值：进一步缩小动态图与静态图的性能差距，同时提升大图的可移植性



2023 年至今：持续优化生态适配

• 核心方向：
  • 与深度学习框架（PyTorch、TensorFlow）更紧密集成，简化动态图的构建接口
  • 增强对稀疏图、异构计算（CPU+GPU）的支持
  • 提供更详细的图性能分析工具（如 NSight 中的 Graph Profiler）
• 价值：降低开发者使用门槛，推动计算图在更多领域的普及（如科学计算、自动驾驶）



演进主线总结

• 静态图阶段（2017-2019）：解决 “重复启动开销” 问题，支持固定流程的高效执行
• 动态图突破（2020）：通过 Graph Segment Table 实现设备端动态分支，解决灵活性问题
• 深度优化阶段（2021 - 至今）：引入图编译器并持续硬件协同，提升复杂图的性能与易用性



3. 为什么需要 CUDA Graph？

在传统的 CUDA 流（Stream）模型中，每个操作（如内核启动、内存拷贝）需要逐个提交到流中，GPU 驱动程序需要逐个处理这些请求。当存在大量小任务时，CPU 与 GPU 之间的通信开销会显著增加。CUDA Graph 通过预录制执行序列，将多个操作合并为一个整体提交。

Graph Launch 一次性提交所有工作，降低 CPU 开销（减少 CPU 调度开销、避免多次启动内核的延迟、提高 GPU 利用率）：



示例：用 CUDA Graph 加速向量加法

• 传统流模式：每次循环需要约 5μs 的 CPU 调度开销，100 次循环总开销约 500μs

#include <cstdio>
#include <cuda_runtime.h>

__global__ void vectorAdd(float* A, float* B, float* C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) C[i] = A[i] + B[i];
}

int main() {
    const int N = 1 << 20; // 1M 元素
    float *A, *B, *C;
    cudaMalloc(&A, N * sizeof(float));
    cudaMalloc(&B, N * sizeof(float));
    cudaMalloc(&C, N * sizeof(float));

    // 执行 100 次向量加法
    for (int i = 0; i < 100; i++) {
        dim3 block(256);
        dim3 grid((N + block.x - 1) / block.x);
        vectorAdd<<<grid, block>>>(A, B, C, N);
        cudaDeviceSynchronize(); // 同步等待完成
    }

    cudaFree(A);
    cudaFree(B);
    cudaFree(C);
    return 0;
}

• CUDA Graph 模式：首次录制图需要约 50μs，后续每次循环仅需 0.5μs，总开销约 50 + 100*0.5 = 100μs，性能提升约 5倍

#include <cstdio>
#include <cuda_runtime.h>

__global__ void vectorAdd(float* A, float* B, float* C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) C[i] = A[i] + B[i];
}

int main() {
    const int N = 1 << 20; // 1M 元素
    float *A, *B, *C;
    cudaMalloc(&A, N * sizeof(float));
    cudaMalloc(&B, N * sizeof(float));
    cudaMalloc(&C, N * sizeof(float));

    // 创建 CUDA 流和事件
    cudaStream_t stream;
    cudaStreamCreate(&stream);

    // 1. 创建并记录 CUDA Graph
    cudaGraph_t graph;
    cudaGraphExec_t instance;
    cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal); // 开始录制

    // 录制内核启动操作
    dim3 block(256);
    dim3 grid((N + block.x - 1) / block.x);
    vectorAdd<<<grid, block, 0, stream>>>(A, B, C, N);

    cudaStreamEndCapture(stream, &graph); // 结束录制
    cudaGraphInstantiate(&instance, graph, NULL, NULL, 0); // 实例化图

    // 2. 执行图（替代传统循环）
    for (int i = 0; i < 100; i++) {
        cudaGraphLaunch(instance, stream); // 提交整个图执行
        cudaStreamSynchronize(stream);
    }

    // 释放资源
    cudaGraphDestroy(graph);
    cudaGraphExecDestroy(instance);
    cudaStreamDestroy(stream);
    cudaFree(A);
    cudaFree(B);
    cudaFree(C);
    return 0;
}



3. CUDA Graph 工作流程

1. 定义阶段：定义 CUDA 图的结构、内容，告诉 CUDA 节点、参数、依赖关系，目标GPU是哪个，要启动到哪个流
2. 实例化阶段：获取CUDA Graph的模板，并执行配置和初始化从而节省启动时所需要执行的操作，生成的CUDA Graph实例即为可执行图形
3. 执行阶段：实例化后的CUDA Graph可以被直接部署到CUDA Stream当中（类似于一个kernel函数）可以重复启动无需多次实例化

示例：

1. 创建CUDA Graph

//创建CUDA Graph
cudaGraphCreate(&graph, 0);

//定义需要的CUDA Graph节点
cudaGraphAddKernelNode(&a, graph, NULL, 0, &nodeParams);
cudaGraphAddKernelNode(&b, graph, NULL, 0, &nodeParams);
cudaGraphAddKernelNode(&c, graph, NULL, 0, &nodeParams);
cudaGraphAddKernelNode(&d, graph, NULL, 0, &nodeParams);

// 定义上一步中CUDA节点之间的连接边
cudaGraphAddDependencies(graph, &a, &b, 1);     // A->B
cudaGraphAddDependencies(graph, &a, &c, 1);     // A->C
cudaGraphAddDependencies(graph, &b, &d, 1);     // B->D
cudaGraphAddDependencies(graph, &c, &d, 1);     // C->D

2. 对执行的CUDA Stream进行捕获——核心API：cudaStreamBeginCapture() 与cudaStreamEndCapture() 控制捕获开始与结束

cudaGraph_t graph;
//开始捕捉
cudaStreamBeginCapture(stream);

//CUDA Stream中执行kernel函数
kernel_A<<< ..., stream >>>(...);
kernel_B<<< ..., stream >>>(...);
libraryCall(stream);
kernel_C<<< ..., stream >>>(...);

//CUDA Stream运行完成后停止捕获，得到一个CUDA Graph
cudaStreamEndCapture(stream, &graph);

捕获CUDA Stream时，CUDA Stream中的kernenl函数或其他操作不会被enqueue到队列中排队等待执行，而是直接被构建到CUDA Graph中，随后通过cudaStreamEndCapture 结束捕获并得到CUDA Graph。



4. CUDA Graph如何提升大模型推理性能

在对延迟极度敏感的大语言模型推理场景中，GPU的强大算力常被CPU的控制开销所限制。

可以将传统架构，比喻为一个“急性子的主厨（CPU）”与“一群技艺精湛但绝对服从的帮厨（GPU流多处理器SMs）”的厨房。

• 主厨 (CPU)：负责阅读菜谱（执行推理代码），并将每一个烹饪步骤（如矩阵乘法、激活函数等）拆解成独立的指令，然后逐条下达。
• 帮厨 (GPU)：拥有海量并行处理能力，能瞬间完成任何单一指令（执行一个CUDA Kernel），但完成一步后，必须停下，等待主厨的下一道命令。

这个流程是严格的命令式（Imperative）：主厨每喊一个指令，帮厨们就迅速完成，然后集体“立正”，等待主'厨的下一个指令。对于一次包含数千个操作的复杂LLM推理，主厨需要不间断地下达数千个指令。问题在于，主厨下达指令本身需要时间，这个时间，就是所谓的控制开销（Control Overhead）。



CPU的控制开销主要来源于四个方面：

1. 内核启动延迟 (Kernel Launch Latency)：CPU通过CUDA驱动向GPU发起一次Kernel启动，这个过程本身就存在微秒级的固定延迟。对于LLM中海量的小型Kernel（如逐元素加法、激活函数），所有启动延迟累加起来，甚至可能超过Kernel的实际执行时间。
2. CUDA API调用开销：除了启动Kernel，每一次内存操作（cudaMalloc, cudaMemcpy）、同步（cudaStreamSynchronize）等API调用，都伴随着CPU与GPU驱动间的上下文切换和验证，这些看似不起眼的操作，积少成多，构成了显著的开销。
3. CPU调度抖动 (Jitter)：CPU作为通用处理器，其上的推理线程会不可避免地受到操作系统（OS）调度的影响。任何微小的中断或上下文切换，都会打乱向GPU发送指令的节奏，引入不确定的延迟“毛刺”。
4. 动态性带来的重复开销：对于每一次新的推理请求，CPU都需要重复地执行几乎完全相同的逻辑判断、参数计算和Kernel启动序列。这在架构上，是一种巨大的冗余浪费。

在训练阶段，较大的批次（Batch Size）和计算密度可以摊销这些微秒级的开销。但在要求低延迟的在线推理服务中（Batch Size通常为1），这些开销被无限放大，CPU成为了那个发号施令最慢的环节，导致GPU宝贵的计算核心被大量浪费在空闲等待中。



CUDA Graph将GPU的执行模型，从依赖CPU实时指挥的“命令式”，转变为一次性定义、可重复执行的声明式（Declarative）模型。

阶段一：捕获（Capture）—— 绘制静态执行图

在此阶段，CPU依然像往常一样，按顺序执行一次完整的推理计算流。CUDA驱动此时将CPU发出的所有CUDA相关操作（Kernel启动、内存拷贝、事件同步等），连同它们的参数、依赖关系，在GPU驱动内部构建成一个有向无环图（Directed Acyclic Graph, DAG）。

这个过程，好比主厨不再直接对帮厨们喊话，而是花时间将一整套复杂菜肴（如“法式酥皮鹅肝鸭肉派”）的全部流程，事无巨细地绘制成一张SOP流程图，这张图，就是CUDA Graph。

# 传统模式 vs. CUDA Graph模式
 
# --- 传统命令式模型 (每次请求都需要CPU逐条发指令) ---
def traditional_inference(input_data):
    cudaMemcpy(d_in, input_data, ...)
    launch_kernel_1(d_in, d_mid, ...)
    launch_kernel_2(d_mid, d_out, ...)
    cudaMemcpy(result, d_out, ...)
    return result
 
# --- CUDA Graph 声明式模型 ---
# 1. 捕获阶段 (执行一次，生成图)
stream = cudaStreamCreate()
cudaStreamBeginCapture(stream)
# ---- 把要录制的操作放入流中 ----
cudaMemcpy(...)
launch_kernel_1(...)
launch_kernel_2(...)
cudaMemcpy(...)
# --------------------------------
graph = cudaStreamEndCapture(stream)
graph_exec = cudaGraphInstantiate(graph) # 实例化可执行图
 
# 2. 执行阶段 (后续所有请求，CPU只需一个指令)
def graph_inference(input_data):
    # (可能需要更新输入/输出数据的指针)
    update_graph_params(graph_exec, input_data, result_buffer)
    cudaGraphLaunch(graph_exec, stream) # CPU只需轻量级启动
    cudaStreamSynchronize(stream)

阶段二：实例化与执行

一旦图捕获完成，CPU的角色就发生了根本性转变。在后续的每一次推理中，CPU不再需要重复上千次的API调用，只需向GPU发出一个极其轻量的指令：“执行这张图（Launch Graph）”

GPU 驱动在接收到这个单一指令后，将接管全部的控制权。它手握完整的计算图，可以在GPU内部，以最高效、最低开销的方式，调度执行图中定义的所有操作，实现了Kernel的 Back-to-Back 执行，彻底绕过了与CPU的反复通信和操作系统的调度抖动。

这相当于厨房拿到了SOP后，帮厨们形成了“肌肉记忆”。主厨每次只需喊一声菜名（“来一份法式酥皮鹅肝鸭肉派！”），整个后厨便能心领神会，行云流水般地完成所有工序，中间无需主厨再做任何干预。

通过这种“一次捕获，多次重放”的架构，CUDA Graph将原本分散在无数次CPU-GPU交互中的控制开销，一次性地摊销在了初始的捕获阶段。在至关重要的执行阶段，实现了近乎“零”的CPU开销，将性能瓶颈重新交还给GPU的算力本身。



5. CUDA Graph在LLM推理中的应用模式

模式一：静态输入的完全图化（最理想）

这是最简单直接的应用模式。当推理请求的输入形状（如Batch Size, Sequence Length）完全固定时，从数据拷贝到计算再到结果回传的整个端到端流程，都可以被完整捕获到一个CUDA Graph中。

但是缺乏灵活性。一旦输入形状改变，整个图就需要废弃并重新捕获，这在动态性强的在线服务中是不可接受的。

模式二：“分段图化”与“动态参数”组合拳（最实用）

• 策略1——“分段图化” (Segmented Graphing)
  • LLM的推理过程可清晰地划分为两个阶段：Prefill（对输入Prompt的并行处理）和Decoding（逐Token的自回归生成）。
  • Prefill阶段的计算图结构与输入序列长度强相关，是动态性的主要来源。而Decoding阶段，由于每次只生成一个Token，其计算模式是固定且高度重复的。
  • 因此，最佳实践是将高度重复的Decoding阶段图化，称之为“Decoding Graph”或“Step Graph”，能够优化掉推理过程中绝大部分（通常是95%以上）的CPU控制开销。

• 策略2——“图更新”与“动态参数” (Graph Update)

  • CUDA提供了Graph Update机制，允许在不重新捕获整个图的情况下，修改图中某些节点（如memcpy或Kernel）的参数，例如指向输入/输出数据的内存地址指针。这在架构上实现了“结构静态，数据动态”。可以捕获一个通用的计算图，在每次执行前，通过Graph Update 将其I/O指针动态地指向当前请求的实际数据缓冲区，极大地提升了灵活性。

• 策略3——“装桶与填充” (Bucketing & Padding)

  • 这是处理动态序列长度的经典工程实践。可以预先为一系列离散的、有代表性的序列长度（如64, 128, 256, 512...）分别捕获并缓存对应的CUDA Graph（目前Sglang就是这么做的）。运行时，当接收到一个请求，我们将其输入序列填充（Pad）到最接近且更大的那个“桶（Bucket）”的长度，然后直接调用该桶预先编译好的Graph。这是一种典型的空间换时间的架构权衡，通过增加内存占用换取了在动态输入下的高性能执行。

模式三：作为“顶层粘合剂”与高性能库集成

在一个设计精良的推理系统中，CUDA Graph应作为顶层调度器，其图中的节点调用的正是那些经过极致优化的计算单元。

• 封装高性能库：将对NVIDIA cuBLAS（矩阵运算）、cuDNN（卷积运算）等官方库的调用序列，封装进Graph中。
• 集成自定义核：将像FlashAttention、vLLM中PagedAttention这样的一系列高度优化的自定义CUDA Kernel调用，捕获到Graph中，实现优化的“强强联合”，消除这些自定义核之间的启动开销。



style="display: none !important;">