1

动态并行示例：

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

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

4

2 回答 2

4

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

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

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

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

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

  • 如果父网格启动子网格但未加入，则所有父块可能在子块被调度之前完成。
  • 如果父网格加入，那么所有子网格很可能在父块完成之前完成。
  • 如果父网格启动超过 GPU 可以同时执行的，那么答案就在中间。

Nsight VSE CUDA Trace 和 Visual Profiler 具有用于跟踪 CDP 网格的额外可视化工具。GTC 2013 演示文稿使用 NVIDIA Nsight Visual Studio 版本分析和优化 CUDA 内核代码的视频（但不是幻灯片）提供了有关 CDP 可视化的最佳文档。17:15 开始观看。

于 2015-06-12T01:41:18.980 回答
0

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

于 2015-06-11T11:35:44.443 回答