Блог

Если я использую динамический параллелизм, в каком потоке будет работать дочерняя сетка?

Например, у меня есть одно ядро с именем A, а другое ядро с именем B. B запускается A.

Если ядро A выполняется в stream_A, а также если I не указывает идентификатор потока для ядра B, в каком потоке будет выполняться ядро? Это поток по умолчанию или он унаследует поток, в котором выполняется A?

Основываясь на том, что я вижу в документации, а также на моих собственных экспериментах, я наблюдаю, что в целом нет (упорядоченной) связи между потоками на хосте и потоками на устройстве.

Мы можем рассмотреть два случая:

В первом случае для созданных потоков я бы сказал, что это явно описано в документации.

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

 $ cat cdp.cu
#include <cstdio>
#include <cassert>
#include <unistd.h>
const unsigned long long dt = 10000000000ULL;
__device__ void delay(){
        unsigned long long start = clock64();
        while (start dt > clock64());
}
__global__ void child(){
        delay();
}
__global__ void parent(){
        child<<<1,1>>>();
        cudaError_t err = cudaGetLastError();
        assert(err == cudaSuccess);
        err = cudaDeviceSynchronize();
        assert(err == cudaSuccess);
}
__global__ void pk(){
        printf("hellon");
}
int main(){
        cudaStream_t s1, s2;
        cudaStreamCreate(amp;s1);
        cudaStreamCreate(amp;s2);
        parent<<<1,1, 0, s1>>>();
        sleep(1);
        pk<<<1,1,0,s2>>>();
        cudaDeviceSynchronize();
}
$ nvcc -o cdp cdp.cu -rdc=true -lcudadevrt -lineinfo
$ time ./cdp                 
hello

real    0m7.276s
user    0m4.193s
sys     0m2.061s
$
 

(Ubuntu 18.04, CUDA 11.4, V100)

Когда я запускаю приведенный выше код, я замечаю, что после запуска cdp исполняемого файла консоль простаивает примерно 1 секунду. После этого консоль печатает «привет», а затем консоль остается бездействующей в течение 6-7 секунд, а затем приложение завершает работу. (Пожалуйста, не ожидайте такого же поведения в графическом процессоре WDDM.)

Если бы нулевой поток на устройстве был «таким же, как» или «унаследован от» НУЛЕВОГО потока на хосте, тогда мы ожидали бы, что распечатка «привет» появится только непосредственно перед выходом из приложения. pk Ядру не будет разрешено запускаться до завершения работы child ядра, если задействована семантика потока NULL хоста. Поэтому мы должны сделать вывод, что нулевой поток устройства не совпадает с нулевым потоком хоста.

Мы можем использовать немного логики, чтобы убедить себя, что это также не совсем то же самое, что наследование основного (созданного) потока, используемого для запуска parent ядра. Мы можем прочитать в документации, что использование потока NULL для запуска дочернего ядра в том же потоковом блоке будет иметь поведение потока NULL. Однако это не могло быть правдой, если поток с нулевым значением устройства был просто таким же, как / унаследованный от потока, созданного хостом. Согласно документации, если у меня есть два дочерних запуска ядра из одного и того же блока потоков, один в поток NULL, а другой в поток, созданный устройством, тогда мы не ожидаем, что они смогут выполняться одновременно — это поведение потока NULL. Но если нулевой поток был просто унаследованным потоком хоста, созданным потоком, нет причин ожидать, что они не смогут выполняться одновременно.

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

Если вы хотите получить разъяснения в документации, обычным предложением является сообщение об ошибке.

Вместо того, чтобы беспокоиться об особенностях поведения нулевого потока, совет, который я даю при обучении CUDA, заключается в том, что если вас беспокоят сложные сценарии параллелизма, выполняйте всю свою работу с созданными потоками. Оставьте нулевой поток позади. Все, что вы хотите сделать, может быть сделано исключительно с созданными потоками. Конечно, делайте, что хотите.