在CUDA中哪些类型的变量会占用寄存器？

Question

__global__ void add( int *c, const int* a, const int* b )
{
    int x = blockIdx.x;
    int y = blockIdx.y;
    int offset = x + y * gridDim.x;
    c[offset] = a[offset] + b[offset];
}

在上面的示例中，我猜想x、y和offset被保存在寄存器中，而nvcc -Xptxas -v显示有4个寄存器，24+16字节的smem，profiler也显示了4个寄存器，ptx文件的开头是：

.reg .u16 %rh<4>;
.reg .u32 %r<9>;    
.reg .u64 %rd<10>;  
.loc    15  21  0   

$LDWbegin__Z3addPiPKiS1_:   
.loc    15  26  0  

能否有人澄清寄存器的使用？在Fermi中，每个线程的最大寄存器数量为63。在我的程序中，我想测试当内核消耗过多寄存器时的情况（因此变量可能必须自动存储在本地内存中，从而导致性能降低）。然后在这一点上，我可以将一个内核拆分成两个，使每个线程都有足够的寄存器。假设SM资源足以支持并发内核。

我不确定我是否正确。

- user1525320

---

您的问题是"为什么这段代码要使用4个寄存器而不是3个?" 如果是这样，答案是：为了将 a[offset] 和 b[offset] 相加，必须获取这两个值。在获取另一个值时，必须将其先获取的值存储在某个地方。因此需要一个额外的寄存器。 - David Schwartz

谢谢您的回答，那么我们可以说中间变量将会保存在寄存器中吗？ - user1525320

如果必要的话，是的。但是很难确定何时需要，并且这甚至可能因硬件目标而异。 - David Schwartz

既然寄存器使用是复杂的，有没有一种方法可以找到内核寄存器使用的边界，因为我想测试寄存器溢出的情况，但当我尝试声明更多变量时，寄存器使用仍然保持不变。 - user1525320

Answer 1

PTX的寄存器分配与内核最终寄存器消耗完全无关。 PTX只是最终机器代码的中间表示，并使用静态单赋值形式，这意味着在PTX中的每个寄存器仅被使用一次。 即使PTX代码有数百个寄存器，也可以编译成只有几个寄存器的内核。

寄存器分配由ptxas作为完全独立的编译步骤完成（静态或由驱动程序进行即时处理，或两者兼备），并且它可以对输入的PTX执行大量的代码重新排序和优化来提高吞吐量并节省寄存器，这意味着原始C语言变量或PTX寄存器与汇编后的内核中的最终寄存器数量之间很少或没有关系。

nvcc提供了一些影响汇编器寄存器分配行为的方法。您可以使用__launch_bounds__来提供启发式提示以影响寄存器分配，并且编译器/汇编器采用-maxrregcount参数（可能会导致寄存器溢出到本地内存，这可能会降低性能）。 volatile关键字在旧版的nvopen64编译器中曾经有所影响，并可以影响局部内存溢出行为。但是您无法任意控制或指导原始C代码或PTX汇编语言代码的寄存器分配。

- talonmies

---

非常感谢，talonmies。所以我猜我们在内核中无法控制寄存器的使用？编译器总是做很多事情。 - user1525320

你可以使用__launch_bounds__来为编译器提供启发式提示，从而影响寄存器分配。编译器/汇编器还可以使用-maxrregcount参数。volatile关键字在早期版本的nvopen64编译器中有所不同，并且可能会影响本地内存溢出行为。但是，你不能在原始C代码中任意控制或引导寄存器分配。 - talonmies