GPU 共享內存bank沖突(shared memory bank conflicts)
時間 2016-11-05 21:47:58 FindSpace
原文
http://www.findspace.name/easycoding/1784
主題 共享內存
Introduction

本文總結了GPU上共享內存的bank conflicts。主要翻譯自Reference和簡單解釋了課件內容。
共享內存(Shared Memory)

因為shared mempory是片上的（ Cache級別 ），所以比局部內存(local memory)和全局內存(global memory)快很多，實際上，shared memory的延遲要比沒有緩存的全局內存延遲小100倍（如果線程之間沒有bank conflicts的話）。在同一個block的線程共享一塊shared memory。線程可以訪問同一個block內的其他線程讓shared memory從全局內存加載的數據。這個功能（結合線程同步，thread synchronization）有很多作用，比如實現用戶管理的數據cache，高性能的并行協作算法（比如并行規約，parallel reduction）等。
什么是bank

bank是一種劃分方式。在cpu中，訪存是訪問某個地址，獲得地址上的數據，但是在這里，是一次性訪問banks數量的地址，獲得這些地址上的所有數據，并邏輯映射到不同的bank上。類似內存讀取的控制。
共享內存bank conflicts

為了實現內存高帶寬的同時訪問，shared memory被劃分成了可以同時訪問的等大小內存塊(banks)。因此，內存讀寫n個地址的行為則可以以b個獨立的bank同時操作的方式進行，這樣有效帶寬就提高到了一個bank的b倍。

然而，如果多個線程請求的內存地址被映射到了同一個bank上，那么這些請求就變成了串行的(serialized)。硬件將把這些請求分成x個沒有沖突的請求序列，帶寬就降成了原來的x分之一。但是如果一個warp內的所有線程都訪問同一個內存地址的話，會產生一次廣播(boardcast)，這些請求會一次完成。計算能力2.0及以上的設備也具有組播(multicast)能力，可以同時響應同一個warp內訪問同一個內存地址的部分線程的請求。

為了最小化bank conflicts，理解內存地址是如何映射到banks是很重要的。shared memory 中連續的32位字被分配到連續的banks，每個clock cycle每個bank的帶寬是32bits。

計算能力1.x的設備上warpsize=32,bank數量是16.一個warp的共享內存請求被分成兩個，一個是前半個warp，一個是后半個warp的請求。

計算能力2.0的設備，warpsize=32，bank的數量也是32.這樣內存請求就不再劃分成前后兩個。

計算能力3.x的設備bank的大小可以自定義配置了， cudaDeviceSetSharedMemConfig() 配置成 cudaSharedMemBankSizeFourByte 四個字節或者 cudaSharedMemBankSizeEightByte 。設置成8字節可以有效避免雙精度數據的bank conflicts。
樣例 1

假設warpsize為8，bank數量為8.

原始代碼：

__global__ void reduce0(int *g_idata, int *g_odata) { extern __shared__ int sdata[];// each thread loads one element from global to shared memunsigned int tid = threadIdx.x;unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;sdata[tid] = g_idata[i];__syncthreads();// do reduction in shared memfor(unsigned int s = 1; s < blockDim.x; s *= 2){int index = 2*s*tid;if(index < blockDim.x){sdata[index] += sdata[index + s];}__syncthreads();// write result for this block to global memif (tid == 0) g_odata[blockIdx.x] = sdata[0]; }

sdata是定義在shared memory上的數組。

s = 1時，所有的線程都執行一次for循環內的語句，那么線程4訪問的sdata[8]和sdata[9]映射到了bank[0]和bank[1]，而本身線程0訪問的地址就被映射到了bank[0]和bank[1]，從而導致同一個warp里的線程訪問的地址映射到了同樣的bank，不得不串行處理，出現了bank conflicts。

改為：

for (unsigned int s = blockDim.x/2; s > 0; s >>= 1){if (tid < s){sdata[tid] += sdata[tid + s];}__syncthreads(); }

由于在一個循環里訪問了兩次sdata，所以不得不分成兩次訪問，但是每次訪問所有的線程訪問地址都映射在了8個bank內，且沒有沖突，因此達到了最高帶寬。
實驗結果

樣例2

warp size = 32, banks = 16,（計算能力1.x的設備）數據映射關系如下：


以2-way bank conflicts為例，s = 2時，16個線程threadIdx.x 從0-15，base = 0假設，則訪問順序如圖所示，thread 0 訪問shared[0],thread1訪問shared[2]..而thread8訪問的數據地址是shared[16]，但是由于index到15，所以映射到了bank0上，而thread8-15和thread0-7都是同一個warp里的線程，但是由于一個bank同時只能喂給一個thread，因此訪問需要變成串行，即thread0-7先訪問一次，再thread8-15訪問。

總結

以上是生活随笔為你收集整理的GPU 共享内存bank冲突(shared memory bank conflicts)的全部內容，希望文章能夠幫你解決所遇到的問題。

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