什么样的变量消耗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 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资源足以用于并发内核.

我不确定我是不对的.

Answer 1

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

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

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