CUDA 线程执行模型分析（一）招兵------ GPU的革命

 

CUDA 线程执行模型分析（一）招兵

------ GPU 的革命

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

下面是得到的结果：

(0,0) (1,0) (2,0) (3,0) (4,0) (5,0) (6,0) (7,0) (8,0) (9,0) (10,0) (11,0) (12,0) (13,0) (14,0) (15,0) (16,0) (17,0) (18,0) (19,0) (20,0) (21,0) (22,0) (23,0) (24,0) (25,0) (26,0) (27,0) (28,0) (29,0) (30,0) (31,0) (32,0) (33,0) (34,0) (35,0) (36,0) (37,0) (38,0) (39,0) (40,0) (41,0) (42,0) (43,0) (44,0) (45,0) (46,0) (47,0) (48,0) (49,0) (50,0) (51,0) (52,0) (53,0) (54,0) (55,0) (56,0) (57,0) (58,0) (59,0) (60,0) (61,0) (62,0) (63,0) (64,0) (65,0) (66,0) (67,0) (68,0) (69,0) (70,0) (71,0) (72,0) (73,0) (74,0) (75,0) (76,0) (77,0) (78,0) (79,0) (80,0) (81,0) (82,0) (83,0) (84,0) (85,0) (86,0) (87,0) (88,0) (89,0) (90,0) (91,0) (92,0) (93,0) (94,0) (95,0) (96,0) (97,0) (98,0) (99,0) (100,0) (101,0) (102,0) (103,0) (104,0) (105,0) (106,0) (107,0) (108,0) (109,0) (110,0) (111,0) (112,0) (113,0) (114,0) (115,0) (116,0) (117,0) (118,0) (119,0) (120,0) (121,0) (122,0) (123,0) (124,0) (125,0) (126,0) (127,0) (128,0) (129,0) (130,0) (131,0) (132,0) (133,0) (134,0) (135,0) (136,0) (137,0) (138,0) (139,0) (140,0) (141,0) (142,0) (143,0) (144,0) (145,0) (146,0) (147,0) (148,0) (149,0) (150,0) (151,0) (152,0) (153,0) (154,0) (155,0) (156,0) (157,0) (158,0) (159,0) (160,0) (161,0) (162,0) (163,0) (164,0) (165,0) (166,0) (167,0) (168,0) (169,0) (170,0) (171,0) (172,0) (173,0) (174,0) (175,0) (176,0) (177,0) (178,0) (179,0) (180,0) (181,0) (182,0) (183,0) (184,0) (185,0) (186,0) (187,0) (188,0) (189,0) (190,0) (191,0) (192,0) (193,0) (194,0) (195,0) (196,0) (197,0) (198,0) (199,0) (200,0) (201,0) (202,0) (203,0) (204,0) (205,0) (206,0) (207,0) (208,0) (209,0) (210,0) (211,0) (212,0) (213,0) (214,0) (215,0) (216,0) (217,0) (218,0) (219,0) (220,0) (221,0) (222,0) (223,0) (224,0) (225,0) (226,0) (227,0) (228,0) (229,0) (230,0) (231,0) (232,0) (233,0) (234,0) (235,0) (236,0) (237,0) (238,0) (239,0) (240,0) (241,0) (242,0) (243,0) (244,0) (245,0) (246,0) (247,0) (248,0) (249,0) (250,0) (251,0) (252,0) (253,0) (254,0) (255,0) (256,0) (257,0) (258,0) (259,0) (260,0) (261,0) (262,0) (263,0) (264,0) (265,0) (266,0) (267,0) (268,0) (269,0) (270,0) (271,0) (272,0) (273,0) (274,0) (275,0) (276,0) (277,0) (278,0) (279,0) (280,0) (281,0) (282,0) (283,0) (284,0) (285,0) (286,0) (287,0) (288,0) (289,0) (290,0) (291,0) (292,0) (293,0) (294,0) (295,0) (296,0) (297,0) (298,0) (299,0) (300,0) (301,0) (302,0) (303,0) (304,0) (305,0) (306,0) (307,0) (308,0) (309,0) (310,0) (311,0) (312,0) (313,0) (314,0) (315,0) (316,0) (317,0) (318,0) (319,0) (320,0) (321,0) (322,0) (323,0) (324,0) (325,0) (326,0) (327,0) (328,0) (329,0) (330,0) (331,0) (332,0) (333,0) (334,0) (335,0) (336,0) (337,0) (338,0) (339,0) (340,0) (341,0) (342,0) (343,0) (344,0) (345,0) (346,0) (347,0) (348,0) (349,0) (350,0) (351,0) (352,0) (353,0) (354,0) (355,0) (356,0) (357,0) (358,0) (359,0) (360,0) (361,0) (362,0) (363,0) (364,0) (365,0) (366,0) (367,0) (368,0) (369,0) (370,0) (371,0) (372,0) (373,0) (374,0) (375,0) (376,0) (377,0) (378,0) (379,0) (380,0) (381,0) (382,0) (383,0) (384,0) (385,0) (386,0) (387,0) (388,0) (389,0) (390,0) (391,0) (392,0) (393,0) (394,0) (395,0) (396,0) (397,0) (398,0) (399,0) (400,0) (401,0) (402,0) (403,0) (404,0) (405,0) (406,0) (407,0) (408,0) (409,0) (410,0) (411,0) (412,0) (413,0) (414,0) (415,0) (416,0) (417,0) (418,0) (419,0) (420,0) (421,0) (422,0) (423,0) (424,0) (425,0) (426,0) (427,0) (428,0) (429,0) (430,0) (431,0) (432,0) (433,0) (434,0) (435,0) (436,0) (437,0) (438,0) (439,0) (440,0) (441,0) (442,0) (443,0) (444,0) (445,0) (446,0) (447,0) (448,0) (449,0) (450,0) (451,0) (452,0) (453,0) (454,0) (455,0) (456,0) (457,0) (458,0) (459,0) (460,0) (461,0) (462,0) (463,0) (464,0) (465,0) (466,0) (467,0) (468,0) (469,0) (470,0) (471,0) (472,0) (473,0) (474,0) (475,0) (476,0) (477,0) (478,0) (479,0) (480,0) (481,0) (482,0) (483,0) (484,0) (485,0) (486,0) (487,0) (488,0) (489,0) (490,0) (491,0) (492,0) (493,0) (494,0) (495,0) (496,0) (497,0) (498,0) (499,0) (500,0) (501,0) (502,0) (503,0) (504,0) (505,0) (506,0) (507,0) (508,0) (509,0) (510,0) (511,0)

一目了然，正是我们定义的 512个线程。

现在大家应该知道 CUDA是怎么运行了？0 – 0 要是还不清楚，跟贴问我吧 - - ！

 

好了，现在招兵的任务已经完成了，兵是招来了，下面就改训练了，不过天时已晚，大家还是吃饱喝足，好好的休息一下吧～

明天继续 GPU的革命。

且听：《沙家浜 》里胡傳魁已經出場了：十幾個人來七八條槍，不是很威風 嗎？ ----我可要说十几个人怎么来分这七八条枪。

士兵们执行任务了，是一哄而上，还是……且听 SIMD- -！

 

PS:上面在输出结果的时候用了一个小技巧，或许有的人知道，就不用看了，有不知道的，接着看下去，或许可以方便你以后的调试：）

再输出的时候，在系统 debugging信息里面的command Arguments里面用一个管道输出到一个1.txt文件里面。管道用法大家自己可以到网上查～我就不用讲解了。也很简单，就是在你的哦xx.exe 后面加一个 “>” 管道，对了还有 “<” 管道，呵呵：） 一个是输出，一个是输入。