序言：有个不会写计算机程序的朋友看了blog，问我，这个GPU也能当故事写吗？我觉得或许GPU真的算是一场革命吧，他的发展或许在酝酿中，不过到08年底，09年初，一定会有一场轰轰烈烈的竞争。那个时候或许从OS层面都会给人带来震撼。如果把CPU的多core看成由几个特种兵组成的，每个特种兵都手里面都拿着8杆枪（SSE）。那么GPU可以看成农民起义……一上来就是成百上千的人，虽然单兵作战能力比不上CPU的单个core，但是毕竟人数众多。就现在GPU的性能，在并行运算上如果不考虑double硬件的成本，已经早早超过CPU的并行运算能力。这或许就是一场革命，这次革命不知是简简单单的GPU和CPU的转变，而是并行算法和串行算法的竞争。并行算法虽然研究到现在已经有很多年，但是真正的实际运用，离我们普通大众还是差很远。但是GPU，并行计算的出现，一下子把我们和并行计算的距离拉近了好多。现在在学校里面学习计算机的时候都是从串行算法开始，养成了很多固定的串行思维。当遇到问题并行划分的时候，就还带着串行的思想，那就不好了：）

　　正文：前面我们已经说到线程的一些概念，但是这些概念都是软环节的。我们常常会听到某某单位说他们的软硬件配置如何如何的好。软件再好，每个士兵都是可造之才，但是如果硬件条件跟不上，也没他们的勇武之地。就像国内做过一些跳槽原因的统计，很多人已经跳槽都是为了高工资，但实际统计结果表明，很多人都觉得在以前那个公司里面学不到东西，或者说得不到发挥自己长处的地方。那就得看公司有没有这样的机会让你发挥你的能力了。看到这里，扯远了～很多人都看得不耐烦了……书接上回《CUDA 线程执行模型分析（二）大军未动粮草先行------GPU的革命》。已经讲到CUDA在线程模型是一个什么样子，经过几天的吸收，也应该在脑子里面有一些印象了。但是你会问一个问题，我们是不是可以开无数个线程来执行啦？或许招兵的人都想找到很多很多的人，但是你得也考虑一下你的粮食有多少，军营有多大。在这么我们得讨论一下现在的支持CUDA的nvidia的显卡现在的硬件情况。

　　<!--[if !vml]--><!--[endif]-->空洞的讲解或许还是没怎么又说服力，下面以G80为例子。

　　1. G80里面有16个Multiprocessor.

　　2. 每个Multiprocessor都有一组（G80里面是8个）32位的Processor（每个Processor都是SIMD架构，什么叫SIMD架构：军训的时候，大家都到了食堂，不是像在学校里面，每个人那自己的碗筷就三三俩俩的去吃饭，那可要讲究纪律，啥叫纪律，一群人站在桌子面前，连长没发话，谁也不敢坐下来……连长一声令下：“坐下”。所有的人才按照敢坐下来，也是同时坐下来- -！要是谁没有同步坐下来，那就惨了- -！再来一次，一定要是同步坐下去的，都能听到声的，咵！恨不得把板凳给坐碎了------还是部队的东西结实，兄弟们怎么坐到军训结束只有坐坏屁股的，没有听说凳子坐坏的- -！所以啊，爱惜公物就是爱惜自己。记住了吧，这就是SIMD：Single Instruction Multiple Data 。 ）还有共用的Instruction Unit（这玩儿就不用翻译是啥了吧～看SIMD，自己理解去）。在G80里面有两个SFU模块。

　　3. 每一个时钟周期内，按照warp（这玩儿咋翻译啦？就理解为运行的时候，一个block里面一起运行的thread，例如block里面有512个thread，但是每次只有32个thread在运行，那么这32个thread就是一个运行的warp组- -! 还好不是rap- -！俺就真没法解释了）

　　4. 每一个warp里面包含的thread数量是有限的，现在的规定是32个。将来不知道会不会有变化？不知道，这个只有CUDA开发人员知道了。

