О пирамидальном сложении на параллельной архитектуре

На параллельных архитектурах часто приходится делать операцию 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)
  1. for(int stride = N / 2; stride > 0; stride >>= 1){
  2.      __syncthreads();
  3.      for(int iAccum = threadIdx.x; iAccum < stride; iAccum += blockDim.x)
  4.                accumResult[iAccum] += accumResult[stride + iAccum];
  5. }
Здесь вложенный for - это на самом деле такой IF, который для рассматриваемого случая, когда N == blockDim.x, удовлетворяется только для части потоков, а выполняется тело цикла для этих потоков только один раз. Свой "умный" код не привожу, слишком уж умный.

Впрочем, есть и специальные случаи. Если известно, что количество складываемых элементов является степенью четверки, то ручное написание кода (без лишних IF) дает еще примерно полуторакратный рост производительности. Вот код:

  1. for(int stride = blockDim.x / 4; stride > 0; stride >>=2){
  2.   __syncthreads();
  3.    for(int iAccum = threadIdx.x; iAccum < stride;
  4.          iAccum += blockDim.x)
  5.         {
  6.                         data[iAccum] += data[stride + iAccum];
  7.                         data[iAccum] += data[stride*2 + iAccum];
  8.                         data[iAccum] += data[stride*3 + iAccum];
  9.         }
  10. }
Да, дальнейший unroll дает падение производительности.

Tags: