Share memory是片上資源，生命周期是整個block中，它的數據讀寫十分快，有1個cycle latency。在Share memory中，經常存在bank conflict問題，如果沒有bank conflict問題，它的數據讀寫可以和片上的寄存器（Register）一樣快。因此，我們需要盡量減少bank conflicts.
首先，什么是bank？我們以capability 1.x為例，Share memory被等分成同等尺寸大小的存儲器模式，即banks，以下圖為例：

每一個bank的帶寬為32-bit，即4 bytes。連續的32-bit字節可以被分到連續的bank中去。我們以G80為例，它具有16個banks，所以每一個bank=4-byte address % 16，此外，bank的數量和半個warp的thread數量一致，所以share memory對一個warp的請求，分成了兩個先后請求來做，不同的半個warp之間沒有bank conflicts。如果同時有兩個請求在同一個bank中，就會出現bank conflicts，如下圖：

對于share memory，有數據讀寫有快慢兩種情況：

快：1)半個warp的所有threads在同一時刻讀寫share memory的不同bank。2)半個warp中的所有threads在同一時刻以廣播的形式讀寫share memory中的同一個bank。

慢：半個warp的所有threads在同一時刻讀寫share memory的多個bank(不是全部)，這時必須串行讀寫。

以G80為例，其share memory具有16個bank，在下面的圖示中，只有S為奇數時，才不會存在bank conflicts。


當s=4時，

同樣，在存儲結構體數據時，也有可能會出現bank conflicts。

例如如下代碼存在三路沖突：

struct vector { float x, y, z; };__shared__ struct vector vectors[32];struct vector v = vectors[baseIndex + threadIdx.x];

因此，對于share memory，最好的讀寫方式每個thread讀取block中連續的元素。

以如下例程為例：

對于一個二維數組，半個warp的threads讀取每一列，這存在了16路bank conflicts。

我們有兩種方式可以解決：

1) 在每一行的后面加一個元素，例如：

2)在處理矩陣之前轉置矩陣

我們對上一節global memory中的coalescing的例程，作進一步優化，分配額外的一列：

代碼如下：

__global__ void transpose_exp(float *odata, float *idata, int width, int height){__shared__ float block[BLOCK_DIM][BLOCK_DIM+1];unsigned int xIndex = blockIdx.x * BLOCK_DIM + threadIdx.x;unsigned int yIndex = blockIdx.y * BLOCK_DIM + threadIdx.y;if( (xIndex < width)&&(yIndex < height) {unsigned int index_in = xIndex + yIndex * width;block[threadIdx.y][threadIdx.x] = idata[index_in];}__syncthreads();xIndex = blockIdx.y * BLOCK_DIM + threadIdx.x;yIndex = blockIdx.x * BLOCK_DIM + threadIdx.y;if( (xIndex < height)&&(yIndex < width) ){unsigned int index_out = yIndex * height + xIndex;odata[index_out] = block[threadIdx.x][threadIdx.y];}}

實驗結果為：

128x128: 0.011ms vs. 0.022ms (2.0X speedup) 512x512: 0.07ms vs. 0.33ms (4.5X speedup) 1024x1024: 0.30ms vs. 1.92ms (6.4X speedup) 1024x2048: 0.79ms vs. 6.6ms (8.4X speedup)

我們再來考慮矩陣相乘時，會不會出現bank conflicts問題，判斷技巧在于固定k，考慮同一時刻thread的讀些問題

1)當在讀一個Ms時，在同一時刻，半個warp的16個threads以廣播的形式讀取同一個bank

for (int k = 0; k < 16; ++k) //ty is constant over a half-warp Csub += Ms[ty][k] * Ns[k][tx];//here k is fixed over a half-warp

2)當在讀一個Ns時，在同一時刻，半個warp的16個threads連續讀取不同的bank，stride=1.

for (int k = 0; k < 16; ++k) //tx is constant over a half-warp Csub += Ms[ty][k] * Ns[k][tx];

所以綜上所述，矩陣相乘的例子并沒有bank conflicts問題。

需要注意的是，我們上述所說的均為capability 1.x的情況，在2.x 3.x中，每個share memory有32個bank，具體還是要看技術手冊的。

總結

以上是生活随笔為你收集整理的Share memory中bank conflict问题的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。