Что происходит с этим кодом CUDA, который возвращает этот неожиданный вывод?

Question

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

Кажется, что-то не так с функцией устройства, которую я определил как arrAdd. Вот таблица результатов до и после каждого дочернего ядра в родительском ядре k2.

Initially    : k1   = { -1   0   0   0   0 }
Post arrInit : temp = { .25 .25 .25 .25 .25}
Post arrMult : temp = {-.25  0   0   0   0 }
post arrAdd  : temp = { -8   0   0   0   0 }
Expected     : temp = {-.50  0   0   0   0 }


__global__ void k2(double* concs, int* maxlength, double* k1s, double* k2s, double * temp, double* tempsum)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    double a21 = .25;

    arrInit<<< 1, *maxlength >>>(temp, a21);                //temp = a21
    cudaDeviceSynchronize();
    arrMult<<< 1, *maxlength >>>(k1s, temp, temp);          //temp = a21*k1
    cudaDeviceSynchronize();
    arrAdd<<< 1, *maxlength >>>(temp, temp, temp);          //temp = 2*a21*k1
    cudaDeviceSynchronize();
}

__global__ void arrAdd(double* a, double* b, double* c)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    c[idx]=a[idx]+b[idx];
}
__global__ void arrMult(double* a, double* b, double* c)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    c[idx]=a[idx]*b[idx];
}
__global__ void arrInit(double* a, double b)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    a[idx]=b;
}

Answer 1

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

Когда вы используете динамический параллелизм, будьте осторожны с этими элементами:

1. Самое глубокое, что вы можете сделать, это 24 (CC=3,5).

2. Количество динамических ядер, ожидающих одновременного запуска, ограничено (по умолчанию 2048 при CC = 3,5), но может быть увеличено.

3. Держите родительское ядро ​​занятым после вызова дочернего ядра, иначе есть большая вероятность того, что вы потратите ресурсы

Я предполагаю, что ваши странные неверные результаты происходят от второго фактора, упомянутого выше. Когда вы достигнете предела, некоторые из динамических ядер просто не будут работать, и если вы не будете проверять наличие ошибок, вы не заметите, потому что механизм создания ошибок для каждого потока.

Вы можете увеличить этот лимит с помощью cudaDeviceSetLimit(), имеющего cudaLimitDevRuntimePendingLaunchCount в качестве лимита. Но чем больше вы указываете, тем больше вы занимаете пространство глобальной памяти. Взгляните на раздел C.4.3.1.3 документации здесь.