GPU动态并行：父子网格的内存共享与同步

目录

背景

概念

嵌套执行

动态并行的hello world

内存

背景

本章先简要介绍动态并行，后期知识面更广之后，会进一步补充相应的知识

概念

动态并行就是允许在GPU核函数内部，再启动新的核函数。
简单来说就是递归，就是我们之前学过的递归，但是之前学的是普通函数调用普通函数

那么核函数是否可以在内部在调用核函数呢？？？完全可以

在此之前，我们写的所有核函数都是由CPU端调用的。动态并行打破了这一限制，让GPU自己也能“生”出新的任务。

• 父网格/父线程块/父线程：在核函数内部启动新核函数的那个网格、线程块、线程。

• 子网格/子线程块/子线程：被父线程启动的新核函数。

动态并行的另一个好处是等到执行的时候再配置创建多少个网格，多少个块，这样就可以动态的利用GPU硬件调度器和加载平衡器，通过动态调整，来适应负载。并且在内核中启动内核可以减少一部分数据传输消耗。（这里我们稍后会讲解）

嵌套执行

在CPU编程的时候，我们提到过父线程，子线程，子线程是父线程pthread_create()出来的，和父线程共享虚拟地址空间

回到GPU，那么也会有相应父子线程的概念，但是不要弄混了

名词相对多了些，比如父网格，父线程块，父线程，对应的子网格，子线程块，子线程。

• 启动者：是父线程块里的一个具体的线程（父线程）。比如代码里常见 if (tid == 0) { kernel<<<...>>>(...); }，就是让线程0去启动子任务。

• 被启动的：是一个完整的子网格（也叫子核函数）。这个子网格可以有自己的网格大小（gridDim）和线程块大小（blockDim），它内部包含了许多子线程块和子线程。不是只启动一个子线程或子Block，而是启动一整个新的核函数。

如上图

主机启动一个网格（也就是一个内核）-> 此网格（父网格）在执行的过程中启动新的网格（子网格们）->所有子网格们都运行结束后-> 父网格才能结束，否则要等待

如果你调用的时候没有显式同步父子，cuda运行时也会保证有个隐式的同步，让子网格结束父网格才能结束，图中显式的同步了父网格和子网格，通过设置栅栏的方法。

• 父线程块是“等待者”：一个父线程块，必须等待它内部所有线程启动的所有子网格都执行完毕。

• 子网格是“被等待者”：父线程块会一直卡在核函数末尾，不会真正结束，直到它启动的所有子网格任务全部完成。

这个过程是隐式的，不需要你写任何代码。之前学的 cudaMemcpy 会自动等待前面的核函数完成，道理相似。在这里，父网格退出时会自动触发一个隐藏的等待动作，确保所有子网格都已结束。

父线程块启动子网格时，需要“显式的同步”，即不同的线程束需要都执行到子网格调用那一句。

举个例子，同一个父线程块中，线程 0（属于 Warp A）要先启动子网格 A，线程 1（属于 Warp B）要先启动子网格 B。如果 Warp A 先跑到启动点，而 Warp B 还没到，整个父线程块不会立刻行动。系统必须等父线程块里所有要启动子网格的线程都到达他们的启动点，完成同步，然后才一次性地去执行这些子网格。

这意味着，在父线程块内部，所有线程的第一次子网格启动调用，构成了一个隐式的全局集合点。

动态并行的hello world

#include <cuda_runtime.h>
#include <stdio.h>
__global__ void nesthelloworld(int iSize,int iDepth)
{
    unsigned int tid=threadIdx.x;
    printf("depth : %d blockIdx: %d,threadIdx: %d\n",iDepth,blockIdx.x,threadIdx.x);
    if (iSize==1)
        return;
    int nthread=(iSize>>1);
    if (tid==0 && nthread>0)
    {
        nesthelloworld<<<1,nthread>>>(nthread,++iDepth);
        printf("-----------> nested execution depth: %d\n",iDepth);
    }

}

int main(int argc,char* argv[])
{
    int size=64;
    int block_x=2;
    dim3 block(block_x,1,1);
    dim3 grid((size+block.x-1)/block.x,1,1);
    nesthelloworld<<<grid,block>>>(size,0);         //grid(32,1,1)  block(2,1,1) 
    cudaGetLastError();
    cudaDeviceReset();
    return 0;
}

• 每个线程先打印自己的深度、Block 索引和线程索引。

• 检查 iSize 是否等于 1：

  • 是 → 直接返回，不再递归。

  • 否 → 计算子网格的线程数 nthread = iSize / 2。

• 只有 tid == 0 的线程（每个 Block 中的第一个线程）才会启动子网格。

• 子网格配置为 <<<1, nthread>>>：只有 1 个 Block，但包含 nthread 个线程。

• iDepth 加 1 传给下一层，同时打印一条嵌套提示。

关键点：每个父 Block 中只有线程 0 启动子网格，但其他线程（线程 1）什么也不做，最终会等待线程 0 启动的子网格完成（隐式同步）。

层级 (depth)	iSize	启动方	子网格配置	线程总数
0 (父)	64	CPU	grid(32), block(2)	64 线程
1	32	32 个父 Block 中的线程 0	<<<1, 32>>> (每个 Block 启动一个)	32×32 = 1024 线程
2	16	第 1 层子网格中的 32 个 Block 的线程 0	<<<1, 16>>> (每个启动一个)	32×16 = 512 线程
3	8	第 2 层子网格中的 32个 Block 的线程 0	<<<1, 8>>>	32*8=256
4	4	第 3 层子网格中的 32 个 Block 的线程 0	<<<1, 4>>>	32*4=128
5	2	第 4层子网格中的 32 个 Block 的线程 0	<<<1, 2>>>	32*2=64
6	1	无，因为iSize==1了	不启动，直接返回	无

