其實這兩天一直不知道什么叫bank conflict沖突，這兩天因為要看那個矩陣轉置優化的問題，里面有講到這些問題，但是沒辦法，為了要看懂那個bank conflict沖突，我不得不去找資料，說句實話我現在不是完全弄明白，但是應該說有點眉目了，現在我就把網上找的整理一下，放在這邊，等哪天完全弄明白了我就在修改里面的錯誤。
Tesla 的每個 SM 擁有 16KB 共享存儲器，用於同一個線程塊內的線程間通信。為了使一個 half-warp 內的線程能夠在一個內核周期中並行訪問，共享存儲器被組織成 16 個 bank，每個 bank 擁有 32bit 的寬度，故每個 bank 可保存 256 個整形或單精度浮點數，或者說目前的bank 組織成了 256 行 16 列的矩陣。如果一個 half-warp 中有一部分線程訪問屬於同一bank 的數據，則會產生 bank conflict，降低訪存效率，在沖突最嚴重的情況下，速度會比全局顯存還慢，但是如果 half-warp 的線程訪問同一地址的時候，會產生一次廣播，其速度反而沒有下降。在不發生 bank conflict 時，訪問共享存儲器的速度與寄存器相同。在不同的塊之間，共享存儲器是毫不相關的。 ------風辰的 CUDA 入門教程
里面說的很清楚就是每個bank有1KB的存儲空間。
Shared memory 是以 4 bytes 為單位分成 banks。因此，假設以下的數據：
__shared__ int data[128];
那么，data[0] 是 bank 0、data[1] 是 bank 1、data[2] 是 bank 2、…、data[15] 是bank 15，而 data[16] 又回到 bank 0。由於 warp 在執行時是以 half-warp 的方式執行，因此分屬於不同的 half warp 的 threads，不會造成 bank conflict。
因此，如果程序在存取 shared memory 的時候，使用以下的方式：
int number = data[base + tid];
那就不會有任何 bank conflict，可以達到最高的效率。但是，如果是以下的方式：
int number = data[base + 4 * tid];
那么，thread 0 和 thread 4 就會存取到同一個 bank，thread 1 和 thread 5 也是同 樣，這樣就會造成 bank conflict。在這個例子中，一個 half warp 的 16 個 threads 會有四個threads 存取同一個 bank，因此存取 share memory 的速度會變成原來的 1/4。
一個重要的例外是，當多個 thread 存取到同一個 shared memory 的地址時，shared memory 可以將這個地址的 32 bits 數據「廣播」到所有讀取的 threads，因此不會造成 bank conflict。例如：
int number = data[3];
這樣不會造成 bank conflict，因為所有的 thread 都讀取同一個地址的數據。
很多時候 shared memory 的 bank conflict 可以透過修改數據存放的方式來解決。例如，以下的程序：
data[tid] = global_data[tid];
...
int number = data[16 * tid];
會造成嚴重的 bank conflict，為了避免這個問題，可以把數據的排列方式稍加修改，把存取方式改成：
int row = tid / 16;
int column = tid % 16;
data[row * 17 + column] = global_data[tid];
...
int number = data[17 * tid];
這樣就不會造成 bank conflict 了。
簡單的說，矩陣中的數據是按照bank存儲的，第i個數據存儲在第i%16個bank中。一個block要訪問shared memory，只要能夠保證以其中相鄰的16個線程一組訪問thread，每個線程與bank是一一對應就不會產生bank conflict。否則會產生bankconflict，訪存時間成倍增加，增加的倍數由一個bank最多被多少個thread同時訪問決定。有一種極端情況，就是所有的16個thread同時訪問同一bank時反而只需要一個訪問周期，此時產生了一次廣播。
下面有一些小技巧可以避免bank conflict 或者提高global存儲器的訪問速度
1. 盡量按行操作，需要按列操作時可以先對矩陣進行轉置
2. 划分子問題時，使每個block處理的問題寬度恰好為16的整數倍，使得訪存可以按照 s_data[tid]=i_data[tid]的形式進行
3. 使用對齊的數據格式，盡量使用nvidia定義的格式如float3,int2等，這些格式本身已經對齊。
4. 當要處理的矩陣寬度不是16的整數倍時，將其補為16的整數倍，或者用malloctopitch而不是malloc。
5. 利用廣播，例如s_odata[tid] = tid%16 < 8 ? s_idata[tid] : s_idata[15];會產生8路的塊訪問沖突而用：s_odata[tid]=s_idata[15];s_odata[tid]= tid%16 < 8 ? s_idata[tid] : s_data[tid]; 則不會產生塊訪問沖突