前面博客中我們說(shuō)到了共享內(nèi)存的使用方法以及一些高級(jí)特性形纺，并簡(jiǎn)單說(shuō)明了一下bank沖突哗蜈，這里我們將會(huì)通過(guò)一些簡(jiǎn)單的例子來(lái)詳細(xì)介紹一下bank沖突谷遂。
為了獲得較高的內(nèi)存帶寬钠糊，共享存儲(chǔ)器被劃分為多個(gè)大小相等的存儲(chǔ)器模塊泻云，稱(chēng)為bank，可以被同時(shí)訪問(wèn)头谜。因此任何跨越b個(gè)不同的內(nèi)存bank的對(duì)n個(gè)地址進(jìn)行讀取和寫(xiě)入的操作可以被同時(shí)進(jìn)行骏掀，這樣就大大提高了整體帶寬 ——可達(dá)到單獨(dú)一個(gè)bank帶寬的b倍。但是很多情況下柱告，我們無(wú)法充分發(fā)揮bank的功能截驮，以致于shared memory的帶寬非常的小，這可能是因?yàn)槲覀冇龅搅薭ank沖突际度。
bank沖突
當(dāng)一個(gè)warp中的不同線程訪問(wèn)一個(gè)bank中的不同的字地址時(shí)葵袭，就會(huì)發(fā)生bank沖突。
如果沒(méi)有bank沖突的話(huà)乖菱，共享內(nèi)存的訪存速度將會(huì)非常的快坡锡，大約比全局內(nèi)存的訪問(wèn)延遲低100多倍，但是速度沒(méi)有寄存器快窒所。然而鹉勒，如果在使用共享內(nèi)存時(shí)發(fā)生了bank沖突的話(huà)，性能將會(huì)降低很多很多吵取。在最壞的情況下禽额，即一個(gè)warp中的所有線程訪問(wèn)了相同bank的32個(gè)不同字地址的話(huà)，那么這32個(gè)訪問(wèn)操作將會(huì)全部被序列化皮官，大大降低了內(nèi)存帶寬脯倒。
NOTE:不同warp中的線程之間不存在什么bank沖突。
共享內(nèi)存的地址映射方式
要解決bank沖突捺氢，首先我們要了解一下共享內(nèi)存的地址映射方式藻丢。
在共享內(nèi)存中，連續(xù)的32-bits字被分配到連續(xù)的32個(gè)bank中摄乒，這就像電影院的座位一樣：一列的座位就相當(dāng)于一個(gè)bank悠反，所以每行有32個(gè)座位残黑，在每個(gè)座位上可以“坐”一個(gè)32-bits的數(shù)據(jù)(或者多個(gè)小于32-bits的數(shù)據(jù)，如4個(gè)char型的數(shù)據(jù)问慎，2個(gè)short型的數(shù)據(jù))萍摊；而正常情況下，我們是按照先坐完一行再坐下一行的順序來(lái)坐座位的如叼，在shared memory中地址映射的方式也是這樣的冰木。下圖中內(nèi)存地址是按照箭頭的方向依次映射的：
上圖中數(shù)字為bank編號(hào)。這樣的話(huà)笼恰，如果你將申請(qǐng)一個(gè)共享內(nèi)存數(shù)組(假設(shè)是int類(lèi)型)的話(huà)踊沸，那么你的每個(gè)元素所對(duì)應(yīng)的bank編號(hào)就是地址偏移量(也就是數(shù)組下標(biāo))對(duì)32取余所得的結(jié)果，比如大小為1024的一維數(shù)組myShMem：
- myShMem[4]: 對(duì)應(yīng)的bank id為#4 (相應(yīng)的行偏移量為0)
- myShMem[31]: 對(duì)應(yīng)的bank id為#31 (相應(yīng)的行偏移量為0)
- myShMem[50]: 對(duì)應(yīng)的bank id為#18 (相應(yīng)的行偏移量為1)
- myShMem[128]: 對(duì)應(yīng)的bank id為#0 (相應(yīng)的行偏移量為4)
- myShMem[178]: 對(duì)應(yīng)的bank id為#18 (相應(yīng)的行偏移量為5)
典型的bank訪問(wèn)方式
下面我介紹幾種典型的bank訪問(wèn)的形式社证。
下面這這種訪問(wèn)方式是典型的線性訪問(wèn)方式(訪問(wèn)步長(zhǎng)(stride)為1)逼龟，由于每個(gè)warp中的線程ID與每個(gè)bank的ID一一對(duì)應(yīng)，因此不會(huì)產(chǎn)生bank沖突追葡。
下面這種訪問(wèn)雖然是交叉的訪問(wèn)腺律，每個(gè)線程并沒(méi)有與bank一一對(duì)應(yīng)，但每個(gè)線程都會(huì)對(duì)應(yīng)一個(gè)唯一的bank宜肉，所以也不會(huì)產(chǎn)生bank沖突匀钧。
下面這種雖然也是線性的訪問(wèn)bank，但這種訪問(wèn)方式與第一種的區(qū)別在于訪問(wèn)的步長(zhǎng)(stride)變?yōu)?谬返，這就造成了線程0與線程28都訪問(wèn)到了bank 0之斯，線程1與線程29都訪問(wèn)到了bank 2...，于是就造成了2路的bank沖突遣铝。我在后面會(huì)對(duì)以不同的步長(zhǎng)(stride)訪問(wèn)bank的情況做進(jìn)一步討論佑刷。
下面這種訪問(wèn)造成了8路的bank沖突，
這里我們需要注意酿炸，下面這兩種情況是兩種特殊情況：
上圖中瘫絮，所有的線程都訪問(wèn)了同一個(gè)bank，貌似產(chǎn)生了32路的bank沖突填硕，但是由于廣播(broadcast)機(jī)制(當(dāng)一個(gè)warp中的所有線程訪問(wèn)一個(gè)bank中的同一個(gè)字(word)地址時(shí)檀何，就會(huì)向所有的線程廣播這個(gè)字(word))，這種情況并不會(huì)發(fā)生bank沖突廷支。
同樣，這種訪問(wèn)方式也不會(huì)產(chǎn)生bank沖突：
這就是所謂的多播機(jī)制(multicast)——當(dāng)一個(gè)warp中的幾個(gè)線程訪問(wèn)同一個(gè)bank中的相同字地址時(shí)栓辜，會(huì)將該字廣播給這些線程恋拍。
NOTE:這里的多播機(jī)制(multicast)只適用于計(jì)算能力2.0及以上的設(shè)備，上篇博客中已經(jīng)提到藕甩。
數(shù)據(jù)類(lèi)型與bank沖突
我們都知道施敢，當(dāng)每個(gè)線程訪問(wèn)一個(gè)32-bits大小的數(shù)據(jù)類(lèi)型的數(shù)據(jù)(如int周荐，float)時(shí)，不會(huì)發(fā)生bank沖突僵娃。
extern __shared__ int shrd[];
foo = shrd[baseIndex + threadIdx.x]
但是如果每個(gè)線程訪問(wèn)一個(gè)字節(jié)(8-bits)的數(shù)據(jù)時(shí)概作，會(huì)不會(huì)發(fā)生bank沖突呢？很明顯這種情況會(huì)發(fā)生bank沖突的默怨，因?yàn)樗膫€(gè)線程訪問(wèn)了同一個(gè)bank讯榕，造成了四路bank沖突。同理匙睹，如果是short類(lèi)型（16-bits）也會(huì)發(fā)生bank沖突愚屁，會(huì)產(chǎn)生兩路的bank沖突，下面是這種情況的兩個(gè)例子：
extern __shared__ char shrd[];
foo = shrd[baseIndex + threadIdx.x];
extern __shared__ short shrd[];
foo = shrd[baseIndex + threadIdx.x];
訪問(wèn)步長(zhǎng)與bank沖突
我們通常這樣來(lái)訪問(wèn)數(shù)組：每個(gè)線程根據(jù)線程編號(hào)tid與s的乘積來(lái)訪問(wèn)數(shù)組的32-bits字(word)：
extern __shared__ float shared[];
float data = shared[baseIndex + s * tid];
如果按照上面的方式痕檬，那么當(dāng)s*n是bank的數(shù)量(即32)的整數(shù)倍時(shí)或者說(shuō)n是32/d的整數(shù)倍(d是32和s的最大公約數(shù))時(shí)霎槐，線程tid和線程tid+n會(huì)訪問(wèn)相同的bank。我們不難知道如果tid與tid+n位于同一個(gè)warp時(shí)梦谜，就會(huì)發(fā)生bank沖突丘跌，相反則不會(huì)。
仔細(xì)思考你會(huì)發(fā)現(xiàn)唁桩，只有warp的大小(即32)小于等于32/d時(shí)闭树，才不會(huì)有bank沖突，而只有當(dāng)d等于1時(shí)才能滿(mǎn)足這個(gè)條件朵夏。要想讓32和s的最大公約數(shù)d為1蔼啦，s必須為奇數(shù)。于是仰猖，這里有一個(gè)顯而易見(jiàn)的結(jié)論：當(dāng)訪問(wèn)步長(zhǎng)s為奇數(shù)時(shí)捏肢，就不會(huì)發(fā)生bank沖突。
bank沖突的例子
既然我們已經(jīng)理解了bank沖突饥侵，那我們就小試牛刀鸵赫，來(lái)練習(xí)下吧！下面我們以并行計(jì)算中的經(jīng)典的歸約算法為例來(lái)做一個(gè)簡(jiǎn)單的練習(xí)躏升。
假設(shè)有一個(gè)大小為2048的向量辩棒，我們想用歸約算法對(duì)該向量求和。于是我們申請(qǐng)了一個(gè)大小為1024的線程塊膨疏，并聲明了一個(gè)大小為2048的共享內(nèi)存數(shù)組一睁，并將數(shù)據(jù)從全局內(nèi)存拷貝到了該共享內(nèi)存數(shù)組。
我們可以有以下兩種方式實(shí)現(xiàn)歸約算法：
不連續(xù)的方式：
連續(xù)的方式：
下面我們用具體的代碼來(lái)實(shí)現(xiàn)上述兩種方法佃却。
// 非連續(xù)的歸約求和
__global__ void BC_addKernel(const int *a, int *r)
{
__shared__ int cache[ThreadsPerBlock];
int tid = blockIdx.x * blockDim.x + threadIdx.x;
int cacheIndex = threadIdx.x;
// copy data to shared memory from global memory
cache[cacheIndex] = a[tid];
__syncthreads();
// add these data using reduce
for (int i = 1; i < blockDim.x; i *= 2)
{
int index = 2 * i * cacheIndex;
if (index < blockDim.x)
{
cache[index] += cache[index + i];
}
__syncthreads();
}
// copy the result of reduce to global memory
if (cacheIndex == 0)
r[blockIdx.x] = cache[cacheIndex];
}
上述代碼實(shí)現(xiàn)的是非連續(xù)的歸約求和者吁，從int index = 2 * i * cacheIndex和cache[index] += cache[index + i];兩條語(yǔ)句，我們可以很容易判斷這種實(shí)現(xiàn)方式會(huì)產(chǎn)生bank沖突饲帅。當(dāng)i=1時(shí)复凳，步長(zhǎng)s=2xi=2瘤泪，會(huì)產(chǎn)生兩路的bank沖突；當(dāng)i=2時(shí)育八，步長(zhǎng)s=2xi=4对途，會(huì)產(chǎn)生四路的bank沖突...當(dāng)i=n時(shí)，步長(zhǎng)s=2xn=2n髓棋∈堤矗可以看出每一次步長(zhǎng)都是偶數(shù)，因此這種方式會(huì)產(chǎn)生嚴(yán)重的bank沖突仲锄。
NOTE:在《GPU高性能運(yùn)算之CUDA》這本書(shū)中對(duì)實(shí)現(xiàn)不連續(xù)的歸約算法有兩種代碼實(shí)現(xiàn)方式劲妙，但筆者發(fā)現(xiàn)書(shū)中的提到(p179)的兩種所謂相同計(jì)算邏輯的函數(shù)reduce0和reduce1，其實(shí)具有本質(zhì)上的不同儒喊。前者不會(huì)發(fā)生bank沖突镣奋，而后者(即本文中所使用的)才會(huì)產(chǎn)生bank沖突。由于前者線程ID要求的條件比較“苛刻”怀愧，只有滿(mǎn)足tid % (2 * s) == 0的線程才會(huì)執(zhí)行求和操作(sdata[tid]+=sdata[tid+i])侨颈；而后者只要滿(mǎn)足index(2 * s * tid，即線程ID的2xs倍)小于線程塊的大小(blockDim.x)即可芯义」福總之，前者在進(jìn)行求和操作(sdata[tid]+=sdata[tid+i])時(shí)扛拨，線程的使用同樣是不連續(xù)的耘分，即當(dāng)s=1時(shí)，線程編號(hào)為0,2,4,...,1022绑警；而后者的線程使用是連續(xù)的求泰，即當(dāng)s=1時(shí)，前512個(gè)線程(0,1,2,...,511)在進(jìn)行求和操作(sdata[tid]+=sdata[tid+i])计盒，而后512個(gè)線程是閑置的渴频。前者不會(huì)出現(xiàn)多個(gè)線程訪問(wèn)同一bank的不同字地址，而后者正如書(shū)中所說(shuō)會(huì)產(chǎn)生嚴(yán)重的bank沖突北启。(書(shū)中用到的s與本文中多次用到的步長(zhǎng)s不是同一個(gè)變量卜朗，注意不要混淆這兩個(gè)變量)當(dāng)然這些只是筆者的想法，如有不同咕村，歡迎來(lái)與我討論场钉，郵箱:chaoyanglius@outlook.com。
// 連續(xù)的歸約求和
__global__ void NBC_addKernel2(const int *a, int *r)
{
__shared__ int cache[ThreadsPerBlock];
int tid = blockIdx.x * blockDim.x + threadIdx.x;
int cacheIndex = threadIdx.x;
// copy data to shared memory from global memory
cache[cacheIndex] = a[tid];
__syncthreads();
// add these data using reduce
for (int i = blockDim.x / 2; i > 0; i /= 2)
{
if (cacheIndex < i)
{
cache[cacheIndex] += cache[cacheIndex + i];
}
__syncthreads();
}
// copy the result of reduce to global memory
if (cacheIndex == 0)
r[blockIdx.x] = cache[cacheIndex];
}
由于每個(gè)線程的ID與操作的數(shù)據(jù)編號(hào)一一對(duì)應(yīng)懈涛，因此上述的代碼很明顯不會(huì)產(chǎn)生bank沖突惹悄。
參考資料
- C語(yǔ)言程序設(shè)計(jì)現(xiàn)代方法，[美]K.N.King著肩钠，人民郵電出版社
- 英偉達(dá)CUDA C programming guide v7.0
- 威斯康星大學(xué)仿真實(shí)驗(yàn)室CUDA課程講義10-07-2013:http://sbel.wisc.edu/Courses/ME964/2013/
- GPU高性能運(yùn)算之CUDA泣港，張舒，褚艷利价匠，中國(guó)水利水電出版社
