为什么存在寄存器溢出机制时，由于寄存器使用过多而无法启动内核？

Question

1) 内核何时开始将寄存器溢出到本地内存？

2) 当没有足够的寄存器时，CUDA 运行时如何决定不启动内核并抛出太多资源请求错误？启动内核需要多少寄存器？

3) 既然有寄存器溢出机制，那么即使没有足够的寄存器，难道不应该启动所有的CUDA内核吗？

Answer 1

1) 内核何时开始将寄存器溢出到本地内存？

这完全在编译器的控制之下。它不是由运行时执行的，并且没有关于它的动态运行时决策。当您的代码到达溢出点时，这意味着编译器已插入如下指令：

STL  [R0], R1

在这种情况下，R1 被存储到本地内存中，本地内存地址在R0 中给出。这将是一个溢出商店。 （在该指令之后，R1 可以用于/加载其他内容。）当然，编译器知道它何时执行此操作，因此它可以报告它选择使用的溢出加载和溢出存储的数量/制作。您可以使用-Xptxas=-v 编译器开关获取此信息（以及寄存器使用情况和其他信息）。

编译器（除非您对其进行限制，请参见下文）主要关注性能上的寄存器使用决策，否则较少关注实际使用了多少寄存器。第一要务是性能。

2) 当没有足够的寄存器时，CUDA 运行时如何决定不启动内核并抛出太多资源请求错误？启动内核需要多少寄存器？

在编译时，当你的内核代码被编译时，编译器不知道它会如何启动。它不知道您的启动配置会是什么样子（块数、每个块的线程数、动态分配的共享内存量等）事实上，编译过程大部分都在进行，就好像正在编译的东西是单个线程一样。

在编译期间，编译器会针对寄存器分配（寄存器的使用方式和使用位置）做出一系列静态决定。 CUDA 有 binary utilities 可以帮助理解这一点。寄存器分配在运行时不会改变，绝不是动态的，因此完全在编译时确定。因此，在完成给定设备代码功能的编译时，通常可以确定需要多少个寄存器。编译器将此信息包含在二进制编译对象中。

在运行时，在内核启动时，CUDA 运行时现在知道：

• 给定内核需要多少个寄存器（每个线程）
• 我们在什么设备上运行，因此总限制是多少
• 启动配置是什么（块、线程）

组合这 3 条信息意味着运行时可以立即知道是否有或将有足够的“寄存器空间”用于启动。粗略地说，通过/失败算法是发射是否满足这个不等式：

 registers_per_thread*threads_per_block <= max_registers_per_multiprocessor

在这个等式中也需要考虑粒度。寄存器通常在运行时以 2 或 4 个为一组分配，即在应用不等式测试之前，registers_per_thread 数量可能需要四舍五入到下一个整数倍数，例如 2 或 4。 registers_per_thread 数量由编译器确定，如前所述。 threads_per_block 数量来自您的内核启动配置。 max_registers_per_multiprocessor 数量是机器可读的（即它是您正在运行的 GPU 的函数）。如果您愿意，可以通过研究deviceQuery CUDA 示例代码了解如何自己检索该数量。

3）既然有寄存器溢出机制，那么即使没有足够的寄存器，难道不应该启动所有的CUDA内核吗？

我重申，寄存器分配（和寄存器溢出决定）完全是一个静态编译时过程。不进行运行时决策或更改。寄存器分配完全可以从编译的代码中检查。因此，由于在运行时无法进行任何调整，因此无法进行任何更改以允许任意启动。任何此类更改都需要重新编译代码。虽然这在理论上可能是可行的，但目前尚未在 CUDA 中实现。此外，它有可能导致可变且可能不可预测的行为（在性能方面），因此可能有理由不这样做。

通过适当地限制编译器对寄存器分配的选择，可以使所有内核“可启动”（关于寄存器限制）。 __launch_bounds__ 和 compiler switch -maxrregcount 是实现这一目标的几种方法。 CUDA 提供occupancy calculator 和occupancy API 来帮助完成此过程。