Синхронизация глубины вложенных ядер

Question

Давайте возьмем следующий код, где есть родительское и дочернее ядро.От указанного родительского ядра мы хотим запустить threadIdx.x дочерних ядер в разных потоках, чтобы максимизировать параллельную пропускную способность.Затем мы ждем тех потомков с cudaDeviceSynchronize(), так как родительское ядро ​​должно увидеть изменения, внесенные в global память.

Теперь предположим, что мы также хотим запустить n родительские ядра с потоками и, междукаждый набор n родительских ядер, которые мы хотим запустить параллельно, мы также должны ждать результатов, используя cudaDeviceSynchronize()

Как это будет вести себя?

Начиная с это официальное введение в динамический параллелизм от Nvidia Я бы подумал, что parent kernel[0] будет ждать только потоков, запущенных в нем.это правильно?Если нет, что происходит?

ПРИМЕЧАНИЕ: я знаю, что одновременно может работать только очень много потоков (32 в моем случае), но это больше для максимального заполнения

РЕДАКТИРОВАТЬ:маленький пример кода

__global__ void child_kernel (void) {}
__global__ void parent_kernel (void) 
{
    if (blockIdx.x == 0)
    {
        cudaStream_t s;
        cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);

        child_kernel <<<1,10,0,s>>> ();
        cudaStreamDestroy(s);
    }
    cudaDeviceSynchronize();
}
for (int i=0; i<10; i++)
{
    cudaStream_t s;
    cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);

    parent_kernel <<<10,10,0,s>>> ();
    cudaStreamDestroy(s);
}
cudaDeviceSynchronize();

Answer 1

Родительские ядра будут ждать завершения всех порожденных дочерних ядер, прежде чем родительское ядро ​​завершит работу.Это описано в документации по динамическому параллелизму :

. Вызов и завершение дочерних сеток правильно вложено, что означает, что родительская сетка не считается завершенной, пока все дочерние сетки не созданыего темы завершены.Даже если вызывающие потоки не синхронизируются явно в запущенных дочерних сетках, среда выполнения гарантирует неявную синхронизацию между родителем и дочерним элементом.

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

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

Не ясно, что вы спрашивали об этом, но обратите внимание, что для кода устройства в вашем примере cudaDeviceSynchronize() гарантирует только то, что этот поток будет ждатьдля дочернего ядра, чтобы завершить, и аналогичным образом только обеспечивает порядок видимости результатов порядка для этого потока .Если вы хотите, чтобы другие потоки в этом же блоке могли наблюдать результаты глобальной памяти из дочернего ядра, порожденного потоком 0 (просто для примера), то вы должны выполнить операцию cudaDeviceSynchronize () в потоке 0 с помощью__syncthreads() операция.После этого __syncthreads() другие потоки в том же блоке будут иметь гарантированную видимость глобальных результатов, создаваемых дочерним ядром, запущенным потоком 0 (или дочерним ядром, запущенным любым потоком, за которым следует вызов cudaDeviceSynchronize (), который предшествуетвышеупомянутый __syncthreads()).

В среде CDP следует учитывать еще несколько ограничений: предел вложенности и ожидаемый предел запуска .