没有合适的资源?快使用搜索试试~ 我知道了~
CUDA学习入门资料整理
5星 · 超过95%的资源 需积分: 32 49 下载量 133 浏览量
2009-03-04
13:25:36
上传
评论 5
收藏 1.49MB PDF 举报
温馨提示
试读
73页
资料内容整理至openhero(赵开勇)的blog,非常好的CUDA入门资料。(thanks to openhero),所有贡献和辛苦属于赵开勇
资源推荐
资源详情
资源评论
CUDA 线程执行模型分析(一)招兵
------ GPU 的革命
序:或许看到下面的内容的时候,你会觉得和传统的讲解线程,和一些讲
解计算机的书的内容不是很相同。我倒觉得有关计算机,编程这些方面的内容,
并不都是深奥难懂的,再深奥难懂的事情,其实本质上也是很简单的。一直以为
计算机编程就像小时候搭建积木一样,只要知道游戏规则,怎么玩就看你自己了。
或许是从小学那会,就喜欢在做数学题的时候用一些简便方法来解题,养成了一
些习惯,喜欢把复杂的问题都会尝试用最简单的额方法来解决,而不喜欢把简单
的问题弄得很复杂。不再多说了,有的朋友已经看得不耐烦了……ps:再罗嗦一
句,如果下面看不明白的,就当小说看了,要是觉得不像小说,那就当故事看,
要是觉得故事不完整,写得太乱,那就当笑话看,在各位学习工作之余能博得大
家一笑,也倍儿感荣幸……ps2:想好再说……突然想到了,确实是了一段时间再
想到的,既然叫 GPU 革命,那就得招集队伍啊,下面我就开始招兵了。
正题:
要真正进入 CUDA 并行化开发,就必须先了解 CUDA 的运行模型,才能在
这个基础上做并行程序的开发。
CUDA 在执行的时候是让 host 里面的一个一个的 kernel 按照线程网格(G
rid)的概念在显卡硬件(GPU)上执行。每一个线程网格又可以包含多个线程
块(block),每一个线程块中又可以包含多个线程(thread)。
在这里我们可以拿古时候的军队作为一个例子来理解这里的程序执行模型。
每一个线程,就相当于我们的每一个士兵,在没有当兵之前,大家都不知道自己
做什么。当要执行某一个大的军事任务的时候,大将军发布命令,大家来要把对
面的敌人部队的 n 个敌人消灭了。然后把队伍分成 M 个部分,每一个部分完成
自己的工作,有的是做侦查的工作,有的是做诱敌的工作,有的是做伏击的工作,
有的是做后备的工作,有的是做后勤的工作……反正把一个大任务按照不同的类
别,不同的流程不同,分别由 M 个部分来完成。
这里我们可以把大将军看着是 Host,它把这次军事行动分解成一个一个的
kernel:kernel_1,kernel_2……kernel_M,每一个 kernel 就交给每一个 G
rid(副将?千户?就看管的人多人少了,如果 GPU 硬件支持少一点,那就是千
户;要是 GPU 硬件高级一些,管理的人多一些,那就副将?戚家军也不过四五
千人,咱也不能太贪心,一下子就想统军百万,再说了,敢问世上韩信一样的将
才又哪有那么多啦?)来完成。当要执行这些任务的时候,每一个 Grid 又把任
务分成一部分一部分的,毕竟人太多,他一个官不过来,他只要管理几个团队中
间的高级军官就可以了。Grid 又把任务划分为一个个的 Block(百户?),这
里每一个 Grid 管理的 Block 也是有限的,(人就那么多……想管多少得看硬件的
支持)。毕竟显卡上的 GPU 硬件还是很少,Thread(线程)相对于真正的军队
来说人还是少了很多。所以到 Block 这个层的时候,就直接管理每一个 Threa
d(士兵)。
由于古代通信不是很方便(从 GPU 的发展史来看,如果按照中国的历史,
现在的 GPU 也就还处在战国时代吧……),所以每一个 Block(百户)内部的 T
hread(士兵)才能方便的通信,按照既定的规则进行同步;而各个 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 叫到他的号了,他得马上回答:我在 Blo
ck(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 就隐藏在 co
py 的代码当中……so,截图……
剩余72页未读,继续阅读
资源评论
- taotao9304182015-04-22适合初学者,看看内容吧。
- leimingjun9009282013-05-16挺好的 对于初学很有用
- kelnawei2013-03-13好东西 适合初学者理解 但是里面的图片看不了 郁闷
- Fuhao19912014-01-01适合初学者,里边的内容还是蛮好用的
jazzwong
- 粉丝: 0
- 资源: 2
上传资源 快速赚钱
- 我的内容管理 展开
- 我的资源 快来上传第一个资源
- 我的收益 登录查看自己的收益
- 我的积分 登录查看自己的积分
- 我的C币 登录后查看C币余额
- 我的收藏
- 我的下载
- 下载帮助
安全验证
文档复制为VIP权益,开通VIP直接复制
信息提交成功