一、简介

1. 综述

动态并行是 CUDA 编程模型的扩展，使 CUDA 内核能够直接在 GPU 上创建新工作并与其同步。 在程序中任何需要的地方动态创建并行性都提供了令人兴奋的功能。

直接从 GPU 创建工作的能力可以减少在主机和设备之间传输执行控制和数据的需要，因为现在可以由在设备上执行的线程在运行时做出启动配置决策。 此外，可以在运行时在内核内内联生成依赖于数据的并行工作，动态利用 GPU 的硬件调度程序和负载平衡器，并根据数据驱动的决策或工作负载进行调整。 以前需要修改以消除递归、不规则循环结构或其他不适合平面、单级并行性的构造的算法和编程模式可以更透明地表达。

本文档描述了支持动态并行性的 CUDA 扩展功能，包括利用这些功能所需的对 CUDA 编程模型的修改和添加，以及利用这种附加功能的指南和最佳实践。

动态并行仅受计算能力 3.5 及更高版本的设备支持。

2. 术语

网格

网格是线程的集合，网格中的线程执行内核函数并被划分为线程块。

线程块

线程块是在同一多处理器 (SM) 上执行的一组线程，线程块中的线程可以访问共享内存并且可以显式同步。

内核函数

内核函数是一个隐式并行子例程，它在 CUDA 执行和内存模型下为网格中的每个线程执行。

主机

主机是指最初调用CUDA的执行环境，通常是在系统的 CPU 处理器上运行的线程。

父

父线程、线程块或网格是启动新网格（子网格）的线程。 在所有启动的子网格也完成之前，父网格才被视为完成。

子

子线程、块或网格是由父网格启动的线程、块或网格。 子网格必须在父线程、线程块或网格被视为完成之前完成。

线程块作用域

具有线程块作用域的对象具有单个线程块的生命周期。 它们仅在由创建对象的线程块中的线程操作时具有定义的行为，并在创建它们的线程块完成时被销毁。

设备运行时

设备运行时是指可用于使内核函数使用动态并行性的运行时系统和 API。

二、执行环境和内存模型

1. 执行环境

CUDA 执行模型基于线程、线程块和网格的原语，其中内核函数定义由线程块和网格内的各个线程执行的程序。 当调用内核函数时，网格的属性由执行配置描述，该执行配置在 CUDA 中具有特殊语法。 CUDA 中对动态并行性的支持扩展了在新网格上配置、启动和隐式同步到设备上运行的线程的能力。

（1）父网格和子网格

配置和启动新网格的设备线程属于父网格，被调用创建的网格是子网格。

父网格在其线程创建的所有子网格全部都完成之后才被视为完成，并且运行时保证父网格和子网格之间的隐式同步。

（2）CUDA原语范围

在主机和设备上，CUDA 运行时都提供了一个 API，用于启动内核并通过流和事件跟踪启动之间的依赖关系。 在主机系统上，启动状态以及引用流和事件的 CUDA 原语由进程内的所有线程共享； 然而进程独立执行并且不能共享 CUDA 对象。

在设备上，启动的内核和 CUDA 对象对网格中的所有线程都是可见的。 例如，这意味着流可以由一个线程创建并由网格中的任何其他线程使用。

（3）同步

CUDA 11.6 中已弃用与父块中的子内核显式同步（即在设备代码中使用 cudaDeviceSynchronize()），并在compute_90+ 编译中删除。 对于计算能力 < 9.0，需要通过指定 -DCUDA_FORCE_CDP1_IF_SUPPORTED 进行编译时选择，才能继续在设备代码中使用 cudaDeviceSynchronize()。 请注意，这预计会在未来的 CUDA 版本中完全删除。