CUDA GPU程式設計如何避免Bank conflict
CUDA GPU程式設計如何避免Bank conflict
強烈推薦參考書：Professional CUDA C Programming，本書第234頁開始對bank conflict有非常詳盡的講解，本文正是從本書的學習和自己的程式設計經驗中總結而來。
1 為什麼要避免Bank conflict
Shared memory在晶片SM（Streaming Multiprocessors）內部，相比片外的Global memory擁有大得多的記憶體頻寬，當對該記憶體資料讀寫操作頻繁，則建議先將Global memory中的資料載入到Shared memory再進行讀寫操作，會大大提高程式計算效能。然而，Shared memory和諸如if else語句引起的warp divergence類似，當操作不當時，會導致程式效能的大大降低
2 什麼是Bank conflict以及Bankconflict的產生
為了提高記憶體讀寫頻寬，共享記憶體被分割成了32個等大小的記憶體塊，即Bank。因為一個Warp有32個執行緒，相當於一個執行緒對應一個記憶體Bank。這個分割方法通常是按每4個位元組一個bank，計算能力3.x的GPU也可以8個位元組一個bank，如圖1所示。使用者建立的共享記憶體就按照地址依次對映到這些bank中。
圖1 共享記憶體bank（Professional CUDA C Programming p237）
理想情況下就是不同的執行緒訪問不同的bank，可能是規則的訪問，如執行緒0讀寫bank0，執行緒1讀寫bank1，也可能是不規則的，如執行緒0讀寫bank1，執行緒1讀寫bank0。這種同一個時刻每個bank只被最多1個執行緒訪問的情況下不會出現Bank conflict。特殊情況如果有多個執行緒同時訪問同一個bank的同一個地址的時候也不會產生Bank conflict，即broadcast
Bank conflict產生後，同一個bank的記憶體讀寫將被序列化，而寫入同一個地址時將只有其中一個執行緒能夠成功寫入（要使得每個執行緒都能成功寫入，需要使用原子操作atomic Instructions）。
3 如何避免產生Bank conflict
Bank conflict主要出現在Global memory與Shared memory的資料交換，以及裝置函式對Shared memory的操作中。
Global memory與Shared memory的資料交換中，最好是每次32個執行緒讀寫32個連續的word。這樣既可以滿足不發生Bank conflict，又滿足了Global memory的Coalesced Access。如shared merory為64個字，則在迴圈中，第一次，thread0 讀0word，thread1 讀1word，... ；第二次，thread0 讀32word，thread1 讀33word，... 。
裝置函式對Shared memory的操作中則有很多注意點，需要仔細分析自己的程式，特別是不規律訪問時。當每個執行緒只訪問自己專屬的資料時，當每個執行緒只有1個word的共享記憶體記憶體時則不會出現衝突。但當每個執行緒保有一個向量或矩陣時，則需要仔細分析。一個M行N列的二維陣列（矩陣）在記憶體中也是連續存放的，因而和一維陣列沒有本質區別，可以看成一個N*M長度的一維陣列。後面都從一維陣列來討論。
例如一個執行緒塊有32個執行緒，每個執行緒有一個長度為6的陣列（向量）。則這個陣列可以有3種宣告的方式：
方式1， 一維陣列方式：__shared__ int Vector1[32*6];
方式2， 二維陣列__shared__ int Vector1[32][6];
方式3， 二維陣列__shared__ int Vector1[6][32];
方式1的情況下，看自己如何對這個一維陣列進行分割，如果每6個連續的字為一個向量則結果和方式2儲存方式相同。方式2，3比較直觀，但是可能具有不同的效能，他們儲存方式如圖2。
圖2 陣列不同的儲存方式
當每個執行緒同時訪問自己向量的第一個元素時，按方式2儲存則每個執行緒訪問的字地址將為：tid*6+0，對應bank為(tid*6+0)%32，tid為執行緒索引threadIdx.x。可以檢查發現將出現Bank conflict。如執行緒0和16，1和17等出現衝突。而按方式3儲存，則顯然每個向量第一個元素都儲存在不同的bank中，不會引起Bank conflict。因而改變資料儲存方式是一種避免Bank conflict的方式。
但按方式3儲存最大的缺點就是每個執行緒保有的陣列元素被離散儲存了，某些情況下對程式設計造成了很大的不便，而方式2的優點正在於每個陣列的元素都是連續儲存的，特別是當執行緒保有的是一個矩陣時，會帶來巨大的便利。在上面的例子中，如果Vector1其實是一個2行3列的矩陣，則對於方式2儲存時，可以用指標將一維陣列轉換為二維陣列的訪問：
int(*pD0)[3]=( int (*)[3])& Vector1 [tid][0];
然後就可以使用諸如pD0[1][2]的方式來替代Vector1進行訪問，這將大大簡化一維陣列的索引計算問題。因此如果通過某種方式使按方式2儲存也能避免Bank conflict就好了。事實上通過一種非常簡單的方式就可以達到這一點：上面的例子中向量的長度為6，是一個偶數，只要長度為偶數，按照方式2儲存就會引入Bank conflict，而只要是奇數，則並不會導致這種衝突。因而當陣列長度為偶數時，只需要將共享記憶體的陣列長度增加1變為奇數，然後只使用前面的偶數個元素即可：
__shared__ int Vector1[32][6+1];
這樣當每個執行緒同時訪問自己向量的第一個元素時，按方式2儲存則每個執行緒訪問的字地址將為：tid*7+0，對應bank為(tid*7+0)%32，就不會出現引入Bank conflict的問題。唯一的一點瑕疵便是浪費了32個字的共享記憶體空間。
以上便是幾種可行的避免Bank conflict衝突的方式，歡迎同行批評指正、提出新的思路，感激不盡。
Luchang Li
2015.12.10 in HUST
--------------------- 本文來自 Luchang-Li 的CSDN 部落格 ，全文地址請點選：https://blog.csdn.net/u013701860/article/details/50253343?utm_source=copy