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保存在寄存器中，同时

• nvcc-Xptxas-v gives 4 registers, 24+16 bytes smem

• profiler显示4个寄存器

• 和头ptx file:

  .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资源足以满足并发内核的需要。

我不确定我是否正确。

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

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

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