c++ - 为什么我的 cuda 程序在 block 上使用 128 个线程后变慢了？

Question

我有一个简单的 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)我使用不同的 block 大小执行了一些测试:

• 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 block 大小，
• 执行时间因 block 大小而异128和 129 ， 尤其:
  • 对于 128执行时间为 280 毫秒，
  • 对于 129执行时间为 386 毫秒。

请问是什么导致 block 大小的执行时间不同128和 129 .

我的 GPU 是特斯拉 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

Answer 1

在其中一条评论中向我们提供了准确的时差后，即:

• 280 毫秒，最多 128 个线程，
• 129 个以上线程为 386 毫秒，

我认为它间接支持了我关于与 warp 调度相关的问题的理论。看GK210 whitepaper ，这是K80中使用的芯片:

• K80 SMX 具有一个 quad warp scheduler，请参阅 Quad Warp Scheduler 部分，
• 这意味着 K80 SMX 能够一次调度多达 128 个线程(4 个 warps == 128 个线程)，然后这些线程同时执行，

因此，对于 129 个线程，调度不能一次发生，因为 SMX 必须调度 5 个 warps，即调度将分两步进行。

如果以上为真，那么我预计:

• block 大小 1 - 128 的执行时间大致相同，
• block 大小为 129 - 192 的执行时间大致相同。

192 是 SMX 上的核心数，请参阅白皮书。提醒一下 - 整个 block 总是安排在一个 SMX 上，因此很明显，如果你生成超过 192 个线程，那么这些线程肯定无法并行执行，并且 193+ 个线程的执行时间应该更长。

您可以通过将内核代码简化到几乎什么都不做的程度来验证上述论点，因此执行时间是否仅由于调度而花费更长的时间应该或多或少是显而易见的(不会有其他限制因素，例如内存吞吐量)。

免责声明:以上只是我的假设，因为我无权访问 K80，也无权访问任何其他具有四元扭曲调度程序的 GPU，因此我无法正确分析您的代码。但无论如何，我相信这就是你的任务 - 为什么不使用 nvprof 并自己分析你的代码？然后你应该能够看到时差在哪里。