ABI如何定义GPU中的寄存器数量？答案

【问题标题】：How does the ABI defines the number of registers in GPU?ABI如何定义GPU中的寄存器数量？
【发布时间】：2013-03-01 15:09:32
【问题描述】：

CUDA Compiler Driver NVCC - Options for steering GPU code generation 中有一行对我来说模棱两可：

小于 ABI 要求的最小寄存器的值将被编译器提高到 ABI 最小限制。

对于__global__ 和__device__ 函数使用的寄存器数量，ABI 是否有任何标准或限制？

【问题讨论】：

我认为（现在找不到参考资料）ABI 至少需要 16 个寄存器。因此，如果您指定较低的寄存器计数（例如，使用 -maxrregcount），编译器会将指定的限制提高到 ABI 要求的最小值，并打印一条建议消息，说明它这样做了。至于每个线程可用的最大 32 位寄存器数量，取决于 GPU 架构：sm_1x 有 124 个寄存器，sm_2x 有 63 个寄存器，sm_3x 有 254 个寄存器。
但是 ABI 与 GPU 寄存器有什么关系？它甚至在 CUDA 诞生之前就已经存在了吗？
一般来说，ABI（应用程序二进制接口）是针对存储布局、向函数传递参数、将函数结果返回给调用者等架构特定的约定。ABI（包括 x86_64、ARM ) 通常为特定任务指定特定寄存器，例如堆栈指针、函数返回值、函数参数等。由于 GPU 架构允许每个线程的寄存器数量可变，因此使用 ABI 需要存在最少数量的寄存器来填充这些寄存器定义的角色。如果我没记错的话，CUDA 引入了 2.0 版的 ABI。
好的，所以 CUDA 有一个 ABI，很好，如果你能找到参考资料（可能是 CUDA ABI），那就太好了，但尽管发布一个答案，我会接受。
谢谢，我提供了上面的扩展版本作为答案，并更正了将 ABI 从 2.0 引入到 3.0 的 CUDA 版本（经过几年和许多 CUDA 版本很容易混淆)。

标签： cuda

【解决方案1】：

我认为（现在找不到参考）CUDA ABI 至少需要 16 个寄存器。因此，如果您指定较低的寄存器计数（例如，使用 -maxrregcount），编译器会将指定的限制提高到 ABI 要求的最小值，并打印一条建议消息，说明它这样做了。至于每个线程可用的最大 32 位寄存器数量，取决于 GPU 架构：sm_1x 有 124 个寄存器，sm_2x 有 63 个寄存器，sm_3x 有 254 个寄存器。

一般来说，ABI（应用程序二进制接口）是针对存储布局、将参数传递给函数、将函数结果返回给调用者等架构的特定约定。ABI（包括 x86_64、ARM）通常指定特定的寄存器对于特定任务，例如堆栈指针、函数返回值、函数参数等。由于 GPU 架构允许每个线程的寄存器数量可变，因此使用 ABI 需要存在最少数量的寄存器来填充这些定义的角色。如果我没记错的话，CUDA 引入了 3.0 版的 ABI，这是第一个支持 Fermi 级 GPU 的版本。

ABI 需要 2.0 或更高的计算能力。较旧的 GPU 架构缺乏 ABI 所需的硬件功能。大多数较新的 CUDA 功能，例如设备端 printf() 和 malloc()，调用函数，单独编译等都依赖并需要使用 ABI，并且默认情况下在编译器生成的代码中使用它为 sm_20 和多于。您可以通过 -Xptxas -abi=no 禁用 ABI。我强烈建议不要这样做。

【讨论】：