compiler-construction - 寄存器和共享内存取决于编译计算能力？

Question

你好，当我用 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)

到目前为止谢谢!

Answer 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 倍。