时间戳

首页 > 解决方案 > CUDA 内核：循环次数增加 10% 时性能下降 10 倍

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

问题描述

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

标签： performanceloopsparallel-processingcudakernel

解决方案

问题似乎来自循环展开。

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

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

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

推荐阅读