标量变量和寄存器：CUDA

Question

我觉得CUDA将会尝试在寄存器分配标量，并在费米级GPU，每个线程有63个寄存器。 我的代码是这样的：

__global__ void test20 (double a) 
{ 
    double i1=1.0; 
    double i2=2.0; 
    double i3=3.0; 
    double i4=4.0; 
    double i5=5.0; 
    double i6=6.0; 
    double i7=7.0; 
    double i8=8.0; 
    double i9=9.0; 
    double i10=10.0; 
    ... 

    a = i1+i2+i3 ... i20 
} 

但是，当我看到使用NVVP每线程寄存器的数量，我只看到2个寄存器每个线程分配的，当我希望更高的数字。即使我将变量减少到10，分配的寄存器数量也保持不变。为什么会发生这种情况，如何确保我有n个变量，CUDA使用n个寄存器（考虑到每个变量可以存储在一个寄存器中）？

编辑：

继建议，我已经修改了代码：

__global__ void test (double *a) 
{ 
    double reg1; 
    double reg2; 
    double reg3; 
    double reg4; 
    double reg5; 
    double reg6; 
    double reg7; 
    double reg8; 
    ....till 40 
    reg1 = log10f(a[0]); 
    reg2 = log10f(a[1]); 
    reg3 = log10f(a[2]); 
    reg4 = log10f(a[3]); 
    reg5 = log10f(a[4]); 
    reg6 = log10f(a[5]); 
    reg7 = log10f(a[6]); 
    reg8 = log10f(a[7]); 
    reg9 = log10f(a[8]); 
    ....till 40 
    a[0] = reg1; 
    a[1] = reg2; 
    a[2] = reg3; 
    a[3] = reg4; 
    a[4] = reg5; 
    a[5] = reg6; 
    a[6] = reg7; 
    a[7] = reg8; 
    } 

我memcpy -ing阵列a回到主机。我现在看到每个线程使用全部63个寄存器：ptxas info : Used 62 registers, 40 bytes cmem[0]。虽然我通过的变量比寄存器中可以容纳的要多得多，但是我没有看到任何溢出到本地内存中;我认为NVCC正在优化代码以仅使用寄存器。

Answer 1

如果按照@talonmies建议，使用不能在运行时计算的表达式，你可能仍然无法获得每一个申报登记（或在这种情况下，2个寄存器来保存双）。您可能还必须保持该变量在整个持续时间内保持活动状态。

__global__ void test20 (double a) 
{ 
    double i1=1.0 * a; 
    double i2=2.0 * i1; 
    double i3=3.0 * i2; 
    double i4=4.0 * i3; 
    double i5=5.0 * i4; 

    a = i1+i2+i3+i4+i5; 

    printf("a = %f = %f + %f + %f + %f + %f\n", a, i1, i2, i3, i4, i5); 
} 

这是在浏览器中编写的示例代码。目标是保留寄存器中的值。此示例没有实际应用，因为编译器的目标是使用最少的寄存器。唯一的价值在于调试使变量在其范围内保持活动状态。

如果您想了解寄存器的使用，你应该使用cuobjump -sass转储为内核的汇编代码。