GPGPU: влияние размера блока на производительность программы, почему моя программа работает быстрее при очень определенных размерах?

Question

Моя программа Cuda получает значительное повышение производительности (в среднем) в зависимости от размера блоков & # блоков; где общее количество «нитей» остается неизменным. (Я не уверен, что поток является правильной терминологией ... но я буду использовать его здесь, где для каждого ядра общее количество потоков (# блоков) * (размер блока)). Я сделал несколько графиков, чтобы проиллюстрировать мою мысль.

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

В основном я просматриваю большие массивы, которые логически рассматриваются как 2D-массивы, где каждый поток добавляет элемент из массива, а также добавляет квадрат этого значения к другой переменной, а затем в конце записывает значение в другой массив , где во время каждого считывания все потоки сдвигаются определенным образом. Вот мой код ядра:

__global__ void MoveoutAndStackCuda(const float* __restrict__ prestackTraces, float* __restrict__ stackTracesOut, 
    float* __restrict__ powerTracesOut, const int* __restrict__ sampleShift, 
    const unsigned int samplesPerT, const unsigned int readIns, 
    const unsigned int readWidth, const unsigned int defaultOffset) { 

    unsigned int globalId = ((blockIdx.x * blockDim.x) + threadIdx.x); // Global ID of this thread, starting from 0 to total # of threads 

    unsigned int jobNum = (globalId/readWidth); // Which array within the overall program this thread works on 
    unsigned int readIndex = (globalId % readWidth) + defaultOffset; // Which sample within the array this thread works on 

    globalId = (jobNum * samplesPerT) + readIndex; // Incorperate default offset (since default offset will also be the offset of 
                // index we will be writing to), actual globalID only needed for above two variables. 

    float stackF = 0.0; 
    float powerF = 0.0; 

    for (unsigned int x = 0; x < readIns; x++) { 

    unsigned int indexRead = x + (jobNum * readIns); 

    float value = prestackTraces[readIndex + (x * samplesPerT) + sampleShift[indexRead]]; 

    stackF += value; 
    powerF += (value * value); 
    } 

    stackTracesOut[globalId] = stackF; 
    powerTracesOut[globalId] = powerF; 
} 

Теперь для мяса этого поста, при вызове этого кода

MoveoutAndStackCuda<<<threadGroups, threadsPerGroup>>>(*prestackTracesCudaPtr, 
    *stackTracesOutCudaPtr, *powerTracesOutCudaPtr, 
    *sampleShiftCudaPtr, samplesPerT, readIns, 
    readWidth, defaultOffset); 

Все, что я сделал, отличаются threadGroups и threadsPerGroup внутри < < < >>>, где threadGroups.x * threadsPerGroup.x остается неизменным. (Как указывалось ранее, это одномерная задача).

Я увеличил размер блока на 64, пока не достиг 1024. Я не ожидал изменений, потому что я понял, что размер блока больше 32, что, по моему мнению, является количеством ALU в ядре, оно будет работать как как можно быстрее. Посмотрите на этом графике я сделал:

Для этого конкретного размера общего числа потоков является 5000 * 5120, так, например, если размер блока равен 64, то есть ((5000 * 5120)/64). По какой-то причине существует значительный прирост производительности при размере блока 896, 768 и 512. Почему?

Я знаю, что это выглядит случайным образом, но каждая точка на этом графике - это 50 тестов, усредненных вместе!

Вот еще один график, на этот раз, когда общее количество потоков будет (8000 * 8192). На этот раз бустер на 768 и 960.

Еще один пример, на этот раз для работы, которая меньше, чем у двух других проблем (всего резьб 2000 * 2048):

Фактически вот альбом, который я сделал из этих графиков, с каждым графиком, представляющим другой размер проблемы: graph album.

Я запускаю этот a Quadro M5000, который имеет 2048 Cuda Cores. Я считаю, что у каждого Cuda Core есть 32 ALU, поэтому я полагаю, что общее количество вычислений, которое может происходить в любой момент времени (2048 * 32)?

