Почему моя программа cuda стала медленнее после использования 128 потоков на блоках?


У меня есть простое приложение cuda со следующим кодом:

#include <stdio.h>
#include <sys/time.h>
#include <stdint.h>
__global__
void daxpy(int n, int a, int *x, int *y) {
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  y[i] = x[i];
  int j;
  for(j = 0; j < 1024*10000; ++j) {
     y[i] += j%10;
  }
}
// debug time
void calc_time(struct timeval *start, const char *msg) {
   struct timeval end;
   gettimeofday(&end, NULL);
   uint64_t us = end.tv_sec * 1000000 + end.tv_usec - (start->tv_sec * 1000000 + start->tv_usec);
   printf("%s cost us = %llun", msg, us);
   memcpy(start, &end, sizeof(struct timeval));
}
void do_test() {
   unsigned long n = 1536;
   int *x, *y, a, *dx, *dy;
   a = 2.0;
   x = (int*)malloc(sizeof(int)*n);
   y = (int*)malloc(sizeof(int)*n);
   for(i = 0; i < n; ++i) {
      x[i] = i;
   }

   cudaMalloc((void**)&dx, n*sizeof(int));
   cudaMalloc((void**)&dy, n*sizeof(int));
   struct timeval start;
   gettimeofday(&start, NULL);
   cudaMemcpy(dx, x, n*sizeof(int), cudaMemcpyHostToDevice);

   daxpy<<<1, 512>>>(n, a, dx, dy); // this line 
   cudaThreadSynchronize();
   cudaMemcpy(y, dy, n*sizeof(int), cudaMemcpyDeviceToHost);
   calc_time(&start, "do_test ");
   cudaFree(dx);
   cudaFree(dy);
   free(x);
   free(y);
}
int main() {
   do_test();
   return 0;
}

Вызов ядра gpu - это daxpy<<<1, 512>>>(n, a, dx, dy) , и я провел несколько тестов, используя различные размеры блоков:

  • daxpy<<<1, 32>>>(n, a, dx, dy)
  • daxpy<<<1, 64>>>(n, a, dx, dy)
  • daxpy<<<1, 128>>>(n, a, dx, dy)
  • daxpy<<<1, 129>>>(n, a, dx, dy)
  • daxpy<<<1, 512>>>(n, a, dx, dy)

... и сделал следующие замечания:

  • время выполнения одинаково для 32, 64, и 128 размеры блоков,
  • время выполнения отличается для размеров блоков 128 и 129, в определенный:
    • для 128 время выполнения составляет 280 МС,
    • для 129 время выполнения составляет 386 МС.
Я хотел бы спросить, в чем причина разницы во времени выполнения для размеров блоков 128 и 129.

Мой графический процессор-tesla K80:

CUDA Driver Version / Runtime Version          6.5 / 6.5
CUDA Capability Major/Minor version number:    3.7
Total amount of global memory:                 11520 MBytes (12079136768 bytes)
(13) Multiprocessors, (192) CUDA Cores/MP:     2496 CUDA Cores
GPU Clock rate:                                824 MHz (0.82 GHz)
Memory Clock rate:                             2505 Mhz
Memory Bus Width:                              384-bit
L2 Cache Size:                                 1572864 bytes
Maximum Texture Dimension Size (x,y,z)         1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
Maximum Layered 1D Texture Size, (num) layers  1D=(16384), 2048 layers
Maximum Layered 2D Texture Size, (num) layers  2D=(16384, 16384), 2048 layers
Total amount of constant memory:               65536 bytes
Total amount of shared memory per block:       49152 bytes
Total number of registers available per block: 65536
Warp size:                                     32
Maximum number of threads per multiprocessor:  2048
Maximum number of threads per block:           1024
Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
Maximum memory pitch:                          2147483647 bytes
Texture alignment:                             512 bytes
Concurrent copy and kernel execution:          Yes with 2 copy engine(s)
Run time limit on kernels:                     No
Integrated GPU sharing Host Memory:            No
Support host page-locked memory mapping:       Yes
Alignment requirement for Surfaces:            Yes
Device has ECC support:                        Enabled
Device supports Unified Addressing (UVA):      Yes
Device PCI Bus ID / PCI location ID:           135 / 0
1 2

1 ответ:

После предоставления нам точных временных различий в одном из комментариев, т. е.:

  • 280 МС Для до 128 потоков,
  • 386мс для 129 + потоков,
Я думаю, что это косвенно поддерживает мою теорию проблемы, связанной с варп-планированием. Посмотрите на gk210 whitepaper , который является чипом, используемым в K80:
  • К80 SMX параметры особенности процессор Quad варп планировщик, см. п. квад варп планировщик,
  • это означает, что K80 SMX способен планировать до 128 потоков сразу (4 деформации = = 128 потоков), они затем выполняются одновременно,

Таким образом, для 129 потоков планирование не может произойти сразу, потому что SMX должен запланировать 5 искажений, т. е. планирование будет происходить в два этапа.

Если вышесказанное верно, то я ожидал бы:

  • время выполнения должно быть примерно одинаковым для блоков размером 1-128,
  • время выполнения должно быть примерно одинаковым для блоков размером 129-192.

192 - это число ядра на SMX, см. белую бумагу. Напомним, что целые блоки всегда планируются для одного SMX, и поэтому очевидно, что если вы создадите более 192 потоков, то они наверняка не смогут выполняться параллельно, а время выполнения должно быть выше для 193+ числа потоков.

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

Отказ от ответственности: вышесказанное-это только мои предположения, поскольку у меня нет доступа к K80, ни к какому другому GPU с Quad warp scheduler, поэтому я не могу правильно профилировать ваш код. Но в любом случае, я считаю, что это задача для вас - почему бы не использовать nvprof и не профилировать свой код самостоятельно? Тогда вы сможете увидеть, где лежит разница во времени.