5 GPU也不允許偏心并行的事情多了，我們作為GPU的指令分配者，不能偏心了——給甲做的事情多，而乙沒(méi)事做，個(gè)么甲肯定不爽的來(lái)。所以，在GPU中，叫做線程網(wǎng)絡(luò)的分配。首先還是來(lái)看下GPU的線程網(wǎng)絡(luò)吧，圖2：
圖2 線程網(wǎng)絡(luò) 我們將具體點(diǎn)的，在主機(jī)函數(shù)中如果我們分配的是這樣的一個(gè)東西： dim3 blocks(32,32); dim3 threads(16,16); dim3是神馬？dim3是一個(gè)內(nèi)置的結(jié)構(gòu)體，和Linux下定義的線程結(jié)構(gòu)體是個(gè)類似的意義的東西，dim3結(jié)構(gòu)變量有x，y，z，表示3維的維度。不理解沒(méi)關(guān)系，慢慢看。 kernelfun<<<blocks, threads>>>(); 我們調(diào)用kernelfun這個(gè)內(nèi)核函數(shù)，將blocks和threads傳到<<<,>>>里去，這句話可牛逼大了——相當(dāng)于發(fā)號(hào)施令，命令那些線程去干活。這里使用了32*32 * 16*16個(gè)線程來(lái)干活。你看明白了嗎？blocks表示用了二維的32*32個(gè)block組，而每個(gè)block中又用了16*16的二維的thread組。好吧，我們這個(gè)施令動(dòng)用了262144個(gè)線程！我們先不管GPU內(nèi)部是如何調(diào)度這些線程的，反正我們這一句話就是用了這么多線程。 那我們的內(nèi)核函數(shù)kernelfun()如何知道自己執(zhí)行的是哪個(gè)線程？這就是線程網(wǎng)絡(luò)的特點(diǎn)啦，為什么叫網(wǎng)絡(luò)，是有講究的，網(wǎng)絡(luò)就可以定格到網(wǎng)點(diǎn)： 比如int tid = threadId.x + blockId.x * 16 這里有一個(gè)講究，block是有維度的，一維、二維、三維。 對(duì)于一維的block，tid = threadId.x 對(duì)于（Dx，Dy）二維的block，tid = threadId.x + Dx*threadId.y 對(duì)于（Dx，Dy，Dz）三維的block，tid = threadId.x + Dx*threadId.y + Dz*Dy*threadId.z 我習(xí)慣的用這樣的模式去分配，比較通用： dim3 dimGrid(); dim3 dimBlock(); kerneladd<<<dimGrid, dimBlock>>>(); 這可是萬(wàn)金油啊，你需要做的事情是填充dimGrid和dimBlock的結(jié)構(gòu)體構(gòu)造函數(shù)變量，比如，dimGrid(16, 16)表示用了16*16的二維的block線程塊。 （0,0）（0,1）（0,2）……（0,15） （1,0）（1,1）（1,2）……（1,15） （2,0）（2,1）（2,2）……（2,15） …… （15,0）（15,1）（15,2）……（15,15） （，）是（dimGrid.x， dimGrid.y）的網(wǎng)格編號(hào)。 我們這么理解吧，現(xiàn)在又一群人，我們分成16*16個(gè)小組（block），排列好，比如第3行第4列就指的是（2,3）這個(gè)小組。 而dimBlock（16,16）表示每個(gè)小組有16*16個(gè)成員，如果你想點(diǎn)名第3行第4列這個(gè)小組的里面的第3行第4列那個(gè)同學(xué)，那么，你就是在（2,3）這個(gè)block中選擇了（2,3）這個(gè)線程。這樣應(yīng)該有那么一點(diǎn)可以理解進(jìn)去的意思了吧？不理解透徹么什么關(guān)系，這個(gè)東西本來(lái)就是cuda中最讓我糾結(jié)的事情。我們且不管如何分配線程，能達(dá)到最優(yōu)化，我們的目標(biāo)是先讓GPU正確地跑起來(lái)，計(jì)算出結(jié)果即可，管他高效不高效，管他環(huán)保不環(huán)保。 嘮叨了這么多，下面我們用一個(gè)最能說(shuō)明問(wèn)題的例子來(lái)進(jìn)一步理解線程網(wǎng)絡(luò)分配機(jī)制來(lái)了解線程網(wǎng)絡(luò)的使用。 一維網(wǎng)絡(luò)線程eg：int arr[1000]，對(duì)每個(gè)數(shù)組元素進(jìn)行加1操作。 idea：我們最直接的想法，是調(diào)度1000個(gè)線程去干這件事情。 first pro：我想用一個(gè)小組的1000個(gè)人員去干活。這里會(huì)存在這樣一個(gè)問(wèn)題——一個(gè)小組是不是有這么多人員呢？是的，這個(gè)事情你必須了解，連自己組內(nèi)多少人都不知道，你也不配作指揮官呀。對(duì)的，這個(gè)參數(shù)叫做maxThreadsPerBlock，如何取得呢？ 好吧，cuda定義了一個(gè)結(jié)構(gòu)體cudaDeviceProp，里面存入了一系列的結(jié)構(gòu)體變量作為GPU的參數(shù)，出了maxThreadsPerBlock，還有很多信息哦，我們用到了再說(shuō)。 maxThreadsPerBlock這個(gè)參數(shù)值是隨著GPU級(jí)別有遞增的，早起的顯卡可能512個(gè)線程，我的GT520可以跑1024個(gè)線程，辦公室的GTX650ti2G可以跑1536個(gè)，無(wú)可非議，當(dāng)然多多益善。一開(kāi)始，我在想，是不是程序?qū)⒚總€(gè)block開(kāi)的線程開(kāi)滿是最好的呢？這個(gè)問(wèn)題留在以后在說(shuō)，一口吃不成胖子啦。 好吧，我們的數(shù)組元素1000個(gè)，是可以在一個(gè)block中干完的。 內(nèi)核函數(shù)： #define N 1000 __gloabl__ void kerneladd(int *dev_arr) { int tid = threadId.x; if (tid < 1000) dev_arr[tid] ++; } int main() { int *arr, *dev_arr;// 習(xí)慣的我喜歡在內(nèi)核函數(shù)參數(shù)變量前加個(gè)dev_作為標(biāo)示 // 開(kāi)辟主機(jī)內(nèi)存，arr = (int*)malloc(N*sizeof(int)); // 開(kāi)辟設(shè)備內(nèi)存 // 主機(jī)拷貝到設(shè)備 kerneladd<<<1, N>>>(dev_arr); // 設(shè)備拷貝到主機(jī) // 打印 // 釋放設(shè)備內(nèi)存 // 釋放主機(jī)內(nèi)存 return 0; }
呀，原來(lái)這么簡(jiǎn)單，個(gè)么CUDA也忒簡(jiǎn)單了哇！這中想法是好的，給自己提高信心，但是這種想法多了是不好的，因?yàn)楹竺娴膯?wèn)題多了去了。 盆友說(shuō)，1000個(gè)元素，還不如CPU來(lái)的快，對(duì)的，很多情況下，數(shù)據(jù)量并行度不是特別大的情況下，可能CPU來(lái)的更快一些，比較設(shè)備與主機(jī)之間互相調(diào)度操作，是會(huì)有額外開(kāi)銷的。有人就問(wèn)了，一個(gè)10000個(gè)元素的數(shù)組是不是上面提供的idea就解決不了啦？對(duì)，一個(gè)block人都沒(méi)怎么多，如何完成！這個(gè)情況下有兩條路可以選擇—— 第一，我就用一個(gè)組的1000人來(lái)干活話，每個(gè)人讓他干10個(gè)元素好了。 這個(gè)解決方案，我們需要修改的是內(nèi)核函數(shù)： __global__ void kernelarr(int *dev_arr) { int tid = threadId.x; if(tid < 1000) // 只用0~999號(hào)線程 { //每個(gè)線程處理10個(gè)元素，比如0號(hào)線程處理0、1001、2001、……9001 for(int i = tid; i<N; i=i+1000) { dev_arr[tid] ++; } } } 第二，我多用幾個(gè)組來(lái)干這件事情，比如我用10個(gè)組，每個(gè)組用1000人。 這個(gè)解決方案就稍微復(fù)雜了一點(diǎn)，注意只是一點(diǎn)點(diǎn)哦~因?yàn)?，組內(nèi)部怎么干活和最原始的做法是一樣的，不同之處是，我們調(diào)遣了10個(gè)組去干這件事情。 首先我們來(lái)修改我們的主機(jī)函數(shù)： int main() { …… kerneladd<<<10, 1000>>>(dev_arr);//我們調(diào)遣了10個(gè)組，每個(gè)組用了1000人 …… } 盆友要問(wèn)了，10個(gè)組每個(gè)組1000人，你怎么點(diǎn)兵呢？很簡(jiǎn)單啊，第1組第3個(gè)線程出列，第9組第9個(gè)線程出列。每個(gè)人用組號(hào)和組內(nèi)的編號(hào)定了位置。在線程網(wǎng)絡(luò)中，blockId.x和threadId.x就是對(duì)應(yīng)的組號(hào)和組內(nèi)編號(hào)啦，我必須要這里開(kāi)始形象點(diǎn)表示這個(gè)對(duì)應(yīng)關(guān)系，如果這個(gè)對(duì)應(yīng)關(guān)系是這樣子的[blockId.x，threadId.x]，那么我們的數(shù)組arr[10000]可以這樣分配給這10個(gè)組去干活： (0,0)——arr[0]，(0,1)——arr[1]，……(0,999)——arr[999] (1,0)——arr[0+1*1000]，(1,1)——arr[1+1*1000]，……(1,999)——arr[999+1*1000] …… (9,0)——arr[0+9*1000]，(9,1)——arr[1+9*1000]，……(9,999)——arr[999+9*1000] 是不是很有規(guī)律呢？對(duì)的，用blockId.x和threadId.x可以很好的知道哪個(gè)線程干哪個(gè)元素，這個(gè)元素的下表就是threadId.x + 1000*blockId.x。 這里我想說(shuō)的是，如果我們哪天糊涂了，畫一畫這個(gè)對(duì)應(yīng)關(guān)系的表，也許，就更加清楚的知道我們分配的線程對(duì)應(yīng)的處理那些東西啦。 一維線程網(wǎng)絡(luò)，就先學(xué)這么多了。 二維網(wǎng)絡(luò)線程eg2：int arr[32][16]二維的數(shù)組自增1。 第一個(gè)念頭，開(kāi)個(gè)32*16個(gè)線程好了哇，萬(wàn)事大吉！好吧。但是，朕現(xiàn)在想用二維線程網(wǎng)絡(luò)來(lái)解決，因?yàn)殡抻X(jué)得一個(gè)二維的網(wǎng)絡(luò)去映射一個(gè)二維的數(shù)組，朕看的更加明了，看不清楚自己的士兵，如何帶兵打仗！ 我還是畫個(gè)映射關(guān)系： 一個(gè)block中，現(xiàn)在是一個(gè)二維的thread網(wǎng)絡(luò)，如果我用了16*16個(gè)線程。 (0,0)，(0,1)，……(0,15) (1,0)，(1,1)，……(1,15) …… (15,0)，(15,1)，……(15,15) 呀，現(xiàn)在一個(gè)組內(nèi)的人稱呼變了嘛，一維網(wǎng)絡(luò)中，你走到一個(gè)小組里，叫3號(hào)出列，就出來(lái)一個(gè)，你現(xiàn)在只是叫3號(hào)，沒(méi)人會(huì)出來(lái)！這個(gè)場(chǎng)景是這樣的，現(xiàn)在你班上有兩個(gè)人同名的人，你只叫名，他們不知道叫誰(shuí)，你必須叫完整點(diǎn)，把他們的姓也叫出來(lái)。所以，二維網(wǎng)絡(luò)中的(0,3)就是原來(lái)一維網(wǎng)絡(luò)中的3，二維中的(i,j)就是一維中的(j+i*16)。不管怎么樣，一個(gè)block里面能處理的線程數(shù)量總和還是不變的。 一個(gè)grid中，block也可以是二維的，一個(gè)block中已經(jīng)用了16*16的thread了，那我們一共就32*16個(gè)元素，我們用2個(gè)block就行了。 先給出一個(gè)代碼清單吧，程序員都喜歡看代碼，這段代碼是我抄襲的。第一次這么完整的放上代碼，因?yàn)槲矣X(jué)得這個(gè)代碼可以讓我說(shuō)明我想說(shuō)的幾個(gè)問(wèn)題： 第一，二維數(shù)組和二維指針的聯(lián)系。 第二，二維線程網(wǎng)絡(luò)。 第三，cuda的一些內(nèi)存操作，和返回值的判斷。
簡(jiǎn)要的來(lái)學(xué)習(xí)一下二維網(wǎng)絡(luò)這個(gè)知識(shí)點(diǎn)， dim3 dimBlock(16,16); //定義block內(nèi)的thread二維網(wǎng)絡(luò)為16*16 dim3 dimGrid((COLS+dimBlock.x-1)/(dimBlock.x), (ROWS+dimBlock.y-1)/(dimBlock.y)); //定義grid內(nèi)的block二維網(wǎng)絡(luò)為1*2 unsigned int row = blockDim.y*blockIdx.y + threadIdx.y; //二維數(shù)組中的行號(hào) unsigned int col = blockDim.x*blockIdx.x + threadIdx.x; //二維線程中的列號(hào)
三維網(wǎng)絡(luò)線程dim3定義了三維的結(jié)構(gòu)，但是，貌似二維之內(nèi)就能處理很多事情啦，所以，我放棄學(xué)習(xí)三維。網(wǎng)上看到的不支持三維網(wǎng)絡(luò)是什么意思呢？先放一放。 給自己充充電同一塊顯卡，不管你是二維和三維或一維，其計(jì)算能力是固定的。比如一個(gè)block能處理1024個(gè)線程，那么，一維和二維線程網(wǎng)絡(luò)是不是處理的線程數(shù)一樣呢？
回答此問(wèn)題，先給出網(wǎng)絡(luò)配置的參數(shù)形式——<<<Dg,Db,Ns,S>>>，各個(gè)參數(shù)含義如下： Dg：定義整個(gè)grid的維度，類型Dim3，但是實(shí)際上目前顯卡支持兩個(gè)維度，所以，dim3<<Dg.x, Dg.y, 1>>>第z維度默認(rèn)只能為1，上面顯示出這個(gè)最大有65536*65536*1，每行有65536個(gè)block，每列有65536個(gè)block，整個(gè)grid中一共有65536*65536*1個(gè)block。 Db：定義了每個(gè)block的維度，類型Dim3，比如512*512*64，這個(gè)可以定義3維尺寸，但是，這個(gè)地方是有講究了，三個(gè)維度的積是有上限的，對(duì)于計(jì)算能力1.0、1.1的GPU，這個(gè)值不能大于768，對(duì)于1.2、1.3的不能大于1024，對(duì)于我們?cè)囈辉嚨倪@塊級(jí)別高點(diǎn)的，不能大于1536。這個(gè)值可以獲取哦——maxThreadsPerBlock Ns：這個(gè)是可選參數(shù)，設(shè)定最多能動(dòng)態(tài)分配的共享內(nèi)存大小，比如16k，單不需要是，這個(gè)值可以省略或?qū)?/span>0。 S：也是可選參數(shù)，表示流號(hào)，默認(rèn)為0。流這個(gè)概念我們這里不說(shuō)。 接著，我想解決幾個(gè)你肯定想問(wèn)的兩個(gè)問(wèn)題，因?yàn)槲铱春芏嗳讼胛疫@樣的問(wèn)這個(gè)問(wèn)題： 1 block內(nèi)的thread我們是都飽和使用嗎？ 答：不要，一般來(lái)說(shuō)，我們開(kāi)128或256個(gè)線程，二維的話就是16*16。 2 grid內(nèi)一般用幾個(gè)block呢？ 答：牛人告訴我，一般來(lái)說(shuō)是你的流處理器的4倍以上，這樣效率最高。 回答這兩個(gè)問(wèn)題的解釋，我想抄襲牛人的一段解釋，解釋的好的東西就要推廣呀： GPU的計(jì)算核心是以一定數(shù)量的Streaming Processor(SP)組成的處理器陣列，NV稱之為Texture Processing Clusters(TPC)，每個(gè)TPC中又包含一定數(shù)量的Streaming Multi-Processor(SM)，每個(gè)SM包含8個(gè)SP。SP的主要結(jié)構(gòu)為一個(gè)ALU（邏輯運(yùn)算單元），一個(gè)FPU（浮點(diǎn)運(yùn)算單元）以及一個(gè)Register File(寄存器堆)。SM內(nèi)包含有一個(gè)Instruction Unit、一個(gè)Constant Memory、一個(gè)Texture Memory，8192個(gè)Register、一個(gè)16KB的Share Memory、8個(gè)Stream Processor(SP)和兩個(gè)Special Function Units（SFU）。（GeForce9300M GS只擁有1個(gè)SM）Thread是CUDA模型中最基本的運(yùn)行單元，執(zhí)行最基本的程序指令。Block是一組協(xié)作Thread，Block內(nèi)部允許共享存儲(chǔ)，每個(gè)Block最多包含512個(gè)Thread。Grid是一組Block，共享全局內(nèi)存。Kernel是在GPU上執(zhí)行的核心程序，每一個(gè)Grid對(duì)應(yīng)一個(gè)Kernel任務(wù)。 在程序運(yùn)行的時(shí)候，實(shí)際上每32個(gè)Thread組成一個(gè)Warp，每個(gè) warp 塊都包含連續(xù)的線程，遞增線程 ID 。Warp是MP的基本調(diào)度單位，每次運(yùn)行的時(shí)候，由于MP數(shù)量不同，所以一個(gè)Block內(nèi)的所有Thread不一定全部同時(shí)運(yùn)行，但是每個(gè)Warp內(nèi)的所有Thread一定同時(shí)運(yùn)行。因此，我們?cè)诙xBlock Size的時(shí)候應(yīng)使其為Warp Size的整數(shù)倍，也就是Block Size應(yīng)為32的整數(shù)倍。理論上Thread越多，就越能彌補(bǔ)單個(gè)Thread讀取數(shù)據(jù)的latency ，但是當(dāng)Thread越多，每個(gè)Thread可用的寄存器也就越少，嚴(yán)重的時(shí)候甚至能造成Kernel無(wú)法啟動(dòng)。因此每個(gè)Block最少應(yīng)包含64個(gè)Thread，一般選擇128或者256，具體視MP數(shù)目而定。一個(gè)MP最多可以同時(shí)運(yùn)行768個(gè)Thread，但每個(gè)MP最多包含8個(gè)Block，因此要保持100%利用率，Block數(shù)目與其Size有如下幾種設(shè)定方式：? 2 blocks x 384 threads ? 3 blocks x 256 threads ? 4 blocks x 192 threads ? 6 blocks x 128 threads ? 8 blocks x 96 threads 這些電很重要啊，必須要充！不然，我就很難理解為什么網(wǎng)絡(luò)線程如何分配的。
|
|
|
