На параллельных архитектурах часто приходится делать операцию 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 - в полтора раза быстрее чем умный складывает по восемь
- 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];
- }
Впрочем, есть и специальные случаи. Если известно, что количество складываемых элементов является степенью четверки, то ручное написание кода (без лишних 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];
- }
- }