　　我们还是按照我们的既定方式学习吧～看图说话- -！接下来又是一张图： <!--[if !vml]--><!--[endif]-->

　　刚才我们已经说到了硬件上的处理器的模型，现在我们来看看内存吧。前面的章节我们已经讲到了内存的register，local，shared，constant，texture，global，现在看到对应的硬件的位置了吧？哦，对了，有人会问，local和global跑那里去了，你看看，是不是多了一个Device Memory嘛。Local和Global 都是在Device Memory上的。只是在线程模型运行的时候，为了方便说明线程的模型，把它单独划分出来好说明。其实Constant和Texture都是在Device Memory里面的，但是人家是只读的Cache，所以比Global和local要快一些。不过Cache和shared Memory比起来，又显得慢一些了。

　　硬件的架构差不多就是这么多。如果还不是很清楚，我们再讲讲军营的生活。还是讲吃食堂吃饭吧（回想起大学的时候和几个好朋友一起成立了一个编程爱好者协会……结果后来变成了好吃协会- -！）。由于我们现在的连队人数众多，一个grid里面就可以有65535个thread，所以一下子让这么多人一起吃饭，是不可能的。我们有好多个grid怎么办？所以每次Device（就是显卡）只处理一个grid。但是每个grid里面又有那么多人。一来是咱们地方小，不能容纳那么多人一起吃饭，而来是也没必要再建立那么多食堂。大家可以轮流吃，早吃的早走，晚来的也晚不了多久。

　　这次是Grid1先到达食堂，然后他把士兵分成了几个block，（我们上次已经讲过的，block和grid的关系，应该还记得把）但是现在桌子还是少，为了大家不乱起来，每个block的人就指定到某一个multiprocessor（一个周期只能容纳32个人一起用餐）哪里。

　　这个block1的人就围着multiprocessor1这张桌子坐，但是这张桌子也不是太大，所以block里面的人又分成32个人一组（warp）来吃饭。那如果这个block没有32个人，就占不完这张桌子，所以一个multiprocessor又可以让别的block的thread一起来用餐。

　　因为一次就把所有的register和shared memory分配给了所有的block的所有的thread。所以怎么样合理的少用register和shared memory就可以同时让更多的thread得到处理。

　　讲了这么多，怎么来看自己的显卡那片军营如何拉？下面给出两种方法，一种是用一个小软件，GPU-Z的工具，这个工具可以得到很多显卡的参数，这个可以上网搜索到。

　　<!--[if !vml]--><!--[endif]-->

　　第二种方法是用CUDA的命令，这个只是针对支持CUDA的显卡的，下面给一段代码：

　　/********************************************************************

　　*　InitCUDA.cu

　　*　This is a init CUDA of the CUDA program.

#include <stdio.h>

#include <stdlib.h>

#include <cuda_runtime.h>

/************************************************************************/

/* Init CUDA　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　*/

/************************************************************************/

bool InitCUDA(void)

