本文，在此對作者表示感謝！

線程網(wǎng)絡

我們將具體點的，在主機函數(shù)中如果我們分配的是這樣的一個東西：

dim3 blocks(32,32);

dim3 threads(16,16);

dim3是神馬？dim3是一個內(nèi)置的結構體，和linux下定義的線程結構體是個類似的意義的東西，dim3結構變量有x，y，z，表示3維的維度。不理解沒關系，慢慢看。

kernelfun<<>>();

我們調(diào)用kernelfun這個內(nèi)核函數(shù)，將blocks和threads傳到<<<,>>>里去，這句話可牛逼大了——相當于發(fā)號施令，命令那些線程去干活。這里使用了32*32 * 16*16個線程來干活。你看明白了嗎？blocks表示用了二維的32*32個block組，而每個block中又用了16*16的二維的thread組。好吧，我們這個施令動用了262144個線程！我們先不管GPU內(nèi)部是如何調(diào)度這些線程的，反正我們這一句話就是用了這么多線程。

那我們的內(nèi)核函數(shù)kernelfun()如何知道自己執(zhí)行的是哪個線程？這就是線程網(wǎng)絡的特點啦，為什么叫網(wǎng)絡，是有講究的，網(wǎng)絡就可以定格到網(wǎng)點：

比如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的結構體構造函數(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)格編號。

我們這么理解吧，現(xiàn)在又一群人，我們分成16*16個小組（block），排列好，比如第3行第4列就指的是（2,3）這個小組。

而dimBlock（16,16）表示每個小組有16*16個成員，如果你想點名第3行第4列這個小組的里面的第3行第4列那個同學，那么，你就是在（2,3）這個block中選擇了（2,3）這個線程。這樣應該有那么一點可以理解進去的意思了吧？不理解透徹么什么關系，這個東西本來就是cuda中最讓我糾結的事情。我們且不管如何分配線程，能達到最優(yōu)化，我們的目標是先讓GPU正確地跑起來，計算出結果即可，管他高效不高效，管他環(huán)保不環(huán)保。

嘮叨了這么多，下面我們用一個最能說明問題的例子來進一步理解線程網(wǎng)絡分配機制來了解線程網(wǎng)絡的使用。

一維網(wǎng)絡線程

eg：int arr[1000]，對每個數(shù)組元素進行加1操作。

idea：我們最直接的想法，是調(diào)度1000個線程去干這件事情。

first PRo：我想用一個小組的1000個人員去干活。這里會存在這樣一個問題——一個小組是不是有這么多人員呢？是的，這個事情你必須了解，連自己組內(nèi)多少人都不知道，你也不配作指揮官呀。對的，這個參數(shù)叫做maxThreadsPerBlock，如何取得呢？

好吧，cuda定義了一個結構體cudaDeviceProp，里面存入了一系列的結構體變量作為GPU的參數(shù)，出了maxThreadsPerBlock，還有很多信息哦，我們用到了再說。

maxThreadsPerBlock這個參數(shù)值是隨著GPU級別有遞增的，早起的顯卡可能512個線程，我的GT520可以跑1024個線程，辦公室的GTX650ti2G可以跑1536個，無可非議，當然多多益善。一開始，我在想，是不是程序?qū)⒚總€block開的線程開滿是最好的呢？這個問題留在以后在說，一口吃不成胖子啦。

好吧，我們的數(shù)組元素1000個，是可以在一個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;// 習慣的我喜歡在內(nèi)核函數(shù)參數(shù)變量前加個dev_作為標示

// 開辟主機內(nèi)存，arr = (int*)malloc(N*sizeof(int));

// 開辟設備內(nèi)存

// 主機拷貝到設備

kerneladd<<<1, N>>>(dev_arr);

// 設備拷貝到主機

// 打印

// 釋放設備內(nèi)存

// 釋放主機內(nèi)存

return 0;

}

