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 在这个 post 上的表现，我发现我的 GPU 实际上每个线程有 255 个寄存器。所以，寄存器不是问题。但是，这让我不确定问题出在哪里。

我认为包括我没有使用任何特定的编译器选项也有好处。我曾尝试过 -ptxas 但变化不大。

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

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

Answer 1

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

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

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

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

CUDA, Copying to shared memory increases number of registers used dramatically

cuda

pycuda