寄存器和共享内存,具体取决于编译计算能力



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

到目前为止谢谢!

正如@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 倍。

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

最新更新