呀，原來這么簡單，個么CUDA也忒簡單了哇！這中想法是好的，給自己提高信心，但是這種想法多了是不好的，因為后面的問題多了去了。

盆友說，1000個元素，還不如CPU來的快，對的，很多情況下，數(shù)據(jù)量并行度不是特別大的情況下，可能CPU來的更快一些，比較設備與主機之間互相調(diào)度操作，是會有額外開銷的。有人就問了，一個10000個元素的數(shù)組是不是上面提供的idea就解決不了啦？對，一個block人都沒怎么多，如何完成！這個情況下有兩條路可以選擇——

第一，我就用一個組的1000人來干活話，每個人讓他干10個元素好了。

這個解決方案，我們需要修改的是內(nèi)核函數(shù)：

__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

{

dev_arr[tid] ++;

}

}

}

第二，我多用幾個組來干這件事情，比如我用10個組，每個組用1000人。

這個解決方案就稍微復雜了一點，注意只是一點點哦~因為，組內(nèi)部怎么干活和最原始的做法是一樣的，不同之處是，我們調(diào)遣了10個組去干這件事情。

首先我們來修改我們的主機函數(shù)：

int main()

{

……

kerneladd<<<10, 1000>>>(dev_arr);//我們調(diào)遣了10個組，每個組用了1000人

……

}

盆友要問了，10個組每個組1000人，你怎么點兵呢？很簡單啊，第1組第3個線程出列，第9組第9個線程出列。每個人用組號和組內(nèi)的編號定了位置。在線程網(wǎng)絡中，blockId.x和threadId.x就是對應的組號和組內(nèi)編號啦，我必須要這里開始形象點表示這個對應關系，如果這個對應關系是這樣子的[blockId.x，threadId.x]，那么我們的數(shù)組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]

是不是很有規(guī)律呢？對的，用blockId.x和threadId.x可以很好的知道哪個線程干哪個元素，這個元素的下表就是threadId.x + 1000*blockId.x。

這里我想說的是，如果我們哪天糊涂了，畫一畫這個對應關系的表，也許，就更加清楚的知道我們分配的線程對應的處理那些東西啦。 

一維線程網(wǎng)絡，就先學這么多了。

二維網(wǎng)絡線程

eg2：int arr[32][16]二維的數(shù)組自增1。

第一個念頭，開個32*16個線程好了哇，萬事大吉！好吧。但是，朕現(xiàn)在想用二維線程網(wǎng)絡來解決，因為朕覺得一個二維的網(wǎng)絡去映射一個二維的數(shù)組，朕看的更加明了，看不清楚自己的士兵，如何帶兵打仗！

我還是畫個映射關系：

一個block中，現(xiàn)在是一個二維的thread網(wǎng)絡，如果我用了16*16個線程。

(0,0)，(0,1)，……(0,15)

(1,0)，(1,1)，……(1,15)

……

(15,0)，(15,1)，……(15,15)

呀，現(xiàn)在一個組內(nèi)的人稱呼變了嘛，一維網(wǎng)絡中，你走到一個小組里，叫3號出列，就出來一個，你現(xiàn)在只是叫3號，沒人會出來！這個場景是這樣的，現(xiàn)在你班上有兩個人同名的人，你只叫名，他們不知道叫誰，你必須叫完整點，把他們的姓也叫出來。所以，二維網(wǎng)絡中的(0,3)就是原來一維網(wǎng)絡中的3，二維中的(i,j)就是一維中的(j+i*16)。不管怎么樣，一個block里面能處理的線程數(shù)量總和還是不變的。

一個grid中，block也可以是二維的，一個block中已經(jīng)用了16*16的thread了，那我們一共就32*16個元素，我們用2個block就行了。

先給出一個代碼清單吧，程序員都喜歡看代碼，這段代碼是我抄襲的。第一次這么完整的放上代碼，因為我覺得這個代碼可以讓我說明我想說的幾個問題：

第一，二維數(shù)組和二維指針的聯(lián)系。

