CUDA Graphs 实战：用 cuda.core 在 cuda-samples 中捕获并回放多阶段内核流水线

CUDA Graphs 实战用 cuda.core 在 cuda-samples 中捕获并回放多阶段内核流水线【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples本文围绕 CUDA Samples 仓库中的 Python 示例 cudaGraphs讲解如何借助cuda.core的 Python 化 CUDA Graphs API把多阶段 kernel 流水线捕获为一张 CUDA graph并用单次驱动调用反复回放。你将掌握图捕获的完整调用链GraphBuilder构建 →complete()生成 →upload()上传 →launch()回放理解图捕获指针而非数据的语义并学会用脚本测量小 kernel 场景下 CPU 侧启动开销的节省幅度。背景CUDA Graphs 要解决什么问题CUDA 应用中每个 kernel 的发起都需要 CPU 侧付出启动开销参数打包、命令下发、依赖调度等。对于运行时间很长的 kernel这笔开销可以忽略但对于大量短小的 kernel尤其是在循环中反复发起的场景CPU 启动开销会显著拖慢整体吞吐。CUDA Graphs 的核心思想是把一组有依赖关系的 kernel或其他操作先录制为一张有向无环图DAG之后以单次驱动调用回放整张图。录制只发生一次回放时驱动不再逐 kernel 走完整的启动路径从而大幅削减 CPU 侧开销。本示例源码见 cudaGraphs.py正是这一思路的最小可运行演示用cuda.core捕获一个三段式元素级流水线并与逐 kernel 单独启动的方式做计时对比。示例总览三段流水线与两种运行模式示例计算的流水线公式为r3 (a b) * c - a即三次逐元素运算先加add再乘mul最后减sub中间结果分别写入临时缓冲区r1、r2最终结果写入r3。整个流程以两种模式各执行 N 次迭代Individual launches单独启动每个 stage 一次launch(stream, ...)循环中每轮迭代发起三次独立 kernel 启动CUDA graph replay图回放把同样的三次 launch 录制进一张Graph之后每轮迭代只调用一次graph.launch(stream)。两种路径都会计时、并与参考计算expected (a b) * c - a做精度校验cp.allclose容差rtol1e-5, atol1e-5。最后示例会修改输入缓冲区的内容并再次回放同一张图验证图捕获的是指针而非数据这一关键语义。环境要求与安装硬件NVIDIA GPUCompute Capability 7.0 或更高Volta 及之后架构最小 GPU 显存512 MB。软件CUDA Toolkit 13.0 或更新与cuda-python13.x 匹配Python 3.10 或更新cuda-python13.0.0、cuda-core1.0.0、cupy-cuda13x14.0.0。示例目录内的 requirements.txt 精确列出了这三项依赖。安装命令cd /path/to/cuda-samples/python/2_CoreConcepts/cudaGraphs pip install -r requirements.txt仓库根目录的 python/requirements.txt 还额外声明了numpy2.3.2本示例用 NumPy 承载标量 kernel 参数与参考计算并注明各示例的额外依赖由各目录自己的 requirements.txt 安装。运行示例与命令行参数基本用法cd cuda-samples/python/2_CoreConcepts/cudaGraphs python cudaGraphs.py自定义参数# 更大的向量 更多迭代次数 python cudaGraphs.py --elements 4096 --iters 2000 # 指定 GPU python cudaGraphs.py --device 1三个命令行参数的定义在源码 cudaGraphs.py 的main()中通过argparse声明参数默认值含义--elements1 124096每个向量的元素个数。源码注释明确说明默认取小尺寸正是为了放大启动开销占比--iters1000计时的流水线迭代次数--device0CUDA device id关于参数选择的技巧README 给出了关键提示短向量会放大图回放的加速收益当向量变大时kernel 本身的执行时间占比上升单次启动开销相对微不足道两种方式的耗时会逐渐趋近。源码剖析一内核定义与 JIT 编译示例没有使用预编译的.cu文件而是把三段 kernel 源码以内嵌字符串PIPELINE_KERNELS的形式直接写在 Python 文件中运行时通过cuda.core的Program接口即时编译JIT。三个 kernel 的签名与实现完全对称以vec_add为例extern C __global__ void vec_add(const float* A, const float* B, float* C, size_t N) { size_t tid blockIdx.x * blockDim.x threadIdx.x; size_t stride (size_t)gridDim.x * blockDim.x; for (size_t i tid; i N; i stride) C[i] A[i] B[i]; }vec_mul与vec_sub结构相同只是把换成*与-。注意它们都采用grid-stride loop带步长的网格循环写法即每个线程用tid、stride循环处理多个元素这样内核可以适配任意N而不要求N恰好是 block 数的整数倍。编译与取内核的调用链如下来自 cudaGraphs.pyprogram_options ProgramOptions(stdc17, archfsm_{device.arch}) program Program(PIPELINE_KERNELS, code_typec, optionsprogram_options) module program.compile(cubin) add_k module.get_kernel(vec_add) mul_k module.get_kernel(vec_mul) sub_k module.get_kernel(vec_sub)ProgramOptions指定 C 标准为c17目标架构由当前设备自动推导sm_{device.arch}Program.compile(cubin)把源码编译为 cubin 模块module.get_kernel(...)取出三个具名内核供后续 launch 使用。启动配置则统一使用LaunchConfigconfig LaunchConfig(grid(N 255) // 256, block256)即每个 block 256 线程block 数量按ceil(N / 256)计算。源码剖析二单独启动路径run_pipeline_individual()实现逐 kernel 启动模式每轮迭代依次发起三次 launchlaunch(stream, config, add_k, a.data.ptr, b.data.ptr, r1.data.ptr, np.uint64(size)) launch(stream, config, mul_k, r1.data.ptr, c.data.ptr, r2.data.ptr, np.uint64(size)) launch(stream, config, sub_k, r2.data.ptr, a.data.ptr, r3.data.ptr, np.uint64(size))要点缓冲区通过a.data.ptr传入裸设备指针N用np.uint64(size)包装为标量 kernel 参数这正是 README 中numpy- scalar kernel arguments的用途数据依赖靠同一 stream 上的顺序保证add 的结果写入r1后mul 才能在下一轮读到它计时使用time.perf_counter()先stream.sync()再开始、结束后再stream.sync()确保测得的是 GPU 实际完成时间而非命令下发时间。正式计时前示例会先以n_iters5跑一遍热身warm up目的是触发编译/缓存等一次性开销避免污染正式测量。源码剖析三图捕获与回放build_graph()演示了用cuda.core捕获 CUDA graph 的完整流程共五步graph_builder stream.create_graph_builder() # 1. 从 stream 创建 GraphBuilder graph_builder.begin_building() # 2. 开始录制 launch(graph_builder, config, add_k, ...) # 把 launch 的目标从 stream 换成 graph_builder launch(graph_builder, config, mul_k, ...) launch(graph_builder, config, sub_k, ...) graph_builder.end_building() # 3. 结束录制 graph graph_builder.complete() # 4. 生成可执行的 Graph graph.upload(stream) # 5. 把图结构上传到设备 return graph_builder, graph对应的cuda.coreAPI 即 README Key APIs 中列出的Stream.create_graph_builder()—— 获得一个与 stream 关联的GraphBuilderGraphBuilder.begin_building()/end_building()—— 界定录制区间GraphBuilder.complete()—— 产出可执行的Graph对象Graph.upload(stream)—— 将图结构上传至设备通常在首次回放前调用一次Graph.launch(stream)—— 回放整张图。值得注意的是录制阶段发起 launch 时第一个参数从stream换成了graph_builder—— 同样的launch()函数签名既支持即时启动也支持把启动记入正在构建的图中这正是cuda.core统一抽象带来的简洁性。回放路径run_pipeline_graph()则极其简单循环内只有一行graph.launch(stream)同样先stream.sync()再计时再stream.sync()收尾。一个容易被忽略的细节是 CuPy 与 stream 的绑定main()中在分配缓冲区前调用了cp.cuda.Stream.from_external(stream).use()其作用源码注释是让 CuPy 的分配发生在我们的 stream 上从而保证缓冲区初始化与后续 kernel 启动在时序上串行化避免数据竞争。指针语义同一张图处理新数据CUDA graph 录制时会把 kernel 参数包括缓冲区指针固化在图中但不会固化指针指向的数据内容。因此只要缓冲区还在图就可以反复回放去处理更新后的数据无需重新录制。示例末尾用一组特殊值验证了这一点见 cudaGraphs.pya[:] cp.ones(N, dtypecp.float32) b[:] cp.full(N, 2.0, dtypecp.float32) c[:] cp.full(N, 3.0, dtypecp.float32) device.sync() # r3 (a b) * c - a (1 2) * 3 - 1 8 graph.launch(stream) stream.sync() assert cp.allclose(r3, 8.0), Graph replay with new data produced wrong result写入新数据后仅回放一次图r3便得到预期的8.0从而证明图绑定的是缓冲区地址而非缓冲区内容。这对实践有直接意义——只要输入输出缓冲区的布局和大小不变就可以复用同一张图处理源源不断的新数据例如每帧图像、每个 mini-batch省去反复录制的高昂成本。预期输出与结果解读README 给出的示例输出如下实际数值随 GPU 与宿主 CPU 而异Device: Your GPU Name Compute Capability: X.Y Individual launches: 1000 iters in 0.0085s (8.49 us/iter) Building CUDA graph... Graph replay: 1000 iters in 0.0034s (3.41 us/iter) Graph speedup: 2.49x Graph replay on updated data verified (same graph, new buffer contents) Done解读要点设备信息来自cuda_samples_utils.print_gpu_info()它打印device.name与compute_capabilitycc.major、cc.minor实现见 cuda_samples_utils.py加速比 单独启动总耗时 ÷ 图回放总耗时。示例中为 2.49x但 README 明确提示设备名、算力与加速比都会随硬件变化请勿将任何具体数值当作固定结论默认--elements 4096的短向量刻意放大了启动开销占比因此图回放优势明显把--elements调大后两种方式会逐渐趋近kernel 执行时间主导时启动开销差异被稀释。资源清理与健壮性设计源码main()的try/finally块展示了正确的资源生命周期管理结束时依次graph.close()、graph_builder.close()、stream.close()并把 CuPy 的默认 stream 恢复为cp.cuda.Stream.null.use()。此外导入阶段用try/except ImportError捕获缺失依赖并提示安装 requirements.txt三段计算单独启动、图回放、更新数据回放均以assert cp.allclose(...)校验正确性任何一步出错都会使脚本以非零状态退出。这套捕获 → 校验 → 计时 → 清理的结构可以当作编写 CUDA Graphs 性能基准脚本的参考模板。文件清单与后续深入示例目录结构如下cudaGraphs.py —— 基于cuda.coreCUDA Graphs 的 Python 实现README.md —— 本示例文档requirements.txt —— 示例依赖声明cuda_samples_utils.py —— 公共工具被本示例 import提供print_gpu_info()等。如果想继续深入可以在本仓库中对照阅读cuda.core的Program/ProgramOptions/LaunchConfig/launch还用于 matrixMulSharedMem、parallelReduction 等示例CUDA Graphs 的进阶主题条件节点、内存节点、图级内存足迹等可参考 C 侧的 3_CUDA_Features 目录下的graphConditionalNodes、graphMemoryFootprint、graphMemoryNodes等示例CUDA Graphs 编程指南的完整语义录制限制、可捕获操作集合、实例化与更新机制则见官方 CUDA C Programming Guide 的 CUDA Graphs 章节。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考