CUDA_寄存器和局部存储器

寄存器

GPU片上高速缓存器，基本单元时寄存器文件，每个寄存器文件大小为32bit。

计算能力1.0/1.1版本硬件中，每个SM中寄存器文件数量为8192；而在1.2/1.3硬件中，每个SM中寄存器文件数量为16384.

一般情况下，内核线程中的简单局部变量都在寄存器存储器中。

局部存储器

对于每个线程，局部存储器也是私有的，如果寄存器被消耗完，数据将被存储在局部存储其中。如果每个线程使用了过多的寄存器，或声明了大型结构体或数组，或者编译器无法确定数组的大小，线程的私有数据就有可能会被分配到local memory中。一个线程的输入和中间变量将被保存在寄存器或者局部存储器中。局部存储器中的数据被数保存在显存中，因此对local memory的访问速度很慢。

如下，mt会被存入local memory中

__global__ void localmemDemo(float* A)
{
  unsigned int mt[3];
}

如果在定义线程私有数组的同时，对其就进行了初始化，那么如果数组尺寸不大，这个数组仍然有可能被分到寄存器中。

__global__ void localmemDemo(float* A)
{
  unsigned int mt[3] = {1, 2, 3};
}

在编译时，输出ptx（parallel thread execution）汇编代码（在编译时加上-ptx或者-keep选项），就能观察变量在编译的第一阶段是否被分配到了local memory中，如果一个变量在ptx中以.local助记符声明，可使用ld.local和st.local助记符访问，这个变量就被存放在local memory中。不过，即使初次编译的变量不位于local memory中，在编译的第二阶段仍然有可能根据目标硬件中存储器的大小将变量存放在local memory中。这时，可以通过加上--ptxas-options=-v编译选项用来观察lmem的使用情况。

如果数组较小，并且一定要分配在寄存器中，用以下方法：

__global__ void localmemDemo(float* A)
{
  unsigned int mt0, mt1, mt2;
}