您的三个问题的简短答案是:
- 是的。
- 是的,如果代码是为 64 位主机操作系统编译的。设备指针大小始终与 CUDA 中的主机应用程序指针大小匹配。
- 不。
为了扩展第 3 点,请考虑以下两个简单的内存复制内核:
__global__
void debunk(float *in, float *out, int n)
{
int i = n * (threadIdx.x + blockIdx.x*blockDim.x);
for(int j=0; j<n; j++) {
out[i+j] = in[i+j];
}
}
__global__
void debunk2(float *in, float *out, int n)
{
int i = n * (threadIdx.x + blockIdx.x*blockDim.x);
float *x = in + i;
float *y = out + i;
for(int j=0; j<n; j++, x++, y++) {
*x = *y;
}
}
根据您的估计,debunk
必须使用较少的寄存器,因为它只有两个局部整数变量,而debunk2
有两个额外的指针。然而,当我使用 CUDA 5 发布工具链编译它们时:
$ nvcc -m64 -arch=sm_20 -c -Xptxas="-v" pointer_size.cu
ptxas info : 0 bytes gmem
ptxas info : Compiling entry function '_Z6debunkPfS_i' for 'sm_20'
ptxas info : Function properties for _Z6debunkPfS_i
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 8 registers, 52 bytes cmem[0]
ptxas info : Compiling entry function '_Z7debunk2PfS_i' for 'sm_20'
ptxas info : Function properties for _Z7debunk2PfS_i
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 8 registers, 52 bytes cmem[0]
它们编译为完全相同的寄存器计数。如果你反汇编工具链输出,你会发现除了设置代码之外,最终的指令流几乎是相同的。造成这种情况的原因有很多,但基本上可以归结为两个简单的规则:
- 试图从 C 代码(甚至 PTX 汇编程序)中确定寄存器计数大多是徒劳的
- 试图重新猜测一个非常复杂的编译器和汇编器也大多是徒劳的。