当我用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
之类的东西?