0

我认为 CUDA 会尝试在寄存器中分配标量变量，而在 Fermi 类 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将阵列a返回给主机。我现在看到每个线程都使用了所有 63 个寄存器：ptxas info : Used 62 registers, 40 bytes cmem[0]. 尽管我传递的变量多于寄存器无法容纳的变量，但我没有看到任何溢出到本地内存；我认为 NVCC 正在优化代码以仅使用寄存器。

4

1 回答 1

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 转储内核的汇编代码。

于 2012-09-22T00:11:43.787 回答