Ядро CUDA: производительность падает в 10 раз при увеличении l oop count на 10%

Question

У меня есть простое ядро ​​CUDA для тестирования l oop развертывания, а затем я обнаружил еще одну вещь: когда l oop count равно 10, ядру требуется 34 миллисекунды для выполнения, когда l oop count 90, оно занимает 59 миллисекунд, но когда количество l oop равно 100, время, которое требуется, составляет 423 миллисекунды! Конфигурация запуска такая же, изменилось только l oop count. Итак, у меня вопрос: в чем может быть причина такого падения производительности?

Вот код, вход представляет собой массив из 128x1024x1024 элементов, и я использую PyCUDA:

__global__ void copy(float *input, float *output) {
  int tidx = blockIdx.y * blockDim.x + threadIdx.x;
  int stride = 1024 * 1024;
  for (int i = 0; i < 128; i++) {
    int idx = i * stride + tidx;
    float x = input[idx];
    float y = 0;

    for (int j = 0; j < 100; j += 10) {
      x = x + sqrt(float(j));
      y = sqrt(abs(x)) + sin(x) + cos(x);

      x = x + sqrt(float(j+1));
      y = sqrt(abs(x)) + sin(x) + cos(x);

      x = x + sqrt(float(j+2));
      y = sqrt(abs(x)) + sin(x) + cos(x);

      x = x + sqrt(float(j+3));
      y = sqrt(abs(x)) + sin(x) + cos(x);

      x = x + sqrt(float(j+4));
      y = sqrt(abs(x)) + sin(x) + cos(x);

      x = x + sqrt(float(j+5));
      y = sqrt(abs(x)) + sin(x) + cos(x);

      x = x + sqrt(float(j+6));
      y = sqrt(abs(x)) + sin(x) + cos(x);

      x = x + sqrt(float(j+7));
      y = sqrt(abs(x)) + sin(x) + cos(x);

      x = x + sqrt(float(j+8));
      y = sqrt(abs(x)) + sin(x) + cos(x);

      x = x + sqrt(float(j+9));
      y = sqrt(abs(x)) + sin(x) + cos(x);
    }

    output[idx] = y;
  }
}

Счетчик l oop, о котором я упоминал, это строка:

for (int j = 0; j < 100; j += 10)

И примеры выходных данных здесь:

10 циклов

griddimx: 1 griddimy: 1024 griddimz: 1
blockdimx: 1024 blockdimy: 1 blockdimz: 1
nthreads: 1048576 blocks: 1024
prefetch.py:82: UserWarning: The CUDA compiler succeeded, but said the following:
ptxas info    : 0 bytes gmem, 24 bytes cmem[3]
ptxas info    : Compiling entry function 'copy' for 'sm_61'
ptxas info    : Function properties for copy
32 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 21 registers, 336 bytes cmem[0], 52 bytes cmem[2]

вычисление занимает 34,24 миллисекунды

90 циклов

griddimx: 1 griddimy: 1024 griddimz: 1
blockdimx: 1024 blockdimy: 1 blockdimz: 1
nthreads: 1048576 blocks: 1024
prefetch.py:82: UserWarning: The CUDA compiler succeeded, but said the following:
ptxas info    : 0 bytes gmem, 24 bytes cmem[3]
ptxas info    : Compiling entry function 'copy' for 'sm_61'
ptxas info    : Function properties for copy
32 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 21 registers, 336 bytes cmem[0], 52 bytes cmem[2]

вычисление занимает 59,33 миллисекунды

100 циклов

griddimx: 1 griddimy: 1024 griddimz: 1
blockdimx: 1024 blockdimy: 1 blockdimz: 1
nthreads: 1048576 blocks: 1024
prefetch.py:82: UserWarning: The CUDA compiler succeeded, but said the following:
ptxas info    : 0 bytes gmem, 24 bytes cmem[3]
ptxas info    : Compiling entry function 'copy' for 'sm_61'
ptxas info    : Function properties for copy
32 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 22 registers, 336 bytes cmem[0], 52 bytes cmem[2]

вычисление занимает 422,96 миллисекунды

Answer 1

Проблема, похоже, исходит из l oop разворачивания .

Действительно, случай 10-loops может быть тривиально развернут с помощью NV CC, поскольку l oop фактически всегда выполняется один раз (таким образом, строку for можно удалить, установив j в 0). Случай 90-loops разворачивается с помощью NV CC (реальных итераций всего 9). Таким образом, результирующий код намного больше, но по-прежнему работает быстро, поскольку ветви не выполняются (графические процессоры ненавидят ветви). Однако случай 100-loops - это не развернуто с помощью NV CC (вы достигли порога оптимизатора компилятора). Результирующий код невелик, но он приводит к тому, что во время выполнения выполняется больше ветвей: ветвление выполняется для каждой выполненной l oop итерации (всего 10). Вы можете увидеть разницу в коде сборки здесь .

Вы можете принудительно развернуть, используя директиву #pragma unroll. Однако имейте в виду, что увеличение размера кода может снизить его производительность.

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