Итак, что объясняет эти магические числа?Я полагал, что это может быть общее количество потоков, разделенных # ядрами cuda, или разделенных на (2048 * 32), но до сих пор я не нашел никакой корреляции ни с чем, что простирается на все графики в моем альбоме. Есть ли еще одно испытание, которое я мог бы сделать, чтобы помочь сузить дело? Я хочу узнать, какой размер блока для запуска этой программы за наилучшие результаты.

Также я не включил его, но я также проверил, где размер блока уменьшился на 1 из 32, а вещи стали экспоненциально медленнее. Это имеет смысл для меня с тех пор, у нас меньше локальных потоков на группу, чем ALU в заданном многопроцессоре.

Answer 1

На основании этого заявления:

I увеличивается размер блока на 64, пока я не достиг 1024. Я не ожидал, что никаких изменений, потому что я понял, до тех пор, как размер блока больше, чем 32, который я считаю, является # из ALU в ядре, он будет работать как можно быстрее.

Я бы сказал, что есть важное понятие о графических процессорах, о которых вы, вероятно, не знаете: графические процессоры - это машина с скрытой задержкой. Они скрывают латентность, главным образом, благодаря наличию большого количества доступных (параллельных) работ. Это можно приблизительно суммировать как «много потоков». Совершенно неверная идея с графическими процессорами заключается в том, что, когда у вас достаточно потоков для покрытия количества «ядер» или блоков исполнения, этого достаточно. Это не так.

В качестве (начинающего) графического процессора вы должны в основном игнорировать количество ядер в вашем GPU. Вы хотите лотов тем. И на уровне ядра, и на GPU SM.

В общем, поскольку вы предоставляете больше потоков каждому SM, тогда способность графического процессора скрывать латентность при выполнении другой полезной работы увеличивается. Это объясняет общую тенденцию во всех ваших графиках, что наклон, как правило, снизу слева направо (т. Е. Средняя производительность увеличивается, как правило, при условии, что вы предоставляете более открытую работу каждому SM).

Это не касается пиков и долин. У графических процессоров большое количество архитектурных проблем, которые могут повлиять на производительность. Здесь я не буду предлагать полное лечение. Но возьмем один случай:

Почему производительность на первом графике увеличивается до 512 потоков, а затем уменьшается на 576 потоков?

Это, скорее всего, Размещение эффект. SM в вашем GPU имеет максимальное дополнение к 2048 потокам. Основываясь на предыдущем обсуждении, SM будет иметь максимальную способность скрывать латентность (и, как правило, обеспечивать максимальную среднюю производительность), когда мы максимизируем дополнение нитей до 2048.

Для размера блока 512 потоков мы может поместиться ровно 4 из этих блоков потока на SM, и тогда он будет иметь дополнение к 2048 потокам, из которых можно выбрать для работы и скрытия времени ожидания.

Но если вы измените размер блока резьбы на 576, 4 * 576> 2048, значит, мы больше не можем поместить 4 резьбовых блока на каждом SM. Это означает, что для этой конфигурации ядра каждый SM будет работать с 3 потоковыми блоками, то есть 1728 потоков из 2048 возможных. Это на самом деле хуже, с точки зрения SM, чем предыдущий случай, который разрешил 2048 потоков, и поэтому он может быть индикатором того, почему производительность уменьшается от 512 до 576 потоков (так же, как она увеличилась с 448 до 512, что связано с аналогичным изменением в мгновенном заполнении).

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

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

Чтобы углубиться в это, я бы посоветовал вам потратить некоторое время на изучение занятости, регистров на поток и возможности анализа производительности различных профилографов. Уже есть много и много информации по этим темам; Google - ваш друг, и обратите внимание на question/answers, связанный в комментариях выше, в качестве разумной отправной точки. Чтобы полностью изучить занятость и ее влияние на производительность, требуется больше информации, чем то, что вы здесь дали. Он требует в основном MCVE, а также точную командную строку компиляции, а также платформу, на которой вы работаете, и версию CUDA. Все эти вещи влияют на использование регистров для каждого потока в компиляторе, большинство из которых вы не предоставили.