取决于编译计算能力的寄存器和共享内存？答案

【问题标题】：Registers and shared memory depending on compiling compute capability?取决于编译计算能力的寄存器和共享内存？
【发布时间】：2011-08-27 15:35:35
【问题描述】：

你好，当我使用nvcc -arch=sm_13 编译时，我得到：

ptxas info    : Used 29 registers, 28+16 bytes smem, 7200 bytes cmem[0], 8 bytes cmem[1]

当我使用nvcc -arch=sm_20 时，我得到：

ptxas info    : Used 34 registers, 60 bytes cmem[0], 7200 bytes cmem[2], 4 bytes cmem[16]

我以为所有内核参数都传递到共享内存，但对于 sm_20，似乎并非如此......？！也许它们也被传递到寄存器中？我的函数的头部如下所示：

__global__ void func(double *, double , double, int)

到目前为止谢谢！

【问题讨论】：

标签： compiler-construction cuda nvcc

【解决方案1】：

正如@talonmies 所说，共享内存差异是由于 SM 2.x 设备通过常量而不是共享内存传递内核参数。

然而，SM 2.x 设备中寄存器使用的主要区别之一是，虽然 SM 1.x 设备具有用于加载和存储指令的专用地址寄存器，但 SM 2.x 使用通用寄存器来存储地址。这往往会增加 SM 2.x 上的套准压力。幸运的是，GF100 (SM 2.0) 上的寄存器文件也比 GT200 (SM 1.3) 大 2 倍。

【讨论】：

马克，虽然寄存器文件确实是 Fermi 上的两倍大小，但每个线程的最大寄存器数也是 sm 1.x 设备的一半。 64 个寄存器限制经常在我使用的代码中发挥作用。我知道 Vasily Volkov 还指出了寄存器压力对代码在 Fermi 卡上实现高水平指令级并行性的能力的影响。
是的，Fermi 将每个 SM 的最大线程数提高了 50% (1536/1024)，删除了地址寄存器，将每个线程的最大寄存器数减半，每个 SM 的寄存器文件大小加倍，并略微增加了共享内存访问延迟。这些都增加了对资源使用的敏感性。将寄存器文件加倍会有所帮助，但你是对的，它仍然很紧。但这是大规模并行处理器的现实。当我们设计一个新的 GPU 时，总是需要考虑这样的权衡。不过总的来说，Fermi 在架构上的效率要比特斯拉架构高得多。

【解决方案2】：

在计算能力 2.x 设备中，内核参数存储在常量内存中。寄存器差异可能归结为版本之间为数学库函数生成的代码的差异。内核中是否有超越函数或sqrt之类的东西？

【讨论】：

首先感谢参数，你有这方面的资料吗？进一步：不，没有数学函数！
作为参考，它可能在 ptx 指南中，但我只能说“一位 NVIDIA 开发人员告诉我是这样的”。寄存器中另一个可能的区别是指针。指针在计算 1.x 设备上使用 1 个寄存器，在 Fermi 上使用 2 个寄存器，内部是 64 位