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;
}