• 父网格（第 0 层）启动后，32 个 Block 中的线程 0 各自启动第 1 层子网格。

• 第 1 层子网格启动后，第 0 层的父 Block 在出口处等待。

• 第 1 层的子网格完成后，第 0 层的父 Block 才真正结束。

• 同样，第 1 层的子网格内部又会启动第 2 层，依此类推。

1. cudaGetLastError()

作用：检查最近一次 CUDA API 调用或核函数启动是否有错误。

• 核函数启动是异步的，<<<>>> 调用本身不会直接返回执行错误。

• cudaGetLastError() 会返回上一个 CUDA 操作（包括核函数启动）的错误码。如果没有错误，返回 cudaSuccess。

• 在这个动态并行的例子中，父核函数内部又启动了子核函数。如果子核函数的启动配置非法（例如线程数超限），错误会被记录下来。cudaGetLastError() 可以捕获这个错误，确保程序不会在不知情的情况下继续运行。

2. cudaDeviceReset()

作用：销毁当前 GPU 上所有已分配的显存资源，重置设备状态，并隐式同步，等待所有未完成的核函数执行完毕。

• 之前学过一个关键点：CPU 启动核函数后是异步的，如果 main 函数直接 return 0，CPU 可能不等 GPU 完成就退出了，导致看不到结果。

• cudaDeviceReset() 会强制 CPU 等待 GPU 上所有任务（包括动态并行的所有父网格和子网格）全部完成后，才会清理资源并返回。

• 因此，它确保了在终端上能看到所有 printf 输出，程序得以优雅结束。

nvcc -arch=sm_89 -lcudadevrt --relocatable-device-code true -o mian main.cu

1. -arch=sm_89

• 作用：指定目标 GPU 架构的计算能力（Compute Capability）。

• 为什么是 sm_89：动态并行是从计算能力 3.5（Kepler 架构）开始支持的硬件特性。低于 3.5 的 GPU 无法使用动态并行。我的 RTX 4060 计算能力是 8.9，远高于 3.5，所以完全支持。写成 -arch=sm_35 是最低兼容写法；如果只在自己显卡上跑，可以写成 -arch=sm_xx 来获得针对你 GPU 的优化。

2. -lcudadevrt

• 作用：链接 CUDA 设备运行时库（libcudadevrt.a）。

• 为什么需要它：常规 CUDA 程序中，核函数从 CPU 端启动，运行时支持（如启动核函数、管理设备内存）由主机端运行时库（libcudart.so）提供。但在动态并行中，你在 GPU 核函数内部调用 <<<>>> 启动子核函数，需要一个运行在 GPU 上的微型运行时来处理子核函数的启动。libcudadevrt.a 就是专门为 GPU 端提供这些基础支持的静态库。不使用动态并行就不需要它。

3. --relocatable-device-code true

• 作用：生成可重定位的设备代码。

• 为什么需要它：常规 CUDA 编译中，每个 .cu 文件的 GPU 代码在编译时就确定了最终地址，不同文件之间不能互相调用设备函数。启用这个选项后，设备代码可以被“重新定位”，支持跨文件的设备端函数调用，以及核函数内部启动其他核函数。这是动态并行的底层基础，让设备代码能够像主机代码一样进行分离编译和链接。

• 可以简写为 -rdc true。

执行结果过长，这里就不粘贴了，感兴趣的可以自己运行一下

内存

最头疼的内存，内存竞争对于普通并行就很麻烦了，现在对于动态并行，更麻烦，主要的有下面几点：

规则一：父子共享全局内存和常量内存

• 全局内存（Global Memory）：通过 cudaMalloc 分配的显存。父网格和它启动的子网格都能读写同一块全局内存。

• 常量内存（Constant Memory）：用 __constant__ 声明的只读内存。父子网格共享同一份常量数据。

这意味着：父子之间可以通过全局内存来传递数据。父网格把数据写入全局内存，子网格去读；或者子网格算完写回，父网格再读。但这也带来了数据竞争的风险，因为两者可以同时访问同一块内存。

---

规则二：父子有不同的局部内存

• 局部内存（Local Memory）：每个线程私有的存储空间，编译器在寄存器不够用时会把一些变量放到这里。

• 父线程和子线程的局部内存是完全独立的，彼此不能访问对方的局部内存。

---

规则三：弱一致性下的并发存取

弱一致性意味着：父子网格对全局内存的读写，不保证一个网格写入的数据立刻对另一个网格可见。这和 CPU 多线程中的内存模型类似——你写了一个值，另一个线程可能暂时看不到最新的值，除非有同步操作。

实用结论：如果你在父子网格之间用全局内存传递数据，不能假设“父写完，子立刻能看到”。必须依靠同步点来保证数据一致性。

---

规则四：两个保证一致的时刻

只有两个时刻，父子网格所见的内存是绝对一致的：

• 子网格启动的时刻：父网格在启动子网格之前写入全局内存的数据，子网格启动后一定能看到。

• 子网格结束的时刻：子网格执行期间写入全局内存的数据，在子网格结束时一定对父网格可见。

在这两个时刻之间（子网格正在运行中），父子网格看到的全局内存内容可能不一致。如果你需要在运行过程中交换数据，必须自己设计同步机制（比如用原子操作或内存栅栏）。

---

规则五：共享内存和局部内存的私有性

这条规则和之前学的普通核函数完全一致：

• 共享内存（__shared__）：同一个线程块内的线程共享。父线程块和子线程块的共享内存是彼此独立的，互不可见。

• 局部内存：每个线程私有，对外不可见。

这意味着：父子之间传递数据的唯一公共通道就是全局内存。不能指望通过共享内存来跨核函数通信。