Динамический параллелизм CUDA. Есть ли способ бесконечно вкладывать запуски ядра?
Теперь я использую динамический параллелизм CUDA для создания ядра в функции ядра.
В документе CUDA функции ядра могут запускаться только с фиксированной глубиной рекурсии из-за ограничений ресурсов.
Но в моем проекте я хочу запустить ядро, но родительскому ядру не нужно ждать выхода дочернего ядра. Другими словами, они полностью независимы.
Итак, есть ли способы запуска ядра в функциях ядра, не ограниченные глубиной рекурсии?
Я использую cudaDeviceSetLimit() для установки cudaLimitDevRuntimeSyncDepth, но ограничения все равно есть.
Пример:
__global__ void do_something(MyQueue* queue, Task* task) {
// do something ...
task->execute();
// If queue is not empty, pop from it and launch a kernel to execute it
Task* t = queue->pop();
if (t) {
do_something<<<t->gridSize, t->blockSize, t->mem, stream>>>(queue, t);
}
}
- Графический процессор: NVIDIA A100
- Версия CUDA: 12.4
- флаг компиляции:
-rdc=true -gencode=arch=compute_61,code=sm_61 -gencode=arch=compute_61,code=compute_61
Есть конкретная причина, по которой вы используете рекурсивный подход? Если нет, то я думаю, что наиболее удобным подходом было бы запустить n ядер для полного использования графического процессора, каждое из которых учитывает sizeOf(queue)/n задач.
CUDA 12 фактически представила новую версию динамического параллелизма (CDP2), которая больше не позволяет синхронизировать код устройства, поскольку это было проблемой производительности. Похоже, вам следует просто запускать ядра, используя новый cudaStreamFireAndForget.
Спасибо, я думаю, что CDP2 может решить мою проблему. Я пишу систему времени выполнения на графическом процессоре. Учитывая производительность и другую стратегию планировщика, я перемещаю очередь на графический процессор, поэтому не могу запустить ядро на процессоре. Теперь я пытаюсь использовать динамическую параллель cuda, чтобы решить эту проблему. :-)
Ответы 1
согласно моему тестированию, используя CDP2, кажется возможным запустить ядро, которое будет вызывать подядра для «длинной» последовательности:
# cat t225.cu
#include <iostream>
#include <cstdio>
#include <cstdlib>
using mt = int;
const mt my_stopc = 32768;
__global__ void k(mt *d, const mt my_stop){
mt c = *d;
if (c < my_stop){
c++;
*d = c;
k<<<1,1,0, cudaStreamTailLaunch>>>(d, my_stop);}
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) printf("device: %lu, %s\n", (unsigned long long)c, cudaGetErrorString(err));
}
int main(int argc, char *argv[]){
mt my_stop = my_stopc;
if (argc > 1) my_stop = atol(argv[1]);
mt *d;
cudaMallocManaged(&d, sizeof(d[0]));
*d = 0;
k<<<1,1>>>(d, my_stop);
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) std::cout << "host 1: " << cudaGetErrorString(err) << std::endl;
err = cudaDeviceSynchronize();
if (err != cudaSuccess) std::cout << "host 2: " << cudaGetErrorString(err) << std::endl;
std::cout << "iter: " << *d << std::endl;
}
# nvcc -o t225 t225.cu -arch=sm_89 -rdc=true -lcudadevrt -lineinfo
# ./t225
iter: 32768
# ./t225 70000
iter: 70000
#
CUDA 12.2, графический процессор L4
Если я увеличу счетчик до 1000000, то (родительское) ядро завершит работу примерно за 12 секунд:
# time ./t225 1000000
iter: 1000000
real 0m12.306s
user 0m10.155s
sys 0m2.080s
#