Поток CUDA блокируется при запуске многих ядер (> 1000)

Question

Я обнаружил, что поток CUDA будет блокироваться при запуске большого количества ядер (более 1000).Мне интересно, есть ли какая-нибудь конфигурация, которую я могу изменить?

В моих экспериментах я запускаю маленькое ядро ​​10000 раз.Это ядро ​​запустилось в ближайшее время (около 190us).Ядро запускается очень быстро при запуске первых 1000 ядер.Требуется 4 ~ 5us, чтобы запустить ядро.Но после этого процесс запуска становится медленным.Чтобы запустить новое ядро, нам понадобится около 190 человек.Поток CUDA, кажется, ожидает завершения предыдущего ядра, а размер буфера составляет около 1000 ядер.Когда я создал 3 потока, каждый поток может запустить 1000 асинхронных ядер.Я хочу сделать этот буфер больше.Я пытаюсь установить cudaLimitDevRuntimePendingLaunchCount, но не работает.Есть ли способ?

#include <stdio.h>
#include "cuda_runtime.h"

#define CUDACHECK(cmd) do {                         \
    cudaError_t e = cmd;                              \
    if( e != cudaSuccess ) {                          \
        printf("Failed: Cuda error %s:%d '%s'\n",             \
        __FILE__,__LINE__,cudaGetErrorString(e));   \
        exit(EXIT_FAILURE);                             \
    }                                                 \
} while(0)

// a dummy kernel for test
__global__ void add(float *a, int n) {
    int id = threadIdx.x + blockIdx.x * blockDim.x;
    for (int i=0; i<n; i++) {
        a[id] = sqrt(a[id] + 1);
    }
}

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

    //managing 1 devices
    int nDev = 1;
    int nStream = 1;
    int size = 32*1024*1024;


    //allocating and initializing device buffers
    float** buffer = (float**)malloc(nDev * sizeof(float*));
    cudaStream_t* s = (cudaStream_t*)malloc(sizeof(cudaStream_t)*nDev*nStream);


    for (int i = 0; i < nDev; ++i) {
        CUDACHECK(cudaSetDevice(i));
        //CUDACHECK(cudaDeviceSetLimit(cudaLimitDevRuntimePendingLaunchCount, 10000));
        CUDACHECK(cudaMalloc(buffer + i, size * sizeof(float)));
        CUDACHECK(cudaMemset(buffer[i], 1, size * sizeof(float)));
        for (int j = 0; j<nStream; j++)
        CUDACHECK(cudaStreamCreate(s+i*nStream+j));
    }

    for (int i = 0; i < nDev; ++i) {
        CUDACHECK(cudaSetDevice(i));
        for (int j=0; j<10000; j++) {
            for (int k=0; k<nStream; k++)
            add<<<32, 1024, 0, s[i*nStream+k]>>>(buffer[i], 1000);
        }
    }

    for (int i = 0; i < nDev; ++i) {
        CUDACHECK(cudaSetDevice(i));
        cudaDeviceSynchronize();
    }


    //free device buffers
    for (int i = 0; i < nDev; ++i) {
        CUDACHECK(cudaSetDevice(i));
        CUDACHECK(cudaFree(buffer[i]));
    }

    printf("Success \n");
    return 0;
}

Вот результаты nvprof:

Когда я создаю 3 потока, первое 3000 ядро ​​запускается быстро, а затем становится медленным

nvprof1.png

Когда я создаю 1 поток, первое 1000 ядер запускается быстро, а затем становится медленным

nvprof1.png

Answer 1

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

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

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

Как уже обсуждалось в комментариях, один из возможных подходов к устранению вашей озабоченности по поводу запусков в сценарии с несколькими устройствами - это запуск в ширину, а нев глубину.Под этим я подразумеваю, что вы должны изменить свои циклы запуска, чтобы запустить ядро ​​на устройстве 0, затем на устройстве 1, затем на устройстве 2 и т. Д., Прежде чем запускать следующее ядро ​​на устройстве 0. Это даст вам оптимальную производительность в смыслечто все графические процессоры будут задействованы в обработке как можно раньше в последовательности запуска.

Если вы хотите увидеть изменения в поведении или документации CUDA, общее предложение - стать зарегистрированным разработчиком на разработчике.nvidia.com, затем войдите в свою учетную запись и зарегистрируйте ошибку, используя процесс регистрации ошибок, нажав на имя своей учетной записи в правом верхнем углу.