{

　int count = 0;

　int i = 0;

　cudaGetDeviceCount(&count);

　if(count == 0) {

　　　　fprintf(stderr, "There is no device.n");

　　　　return false;

　}

　for(i = 0; i < count; i++) {

　　　　cudaDeviceProp prop;

　　　　if(cudaGetDeviceProperties(&prop, i) == cudaSuccess)

　　　　{　　　　　　　　

　　　　　　　printf("name:%sn",　　　　　　　　　　　　　　　　　　　　 prop.name);

　　　　　　　printf("totalGlobalMem:%un",　　　　　　　　　 prop.totalGlobalMem);

　　　　　　　printf("sharedMemPerBlock:%un",　　　　　　　　prop.sharedMemPerBlock);

　　　　　　　printf("regsPerBlock:%dn",　　　　　　　　　　　　　 prop.regsPerBlock);

　　　　　　　printf("warpSize:%dn",　　　　　　　　　　　　　　　 prop.warpSize);

　　　　　　　printf("memPitch:%un",　　　　　　　　　　　　　　　 prop.memPitch);

　　　　　　　printf("maxThreadsPerBlock:%dn",　　　　　　　 prop.maxThreadsPerBlock);

　　　　　　　printf("maxThreadsDim:x %d, y %d, z %dn",　　　prop.maxThreadsDim[0],prop.maxThreadsDim[1],prop.maxThreadsDim[2]);

　　　　　　　printf("maxGridSize:x %d, y %d, z %dn",　　　prop.maxGridSize[0],prop.maxGridSize[0],prop.maxGridSize[0]);

　　　　　　　printf("totalConstMem:%un",　　　　　　　　　　prop.totalConstMem);

　　　　　　　printf("major:%dn",　　　　　　　　　　　　　　　　　prop.major);

　　　　　　　printf("minor:%dn",　　　　　　　　　　　　　　　　　prop.minor);

　　　　　　　printf("clockRate:%dn",　　　　　　　　　　　　　　　prop.clockRate);

　　　　　　　printf("textureAlignment:%un",　　　　　　　　 　　　prop.textureAlignment);

　　　　　　　if(prop.major >= 1) {

　　　　　　　　　　break;

　　　　　　　}

　　　　}

　}

　if(i == count) {

　　　　fprintf(stderr, "There is no device supporting CUDA 1.x.n");

　　　　return false;

　}

　cudaSetDevice(i);

　printf("CUDA initialized.n");

　return true;

}

　　序：或许看到下面的内容的时候，你会觉得和传统的讲解线程，和一些讲解计算机的书的内容不是很相同。我倒觉得有关计算机，编程这些方面的内容，并不都是深奥难懂的，再深奥难懂的事情，其实本质上也是很简单的。一直以为计算机编程就像小时候搭建积木一样，只要知道游戏规则，怎么玩就看你自己了。或许是从小学那会，就喜欢在做数学题的时候用一些简便方法来解题，养成了一些习惯，喜欢把复杂的问题都会尝试用最简单的额方法来解决，而不喜欢把简单的问题弄得很复杂。不再多说了，有的朋友已经看得不耐烦了……ps：再罗嗦一句，如果下面看不明白的，就当小说看了，要是觉得不像小说，那就当故事看，要是觉得故事不完整，写得太乱，那就当笑话看，在各位学习工作之余能博得大家一笑，也倍儿感荣幸……ps2：想好再说……突然想到了，确实是了一段时间再想到的，既然叫GPU革命，那就得招集队伍啊，下面我就开始招兵了。

　　正题：

　　要真正进入CUDA并行化开发，就必须先了解CUDA的运行模型，才能在这个基础上做并行程序的开发。

　　CUDA在执行的时候是让host里面的一个一个的kernel按照线程网格（Grid）的概念在显卡硬件（GPU）上执行。每一个线程网格又可以包含多个线程块（block），每一个线程块中又可以包含多个线程（thread）。

　　在这里我们可以拿古时候的军队作为一个例子来理解这里的程序执行模型。每一个线程，就相当于我们的每一个士兵，在没有当兵之前，大家都不知道自己做什么。当要执行某一个大的军事任务的时候，大将军发布命令，大家来要把对面的敌人部队的n个敌人消灭了。然后把队伍分成M个部分，每一个部分完成自己的工作，有的是做侦查的工作，有的是做诱敌的工作，有的是做伏击的工作，有的是做后备的工作，有的是做后勤的工作……反正把一个大任务按照不同的类别，不同的流程不同，分别由M个部分来完成。

