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