CUDA 寄存器使用

Question

CUDA手册指定数量 32 位 每个多处理器的寄存器。是否意味着:

双变量需要两个寄存器？

指针变量需要两个寄存器？ - 在 Fermi 上必须有 6 GB 内存的多个寄存器，对吗？

如果问题 2 的答案是肯定的，那么使用更少的指针变量和更多的指针变量肯定更好int指数。

例如，这个内核代码:

float* p1;               // two regs
float* p2 = p1 + 1000;   // two regs
int i;                   // one reg
for ( i = 0; i < n; i++ )
{
    CODE THAT USES p1[i] and p2[i]
}

理论上需要比这个内核代码更多的寄存器:

float* p1;               // two regs
int i;                   // one reg
int j;                   // one reg
for ( i = 0, j = 1000; i < n; i++, j++ )
{
    CODE THAT USES p1[i] and p1[j]
}

Answer 1

对您的三个问题的简短回答是:

是的。

是的，如果代码是为 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 汇编程序)中确定寄存器数量几乎是徒劳的

试图对非常复杂的编译器和汇编器进行二次猜测也是徒劳的。