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