CUDA 内核：循环次数增加 10% 时性能下降 10 倍

Posted 2023-03-23

【中文标题】CUDA 内核：循环次数增加 10% 时性能下降 10 倍

【英文标题】：CUDA kernel: performance drops by 10x when increased loop count by 10%

【发布时间】：2020-10-08 06:24:59

【问题描述】：

我有一个简单的CUDA内核来测试循环展开，然后发现另一件事：当循环计数为10时，内核需要34毫秒才能执行，当循环计数为90时，需要59毫秒，但是当循环计数是100，耗时423毫秒！ 启动配置是一样的，只是循环次数改变了。 那么，我的问题是，这种性能下降的原因可能是什么？

这里是代码，输入是一个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;
  

我提到的循环计数是这一行：

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 毫秒

【参考方案1】：

问题似乎来自循环展开。

事实上，10-loops 的情况可以通过 NVCC 轻松展开，因为循环实际上总是执行一次（因此可以在 j 设置为 0 的情况下删除for 行）。 90-loops 案例由 NVCC 展开（只有 9 次实际迭代）。因此，生成的代码要大得多，但仍然很快，因为没有执行任何分支（GPU 讨厌分支）。但是，100-loops 案例没有被 NVCC 展开（您达到了编译器优化器的阈值）。生成的代码很小，但会导致在运行时执行更多的分支：每次执行的循环迭代都会执行分支（总共 10 次）。 可以看到汇编代码差异here。

您可以使用指令#pragma unroll 强制展开。但是，请记住，增加代码的大小会降低其性能。

PS：上一版本使用的寄存器数量稍多可能会降低性能，但simulations表明在这种情况下应该没问题。

【讨论】：

我用 #pragma unroll 和增量 j = 1 (j++) 进行了测试，然后计算时间约为 200 毫秒。没有展开它是 550 毫秒。手动展开，只需 41 毫秒！