CUDA 动态并行CDP实战入门解析 cdpSimplePrint 示例的 GPU 递归内核启动【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samplesCUDA Dynamic ParallelismCDP允许运行在 GPU 上的线程直接启动新的内核而无需 CPU 参与是构建自适应、递归型并行算法如快速排序、四叉树、体素细分的核心机制。本文以 cuda-samples 仓库中的 cdpSimplePrint 示例为主线从源码逐行剖析 CDP 的递归内核启动、块唯一标识生成与深度控制逻辑并结合 CMake 构建配置说明如何在满足 SM 3.5 的 GPU 上编译运行该示例帮助读者快速掌握 CDP 的基本编程范式。示例概述与核心技术点根据 cdpSimplePrint 官方 README该示例通过CUDA Dynamic Parallelism实现了一个基于 printf 的简单递归输出程序CPU 首先启动 2 个块、每块 2 个线程的内核随后 GPU 上的每个线程再递归地启动 2 个块 × 2 个线程的子内核直到达到用户指定的最大递归深度max_depth。最终程序会以缩进树的形式打印出每一层块的编号及其父块信息直观展示 CDP 下内核层级启动的调用关系。本示例所属的3_CUDA_Features目录在主 README.md 中被定位为 Samples that demonstrate CUDA Features与 Cooperative Groups、CUDA Graphs 等特性并列属于 CUDA 高级特性演示序列。示例涉及的核心概念与 API 如下表所示类别具体内容关键技术CUDA Dynamic ParallelismCDPCUDA Runtime APIcudaDeviceSynchronize、cudaGetLastError、cudaGetDeviceProperties、cudaDeviceSetLimit辅助工具头文件helper_cuda.h、helper_string.h其中cudaDeviceSetLimit用于调整 CDP 运行时设备端资源限制如cudaLimitDevRuntimeSyncDepth同步深度、cudaLimitDevRuntimePendingLaunchCount待处理启动数量在同目录的 cdpSimpleQuicksort/README.md 和 cdpQuadtree/README.md 中同样被列为关键 API。硬件与软件前置条件计算能力要求原文档明确说明该示例要求计算能力Compute Capability3.5 或更高这也是 CUDA Dynamic Parallelism 的最低硬件门槛。主仓库 README.md 对此特性的描述为CDP (CUDA Dynamic Parallelism) allows kernels to be launched from threads running on the GPU. CDP is only available on GPUs with SM architecture of 3.5 or above.这一约束在源码中也做了显式校验。cdpSimplePrint.cu 中程序在findCudaDevice选定设备后通过cudaGetDeviceProperties查询设备属性并检查if (!(deviceProp.major 3 || (deviceProp.major 3 deviceProp.minor 5))) { printf(GPU %d - %s does not support CUDA Dynamic Parallelism\n Exiting., device, deviceProp.name); exit(EXIT_WAIVED); }若不满足 SM 3.5 条件程序会以EXIT_WAIVED状态退出即测试豁免这是 cuda-samples 中针对特性缺失的统一处理模式。支持的平台维度支持范围操作系统Linux、WindowsCPU 架构x86_64、armv7l依赖与前置安装示例的 README.md 将其运行依赖标注为 CDP 特性该特性由 CUDA Toolkit / CUDA Driver 提供。前置条件是下载并安装与平台对应的 CUDA Toolkit并确认系统满足上述 CDP 依赖。源码级解析CDP 递归打印的实现原理完整的实现位于 cdpSimplePrint.cu全文件约 160 行核心逻辑可分为三个部分全局唯一 ID 生成、块信息打印、递归内核启动。1. 全局块 ID 生成器源码第 37 行定义了一个设备端全局计数器用于为每个块生成唯一标识__device__ int g_uids 0;在内核中每个块由 thread 0 通过原子自增获取唯一 ID并存入共享内存供块内所有线程使用__shared__ int s_uid; if (threadIdx.x 0) { s_uid atomicAdd(g_uids, 1); } __syncthreads();这里体现了两个 CDP 编程要点设备端全局变量g_uids的生命周期跨越整个程序可被各级递归内核共享访问利用atomicAdd保证并发生成的 ID 互不冲突__syncthreads()确保其余线程在打印前已同步拿到s_uid。2. 层级化信息打印print_infoprint_info辅助函数第 42-62 行仅由每个块的 thread 0 执行输出按递归深度depth生成缩进前缀if (depth 0) printf(BLOCK %d launched by the host\n, uid); else { char buffer[32]; for (int i 0; i depth; i) { buffer[3 * i 0] |; buffer[3 * i 1] ; buffer[3 * i 2] ; } buffer[3 * depth] \0; printf(%sBLOCK %d launched by thread %d of block %d\n, buffer, uid, thread, parent_uid); }输出结果以|缩进直观呈现递归层级顶层块由 host 启动深层块则标明由哪个线程thread从哪个父块parent_uid启动这正是 CDP 嵌套启动关系的可视化表达。3. 递归内核启动CDP 核心主内核cdp_kernel第 71-92 行在完成 ID 分配与打印后若未达到max_depth则直接从 GPU 上启动新的子内核__global__ void cdp_kernel(int max_depth, int depth, int thread, int parent_uid) { ... if (depth max_depth) { return; } cdp_kernelgridDim.x, blockDim.x(max_depth, depth, threadIdx.x, s_uid); }这段代码浓缩了 CDP 的两个关键特征设备端启动cdp_kernel...出现在设备函数内部编译与运行时均由 CUDA 动态并行机制处理层级参数传递子内核携带当前线程号threadIdx.x与父块 IDs_uid作为thread与parent_uid形成完整的家族树信息depth每层递增控制递归深度网格与块维度沿用父内核的gridDim.x、blockDim.x。4. 主机端启动与收尾main函数第 97-160 行完成命令行解析、设备选择、内核启动与错误检查cdp_kernel2, 2(max_depth, 0, 0, -1); checkCudaErrors(cudaGetLastError()); checkCudaErrors(cudaDeviceSynchronize());初始启动配置为2 个块 × 2 个线程parent_uid传-1表示没有父块cudaGetLastError()捕获启动阶段的异步错误cudaDeviceSynchronize()等待包括所有设备端递归子内核在内的全部工作完成——这是 CDP 程序的必要收尾步骤因为设备端启动的内核同样归属当前流。程序启动时会打印预期块数量。以max_depth2为例CPU 启动 2 块每块 2 线程每个线程再启动 2 块共246块递归更深时块数按×4增长2, 8, 32, 128...。命令行参数与运行示例程序支持两个命令行参数参数含义默认值取值范围depthmax_depth递归最大深度21 ~ 8传入-h/-help时打印用法说明并退出源码第 104-109 行depth超出 [1, 8] 范围时程序提示depth parameter has to be between 1 and 8并以失败码退出第 111-118 行深度上限 8 意味着最多启动2 8 32 ... 2^7级别的块避免递归过度膨胀。典型运行方式假设已构建生成可执行文件# 默认深度 2 ./cdpSimplePrint # 指定递归深度为 3 ./cdpSimplePrint depth3运行后将看到以|缩进输出的块树例如顶层输出BLOCK 0 launched by the host深层输出| BLOCK 1 launched by thread 0 of block 0等直观展示 CDP 的多级启动结构。构建方式示例目录 cdpSimplePrint 使用 CMake 构建其 CMakeLists.txt 要点如下project(cdpSimplePrint LANGUAGES C CXX CUDA)声明 CUDA 语言支持并通过find_package(CUDAToolkit REQUIRED)定位 Toolkit根据目标架构自动设置CMAKE_CUDA_ARCHITECTURESaarch64TegraToolkit 13.0下为87 110其余平台为75 80 86 89 90 100 110 120开启CMAKE_CUDA_FLAGS -Wno-deprecated-gpu-targets调试构建ENABLE_CUDA_DEBUG附加-G否则附加-lineinfo关键点通过set_target_properties(cdpSimplePrint PROPERTIES CUDA_SEPARABLE_COMPILATION ON)开启可分离编译并在编译选项中加入--extended-lambda、C17 / CUDA 17 标准——CDP 的设备端启动依赖于可分离编译与设备链接device-link阶段通过include_directories(../../../Common)引入仓库公共头文件如helper_cuda.h、helper_string.h最终由InstallSamples.cmake生成安装规则。从仓库根目录构建整个 CUDA Features 集或单独构建该示例时可使用标准的 CMake 流程cmake -S . -B build cmake --build build --target cdpSimplePrint延伸阅读仓库中的其他 CDP 示例3_CUDA_Features目录下还包含多个以 CDP 为技术基础的进阶示例可作为本示例的延伸学习材料cdpSimpleQuicksort基于 CDP 的 GPU 快速排序涉及cudaDeviceSetLimit调整运行时资源上限cdpAdvancedQuicksort高级版快速排序其源码中使用cudaDeviceSetLimit(cudaLimitDevRuntimePendingLaunchCount, 4096)提升设备端待启动内核容量见 cdpAdvancedQuicksort.cucdpQuadtree基于 CDP 的四叉树构建cdpBezierTessellation贝塞尔曲面细分体现 CDP 在图形学中的应用。小结cdpSimplePrint 以最小可读代码演示了 CDP 的完整编程闭环设备端原子计数生成唯一块 ID、层级化 printf 输出调用树、以及通过...在 GPU 上递归启动子内核。掌握这一模式后读者可以进一步阅读 cdpSimpleQuicksort、cdpAdvancedQuicksort 等示例将 CDP 应用于真正的分治算法与自适应计算场景。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
CUDA 动态并行（CDP）实战入门：解析 cdpSimplePrint 示例的 GPU 递归内核启动
CUDA 动态并行CDP实战入门解析 cdpSimplePrint 示例的 GPU 递归内核启动【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samplesCUDA Dynamic ParallelismCDP允许运行在 GPU 上的线程直接启动新的内核而无需 CPU 参与是构建自适应、递归型并行算法如快速排序、四叉树、体素细分的核心机制。本文以 cuda-samples 仓库中的 cdpSimplePrint 示例为主线从源码逐行剖析 CDP 的递归内核启动、块唯一标识生成与深度控制逻辑并结合 CMake 构建配置说明如何在满足 SM 3.5 的 GPU 上编译运行该示例帮助读者快速掌握 CDP 的基本编程范式。示例概述与核心技术点根据 cdpSimplePrint 官方 README该示例通过CUDA Dynamic Parallelism实现了一个基于 printf 的简单递归输出程序CPU 首先启动 2 个块、每块 2 个线程的内核随后 GPU 上的每个线程再递归地启动 2 个块 × 2 个线程的子内核直到达到用户指定的最大递归深度max_depth。最终程序会以缩进树的形式打印出每一层块的编号及其父块信息直观展示 CDP 下内核层级启动的调用关系。本示例所属的3_CUDA_Features目录在主 README.md 中被定位为 Samples that demonstrate CUDA Features与 Cooperative Groups、CUDA Graphs 等特性并列属于 CUDA 高级特性演示序列。示例涉及的核心概念与 API 如下表所示类别具体内容关键技术CUDA Dynamic ParallelismCDPCUDA Runtime APIcudaDeviceSynchronize、cudaGetLastError、cudaGetDeviceProperties、cudaDeviceSetLimit辅助工具头文件helper_cuda.h、helper_string.h其中cudaDeviceSetLimit用于调整 CDP 运行时设备端资源限制如cudaLimitDevRuntimeSyncDepth同步深度、cudaLimitDevRuntimePendingLaunchCount待处理启动数量在同目录的 cdpSimpleQuicksort/README.md 和 cdpQuadtree/README.md 中同样被列为关键 API。硬件与软件前置条件计算能力要求原文档明确说明该示例要求计算能力Compute Capability3.5 或更高这也是 CUDA Dynamic Parallelism 的最低硬件门槛。主仓库 README.md 对此特性的描述为CDP (CUDA Dynamic Parallelism) allows kernels to be launched from threads running on the GPU. CDP is only available on GPUs with SM architecture of 3.5 or above.这一约束在源码中也做了显式校验。cdpSimplePrint.cu 中程序在findCudaDevice选定设备后通过cudaGetDeviceProperties查询设备属性并检查if (!(deviceProp.major 3 || (deviceProp.major 3 deviceProp.minor 5))) { printf(GPU %d - %s does not support CUDA Dynamic Parallelism\n Exiting., device, deviceProp.name); exit(EXIT_WAIVED); }若不满足 SM 3.5 条件程序会以EXIT_WAIVED状态退出即测试豁免这是 cuda-samples 中针对特性缺失的统一处理模式。支持的平台维度支持范围操作系统Linux、WindowsCPU 架构x86_64、armv7l依赖与前置安装示例的 README.md 将其运行依赖标注为 CDP 特性该特性由 CUDA Toolkit / CUDA Driver 提供。前置条件是下载并安装与平台对应的 CUDA Toolkit并确认系统满足上述 CDP 依赖。源码级解析CDP 递归打印的实现原理完整的实现位于 cdpSimplePrint.cu全文件约 160 行核心逻辑可分为三个部分全局唯一 ID 生成、块信息打印、递归内核启动。1. 全局块 ID 生成器源码第 37 行定义了一个设备端全局计数器用于为每个块生成唯一标识__device__ int g_uids 0;在内核中每个块由 thread 0 通过原子自增获取唯一 ID并存入共享内存供块内所有线程使用__shared__ int s_uid; if (threadIdx.x 0) { s_uid atomicAdd(g_uids, 1); } __syncthreads();这里体现了两个 CDP 编程要点设备端全局变量g_uids的生命周期跨越整个程序可被各级递归内核共享访问利用atomicAdd保证并发生成的 ID 互不冲突__syncthreads()确保其余线程在打印前已同步拿到s_uid。2. 层级化信息打印print_infoprint_info辅助函数第 42-62 行仅由每个块的 thread 0 执行输出按递归深度depth生成缩进前缀if (depth 0) printf(BLOCK %d launched by the host\n, uid); else { char buffer[32]; for (int i 0; i depth; i) { buffer[3 * i 0] |; buffer[3 * i 1] ; buffer[3 * i 2] ; } buffer[3 * depth] \0; printf(%sBLOCK %d launched by thread %d of block %d\n, buffer, uid, thread, parent_uid); }输出结果以|缩进直观呈现递归层级顶层块由 host 启动深层块则标明由哪个线程thread从哪个父块parent_uid启动这正是 CDP 嵌套启动关系的可视化表达。3. 递归内核启动CDP 核心主内核cdp_kernel第 71-92 行在完成 ID 分配与打印后若未达到max_depth则直接从 GPU 上启动新的子内核__global__ void cdp_kernel(int max_depth, int depth, int thread, int parent_uid) { ... if (depth max_depth) { return; } cdp_kernelgridDim.x, blockDim.x(max_depth, depth, threadIdx.x, s_uid); }这段代码浓缩了 CDP 的两个关键特征设备端启动cdp_kernel...出现在设备函数内部编译与运行时均由 CUDA 动态并行机制处理层级参数传递子内核携带当前线程号threadIdx.x与父块 IDs_uid作为thread与parent_uid形成完整的家族树信息depth每层递增控制递归深度网格与块维度沿用父内核的gridDim.x、blockDim.x。4. 主机端启动与收尾main函数第 97-160 行完成命令行解析、设备选择、内核启动与错误检查cdp_kernel2, 2(max_depth, 0, 0, -1); checkCudaErrors(cudaGetLastError()); checkCudaErrors(cudaDeviceSynchronize());初始启动配置为2 个块 × 2 个线程parent_uid传-1表示没有父块cudaGetLastError()捕获启动阶段的异步错误cudaDeviceSynchronize()等待包括所有设备端递归子内核在内的全部工作完成——这是 CDP 程序的必要收尾步骤因为设备端启动的内核同样归属当前流。程序启动时会打印预期块数量。以max_depth2为例CPU 启动 2 块每块 2 线程每个线程再启动 2 块共246块递归更深时块数按×4增长2, 8, 32, 128...。命令行参数与运行示例程序支持两个命令行参数参数含义默认值取值范围depthmax_depth递归最大深度21 ~ 8传入-h/-help时打印用法说明并退出源码第 104-109 行depth超出 [1, 8] 范围时程序提示depth parameter has to be between 1 and 8并以失败码退出第 111-118 行深度上限 8 意味着最多启动2 8 32 ... 2^7级别的块避免递归过度膨胀。典型运行方式假设已构建生成可执行文件# 默认深度 2 ./cdpSimplePrint # 指定递归深度为 3 ./cdpSimplePrint depth3运行后将看到以|缩进输出的块树例如顶层输出BLOCK 0 launched by the host深层输出| BLOCK 1 launched by thread 0 of block 0等直观展示 CDP 的多级启动结构。构建方式示例目录 cdpSimplePrint 使用 CMake 构建其 CMakeLists.txt 要点如下project(cdpSimplePrint LANGUAGES C CXX CUDA)声明 CUDA 语言支持并通过find_package(CUDAToolkit REQUIRED)定位 Toolkit根据目标架构自动设置CMAKE_CUDA_ARCHITECTURESaarch64TegraToolkit 13.0下为87 110其余平台为75 80 86 89 90 100 110 120开启CMAKE_CUDA_FLAGS -Wno-deprecated-gpu-targets调试构建ENABLE_CUDA_DEBUG附加-G否则附加-lineinfo关键点通过set_target_properties(cdpSimplePrint PROPERTIES CUDA_SEPARABLE_COMPILATION ON)开启可分离编译并在编译选项中加入--extended-lambda、C17 / CUDA 17 标准——CDP 的设备端启动依赖于可分离编译与设备链接device-link阶段通过include_directories(../../../Common)引入仓库公共头文件如helper_cuda.h、helper_string.h最终由InstallSamples.cmake生成安装规则。从仓库根目录构建整个 CUDA Features 集或单独构建该示例时可使用标准的 CMake 流程cmake -S . -B build cmake --build build --target cdpSimplePrint延伸阅读仓库中的其他 CDP 示例3_CUDA_Features目录下还包含多个以 CDP 为技术基础的进阶示例可作为本示例的延伸学习材料cdpSimpleQuicksort基于 CDP 的 GPU 快速排序涉及cudaDeviceSetLimit调整运行时资源上限cdpAdvancedQuicksort高级版快速排序其源码中使用cudaDeviceSetLimit(cudaLimitDevRuntimePendingLaunchCount, 4096)提升设备端待启动内核容量见 cdpAdvancedQuicksort.cucdpQuadtree基于 CDP 的四叉树构建cdpBezierTessellation贝塞尔曲面细分体现 CDP 在图形学中的应用。小结cdpSimplePrint 以最小可读代码演示了 CDP 的完整编程闭环设备端原子计数生成唯一块 ID、层级化 printf 输出调用树、以及通过...在 GPU 上递归启动子内核。掌握这一模式后读者可以进一步阅读 cdpSimpleQuicksort、cdpAdvancedQuicksort 等示例将 CDP 应用于真正的分治算法与自适应计算场景。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
