序：或许看到下面的内容的时候，你会觉得和传统的讲解线程，和一些讲解计算机的书的内容不是很相同。我倒觉得有关计算机，编程这些方面的内容，并不都是深奥难懂的，再深奥难懂的事情，其实本质上也是很简单的。一直以为计算机编程就像小时候搭建积木一样，只要知道游戏规则，怎么玩就看你自己了。或许是从小学那会，就喜欢在做数学题的时候用一些简便方法来解题，养成了一些习惯，喜欢把复杂的问题都会尝试用最简单的额方法来解决，而不喜欢把简单的问题弄得很复杂。不再多说了，有的朋友已经看得不耐烦了……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;