cuda - 寄存器溢出是否可能导致 CUDA_EXCEPTION_5，Warp Out-Of-Range Address 错误？

Question

我收到 CUDA_EXCEPTION_5，Warp Out-of-range Address 错误，我正在尝试找出可能导致该错误的各种情况。

我正在努力将一个 C 项目(由其他人编写)移植到 CUDA。 C 代码非常依赖寄存器，在堆栈中实例化了许多数组。我假设很可能会发生寄存器溢出，这可能会触发 warp out-of-range 错误。

请注意，我想先让它运行起来，然后再开始优化代码。

我正在使用 Compute Capable 3.0 硬件，根据维基百科，它具有 512KB 的“每个线程本地内存”。我在别处读到，每个 SM 有 512KB 的寄存器空间。每个运行线程有可能有 512KB 的寄存器空间吗？

我目前正在按如下方式执行我的内核(是的，我知道它非常慢):

dim3 grid(28800,1);
cuPlotLRMap<<<grid,1>>>(...)

一些细节(我不知道这会有多大帮助):

我的硬件有 7 个 SM。有 112 个运行 block ，那么这是否意味着每个 block 获得 512k 寄存器空间的 1/16？

我也明白，如果一个线程超出寄存器空间，它可能会溢出到全局内存中。发生这种情况时是否有可能并发线程溢出到同一个全局内存空间？

Answer 1

512KB of "local memory per thread". I read elsewhere it has 512KB of register space per SM. Is it possible to have 512KB of register space per running thread?

参见 Compute Capabilities CUDA C 编程指南中的表格。计算能力 2.x 及更高版本的设备支持每个线程最大 512KB 的本地内存。函数 cudaDeviceSetLimit(cudaLimitStackSize, bytesPerThread) 可用于设置该值。我相信默认值是每个线程 2 KB。

My hardware has 7 SMs. There are 112 running blocks, so does this mean each block gets 1/16th of 512k worth of register space?

计算能力 3.x 设备每个多处理器最多可以有 16 个驻留 block 。这假定您的寄存器/线程、线程/ block 或共享内存/ block 不会将内核限制为小于设备最大值。 Visual Profiler 和 Nsight VSE CUDA Profiler 内核使用的配置。

目前，您只启动了 1 个线程/ block 。您应该为每个 block (32) 启动 WARP_SIZE 的倍数。

I also understand if a thread exceeds the register space it can overflow into global memory. Is it possible for concurrent threads to overflow into the same global memory space when this occurs?

在编译或 JIT 时，编译器将执行寄存器分配。如果每个线程没有足够的寄存器，那么编译器将溢出到本地内存。此操作是确定性的，在运行时不确定。

计算能力 3.0 设备限制为 63 个寄存器/线程。计算能力 3.5 设备限制为每个线程 255 个寄存器。