动态并行（Dynamic Parallelism）：在GPU上启动新内核-平芜编程栈

动态并行（Dynamic Parallelism）是 CUDA 5.0 引入的一项功能，它允许一个正在 GPU 上运行的 Kernel（称为父 Kernel）启动一个新的 Kernel（称为子 Kernel）。

1. 动态并行的核心概念

1.1 传统的 CUDA 启动模式（Host-Side Launch）

在传统的 CUDA 模型中，所有 Kernel 启动都必须由 **Host（CPU）**执行。CPU 负责配置参数、调用启动语法<<<>>>，并将任务推送到 GPU。这种模式是：

Host→Device\text{Host} \rightarrow \text{Device}Host→Device

1.2 动态并行的启动模式（Device-Side Launch）

动态并行打破了这一限制，允许 Kernel 在 GPU 上启动其他 Kernel。这创建了一个 Kernel 启动的层次结构：

Host→Device (Parent Kernel)→Device (Child Kernel)\text{Host} \rightarrow \text{Device (Parent Kernel)} \rightarrow \text{Device (Child Kernel)}Host→Device (Parent Kernel)→Device (Child Kernel)

Parent Kernel (父 Kernel)：在 GPU 上执行，并包含启动子 Kernel 的代码。
Child Kernel (子 Kernel)：被父 Kernel 启动，也在 GPU 上执行。

1.3 优势：灵活的并行与工作分解

动态并行的主要优势在于：

自适应工作分配：父 Kernel 可以根据运行时的数据和中间结果，动态地决定如何分解后续工作，以及启动多少个线程块和线程。
深层次递归：它可以方便地实现并行递归算法，如树遍历、分治算法（Divide-and-Conquer），而无需频繁地回到 CPU 进行调度。
减少 CPU/GPU 交互开销：避免了父 Kernel 必须等待 CPU 同步并重新启动子任务的开销，从而减少了 PCI Express 总线的流量和 CPU 的负担。

2. 编程实践与语法

使用动态并行在 Device 上启动子 Kernel 的语法与在 Host 上启动 Kernel 的语法完全相同。

2.1 启用与编译

要在 Kernel 中使用动态并行，你必须在编译时启用它：

# 必须使用 -rdc=true (Relocatable Device Code) 链接选项# 以及 -lcudadevrt 库 (Device Runtime Library)nvcc my_file.cu -arch=sm_xx -rdc=true -lcudadevrt -o my_app

2.2 Device-Side 启动 Kernel 代码

在父 Kernel 内部，你可以像在主机代码中一样使用<<<>>>语法：

// 子 Kernel 函数 (也需要使用 __global__ 修饰符) __global__ void childKernel(float* data, int size) { int i = threadIdx.x + blockIdx.x * blockDim.x; if (i < size) { data[i] *= 2.0f; } } // 父 Kernel 函数 __global__ void parentKernel(float* data, int N) { int i = threadIdx.x + blockIdx.x * blockDim.x; if (i == 0) { // 假设只有线程 0 启动子 Kernel // 1. 动态配置启动参数 int threadsPerBlock = 256; int numBlocks = (N + threadsPerBlock - 1) / threadsPerBlock; // 2. Device-Side 启动子 Kernel (语法与 Host 端启动完全相同) printf("Parent Kernel (Thread %d) is launching Child Kernel...\n", i); childKernel<<<numBlocks, threadsPerBlock>>>(data, N); // 3. 隐式同步点：设备端启动默认是同步的 // 在不使用流的情况下，子 Kernel 必须在父 Kernel 继续执行之前完成。 } } // Host 端代码：只需要启动父 Kernel int main() { // ... 分配内存，设置 N ... parentKernel<<<numBlocks, threadsPerBlock>>>(d_data, N); // ... 拷贝结果回 Host ... }

3. 同步与执行模型

动态并行在同步方面与 Host 启动有显著差异，主要是为了确保父子任务的正确执行顺序。

3.1 默认同步行为

关键点：由父 Kernel 启动的子 Kernel 默认是同步的。

父 Kernel 在启动子 Kernel 后会暂停执行，直到子 Kernel 及其所有操作完成。
一旦子 Kernel 完成，控制权才会返回给父 Kernel，父 Kernel 继续执行后续代码。

这种默认的同步行为简化了编程，因为它避免了复杂的父子任务依赖管理。

3.2 异步启动与设备流（Device Streams）

父 Kernel 也可以使用 **设备流（Device Streams）**来实现子 Kernel 的异步启动，类似于 Host Streams：

创建流：使用cudaStream_t child_stream;和cudaStreamCreateWithFlags(&child_stream, cudaStreamNonBlocking);在设备端创建流。
异步启动：在启动子 Kernel 时将流作为第四个参数传入：childKernel<<<numBlocks, threadsPerBlock, 0, child_stream>>> (d_data, N);
同步：父 Kernel 可以使用cudaStreamSynchronize(child_stream);来显式等待子 Kernel 完成。

3.3 限制

嵌套深度：动态并行允许递归或嵌套启动，但为了防止无限递归和资源耗尽，通常建议限制嵌套深度。
设备内存：子 Kernel 共享父 Kernel 的全局内存空间。
不支持的操作：父 Kernel 不能启动新的 Host 线程或执行某些 Host 端的 API 调用。

4. 总结与适用场景

动态并行提供了一种强大的方式来表达那些在运行时才能确定其并行度或工作分解方式的算法。

特性	动态并行 (Device Launch)	传统并行 (Host Launch)
启动者	GPU 线程 (Parent Kernel)	CPU 线程 (Host Code)
调度依赖	依赖于运行时数据和 Kernel 逻辑。	依赖于 CPU 侧的程序控制流。
开销	低（无需 PCI-E 传输或 Host API 开销）。	高（每次启动都需要 Host CPU 参与）。
同步	默认是同步的，易于管理。	默认是异步的，需要`cudaDeviceSynchronize`。
最佳用途	并行递归、树结构遍历、自适应网格细化等。	大规模、规则的矩阵/向量运算。