轉載自 從0開始學習《GPU高性能運算之CUDA》——2
5 GPU也不允許偏心
并行的事情多了，我們作為 GPU 的指令分配者，不能偏心了——給甲做的事情多，而乙沒事做，個么甲肯定不爽的來。所以，在 GPU 中，叫做線程網絡的分配。
我們將具體點的，在主機函數中如果我們分配的是這樣的一個東西：
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 = threadIdx.x + blockIdx.x * 16
這里有一個講究，block 是有維度的，一維、二維、三維。
對于一維的 block : tid = threadIdx.x 對于（Dx，Dy）二維的 block : tid = threadIdx.x + Dx*threadIdx.y 對于（Dx，Dy，Dz）三維的 block : tid = threadIdx.x + Dx*threadIdx.y + Dz*Dy*threadIdx.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 = threadIdx.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 = threadIdx.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 個線程出列。每個人用組號和組內的編號定了位置。在線程網絡中，blockIdx.x 和 threadIdx.x 就是對應的組號和組內編號啦，我必須要這里開始形象點表示這個對應關系，如果這個對應關系是這樣子的[blockIdx.x，threadIdx.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]
是不是很有規律呢？對的，用 blockIdx.x 和 threadIdx.x 可以很好的知道哪個線程干哪個元素，這個元素的下表就是 threadIdx.x + 1000*blockIdx.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*16dim3 dimGrid((COLS+dimBlock.x-1)/(dimBlock.x), (ROWS+dimBlock.y-1)/(dimBlock.y)); //定義grid內的block二維網絡為1*2unsigned 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。流這個概念我們這里不說。
接著，我想解決幾個你肯定想問的兩個問題，因為我看很多人想我這樣的問這個問題：
block 內的 thread 我們是都飽和使用嗎？ 答：不要，一般來說，我們開 128 或 256 個線程，二維的話就是 16*16。
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 個T hread 組成一個 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
這些電很重要啊，必須要充！不然，我就很難理解為什么網絡線程如何分配的。
新聞熱點
疑難解答
