5 GPU也不允许偏心

并行的事情多了，我们作为GPU的指令分配者，不能偏心了——给甲做的事情多，而乙没事做，个么甲肯定不爽的来。所以，在GPU中，叫做线程网络的分配。首先还是来看下GPU的线程网络吧，图2：

file:///C:/Users/HuangFX/AppData/Local/Temp/ksohtml/wps_clip_image-19927.png

图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有如下几种设定方式： &#216; 2 blocks x 384 threads &#216; 3 blocks x 256 threads &#216; 4 blocks x 192 threads &#216; 6 blocks x 128 threads &#216; 8 blocks x 96 threads

这些电很重要啊，必须要充！不然，我就很难理解为什么网络线程如何分配的。

出处: http://www.cnblogs.com/viviman/archive/2012/11/28/2792521.html