下面簡單介紹一些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路訪問衝突