第二，二維線程網(wǎng)絡。

第三，cuda的一些內(nèi)存操作，和返回值的判斷。

簡要的來學習一下二維網(wǎng)絡這個知識點，

dim3 dimBlock(16,16); 

//定義block內(nèi)的thread二維網(wǎng)絡為16*16

dim3 dimGrid((COLS+dimBlock.x-1)/(dimBlock.x), (ROWS+dimBlock.y-1)/(dimBlock.y)); 

//定義grid內(nèi)的block二維網(wǎng)絡為1*2

unsigned int row = blockDim.y*blockIdx.y + threadIdx.y; 

//二維數(shù)組中的行號

unsigned int col = blockDim.x*blockIdx.x + threadIdx.x; 

//二維線程中的列號

三維網(wǎng)絡線程

dim3定義了三維的結構，但是，貌似二維之內(nèi)就能處理很多事情啦，所以，我放棄學習三維。網(wǎng)上看到的不支持三維網(wǎng)絡是什么意思呢？先放一放。

給自己充充電

同一塊顯卡，不管你是二維和三維或一維，其計算能力是固定的。比如一個block能處理1024個線程，那么，一維和二維線程網(wǎng)絡是不是處理的線程數(shù)一樣呢？

回答此問題，先給出網(wǎng)絡配置的參數(shù)形式——<<>>，各個參數(shù)含義如下：

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：這個是可選參數(shù)，設定最多能動態(tài)分配的共享內(nèi)存大小，比如16k，單不需要是，這個值可以省略或?qū)?。

S：也是可選參數(shù)，表示流號，默認為0。流這個概念我們這里不說。

接著，我想解決幾個你肯定想問的兩個問題，因為我看很多人想我這樣的問這個問題：

1 block內(nèi)的thread我們是都飽和使用嗎？

答：不要，一般來說，我們開128或256個線程，二維的話就是16*16。

2 grid內(nèi)一般用幾個block呢？

答：牛人告訴我，一般來說是你的流處理器的4倍以上，這樣效率最高。

回答這兩個問題的解釋，我想抄襲牛人的一段解釋，解釋的好的東西就要推廣呀：

GPU的計算核心是以一定數(shù)量的Streaming Processor(SP)組成的處理器陣列，NV稱之為Texture Processing Clusters(TPC)，每個TPC中又包含一定數(shù)量的Streaming Multi-Processor(SM)，每個SM包含8個SP。SP的主要結構為一個ALU（邏輯運算單元），一個FPU（浮點運算單元）以及一個Register File(寄存器堆)。SM內(nèi)包含有一個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模型中最基本的運行單元，執(zhí)行最基本的程序指令。Block是一組協(xié)作Thread，Block內(nèi)部允許共享存儲，每個Block最多包含512個Thread。Grid是一組Block，共享全局內(nèi)存。Kernel是在GPU上執(zhí)行的核心程序，每一個Grid對應一個Kernel任務。 在程序運行的時候，實際上每32個Thread組成一個Warp，每個 warp 塊都包含連續(xù)的線程，遞增線程 ID 。Warp是MP的基本調(diào)度單位，每次運行的時候，由于MP數(shù)量不同，所以一個Block內(nèi)的所有Thread不一定全部同時運行，但是每個Warp內(nèi)的所有Thread一定同時運行。因此，我們在定義Block Size的時候應使其為Warp Size的整數(shù)倍，也就是Block Size應為32的整數(shù)倍。理論上Thread越多，就越能彌補單個Thread讀取數(shù)據(jù)的latency ，但是當Thread越多，每個Thread可用的寄存器也就越少，嚴重的時候甚至能造成Kernel無法啟動。因此每個Block最少應包含64個Thread，一般選擇128或者256，具體視MP數(shù)目而定。一個MP最多可以同時運行768個Thread，但每個MP最多包含8個Block，因此要保持100%利用率，Block數(shù)目與其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