共享儲存器優化

yyfn風辰發表於2010-06-24
優化
3.4  共享儲存器優化
        共享儲存器比較靠近ＳＭ的，因此其速度相當的快，一般而言可以在一到二個時鐘內讀寫，因此使用共享儲存器取代全域性儲存器會極大的節約頻寬，但是共享儲存器 的使用是有要求，它要求資料有區域性性（一個ｂｌｏｃｋ內共享）重用。當然有時也可用共享儲存器取代暫存器，這會在後面說到。
        一般的使用共享儲存器的方式是：s_x[tid] = x[id];其中tid是塊內執行緒的塊內索引，也就是類似threadIdx，id是執行緒的網格內索引。然後可以訪問ｓ＿ｘ取代訪問ｘ。這樣一般可以提升 程式的效率。
        要注意的是，使用共享儲存器往往伴隨著儲存器的一致性問題，此時可求助於__syncthreads()和memory fence大神。

3.4.1  儲存體衝突
        類似於合併訪問，共享儲存器一次也能滿足　半束執行緒的要求，條件是這半束執行緒訪問的資料在不同的儲存體中。CUDA將共享儲存器組織成16列，每一列稱為 一個儲存體。如果有兩個或以上的執行緒訪問的資料在同一個儲存體中，此時不能在一次儲存體訪問中，滿足半束執行緒的要求，這稱為儲存體衝突。一般而言，儲存體 衝突對效能的影響還不是太大，當然具體的影響要看半束內最多有多少執行緒落入同一儲存體內。要提到的程式設計時要注意的是對於一個ＳＭ來說，共享儲存器的數量是有限的，如果超出其使用限制，其結果未知，我的經驗告訴我，有時沒有問題，有時會有大問題。另外， 使用共享儲存器有時可能會影響程式的可擴充套件性，這個在編寫庫程式碼的時候要特別注意。

來自 “ ITPUB部落格 ” ，連結：http://blog.itpub.net/23057064/viewspace-666187/，如需轉載，請註明出處，否則將追究法律責任。

相關文章