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

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

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

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

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

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

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

計(jì)算能力1.x的設(shè)備上warpsize=32,bank數(shù)量是16.一個(gè)warp的共享內(nèi)存請(qǐng)求被分成兩個(gè)，一個(gè)是前半個(gè)warp，一個(gè)是后半個(gè)warp的請(qǐng)求。

計(jì)算能力2.0的設(shè)備，warpsize=32，bank的數(shù)量也是32.這樣內(nèi)存請(qǐng)求就不再劃分成前后兩個(gè)。

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

假設(shè)warpsize為8，bank數(shù)量為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上的數(shù)組。

s = 1時(shí)，所有的線程都執(zhí)行一次for循環(huán)內(nèi)的語(yǔ)句，那么線程4訪問(wèn)的sdata[8]和sdata[9]映射到了bank[0]和bank[1]，而本身線程0訪問(wèn)的地址就被映射到了bank[0]和bank[1]，從而導(dǎo)致同一個(gè)warp里的線程訪問(wèn)的地址映射到了同樣的bank，不得不串行處理，出現(xiàn)了bank conflicts。

改為：

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

由于在一個(gè)循環(huán)里訪問(wèn)了兩次sdata，所以不得不分成兩次訪問(wèn)，但是每次訪問(wèn)所有的線程訪問(wèn)地址都映射在了8個(gè)bank內(nèi)，且沒(méi)有沖突，因此達(dá)到了最高帶寬。
實(shí)驗(yàn)結(jié)果

樣例2

warp size = 32, banks = 16,（計(jì)算能力1.x的設(shè)備）數(shù)據(jù)映射關(guān)系如下：


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

總結(jié)

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

如果覺(jué)得生活随笔網(wǎng)站內(nèi)容還不錯(cuò)，歡迎將生活随笔推薦給好友。