【CUDA學習】共享存儲器

下面簡單介紹一些cuda中的共享存儲器和全局存儲器　

共享存儲器，shared memory，能夠被同一塊中的全部線程訪問的可讀寫存儲器，生存期是塊的生命期。

Tesla的每一個SM擁有16KB共享存儲器。

在編程過程當中，有靜態的shared memory 動態的shared memory

靜態的shared memory 在程序中定義　　　__shared__ type shared[SIZE];

動態的shared memory 經過內核函數的每三個參數設置大小 extern __shared__ type shared[];

共享存儲器被組織爲16個bank，每一個bank擁有32bit的寬度。

無bank conflict時，一個half-warp內的線程能夠在一個內核週期中並行訪問

對同一bank的同時訪問致使bank conflict 　　只能順序處理　訪存效率下降

若是half-warp的線程訪問同一地址時，會產生一次廣播，不會產生bank conflict

__shared__ float shared[256];

float foo = shared[threadIdx.x];

沒有訪問衝突

__shared__ float shared[256];

float foo = shared[threadIdx.x * 2];

產生2路訪問衝突

__shared__ float shared[256];

float foo = shared[threadIdx.x*8];

產生8路訪問衝突