Почему некоторые дочерние сетки ядра не возвращаются в родительскую сетку мгновенно?

Question

Я пытаюсь использовать динамический параллелизм, запуская ядро ​​<<<1,1>>> ("Контроллер"), которое будет запускать другие вычислительные сетки в зависимости от их результата. Контроллеру нужны результаты вычислений для определения следующего ядра, которое будет запущено.

Я жду результатов ядра с cudaDeviceSynchronize() и Контроллер делает паузы, пока они все не будут сделаны.

Проблема в том, что в большинстве случаев дочерняя сетка долго заканчивалась, но родительский поток не будет продолжаться в течение достаточно долгого времени. Я видел это на временной шкале, созданной Nsight.

Это тестовая программа, которую я сделал, чтобы показать это поведение: (Вам нужно только понять "Контроллер" ядра, остальные два просто вращаются вхолостую, чтобы симулировать вычисления)

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

__global__ void ComputeKernel0();
__global__ void ComputeKernel1();

__device__ int control = 0;

__global__ void Controller()
{
    while (true)
    {
        switch (control)
        {
        case 0:
            ComputeKernel0 << <100, 32 >> > ();
            break;
        case 1:
            ComputeKernel1 << <500, 32 >> > ();
            break;
        default:
            return;
        }
        cudaDeviceSynchronize();
    }
}

int main()
{
    Controller << <1, 1 >> > ();
}

__device__ int counter = 0;
__device__ unsigned data = 0;

__global__ void ComputeKernel0()
{
    while (atomicInc(&data, 1000) != 0);
    if (blockIdx.x == 0 && threadIdx.x == 0)
    {
        counter++;
        if (counter > 50)
            atomicExch(&control, 2);
        else
            atomicExch(&control, 1);
        printf("%i ", counter);
    }
}

__global__ void ComputeKernel1()
{
    while (atomicDec(&data, 10000) != 1);
    if (blockIdx.x == 0 && threadIdx.x == 0)
        atomicExch(&control, 0);
}

Я запускаю это на 1050Ti в режиме WDDM.

На рисунке выше вы видите, что у него очень тонкая линия, представляющая контроллер, затем длинная полоса для ComputeKernel0, затем короткая строка для контроллера и длинная полоса для ComputeKernel1. Затем возобновление работы контроллера занимает очень много времени.

Так что в этом примере он переключается между продолжением мгновенно или очень долго. В моей настоящей программе это даже более эпизодически. Мгновенное продолжение может произойти после обоих ядер.

Есть ли способ, чтобы это всегда продолжалось мгновенно?