【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路訪問沖突

 


免責聲明!

本站轉載的文章為個人學習借鑒使用,本站對版權不負任何法律責任。如果侵犯了您的隱私權益,請聯系本站郵箱yoyou2525@163.com刪除。



 
粵ICP備18138465號   © 2018-2025 CODEPRJ.COM