Dynamic Parallelism: Execute subsequent kernel after parent kernel has finished

Question

I'm developing a cuda kernel which calls itself multiple times (dynamic parallelism). It's necessary that all subsequent kernel calls are executed after the parent kernel has finished. So the algorithm works in turns.

Here's a minimal example: 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

The first subcall is done by putting the kernel in the parent stream. It is not executed at all. The second subcall is done by creating a new stream and waiting for an event e which is recorded when the parent kernel has finished. This subcall is executed immediately.

EDIT:

Streams and events created on the host have undefined behavior when used within any kernel, just as streams and events created by a parent grid have undefined behavior if used within a child grid.

Read more at: http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#ixzz4KngEXdsC

Nevertheless: Is there a way to achieve turn-based dynamic parallelism?

Robert Crovella Robert Crovella · Accepted Answer · 2016-09-20T13:59:24

With CUDA Dynamic Parallelism, strictly speaking, it is not possible for a parent kernel to finish before all of its child kernels have finished.

In your case, the evidence of this is the turnBasedAlgo bar in the profiler which is partially solid (blue) and partially hollow (white). This represents the entire duration of the parent (host-launched) kernel, and the hollow bar portion represents the time during which the parent kernel is nominally waiting while the child kernels spawned from it finish.

The reason for this is documented in the CDP section of the programming guide:

the parent grid is not considered complete until all child grids created by its threads have completed

In practice,however, it is entirely possible for a parent grid to have completed all its "work" (i.e. it is possible for all parent grid threads to have no more instructions to process and therefore effectively "retire") before child grids are finished.

Dynamic Parallelism: Execute subsequent kernel after parent kernel has finished

1 Answers