CUDA寄存器使用情况

Question

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

1. Double variable takes two registers?

2. Pointer variable takes two registers? - It has to be more than one register on Fermi with 6 GB memory, right?

3. If answer to question 2 is yes, it must be better to use less pointer variables and more int indices.

   E. g., this kernel code:

   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]
   }
   

   theoretically requires more registers than this kernel code:

   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]
   }
   

- user2052436

Answer 1

针对你的三个问题，简短回答如下:

1. 是的。
2. 是的，如果代码针对64位主机操作系统进行编译。在CUDA中，设备指针大小始终与主机应用程序指针大小匹配。
3. 不是。

为了扩展第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]

它们编译成完全相同的寄存器数。如果您反汇编工具链输出，您会发现除了设置代码外，最终指令流几乎相同。有许多原因，但基本上归结为两个简单的规则：

1. 从C代码（甚至是PTX汇编程序）中确定寄存器计数大多是徒劳的
2. 试图猜测一个非常复杂的编译器和汇编器也大多是徒劳的。

- talonmies

---

请问您能否解释一下，为什么第三个是“不”？ - user2052436

在简单情况下，优化器是否会生成相同/类似的代码？实际上，我有一个处理3D数组并迭代i0、i1、i2索引的有限差分代码。通常情况下，我需要从当前点p[i]向三个方向中的一个方向移动，其中i=i0+i1*stride1+i2*stride2。因此，如果我引入指针px1=p+1、py1=p+stride1、pz1=p+stride2（可能还有更多-px2=p+2等），并使用p[i]、px1[i]等进行操作，则代码会更清晰。如果优化器无法优化掉所有这些额外的指针，那么这样做会增加寄存器使用量吗？ - user2052436

所以我的问题是：在复杂的内核代码中，引入额外的指针而不是整数索引是否安全？也就是说，是否有可能会导致更高的寄存器使用率，因为对于复杂的代码，优化器实际上将使用具有2个寄存器（用于指针变量）而不是1个寄存器（用于整数索引）的二进制代码？ - user2052436

我建议你编译和反汇编演示内核，并学习它们，直到你理解它们的作用。这个简单的操作基本上需要一个源地址和目标地址寄存器，一个数据寄存器用于值（这里没有间接寻址），一个数据寄存器用于循环计数以及一对寄存器用于评估循环条件。那就是8个寄存器。整数索引在这里是无关紧要的，也不能神奇地保存寄存器：*p+i和p[i]生成相同的代码。 - talonmies