了解 CUDA 中的动态并行性

Posted 2023-04-15

【中文标题】了解 CUDA 中的动态并行性

【英文标题】：Understanding Dynamic Parallelism in CUDA

【发布时间】：2015-08-27 01:40:08

【问题描述】：

动态并行示例：

__global__ void nestedHelloWorld(int const iSize,int iDepth) 
    int tid = threadIdx.x;
    printf("Recursion=%d: Hello World from thread %d" "block %d\n",iDepth,tid,blockIdx.x);
    // condition to stop recursive execution
    if (iSize == 1) return;
    // reduce block size to half
    int nthreads = iSize>>1;
    // thread 0 launches child grid recursively
    if(tid == 0 && nthreads > 0) 
        nestedHelloWorld<<<1, nthreads>>>(nthreads,++iDepth);
        printf("-------> nested execution depth: %d\n",iDepth);
    

用一个块打印，用两个块整个父网格已经完成：

./nestedHelloWorld Execution Configuration: grid 1 block 8
Recursion=0: Hello World from thread 0 block 0
Recursion=0: Hello World from thread 1 block 0
Recursion=0: Hello World from thread 2 block 0
Recursion=0: Hello World from thread 3 block 0
Recursion=0: Hello World from thread 4 block 0
Recursion=0: Hello World from thread 5 block 0
Recursion=0: Hello World from thread 6 block 0
Recursion=0: Hello World from thread 7 block 0
-------> nested execution depth: 1
Recursion=1: Hello World from thread 0 block 0
Recursion=1: Hello World from thread 1 block 0
Recursion=1: Hello World from thread 2 block 0
Recursion=1: Hello World from thread 3 block 0
-------> nested execution depth: 2
Recursion=2: Hello World from thread 0 block 0
Recursion=2: Hello World from thread 1 block 0
-------> nested execution depth: 3
Recursion=3: Hello World from thread 0 block 0

假设我从 threadIdx.x==0 的块中的一个线程启动子网格。我可以假设父网格中的所有其他线程在我启动子网格之前都已完成执行吗？

如果是这样，它是如何工作的？我正在阅读的是，从技术上讲，父网格在子网格之前并未完成。没有关于没有启动子线程的其他父线程的保证。

【参考方案1】：

假设我从一个块中的一个线程启动一个子网格 threadIdx.x==0。我可以假设父网格中的所有其他线程 已经完成执行，直到我启动子网格为 好吗？

没有。您不能对父块中的其他线程或父网格中的其他块的状态做出任何假设。

如果是这样，它是如何工作的？我正在阅读的是父网格是 在子网格之前没有在技术上完成。与保证无关 其他尚未启动子线程的父线程。

当父线程启动子网格时，它会将工作以比自身更高的优先级推送到 GPU。在计算能力 3.5 - 5.x 上，GPU 将安排最高优先级的工作，但不会抢占任何正在运行的块。如果 GPU 已满，则计算工作分配将无法调度子块。当父块完成时，子块将在任何新的父块之前分发。此时设计仍可能死锁。如果启动工作的块执行连接操作（cudaDeviceSynchronize）并且如果子工作由于没有足够的空间来安排子工作或仍在运行而尚未完成，则父块（不是网格）将预先清空自己。这允许子网格向前发展。当子网格完成后，CDP 调度程序将恢复父块。

在父网格的所有块都完成并且所有子网格都完成之前，父网格不会被标记为已完成。

如果父网格启动子网格但未加入，则 可能所有父块都在子块之前完成 预定。 如果父网格加入，那么很可能所有 子网格在父块完成之前完成。 如果父网格启动超过了 GPU 那么答案就在中间。

Nsight VSE CUDA Trace 和 Visual Profiler 具有用于跟踪 CDP 网格的额外可视化工具。 GTC 2013 演示文稿Profiling and Optimizing CUDA Kernel Code with NVIDIA Nsight Visual Studio Edition 中的视频（但不是幻灯片）提供了有关 CDP 可视化的最佳文档。 17:15 开始观看。

【讨论】：

感谢您的详细解释。您能否详细说明父网格“加入”子网格的含义以及为此调用了哪些代码？我在某处读到，如果没有明确说明，它是隐含的，但一切仍然不是很清楚。

Fork-join 是并行编程中的常用模型。在 CDP 中，fork 是子启动，join 是子同步原语。在上述情况下，可以使用 cudaDeviceSynchronize() 来等待孩子的工作完成。

【参考方案2】：

没有。 warp 中的所有线程都以锁步执行，因此如果线程 0 尚未完成，则线程 [1..31] 也没有。块中的其他线程（或warp）可能已完成执行，也可能尚未完成。

【讨论】：

我明白这一点。但是当我将它增加到 2 个或更多时，两者都会在孩子之前打印。这只是随机的吗？见这里：turkpaylasim.com/cevahir/2015/04/20/cudada-kernel-icinde-kernel

我的土耳其语不是很好，但据我所知，您的输出与交错的每个块中的单个块相同。在每个“根”块内，线程按预期运行。在每个子块中（所有子块的 id 均为 0），线程仍处于锁步状态。

哈哈。我也不会说土耳其语，只是想展示 2 个街区的例子。不知道我是否理解你所说的正确。这是否意味着整个父网格将在子网格启动之前完成执行？根据你说的，没有。我可以期望根中的块 0 可能在块 1 中的子 0 之后执行吗？

我应该引用talonmies 并说“您还应该知道编译器和汇编器确实执行指令重新排序......”，再加上我不相信顺序的事实printfs 线程之间。