cuda - 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

• 分析器显示 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资源足以满足并发内核的需要。

我不确定我是否正确。

Answer 1

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

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

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