当前位置： 首页 > news >正文

CUDA graph 简析

news 2026/3/26 17:31:51

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;">

http://www.jsqmd.com/news/466396/

相关文章：

  • 基于微信小程序的课程作业管理系统[小程序]-计算机毕业设计源码+LW文档
  • 别死记硬背！Java的CountDownLatch 核心原理：AQS state 才是关键
  • 知识体系——MCP（四）demo（2）开发mcp client
  • OWASP Top10 2021 完整版：与 SAST 适配的深度解析
  • Rocky Linux 10 上搭建 社区版 GitLab CE
  • 2026年 智能制造实训设备厂家推荐排行榜：高校教学、模拟药厂、生产线实训平台与系统装置一站式解决方案 - 品牌企业推荐师（官方）
  • g更改linux root密码
  • LeetCode 76. 最小覆盖子串（详细技术解析）
  • 虚拟同步发电机（VSG）孤岛与并网的Simulink（2019a）仿真模型搭建与探索
  • 对于【LSTM与GRU在水文预测中的对比分析】的未来改进和建议
  • 工业清洁设备优质品牌推荐榜：驾驶式洗地机/1000公斤高压清洗机/商用洗地机/工业吸尘器/工业洗地机/工业清洗机厂家/选择指南 - 优质品牌商家
  • 2026年比较好的西安租赁洗地机工厂推荐：西安洗地机租赁稳定供应商推荐 - 行业平台推荐
  • visual studio编译wxWidgets
  • 防疫站疫苗预约管理系统_Python django flask
  • 2026宁波好用的芯轴品牌生产厂盘点，如何选择靠谱厂家 - 工业推荐榜
  • 2026河北新河优质MC浇筑尼龙加工件推荐榜：pa66尼龙棒/pp尼龙棒/尼龙加工件源头厂家/浇筑尼龙棒/玻纤mc尼龙浇铸棒/选择指南 - 优质品牌商家
  • Django + Vue3 + YOLO 实现车辆检测、测速预警与违章分析平台
  • 互联网大厂Java面试：谢飞机与严肃面试官的搞笑对决
  • 救命神器! 降AIGC工具 千笔·降AIGC助手 VS WPS AI 专科生专属
  • Gemini认证工具创意开发（技术深度解析）
  • 探讨江苏地区井口装置涂装厂家排名，江苏万和涂装排第几？ - myqiye
  • 论文查重「避坑」全攻略：Paperxie 四大检测体系如何守住你的毕业「生命线」
  • 2026年胰岛素泵市场盘点：哪款专业品牌更胜一筹？
  • leetcode 3600. 升级后最大生成树稳定性 困难
  • 北京/上海/深圳/杭州/南京/无锡高端腕表维修指南：豪爵/库尔沃/蕾蒙威/播威故障养护与维修全解析 - 时光修表匠
  • 收藏备用！程序员转型AI的三个核心赛道（小白/进阶通用）
  • 产品推荐|八戒光度成像系统全新小型化款来了！
  • word打字输入及删除 很卡，延迟几秒钟
  • 《OpenClaw 实战：从 0 到 1 快速入门到进阶实战》一本全面掌握 OpenClaw 云桌面助理的实战指南：第二部分《进阶篇》
  • 《投资-407》长期价值投资考验的是眼光与格局, 考验的是战略方向的能力，其难度远大于战术上勤奋的能力，如何提升这方面的能力？