Некоторые дочерние решетки не выполняются с динамическим параллелизмом CUDA

Question

Я экспериментирую с новой функцией динамического параллелизма в CUDA 5.0 (GTK 110). Я сталкиваюсь с странным поведением, что моя программа не возвращает ожидаемый результат для некоторых конфигураций - не только неожиданный, но и другой результат с каждым запуском.

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

я написал небольшую тестовую программу, чтобы проиллюстрировать это поведение:

#include <stdio.h> 

__global__ void out_kernel(char* d_out, int index) 
{ 
    d_out[index] = 1; 
} 

__global__ void kernel(char* d_out) 
{ 
    int index = blockIdx.x * blockDim.x + threadIdx.x; 
    out_kernel<<<1, 1>>>(d_out, index); 
} 

int main(int argc, char** argv) { 

    int griddim = 10, blockdim = 210; 
    // optional: read griddim and blockdim from command line 
    if(argc > 1) griddim = atoi(argv[1]); 
    if(argc > 2) blockdim = atoi(argv[2]); 

    const int numLaunches = griddim * blockdim; 
    const int memsize = numLaunches * sizeof(char); 

    // allocate device memory, set to 0 
    char* d_out; cudaMalloc(&d_out, memsize); 
    cudaMemset(d_out, 0, memsize); 

    // launch outer kernel 
    kernel<<<griddim, blockdim>>>(d_out); 
    cudaDeviceSynchronize(); 

    // dowload results 
    char* h_out = new char[numLaunches]; 
    cudaMemcpy(h_out, d_out, memsize, cudaMemcpyDeviceToHost); 

    // check results, reduce output to 10 errors 
    int maxErrors = 10; 
    for (int i = 0; i < numLaunches; ++i) { 
     if (h_out[i] != 1) { 
      printf("Value at index %d is %d, should be 1.\n", i, h_out[i]); 
      if(maxErrors-- == 0) break; 
     } 
    } 

    // clean up 
    delete[] h_out; 
    cudaFree(d_out); 
    cudaDeviceReset(); 
    return maxErrors < 10 ? 1 : 0; 
} 

Программа запускает ядро ​​в заданном количестве блоков (первый параметр) с заданным числом нитей каждый (второй параметр). Каждый поток в этом ядре запускает другое ядро ​​с одним потоком. Это дочернее ядро ​​напишет 1 в своей части выходного массива (который был инициализирован 0s).

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

Это происходит только в том случае, если многие дочерние решетки появляются одновременно. В моей тестовой системе (Tesla K20x) это имеет место для 10 блоков, содержащих по 210 потоков. Однако 10 блоков с 200 потоками обеспечивают правильный результат. Но также 3 блока с 1024 потоками, каждая из которых вызывает ошибку. Странно, об ошибке не сообщается во время выполнения. Резервные гриды просто игнорируются планировщиком.

Есть ли у кого-то еще такая же проблема? Является ли это поведение документированным где-то (я ничего не нашел), или это действительно ошибка в среде выполнения устройства?

Answer 1

Вы не делаете error checking любого вида, который я вижу. Вы можете и должны выполнять аналогичную проверку ошибок при запуске ядра устройства. Обратитесь к documentation Эти ошибки не обязательно будут пузырилась до хоста:

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

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

Если бы вы сделали правильную проверку ошибок, вы обнаружите, что в каждом случае, когда ядро ​​не запускалось, API-интерфейс среды cuda возвращал ошибку 69, cudaErrorLaunchPendingCountExceeded.

При сканировании documentation этой ошибки, вы найдете это:

cudaLimitDevRuntimePendingLaunchCount

Управляет объем памяти, отведенные для запусков ядра буферных, которые еще не начали выполнять, из-за либо к нерешенным зависимостям, либо к недостатку ресурсов выполнения. Когда буфер заполнен, запуски установят последнюю ошибку потока до cudaErrorLaunchPendingCountExceeded. Ожидаемый запуск запуска - 2048 запусков.

В 10 блоках * 200 потоков вы запускаете 2000 ядер, и все работает.

В 10 блоках * 210 потоков вы запускаете 2100 ядер, что превышает предел 2048, упомянутый выше.

Обратите внимание, что это несколько динамичный характер; в зависимости от того, как ваше приложение запускает дочерние ядра, вы можете запустить более 2048 ядер без ущерба для этого предела. Но так как ваше приложение запускает все ядра примерно одновременно, вы достигаете предела.

Правильная проверка ошибок cuda рекомендуется в любое время, когда ваш код CUDA не ведет себя так, как вы ожидаете.

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

__global__ void kernel(char* d_out) 
{ 
    int index = blockIdx.x * blockDim.x + threadIdx.x; 
    out_kernel<<<1, 1>>>(d_out, index); 
// cudaDeviceSynchronize(); // not necessary since error 69 is returned immediately 
    cudaError_t err = cudaGetLastError(); 
    if (err != cudaSuccess) d_out[index] = (char)err; 
} 

Предстоящего предел количества запуска является изменяемым. См. Документацию по адресу cudaLimitDevRuntimePendingLaunchCount