动态并行：父内核完成后执行后续内核答案

【问题标题】：Dynamic Parallelism: Execute subsequent kernel after parent kernel has finished动态并行：父内核完成后执行后续内核
【发布时间】：2017-01-28 04:55:54
【问题描述】：

我正在开发一个可多次调用自身的 cuda 内核（动态并行）。所有后续内核调用都必须在父内核完成后执行。所以算法是轮流工作的。

这是一个最小的例子：EDIT

$ cat turn.cu 
#include <stdio.h>

__global__ void turnBasedAlgo(int depth, cudaStream_t stream) {
    if (depth < 3) {
        printf("depth: %d\n", depth);
        cudaEvent_t e;
        cudaEventCreateWithFlags(&e, cudaEventDisableTiming);
        cudaEventRecord(e, stream);

        turnBasedAlgo<<<1,1,0,stream>>>(depth+1, stream);

        cudaStream_t s2;
        cudaStreamCreateWithFlags(&s2, cudaStreamNonBlocking);
        cudaStreamWaitEvent(s2, e, 0);
        turnBasedAlgo<<<1,1,0,s2>>>(depth+1, s2);

        // some work
        clock_t start = clock();
        clock_t end = clock();
        while (end - start < 100000) {
            end = clock();
        }
    }
}

int main(int argc, char **argv) {
    cudaStream_t s;
    cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);
    turnBasedAlgo<<<1,1,0,s>>>(0, s);
    cudaDeviceSynchronize();
    cudaDeviceReset();

    return 0;
}
$ nvcc -arch=sm_35 -dc turn.cu && nvcc -arch=sm_35 -link -o turn turn.o
$ ./turn
depth: 0
depth: 1
depth: 2
$ nvvp ./turn

第一个子调用是通过将内核放入父流中来完成的。它根本不执行。第二个子调用是通过创建一个新流并等待父内核完成时记录的事件e 来完成的。这个子调用会立即执行。

编辑：

在主机上创建的流和事件在任何内核中使用时具有未定义的行为，就像在子网格中使用时由父网格创建的流和事件具有未定义的行为一样。

不过：有没有办法实现回合制动态并行？

【问题讨论】：

请始终提供minimal reproducible example

标签： cuda

【解决方案1】：

对于 CUDA 动态并行，严格来说，父内核不可能在其所有子内核完成之前完成。

在您的情况下，这方面的证据是分析器中的turnBasedAlgo 条，它部分实心（蓝色）和部分空心（白色）。这表示父（主机启动）内核的整个持续时间，而空心条部分表示父内核名义上等待的时间，而子内核从它产生的子内核完成。

原因记录在the programming guide 的 CDP 部分：

在其线程创建的所有子网格都完成之前，父网格不会被认为是完整的

然而，在实践中，父网格完全有可能在子网格之前完成其所有“工作”（即所有父网格线程可能没有更多的指令来处理，因此有效地“退休”）网格已完成。

【讨论】：