О пирамидальном сложении на параллельной архитектуре
На параллельных архитектурах часто приходится делать операцию reduce (складывать и умножать вектора, считать среднее и так далее). В отличие от однопоточной конструкции, где все тривиально, параллельная reduce разбивается на два этапа: сначала мы всеми исполняющими юнитами обрабатываем куски данных, а потом должны сложить (усреднить, поделить) результаты уже меньшим числом процессоров.
Для второго шага reduce обычно используется пирамидальная схема: сначала в N/2 потоков сложим N результатов попарно, затем сложим N/2 в N/4 и так далее. Число итераций равно, очевидно, log2N.
Возникает вопрос, «сколько данных складывать на каждой итерации?» Ведь можно складывать в N/4-N/16-N/256 кучек, можно по 1/8-64-512 и так далее. Из общих соображений, складывать по несколько лучше чем по два. Конечно, потоков получается меньше, но меньше и оверхед на создание-завершение потока.
Для NVidia CUDA идея "делать не по 2", выбирая все динамически, оказалась очень плохой. Да, с одной стороны мы действительно имеем оптимум при сложении по 8 или по 16. С другой стороны, код для вычисления содержит больше условного исполнения, отчего все ухудшается:
- С одной стороны, "умный" код для сложения по 8 примерно втрое быстрее, чем при сложении им же "по 2".
- С другой стороны, код рассчитанный только на сложение по 2 - в полтора раза быстрее чем умный складывает по восемь
Приведу на всякий случай быстрый код, взятый из примера scalarprod (CUDA SDK)
Код:
for(int stride = N / 2; stride > 0; stride >>= 1){
__syncthreads();
for(int iAccum = threadIdx.x; iAccum < stride; iAccum += blockDim.x)
accumResult[iAccum] += accumResult[stride + iAccum];
}
Здесь вложенный for - это на самом деле такой IF, который для рассматриваемого случая, когда N == blockDim.x, удовлетворяется только для части потоков, а выполняется тело цикла для этих потоков только один раз. Свой "умный" код не привожу, слишком уж умный.
Впрочем, есть и специальные случаи. Если известно, что количество складываемых элементов является степенью четверки, то ручное написание кода (без лишних IF) дает еще примерно полуторакратный рост производительности.
Вот код:
Код:
for(int stride = blockDim.x / 4; stride > 0; stride >>=2){
__syncthreads();
for(int iAccum = threadIdx.x; iAccum < stride;
iAccum += blockDim.x)
{
data[iAccum] += data[stride + iAccum];
data[iAccum] += data[stride*2 + iAccum];
data[iAccum] += data[stride*3 + iAccum];
}
}
Да, дальнейший unroll дает падение производительности.
(c) Алексей Тутубалин