CUDA Compiler Driver NVCC - Options for steering GPU code generation中有一行这对我来说模棱两可:
Value less than the minimum registers required by ABI will be bumped up by the compiler to ABI minimum limit.
ABI 对于 __global__ 和 __device__ 函数使用的寄存器数量有任何标准或限制吗？
最佳答案
我认为(现在找不到引用)CUDA ABI 至少需要 16 个寄存器。因此，如果您指定较低的寄存器计数(例如使用 -maxrregcount)，编译器会将指定的限制提高到 ABI 所需的最小值，并打印一条建议消息，说明它已这样做。至于每个线程可用的 32 位寄存器的最大数量，取决于 GPU 架构:sm_1x 为 124 个寄存器，sm_2x 为 63 个寄存器，sm_3x 为 254 个寄存器。
一般来说，ABI(应用程序二进制接口(interface))是一种特定于体系结构的约定，用于存储布局、向函数传递参数、将函数结果传递回调用者等。ABI(包括 x86_64、ARM)通常指定特定寄存器用于特定任务，例如堆栈指针、函数返回值、函数参数等。由于 GPU 架构允许每个线程使用可变数量的寄存器，因此使用 ABI 需要使用最少数量的寄存器来填充这些定义的角色。如果我没记错的话，CUDA 在 3.0 版本中引入了 ABI，这是第一个支持 Fermi 级 GPU 的版本。
ABI 需要计算能力 2.0 或更高。较旧的 GPU 架构缺乏 ABI 所需的硬件功能。大多数较新的 CUDA 功能，例如设备端 printf() 和 malloc()、调用函数、单独编译等都依赖并要求使用 ABI，并且默认情况下在编译器生成的 sm_20 和 sm_20 代码中使用它多于。您可以通过 -Xptxas -abi=no 禁用 ABI。我强烈建议不要这样做。
关于cuda - ABI如何定义GPU中的寄存器数量？，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/15160859/