CUDA，复制到共享内存会显着增加使用的寄存器数量

Question

由于请求的资源过多，我遇到了内核启动失败的问题。我理解这个错误，我可以减少我的块大小来避免它，但我正在努力解决这个问题。

我正在使用 Nvidia Tesla K40c GPU。 我正在使用 pycuda 来解决 PDE 系统。所以，我的目标是对每个线程进行一些本地计算，然后写入共享内存数组。我对 GPU 计算相当陌生，但我所知道的手头问题写在下面。 此问题与下面代码片段中注释掉的代码行有关。我知道共享内存非常适合块中的线程间通信，并且我的共享内存正常工作，直到我尝试从我假设存储在寄存器中的局部变量写入它。我假设这是因为我读取的数组小于特定大小，如果我没记错的话，16 个浮点数，可以存储在寄存器中。我的大小为 4。无论如何，这是避免存储在全局中的目标。

__device__
void step(float *shared_state, int idx)
{
  float dfdxy[NVC]={0};
  get_dfdx(dfdxy,shared_state,idx);
  get_dfdy(dfdxy,shared_state,idx);
  __syncthreads();
  //shared_state[idx+0*SGIDS] += dfdxy[0];
}

这是踪迹。正如我所提到的，我对错误很熟悉。

Traceback (most recent call last):
  File "./src/sweep/sweep.py", line 325, in <module>
    sweep(arr0,targs,order,block_size,euler.step,0)
  File "./src/sweep/sweep.py", line 109, in sweep
    gpu_speed(arr, source_mod, cpu_fcn, block_size,ops,num_tries=20)
  File "./src/sweep/sweep.py", line 175, in gpu_speed
    gpu_fcn(arr_gpu,grid=grid_size, block=block_size,shared=shared_size)
  File "/home/walkanth/.conda/envs/pysweep/lib/python3.6/site-packages/pycuda/driver.py", line 402, in function_call
    func._launch_kernel(grid, block, arg_buf, shared, None)
pycuda._driver.LaunchError: cuLaunchKernel failed: too many resources requested for launch

当我运行带有注释行的代码时，问题就是这个。它说我正在使用 32 个寄存器。这很好，一切正常，因为我低于 63 的限制。

但是，当我取消注释该行时，使用的寄存器数量会上升到 70 个，我怀疑这就是内核启动失败的原因。

所以，有几个问题。

首先，谁能解释一下为什么会这样？我一直在寻找一段时间，但都失败了。

其次，如果没有办法解决这个问题。除了减少block_size之外，有谁知道一些减少我的寄存器使用的技巧？ 我在 nvidia dev 上看到过一些较早的帖子讨论过这个问题，但它们似乎已经过时了。

编辑：

感谢 Michael 在这篇文章中的介绍，我发现我拥有的 GPU 实际上每个线程有 255 个寄存器。所以，寄存器不是问题。但是，这让我不确定问题出在哪里。

我认为包含我没有使用任何特定的编译器选项也是有益的。我曾经尝试过 -ptxas，但它并没有太大变化。

我不想减小 blocksize 块大小，因为在需要外部信息之前我可以进行的计算次数取决于 blocksize 的最小尺寸（x 或 y）。块大小越大，可能的计算就越多。

编辑： 所以，据我了解，我仍然超过了每个 SM 的寄存器总数，这是导致问题的原因。我需要减少使用的寄存器或块大小。

Answer 1

编译器将尝试自动优化寄存器指令的数量；如果您编写的代码最终不会在线程之外的任何地方存储信息，那么根本不应该生成这些指令。这可能就是您在取消注释写入共享内存的行时看到寄存器数量发生巨大变化的原因。

但是，根据https://developer.nvidia.com/cuda-gpus，K40c 的计算能力为 3.5，而根据https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capabilities，计算能力为 3.5 的设备每个线程最多可以有 255 个寄存器，而不是 63 个。因此，如果您仍然只使用每个线程 70 个寄存器，那么这可能不是问题。如果您通过减小块大小不再收到错误，则可以确认这一点；块大小的减少会减少块中的线程数，但不应改变每个线程使用的寄存器数量，因此如果您实际上每个线程的寄存器都用完了，它不应该解决您的问题。

如果不进一步了解您的编译器选项、内核的其余部分以及您如何启动它，我们无法轻易确定资源问题是什么。每个块的寄存器数量和每个多处理器的寄存器数量也有限制；如果减小块大小可以解决问题，那么您很可能超出了这些阈值......并且需要减小块大小。目前尚不清楚您为什么不想减小块大小，但似乎您只是遇到了硬件限制。