　　这里我们可以把大将军看着是Host，它把这次军事行动分解成一个一个的kernel：kernel_1，kernel_2……kernel_M，每一个kernel就交给每一个Grid（副将？千户？就看管的人多人少了，如果GPU硬件支持少一点，那就是千户；要是GPU硬件高级一些，管理的人多一些，那就副将？戚家军也不过四五千人，咱也不能太贪心，一下子就想统军百万，再说了，敢问世上韩信一样的将才又哪有那么多啦？）来完成。当要执行这些任务的时候，每一个Grid又把任务分成一部分一部分的，毕竟人太多，他一个官不过来，他只要管理几个团队中间的高级军官就可以了。Grid又把任务划分为一个个的Block（百户？），这里每一个Grid管理的Block也是有限的，(人就那么多……想管多少得看硬件的支持)。毕竟显卡上的GPU硬件还是很少，Thread（线程）相对于真正的军队来说人还是少了很多。所以到Block这个层的时候，就直接管理每一个Thread（士兵）。

　　由于古代通信不是很方便（从GPU的发展史来看，如果按照中国的历史，现在的GPU也就还处在战国时代吧……），所以每一个Block（百户）内部的Thread（士兵）才能方便的通信，按照既定的规则进行同步；而各个block之间就没那么方便了，大家不能互相通讯。不过同一个（千户）Grid管理的block之间是共享同一个任务分配的资源的。每一个Grid都可以从大将军那里分配到一些任务，和一些粮食，同一个Grid的block都可以分到这个Grid分配到的粮食。而每一个（千户）Grid本身的任务就不一样，所以Grid除了知道自己做的事情外，其他的Grid他都不会知道了。----这差不多就是一个运行模型。下面让我们来看看在GPU中东图例说明：

　　看到这张图，我们可以对应来讲解我们的Thread部队。一个大将军Host，分配了任务中的两个任务（Kernel1, Kernel2）给了千户（Grid1，Grid2）来完成。千户Grid1里面把自己的队伍分成了6个百户Block，然后每一个百户又把任务分配给了自己的士兵（Thread）来具体完成。这里得说明的是，由于千户拿到的任务Kernel是定了的，所以到每个士兵（Thread）也就那里就只会埋头做同样的事情（就像戚继光招的兵：在胡宗宪的幕僚郑若曾所著的《江南经略》中，有着这样一份详细的招生简章，如果不服气，大可以去对照一下：凡选入军中之人，以下几等人不可用，在市井里混过的人不能用，喜欢花拳绣腿的人不能用，年纪过四十的人不能用，在政府机关干过的人不能用。以上尚在其次，更神奇的要求还在下面：喜欢吹牛、高谈阔论的人不能用，胆子小的人不能用，长得白的人不能用，为保证队伍的心理健康，性格偏激（偏见执拗）的人也不能用。……概括起来，戚继光要找的是这样一群人：四肢发达，头脑简单，为人老实，遵纪守法服从政府，敢打硬仗，敢冲锋不怕死，具备二愣子性格的肌肉男。----《明朝那些事儿》）。

　　为了方便统一管理，大家都去掉了自己的名字，按照Grid1，Block（x，y），Thread（x，y）这样的编号来称呼每一个Thread士兵。如果你要找到某一个Thread，你就跑到军营里面大叫：喂，Grid1手下的Block1管理的三排第二个Thread（1，2）出来。对于每个士兵自己来说，他要知道自己的位置，就得知道自己的长官都是谁。Thread（1，2）要知道自己再整个Grid手下算第几个兵（钢七年第……个兵），当Grid1叫到他的号了，他得马上回答：我在Block（1，1）的编号是：　　

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

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

　　如果要得到线性的编号，我们可以自己算一下：这个士兵是在Grid部队下的第xIndex行，yIndex列站着（这里我们必须注意：blockDim.x （这里为5），blockDim.y（这里为3）不是block的坐标，这里是block的size，切记！）。如果从第一个士兵哪里算是线性编号0，那他的线性编号就是

