了解 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

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

Answer 2

Say I launch a child grid from one thread in a block at threadIdx.x==0. Can I assume that all other threads in the parent grid have finished executing up to the point I launched the child grid as well?

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

If so, how does this work? All I'm reading is that a parent grid is not technically finished before a child grid. Nothing about guarantees of other parent threads that have not launched children.

当父线程启动子网格时，它会将工作推送到优先级高于自身的 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.

开始观看

了解 CUDA 中的动态并行性

Understanding Dynamic Parallelism in CUDA

cuda

dynamic-parallelism