了解 CUDA 中的动态并行性

Question

动态并行示例：

__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 的块中的一个线程启动子网格。我可以假设父网格中的所有其他线程在我启动子网格之前都已完成执行吗？

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

Answer 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 开始观看。

Answer 2

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