Почему выполнения ядра в разных потоках не параллельны?

Question

Я только что изучил потоковую технику в CUDA и попробовал ее. Hoverver возвращает нежелательный результат, а именно, потоки не параллельны. (На GPU Tesla M6, ОС Red Hat Enterprise Linux 8)

У меня есть матрица данных с размером (5,2048) и ядро ​​для обработки матрицы.

Мой план состоит в том, чтобы разложить данные по секторам 'nStreams = 4' и использовать 4 потока для параллельного выполнения ядра.

Часть моего кода выглядит следующим образом:

int rows = 5;
int cols = 2048;

int blockSize = 32;
int gridSize = (rows*cols) / blockSize;
dim3 block(blockSize);
dim3 grid(gridSize);

int nStreams = 4;    // preparation for streams
cudaStream_t *streams = (cudaStream_t *)malloc(nStreams * sizeof(cudaStream_t));
for(int ii=0;ii<nStreams;ii++){
    checkCudaErrors(cudaStreamCreate(&streams[ii]));
}

int streamSize = rows * cols / nStreams;
dim3 streamGrid = streamSize/blockSize;

for(int jj=0;jj<nStreams;jj++){
    int offset = jj * streamSize;
    Mykernel<<<streamGrid,block,0,streams[jj]>>>(&d_Data[offset],streamSize);
}    // d_Data is the matrix on gpu

Результат Visual Profiler показывает, что 4 разных потока не параллельны. Поток 13 работает первым, а поток 16 - последним. Между потоком 13 и потоком 12,378us. Каждое выполнение ядра длится около 5us. В строке «Runtime API» выше написано «cudaLaunch».

Не могли бы вы дать мне совет? Спасибо!

(Я не знаю, как загрузить изображения в стеке, поэтому я просто описываю результат словами).

Answer 1

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

Кроме того, ваш Tesla M6 имеет 12 мультипроцессоров, если я не ошибаюсь.Каждый из этих 12 мультипроцессоров Maxwell может содержать до 32 резидентных блоков.Таким образом, общее максимальное количество блоков, находящихся на всем устройстве, достигает 384. Вы запускаете 320 блоков по 32 потока в каждом.Само по себе это не оставляет много места, и вы, вероятно, используете более 32 регистров на поток, поэтому графический процессор будет заполнен одним из этих запусков, что, скорее всего, приводит к тому, что драйвер решает не запускать другое ядро.параллельно.

Параллельный запуск ядра в основном имеет смысл, когда у вас есть, например, кучка небольших ядер, которые делают разные вещи, которые могут работать рядом друг с другом на отдельных мультипроцессорах.Кажется, что ваша рабочая нагрузка может легко заполнить все устройство.Что именно вы надеетесь достичь, запустив несколько ядер параллельно?Почему вы работаете с такими крошечными блоками?Разве не имеет смысла запускать все это как одно большое ядро ​​с большими блоками?Обычно вы хотите иметь как минимум пару перекосов на блок.См., Например, этот вопрос для получения дополнительной информации: Как выбрать размеры сетки и блока для ядер CUDA? Если вы используете разделяемую память, вам также понадобится не менее двух блоков на многопроцессорную систему, поскольку вы выиграли в противном случаеВы даже не сможете использовать все это на некоторых графических процессорах (которые, например, предлагают общую память 96 КиБ на каждый многопроцессорный процессор, но каждый блок может иметь только максимум 48 КиБ)…

Answer 2

Чтобы добавить к существующему ответу (который является полностью правильным), рассмотрите следующую тривиально полную версию кода, который вы разместили в своем вопросе:

__global__
void Mykernel(float* data, int size)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    for(; tid < size; tid+= blockDim.x * gridDim.x) data[tid] = 54321.f;
}

int main()
{
    int rows = 2048;
    int cols = 2048;

    int blockSize = 32;
    dim3 block(blockSize);

    int nStreams = 4;    // preparation for streams
    cudaStream_t *streams = (cudaStream_t *)malloc(nStreams * sizeof(cudaStream_t));
    for(int ii=0;ii<nStreams;ii++){
        cudaStreamCreate(&streams[ii]);
    }

    float* d_Data;
    cudaMalloc(&d_Data, sizeof(float) * rows * cols);
    int streamSize = rows * cols / nStreams;
    dim3 streamGrid = dim3(4);

    for(int jj=0;jj<nStreams;jj++){
        int offset = jj * streamSize;
        Mykernel<<<streamGrid,block,0,streams[jj]>>>(&d_Data[offset],streamSize);
    }    // d_Data is the matrix on gpu


    cudaDeviceSynchronize();
    cudaDeviceReset();
}

Обратите внимание на два отличия - количество блоковколичество запускаемых для каждого ядра уменьшается, а общее количество вычислений для каждого потока увеличивается путем установки rows в 2048. Само ядро ​​содержит цикл с шагом сетки, который позволяет каждому потоку обрабатывать несколько входных данных, обеспечивая обработку всего входного набора данных.независимо от того, сколько всего блоков / потоков запущено.

Профилирование на аналогичном графическом процессоре Maxwell для вашего устройства показывает это:

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

Ядра больше не перекрываются.Зачем?Поскольку задержка драйвера и устройства достаточно высока, а время выполнения каждого ядра достаточно короткое, чтобы не было времени для перекрытия выполнения, даже если ресурсы устройства в противном случае позволили бы это.Таким образом, помимо ограничений требований к ресурсам, описанных в другом ответе, объем вычислений должен быть достаточно большим, чтобы компенсировать фиксированную задержку, связанную с планированием запуска ядра в потоке.

Наконец, я хотел бы предложить правильный подход кустановка схемы параллельного выполнения на основе потоков должна выглядеть примерно так:

int blockSize = 32;
dim3 block(blockSize);
int blocksperSM, SMperGPU = 13; // GPU specific
cudaOccupancyMaxActiveBlocksPerMultiprocessor(&blocksperSM, Mykernel, blockSize, 0); // kernel specific
dim3 streamGrid = blocksperSM * (SMperGPU / nStreams); // assume SMperGPU >> nstreams   

Здесь идея состоит в том, что количество доступных SM (приблизительно) поровну поделено между потоками и количеством блоков, которыеМаксимально занимаемая каждым SM для выбранного размера блока получается для ядра через оккупационный API.

Этот профиль выглядит следующим образом:

, что дает как перекрытие, так и короткое время выполнения, правильно сопоставляя требования к ресурсам ядрадо емкости графического процессора для случая с rows = 2048.