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

Question

我有一个简单的 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 毫秒

Answer 1

问题似乎来自循环展开。

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

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

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

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

CUDA kernel: performance drops by 10x when increased loop count by 10%

parallel-processing

performance

kernel

loops

cuda