Почему моя программа 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 = %llu\n", 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

2

Решение

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

  • 280 мс до 128 потоков,
  • 386мс для 129+ потоков,

Я думаю, что это косвенно поддерживает мою теорию вопроса, связанную с планированием варпа. Посмотрите на Технический документ GK210, который является чипом, используемым в K80:

  • K80 SMX оснащен четырехпозиционным планировщиком, см. Раздел Quad Warp Scheduler,
  • Это означает, что K80 SMX может планировать до 128 потоков одновременно (4 перекоса == 128 потоков), затем они выполняются одновременно,

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

Если вышеупомянутое верно, тогда я ожидал бы:

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

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

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

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

3

Другие решения


По вопросам рекламы [email protected]