CUDA 条件图节点Conditional Graph Nodes实战指南基于 cuda-samples 的 graphConditionalNodes 示例全解析【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples导读本文以 cuda-samples 仓库中的 graphConditionalNodes 示例 为核心系统讲解 CUDA Graphs 条件节点Conditional Graph Nodes的完整使用方式。条件节点允许在已实例化的 CUDA Graph 内部根据运行时条件动态选择执行路径实现图内的 if / if-else / do-while / while / switch 语义从而把以往需要在 Host 侧反复启动图才能实现的动态逻辑下沉到 GPU 端。读完本文你将掌握cudaGraphConditionalHandleCreate、cudaGraphSetConditional、cudaGraphCondTypeIf/While/Switch等核心 API 的用法理解显式构图 流捕获stream capture填充条件体两种混合编程模式并了解驱动版本对条件节点特性的门控规则。一、示例定位在 CUDA Graphs 中引入运行时分支在传统 CUDA Graphs 使用中图一旦实例化cudaGraphInstantiate其节点拓扑与依赖关系便是固定的若要依据设备内存中的数据动态改变计算路径通常只能在 Host 端反复修改图或启动多个图带来额外的启动开销与同步成本。条件节点Conditional Node改变了这一局面它作为图中的一个节点内嵌一个或多个子图body graph并在图执行期间根据条件句柄conditional handle的取值动态决定执行哪个子图甚至循环执行同一子图。正如示例源文件 graphConditionalNodes.cu 注释所述This file demonstrates the usage of conditional graph nodes with a series ofsimpleexample graphs.该示例通过5 个自包含的小型示例覆盖了条件节点的全部核心形态示例函数图结构条件类型特性simpleIfGraphA - B [ C ]IF条件为真才执行体内 kernelsimpleDoWhileGraphB [ A - B - C ]WHILE借助默认值实现 do-while 语义capturedWhileGraphA - B [ C ] - DWHILE流捕获中手动插入条件节点simpleIfElseGraphA - B [ C \| D ]IF双子图真/假各对应一个子图需 CUDA ≥ 12.8simpleSwitchGraphA - B [ C \| D \| E \| F ]SWITCH按整数值选择 4 个子图之一需 CUDA ≥ 12.8下面逐一展开每个示例的实现细节。二、环境与依赖要求2.1 支持的硬件与平台依据示例 README.md 的声明支持的 SM 架构SM 6.0、6.1、7.0、7.2、7.5、8.0、8.6、8.7、8.9、9.0即 Pascal 至 Hopper/Blackwell 一代的主流计算能力支持的操作系统Linux、Windows支持的 CPU 架构x86_64、armv7l前置条件下载并安装与平台对应的 CUDA Toolkit。需要注意的是条件节点本身需要较新的驱动支持源码在main()中做了显式门控详见第六节基础条件节点IF/WHILE要求驱动 ≥ 120.30CUDA 12.3if-else 与 switch 类型要求驱动 ≥ 120.80CUDA 12.8。2.2 构建方式CMake示例通过标准 CMake 集成到整个仓库的构建体系中。其在 3_CUDA_Features/CMakeLists.txt 中以add_subdirectory(graphConditionalNodes)被引入子目录自身的 CMakeLists.txt 定义了构建细节使用find_package(CUDAToolkit REQUIRED)定位 CUDA 工具包编译目标为可执行文件graphConditionalNodes源文件即 graphConditionalNodes.cu默认编译架构覆盖75 80 86 87 89 90 100 110 120并开启CUDA_SEPARABLE_COMPILATION引入仓库公共头文件目录 Common其中包含helper_cuda.h、helper_functions.h等工具头。在仓库根目录执行标准 CMake 流程即可构建整个cpp子树根 CMakeLists.txt 已通过add_subdirectory(cpp)纳入全部 C 示例。三、示例用到的 CUDA Runtime API 全景原文档 README.md 列出了本示例涉及的全部 CUDA Runtime API按其职责可归类如下设备与内存管理cudaMalloc、cudaFree、cudaMemset、cudaDeviceSynchronize图对象管理cudaGraphCreate、cudaGraphDestroy、cudaGraphAddNode、cudaGraphInstantiate、cudaGraphLaunch、cudaGraphExecDestroy条件节点专属cudaGraphConditionalHandleCreate创建条件句柄、cudaGraphSetConditional在 kernel 内部设置条件值流捕获cudaStreamBeginCapture、cudaStreamEndCapture、cudaStreamBeginCaptureToGraph、cudaStreamGetCaptureInfo、cudaStreamCreate、cudaStreamDestroy版本查询cudaDriverGetVersion。此外源码还使用了 README 未列入的cudaStreamUpdateCaptureDependencies用于在流捕获过程中手动修正依赖关系见capturedWhileGraph小节。所有 API 调用都通过公共头文件 Common/helper_cuda.h 中的checkCudaErrors宏进行错误检查任何一步失败都会打印出错位置并终止程序这也是 cuda-samples 系列的一贯风格。四、示例一simpleIfGraph —— 最基本的 IF 条件节点simpleIfGraph构建的图结构为A - B [ C ]节点 A 是上游 kernel负责设置条件节点 B 是 IF 条件节点其条件体内只有一个 kernel C。4.1 条件设置 Kernel节点 A 的 kernelifGraphKernelA读取设备内存中的字节取其最低位作为条件值并通过cudaGraphSetConditional写入条件句柄__global__ void ifGraphKernelA(char *dPtr, cudaGraphConditionalHandle handle) { // In this example, condition is set if *dPtr is odd unsigned int value *dPtr 0x01; cudaGraphSetConditional(handle, value); printf(GPU: Handle set to %d\n, value); }cudaGraphSetConditional是可以在设备端 kernel 中调用的函数它把条件值写入手柄图运行时将据此判定后续条件节点的走向。节点 C 的 kernelifGraphKernelC仅在条件为真时执行// This kernel will only be executed if the condition is true __global__ void ifGraphKernelC(void) { printf(GPU: Hello from the GPU! The condition was true.\n); }4.2 显式构图流程整个构图过程是显式构图的典型范式分为五步第一步创建图与条件句柄。cudaGraphCreate(graph, 0); cudaGraphConditionalHandleCreate(handle, graph);cudaGraphConditionalHandleCreate在指定图中创建条件句柄句柄与图绑定后续由条件节点消费。第二步用cudaGraphNodeParams添加 kernel 节点 A。通过cudaGraphNodeParams结构体描述节点类型与参数params.kernel.func指定 kernel 函数指针gridDim/blockDim设置执行规模kernelParams传递参数数组cudaGraphNodeParams params {cudaGraphNodeTypeKernel}; params.kernel.func (void *)ifGraphKernelA; params.kernel.blockDim.x params.kernel.blockDim.y params.kernel.blockDim.z 1; params.kernel.gridDim.x params.kernel.gridDim.y params.kernel.gridDim.z 1; params.kernel.kernelParams kernelArgs; kernelArgs[0] dPtr; kernelArgs[1] handle; checkCudaErrors(cudaGraphAddNode(kernelNode, graph, NULL, NULL, 0, params));cudaGraphAddNode的第 3~5 个参数为前置依赖节点数组、数组长度等这里节点 A 是首节点因此依赖为空。第三步添加条件节点 B。关键在cudaGraphNodeParams的conditional子结构cudaGraphNodeParams cParams {cudaGraphNodeTypeConditional}; cParams.conditional.handle handle; cParams.conditional.type cudaGraphCondTypeIf; cParams.conditional.size 1; checkCudaErrors(cudaGraphAddNode(conditionalNode, graph, kernelNode, NULL, 0, cParams));type cudaGraphCondTypeIf声明这是一个 IF 类型条件节点size 1该条件节点内嵌1 个子图即仅真分支依赖为上游的kernelNode保证节点 A 先于条件判定执行。第四步填充条件体子图。cudaGraphAddNode返回后cParams.conditional.phGraph_out[0]即为条件节点内部第 0 个子图的句柄直接向其中添加节点即可cudaGraph_t bodyGraph cParams.conditional.phGraph_out[0]; cudaGraphNode_t bodyNode; params.kernel.func (void *)ifGraphKernelC; params.kernel.kernelParams nullptr; checkCudaErrors(cudaGraphAddNode(bodyNode, bodyGraph, NULL, NULL, 0, params));第五步实例化并两次启动验证不同分支。实例化cudaGraphInstantiate后通过cudaMemset把设备内存置为 0 或 1再分别cudaGraphLaunch两次checkCudaErrors(cudaMemset(dPtr, 0, 1)); // Set dPtr to 0 printf(Host: Launching graph with device memory set to 0\n); checkCudaErrors(cudaGraphLaunch(graphExec, 0)); checkCudaErrors(cudaDeviceSynchronize()); checkCudaErrors(cudaMemset(dPtr, 1, 1)); // Set dPtr to 1 checkCudaErrors(cudaGraphLaunch(graphExec, 0)); checkCudaErrors(cudaDeviceSynchronize());运行效果第一次dPtr0偶数只有节点 A 打印 Handle set to 0kernel C 不执行第二次dPtr1奇数除节点 A 外还会看到 kernel C 打印 Hello from the GPU! The condition was true.。五、示例二simpleDoWhileGraph —— 用默认条件值实现 do-whilesimpleDoWhileGraph演示如何用 WHILE 类型条件节点构造do-while 循环循环体至少执行一次直到计数递减到 0 才终止。5.1 关键差异句柄默认值与 IF 示例不同创建条件句柄时传入了默认值与赋值策略cudaGraphConditionalHandleCreate(handle, graph, 1, cudaGraphCondAssignDefault);第三个参数1表示默认条件值为真第四个参数cudaGraphCondAssignDefault表示该默认值只用于首次评估条件之前若后续条件被显式设置则以显式值为准。WHILE 节点在每次迭代开始前评估条件因此默认值为真就意味着循环体至少执行一次——这正是 do-while 语义。5.2 条件体流捕获填充条件节点的type为cudaGraphCondTypeWhilesize 1。与示例一的显式cudaGraphAddNode不同这里示范了第二种填充方式使用流捕获stream capture向已存在的 body graph 追加节点。cudaStream_t captureStream; checkCudaErrors(cudaStreamCreate(captureStream)); checkCudaErrors( cudaStreamBeginCaptureToGraph(captureStream, bodyGraph, nullptr, nullptr, 0, cudaStreamCaptureModeGlobal)); doWhileEmptyKernel1, 1, 0, captureStream(); doWhileEmptyKernel1, 1, 0, captureStream(); doWhileLoopKernel1, 1, 0, captureStream(dPtr, handle); checkCudaErrors(cudaStreamEndCapture(captureStream, nullptr)); checkCudaErrors(cudaStreamDestroy(captureStream));cudaStreamBeginCaptureToGraph允许把流捕获的目标指定为已有的 body graph捕获期间在流上发起的 kernel 启动会被记录为图节点cudaStreamEndCapture后这些节点即落入条件体。循环体包含 3 个节点两个doWhileEmptyKernel仅用于演示的占位节点和一个doWhileLoopKernel。5.3 终止条件递减计数doWhileLoopKernel每次迭代把设备内存中的计数器减 1归零时把条件设为假__global__ void doWhileLoopKernel(char *dPtr, cudaGraphConditionalHandle handle) { if (--(*dPtr) 0) { cudaGraphSetConditional(handle, 0); } printf(GPU: counter %d\n, *dPtr); }Host 侧把计数器初始化为 10 后启动图预期输出 counter 从 10 递减到 1第 10 次迭代减到 0 并置条件为假循环终止且循环体在首次迭代前无需额外判定。六、示例三capturedWhileGraph —— 流捕获中手动插入条件节点第三个示例展示了最复杂的场景在图处于流捕获过程中手动把条件节点嫁接进正在被捕获的图构造A - B [ C ] - D的 while 循环。6.1 在捕获中获取图句柄并创建条件节点先用cudaStreamBeginCapture开启流捕获再通过cudaStreamGetCaptureInfo取得当前正在捕获的图句柄cudaStreamBeginCapture(captureStream, cudaStreamCaptureModeGlobal); // Obtain the handle of the graph cudaStreamGetCaptureInfo(captureStream, status, NULL, graph, dependencies, NULL, numDependencies); // Create the conditional handle cudaGraphConditionalHandle handle; checkCudaErrors(cudaGraphConditionalHandleCreate(handle, graph));随后在捕获流上启动 kernel AcapturedWhileKernel再第二次调用cudaStreamGetCaptureInfo取得节点 A 的依赖信息用于把条件节点 B 正确地接在 A 之后capturedWhileKernel1, 1, 0, captureStream(dPtr, handle); cudaStreamGetCaptureInfo(captureStream, status, NULL, graph, dependencies, NULL, numDependencies); cudaGraphNodeParams cParams {cudaGraphNodeTypeConditional}; cParams.conditional.handle handle; cParams.conditional.type cudaGraphCondTypeWhile; cParams.conditional.size 1; checkCudaErrors(cudaGraphAddNode(conditionalNode, graph, dependencies, NULL, numDependencies, cParams));这里cudaGraphAddNode的依赖取自cudaStreamGetCaptureInfo返回的dependencies即捕获流中最后一个节点节点 A。6.2 修正捕获依赖继续插入节点 D由于条件节点是手动加入捕获图的捕获流自身并不知道它的存在因此需要调用cudaStreamUpdateCaptureDependencies把条件节点声明为捕获流的依赖之后在流上启动的 kernel D 才会正确地排在条件节点之后checkCudaErrors(cudaStreamUpdateCaptureDependencies( captureStream, conditionalNode, NULL, 1, cudaStreamSetCaptureDependencies)); // Insert kernel node D capturedWhileEmptyKernel1, 1, 0, captureStream(); checkCudaErrors(cudaStreamEndCapture(captureStream, graph));6.3 条件体用流捕获填充并解释 while 与 do-while 的差别条件体 C 同样通过cudaStreamBeginCaptureToGraph 捕获 kernel 完成。值得注意源码注释中特别强调的设计考量graphConditionalNodes.cuWe must run the kernel before the loop as well as inside the loop in order to behave like a while loop as opposed to a do-while loop.即WHILE 节点会在进入循环体之前评估条件所以必须在条件节点上游再放一个相同的 kernel节点 A来提前计算条件值否则就会退化成 do-while 行为。由于有上游 kernel 提前设值这里无需使用句柄默认值来初始化条件与示例二形成对照。capturedWhileKernel的逻辑是先打印当前计数若非零则减 1并把新值写回条件句柄__global__ void capturedWhileKernel(char *dPtr, cudaGraphConditionalHandle handle) { printf(GPU: counter %d\n, *dPtr); if (*dPtr) { (*dPtr)--; } cudaGraphSetConditional(handle, *dPtr); }Host 侧分别以计数器 0 与 10 启动图前者因条件为假条件体一次都不执行体现先判定再进入的 while 语义后者循环执行 10 次直至计数归零。七、示例四simpleIfElseGraph —— 真/假双分支CUDA ≥ 12.8simpleIfElseGraph把 IF 条件节点扩展为双分支形态A - B [ C | D ]条件为真执行子图 0kernel C条件为假执行子图 1kernel D。与示例一唯一的实质区别在于创建条件节点时把size设为 2cParams.conditional.handle handle; cParams.conditional.type cudaGraphCondTypeIf; cParams.conditional.size 2; // Set size to 2 to indicate an ELSE graph will be usedsize 2表示该条件节点内嵌2 个子图phGraph_out[0]为真分支、phGraph_out[1]为假分支分别填充即可cudaGraph_t bodyGraph cParams.conditional.phGraph_out[0]; // ... add ifGraphKernelC (true branch) ... bodyGraph cParams.conditional.phGraph_out[1]; // ... add ifGraphKernelD (false branch) ...运行两次dPtr0 与 dPtr1可分别观察到 The condition was false. 与 The condition was true. 两条输出。源文件注释明确标注该特性要求 CUDA ≥ 12.8graphConditionalNodes.cu。八、示例五simpleSwitchGraph —— 多路选择CUDA ≥ 12.8simpleSwitchGraph把条件节点扩展为SWITCH类型A - B [ C | D | E | F ]条件值是一个整数运行时选择第 n 个子图执行若条件值 ≥ 子图数量则任何子图都不执行。上游 kernelswitchGraphKernelA直接读取设备内存中的整数值写入句柄__global__ void switchGraphKernelA(char *dPtr, cudaGraphConditionalHandle handle) { unsigned int value *dPtr; cudaGraphSetConditional(handle, value); printf(GPU: Handle set to %d\n, value); }条件节点声明为cudaGraphCondTypeSwitchsize 4随后向phGraph_out[0]~phGraph_out[3]四个子图分别填充 kernel C、D、E、FcParams.conditional.handle handle; cParams.conditional.type cudaGraphCondTypeSwitch; cParams.conditional.size 4; checkCudaErrors(cudaGraphAddNode(conditionalNode, graph, kernelNode, NULL, 0, cParams)); // Populate the four graph bodies within the SWITCH conditional graph params.kernel.func (void *)switchGraphKernelC; checkCudaErrors(cudaGraphAddNode(bodyNode, cParams.conditional.phGraph_out[0], NULL, NULL, 0, params)); params.kernel.func (void *)switchGraphKernelD; checkCudaErrors(cudaGraphAddNode(bodyNode, cParams.conditional.phGraph_out[1], NULL, NULL, 0, params)); params.kernel.func (void *)switchGraphKernelE; checkCudaErrors(cudaGraphAddNode(bodyNode, cParams.conditional.phGraph_out[2], NULL, NULL, 0, params)); params.kernel.func (void *)switchGraphKernelF; checkCudaErrors(cudaGraphAddNode(bodyNode, cParams.conditional.phGraph_out[3], NULL, NULL, 0, params));Host 侧用一个循环把设备内存依次设为 0~4 并分别启动图条件值 0~3 时依次触发 C/D/E/F 四个 kernel 之一条件值 4≥ size时则无任何分支执行。同样地SWITCH 类型也要求CUDA ≥ 12.8。九、驱动版本门控main() 的兼容性处理示例在main()中通过cudaDriverGetVersion获取驱动版本返回值单位为百位数表示主版本据此决定执行哪些示例graphConditionalNodes.cuint driverVersion 0; cudaDriverGetVersion(driverVersion); printf(Driver version is: %d.%d\n, driverVersion / 1000, (driverVersion % 100) / 10); if (driverVersion 12030) { printf(Skipping execution as driver does not support Graph Conditional Nodes\n); return 0; } simpleIfGraph(); simpleDoWhileGraph(); capturedWhileGraph(); if (driverVersion 12080) { printf(Skipping execution as driver does not support if/else and switch type Graph Conditional Nodes\n); return 0; } simpleIfElseGraph(); simpleSwitchGraph();由此可以明确版本与特性的对应关系驱动 120.30CUDA 12.3 之前完全不支持条件节点程序直接跳过全部示例驱动 ≥ 120.30支持 IF / WHILE 类型前三个示例可运行驱动 ≥ 120.80CUDA 12.8额外支持 if-else 双分支与 SWITCH 多路分支后两个示例可运行。main开头还通过findCudaDevice定义于 Common/helper_cuda.h选择设备若命令行提供-devicen则使用指定设备否则自动选择 GFlops 最高的设备。十、从源码结构看条件节点的编程范式总结综合五个示例可以把 CUDA 条件节点的使用归纳为两条互补的编程路径它们在实际项目中可以混合使用路径一显式构图cudaGraphCreatecudaGraphAddNode。适合条件体规模已知、依赖关系清晰的场景。要点是先cudaGraphConditionalHandleCreate创建句柄再通过cudaGraphNodeParams的conditional子结构handle、type、size添加条件节点随后向phGraph_out[i]返回的子图句柄填充内容。示例一、四、五采用此路径。路径二流捕获stream capture。适合希望用常规启动语法快速生成节点的场景。cudaStreamBeginCaptureToGraph可以把捕获目标直接指定为条件体的子图而在捕获进行中手动插入条件节点时示例三需要配合cudaStreamGetCaptureInfo获取当前依赖、用cudaStreamUpdateCaptureDependencies修正捕获流依赖确保后续节点正确排在条件节点之后。示例二、三采用此路径。两条路径共享同一套关键 API 约定cudaGraphSetConditional必须在设备端 kernel中调用用于把条件值写入句柄IF 类型size 1时仅真分支size 2时扩展出 else 分支WHILE 类型在每次迭代前评估条件配合cudaGraphCondAssignDefault默认值即可实现 do-while若想要纯 while 语义必须在条件节点上游额外放置一个 kernel 预先设值SWITCH 类型按整数值多路选择条件值超出子图数量则不执行任何分支IF-ELSE 与 SWITCH 两个增强特性依赖较新的驱动≥ 120.80。如果你希望继续深入可以在本仓库内对照阅读同属 3_CUDA_Features 目录的其他图相关示例如simpleCudaGraphs、graphMemoryNodes、graphMemoryFootprint、cudaGraphsPerfScaling它们分别从基础构图、图内存管理与性能伸缩等角度补全 CUDA Graphs 的知识拼图。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考