Notes about CUDA Dynamic Parallelism
CUDA Dynamic Parallelism(也简称DP，动态并行)是CUDA编程模型的一个高级特性，它允许GPU内核（Kernel）在设备端自行创建和启动新的内核，而无需CPU的干预。
为什么引入Dynamic Parallelism？
传统 CUDA 模型要求所有内核（Kernel）必须由 CPU 预先启动，在面对递归算法、自适应计算、不规则数据结构等场景时存在显著局限性。
CUDA Dynamic Parallelism就是来解决这个问题的。
DP基本概念
CUDA DP是CUDA编程模型的一个扩展，CUDA5.0引入的关键特性。允许CUDA内核直接在GPU上创建和同步新的工作（即启动新的内核），实现运行时动态任务生成与调度，从而实现更灵活的并行计算模式。
与静态并行相比，DP 的核心创新在于：
- 任务生成的本地化：任务调度逻辑完全在 GPU 上执行，无需 CPU 介入
- 动态资源分配：根据运行时数据特征实时调整并行粒度
- 算法自然映射：支持递归、条件分支等不规则控制流的并行化
工作原理
- 父内核启动：整个过程仍然始于CPU。CPU启动一个普通的内核（称为父内核）
- 设备端启动：在父内核的执行过程中，GPU上的线程可以根据其计算得到的结果，使用标准的语法来启动一个新的子内核
- 设备端同步：父内核中的线程可以使用
cudaDeviceSynchronize()来等待其启动的所有子内核执行完毕，然后再继续自己的工作 - 层次结构：子内核还可以继续启动自己的子内核，从而形成一个嵌套的、层次化的内核启动结构。CUDA确保了这种嵌套执行的正确性
DP核心优势
- 消除主机端瓶颈: 在传统模式中，每次内核启动需经历 “CPU 准备参数→PCIe 传输→GPU 执行” 的固定流程。DP 将任务启动逻辑迁移至设备端，避免了频繁的主机 - 设备同步。例如，在递归计算中，子内核启动延迟从 PCIe 级（微秒级）降至 GPU 内核内函数调用级（纳秒级）。
- 细粒度任务调度: 线程可根据局部数据特征动态决定是否启动子任务。如在自适应网格计算中，仅当网格单元误差超过阈值时才启动子内核进行细化，实现计算资源的精准分配。
- 支持复杂算法结构: 传统CUDA难以高效实现的递归算法（如快速傅里叶变换、光线追踪的光线分叉），可通过DP直接映射为内核层级调用，避免繁琐的迭代化改造。
总结
CUDA Dynamic Parallelism 是一项强大的技术，它通过允许GPU内核自主启动新工作，深化了GPU的独立计算能力。它主要解决了减少CPU-GPU通信瓶颈和简化复杂算法实现的问题，但随之而来的是需要开发者对开销和复杂性进行仔细权衡。在合适的场景下使用，它能带来显著的性能提升和编程便利性。
参考资料:
- NVIDIA | BaM与CUDA Dynamic Parallelism的核心区别
- deepseek prompt:简要介绍下cuda的Dynamic Parallelism技术