1) 内核什么时候开始将寄存器溢出到本地内存？
这完全在编译器的控制之下。它不是由运行时执行的，并且没有关于它的动态运行时决策。当您的代码达到溢出点时，这意味着编译器已插入如下指令：
STL [R0], R1
在这种情况下，R1被存储到本地内存，本地内存地址给出R0。这将是一个溢出商店。 （在那条指令之后，R1当然，编译器知道它何时完成此操作，因此它可以报告它选择使用/制作的溢出加载和溢出存储的数量。您可以使用以下方式获取此信息（以及寄存器使用情况和其他信息）：-Xptxas=-v编译器开关。
编译器（除非您对其进行限制，请参见下文）对寄存器使用做出的决定主要关注性能，否则较少关注实际使用了多少寄存器。第一要务是性能。
2）当没有足够的寄存器时，CUDA运行时如何决定不启动内核并抛出太多资源请求错误？多少个寄存器足以启动内核？
在编译时，当编译内核代码时，编译器不知道它将如何启动。它不知道您的启动配置是什么样的（块数、每个块的线程数、动态分配的共享内存量等）。事实上，编译过程大部分都在进行，就好像正在编译的东西是单个线程一样。
在编译期间，编译器会做出一系列有关寄存器分配的静态决策（如何以及在何处使用寄存器）。 CUDA有二进制实用程序这有助于理解这一点。寄存器分配在运行时不会改变，也不是动态的，因此完全在编译时确定。因此，在完成给定设备代码函数的编译时，通常可以确定需要多少个寄存器。编译器将此信息包含在二进制编译对象中。
在运行时，在内核启动时，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 的函数）。如果您愿意，您可以通过研究如何自行检索该数量deviceQueryCUDA 示例代码。
3）既然存在寄存器溢出机制，即使没有足够的寄存器，难道不应该启动所有CUDA内核吗？
我重申，寄存器分配（和寄存器溢出决策）完全是一个静态编译时过程。不进行运行时决策或更改。寄存器分配完全可以从编译的代码中检查。因此，由于无法在运行时进行调整，因此无法进行任何更改以允许任意启动。任何此类更改都需要重新编译代码。虽然这在理论上是可能的，但目前尚未在 CUDA 中实现。此外，它有可能导致可变且可能不可预测的行为（在性能方面），因此可能有理由不这样做。
通过适当限制编译器对寄存器分配的选择，可以使所有内核“可启动”（相对于寄存器限制）。__launch_bounds__和编译器开关 -maxrregcount有几种方法可以实现这一目标。 CUDA 提供了占用计算器以及一个占用API来帮助完成这个过程。