　　unsigned int index = xIndex（6） + size_x * yIndex（5）; 这里的size_x就是一行一共有多少个士兵（Thread），例如上图，这里一行有3个block每一个block里面的每一行有5个Thread，所以size_x就应该为3×5=15，一个Grid的一行有15个士兵，那刚才叫道的那个人的线性编号就应该是……还要我算吗？如果不得81，自己再算一下……（编号是从0，开始的）计算过程：index = 6+5×15=81；

　　说了那么多，咱也不能光说不练假把式。下面给一个简单的Thread 测试的Demo。本来打算把整个代码都copy过来，但是考虑到又会被别人copy出去，这样copy来copy去，在编程中很容易出现错误，所以作者也不提倡在做编程中直接copy代码，这样很危险的，很多自己都不知道的bug就隐藏在copy的代码当中……so，截图……

　　作者这里做了一个简单的测试，测试了512个线程。这里只有一个grid，从这一个grid里面也只分了一个bock：

　　dim3　　grid(size_x / BLOCK_DIM, 1);--》　　dim3grid(1, 1);一个Grid里面一个blcok

　　dim3　　block(BLOCK_DIM, 1, 1);--》dim3 block(512, 1, 1);一个Block里面分配512个Thread；

　　这里的每一个任务kernel就是：

__global__ static void ThreadDemo1(unsigned int* ret)

{

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

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

if(xIndex < size_x && yIndex < size_y)

{

　　　　unsigned int index = xIndex + size_x * yIndex;

　　　　ret[index] = xIndex;

　　　　ret[index + size_x*size_y] = yIndex;

}

}

　　计算自己的线性id然后把自己的坐标写入到线性id对应的数组里面。Ps：说明一下，这个记录id坐标的数组ret[],ret的前一半记录的是线程的x坐标，后一般是记录的y坐标。PS2：题外话，cu是C的扩展，这里的const定义的常量的用法在ANSIC C里面是行不通的，但是在C++中是可用的。

　　每个任务kernel都说好了，然后就是host下达命令

　　ThreadDemo1<<<grid,block>>>(ret);

　　来运行程序。所有的士兵都开始工作了，把自己的坐标x，y写入到ret数组里面。

　　<!--[if !vml]--><!--[endif]-->空洞的讲解或许还是没怎么又说服力，下面以G80为例子。

　　1. G80里面有16个Multiprocessor.

　　2. 每个Multiprocessor都有一组（G80里面是8个）32位的Processor（每个Processor都是SIMD架构，什么叫SIMD架构：军训的时候，大家都到了食堂，不是像在学校里面，每个人那自己的碗筷就三三俩俩的去吃饭，那可要讲究纪律，啥叫纪律，一群人站在桌子面前，连长没发话，谁也不敢坐下来……连长一声令下：“坐下”。所有的人才按照敢坐下来，也是同时坐下来- -！要是谁没有同步坐下来，那就惨了- -！再来一次，一定要是同步坐下去的，都能听到声的，咵！恨不得把板凳给坐碎了------还是部队的东西结实，兄弟们怎么坐到军训结束只有坐坏屁股的，没有听说凳子坐坏的- -！所以啊，爱惜公物就是爱惜自己。记住了吧，这就是SIMD：Single Instruction Multiple Data 。 ）还有共用的Instruction Unit（这玩儿就不用翻译是啥了吧～看SIMD，自己理解去）。在G80里面有两个SFU模块。

　　3. 每一个时钟周期内，按照warp（这玩儿咋翻译啦？就理解为运行的时候，一个block里面一起运行的thread，例如block里面有512个thread，但是每次只有32个thread在运行，那么这32个thread就是一个运行的warp组- -! 还好不是rap- -！俺就真没法解释了）

　　4. 每一个warp里面包含的thread数量是有限的，现在的规定是32个。将来不知道会不会有变化？不知道，这个只有CUDA开发人员知道了。

　　我们还是按照我们的既定方式学习吧～看图说话- -！接下来又是一张图： <!--[if !vml]--><!--[endif]-->