哪些变量会消耗 CUDA 中的寄存器？

Posted 2023-02-16

【中文标题】哪些变量会消耗 CUDA 中的寄存器？

【英文标题】：What kind of variables consume registers in CUDA?

【发布时间】：2012-07-14 11:18:11

【问题描述】：

__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 registers, 24+16 bytes 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 资源足够并发内核。

我不确定我是否正确。

【问题讨论】：

你的问题是“为什么这段代码使用 4 个寄存器而不是 3 个？”如果是这样，答案是：为了添加a[offset] 和b[offset]，必须获取这两个值。它必须将它首先获取的任何一个存储在某个地方，同时它正在获取另一个。所以还需要一个寄存器。

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

必要时可以。判断何时需要这样做并不总是那么容易，甚至会因硬件目标而异。

知道了：-P 因为寄存器的使用很复杂，有没有办法找到内核寄存器使用的边界，因为我想测试寄存器溢出的情况，但是当我尝试声明更多变量，寄存器用法保持不变。

【参考方案1】：

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

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

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

【讨论】：

非常感谢，talonmies。所以我想我们对内核中的寄存器使用控制无能为力吗？编译器总是做很多事情。

您有__launch_bounds__ 向编译器提供启发式提示，这些提示可能会影响寄存器分配，并且编译器/汇编器采用-maxrregcount 参数。 volatile 关键字用于对旧版本的 nvopen64 编译器产生影响，并可能影响本地内存溢出行为。但是你不能在原始的 C 代码中任意控制或引导寄存器分配。