什么样的变量消耗CUDA中的寄存器？

__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被保存在寄存器中而什么样的变量消耗CUDA中的寄存器？

NVCC -Xptxas -v给出4 registers, 24+16 bytes smem
分析器显示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资源足够用于并发内核。

我不知道我是否正确。

来源

2012-07-14 user1525320

是你的问题“为什么这个代码使用4个寄存器而不是3个？”如果是这样，答案是这样的：为了添加'a [offset]'和'b [offset]'，必须获取这两个值。它必须存储它在第一次获取到的第一个地方，而它正在获取另一个地方。所以还需要一个寄存器。 – 2012-07-14 12:20:43

谢谢您的回答，那么我们可以说中间变量将被保存在寄存器中吗？ – user1525320 2012-07-14 12:37:46

必要时，是的。在需要的时候并不总是很容易辨别，甚至可能因硬件目标而异。 – 2012-07-14 12:38:15

PTX中的寄存器分配与内核的最终寄存器消耗完全无关。 PTX仅是最终机器代码的中间代表，并使用static single assignment form，这意味着PTX中的每个寄存器仅使用一次。一块带有数百个寄存器的PTX可以编译成只有少数寄存器的内核。

寄存器分配由ptxas完成，作为完全独立的编译过程（驱动程序静态地或即时地执行），并且可以对输入PTX执行大量代码重新排序和优化以提高吞吐量并保存寄存器，这意味着PTX中的原始C或寄存器中的变量与组装好的内核的最终寄存器计数之间几乎没有关系。

nvcc确实提供了一些方法来影响汇编程序的寄存器分配行为。您有__launch_bounds__向编译器提供了启发式提示，这会影响寄存器分配，编译器/汇编程序采用-maxrregcount参数（可能会导致寄存器溢出到本地内存，从而降低性能）。 volatile关键字用于对旧版本的基于nvopen64的编译器产生影响，并可能影响本地内存溢出行为。但是，您不能任意控制或引导原始C代码或PTX汇编语言代码中的寄存器分配。

来源

2012-07-14 13:15:23 talonmies

非常感谢，talonmies。所以我想我们的内核中没有任何关于寄存器使用控制的事情？编译器总是做很多事情。 – user1525320 2012-07-14 13:38:41

你有'__launch_bounds__'向编译器提供启发式提示，它可以影响寄存器分配，编译器/汇编器则采用'-maxrregcount'参数。 'volatile'关键字用于与旧版本的nvopen64编译器有所不同，并可能影响本地内存溢出行为。但是你不能任意控制或引导原始C代码中的寄存器分配。 – talonmies 2012-07-14 14:23:34

这有助于很多！再次感谢男士。 – user1525320 2012-07-14 14:48:29

什么样的变量消耗CUDA中的寄存器？

回答

相关问题