0 序言
学习CUDA已经有个把月了,感觉自己学习一门新技术的第一个阶段已经接近尾声,对于一些基本的东西,学习的收获应该作一个总结,我是一个喜欢总结的人。
CUDA是异构编程的一个大头,洋洋洒洒的看了写资料,但是,感觉这个技术没有像C++或者Java那样有自己的权威的《编程思想》来指导系统学习,总是感觉心里不踏实,是不是自己还没掌握深入、或者说心里没底气说自己已经入门了、已经熟悉了、已经精通了。站在一个初学者的角度,作为一个笔记式的记录,讲解自己学习和理解CUDA过程中的一些列想到的、碰到的问题。享受一个东西不一定是结果,可以是从无知到了解到精通的这个整个过程。
1 给自己提几个问题
对的,我想要做什么事情的时候,习惯性的给自己提如下问题:
问题1:GPU高性能运算之CUDA(下午简称CUDA技术)是干嘛用的,我为什么要学它?
这个问题如果我回答是,纯粹为了掌握一门新的技术,如果你还是学生则可以,如果你是一个手里有项目的工程师,那么,我觉得没什么必要去学这个东西。我个人理解CUDA是计算机里面的边缘技术,是对程序执行性能的提高的一种方式,可以理解成一个工具。工具这种东西,你需要用的时候再去学,你不需要用的时候,你知道有这回事情就可以了。如果你用不到的东西,你也很贪心什么都去学,第一学不精,第二计算机的技术太多太广,对应我这样的智商一般的人来说是不靠谱的。
我要学习CUDA,因为项目的任务可分解性较大,粒度相关性小,某项目大概一个计算任务可以分解成3000*3000*300规模的计算,而且这些元素之间完全独立。又加入,你需要对全市100万人,每个人计算下该个体的每年的收入与支出的净值(从1900年计算到2000年),淡然有人说放excel中计算不就得了,好吧……我只是讲个通俗点的例子。那么这个计算规模是1000000*10,倘若要对每个人每年在做点其他什么高级的算法得出一个什么指数,恐怕excel还是实现起来不太容易。所以,这种并行度很大的问题,我们可以考虑用CUDA来解决。也许你也知道,搞图像的,CUDA就显得很重要啦。
书上说:
问题2:我的基础是什么?
学习新的技术,我喜欢和之前学过的某个还熟悉的东西做为比对,这样理解起来可能会快一些。比如学Java的时候,我想着C++,学UML的时候,我想着面向对象编程,学CUDA呢?我有什么么?可能很多人会没有异构编程的历史。我很幸运,之前弄过半年的OpenMP,对的——多核编程技术。后来我了解到,如果是多个GPU,是需要用到OpenMP来做的。OpenMP本身和CUDA好像并没什么关系,但是,片上多核的并行算法是很想通的。很好的一个例子是:奇偶排序。哈哈,给自己提了点学好CUDA的信心。
好了,明白自己的需求和基础,给自己一个学习的定位,什么地方该花时间去琢磨,相比应该很清楚。
2 你好,GPU!
GPU有两大开发商——英伟达和AMD,支持CUDA的是英伟达,很好啊,之前买电脑是有先见之明的,买了英伟达的显卡——GT520,不是特别高端,和单位的GTX650ti2G比起来,有那么一点逊色,但是,好歹可以跑CUDA!
英伟达支持CUDA编程的显卡型号从G8800开始,都是可以的。一开始作为图像处理用,而今,天文地理、数学金融、医疗军事等等,都开始尝试发挥GPU的优势。
GPU的计算核心也是隔年换代,现在已经倒GK10X了,计算能力也是逐渐提升,目前已经最好的有3.0。GPU的架构从原始的到费米的,再到开普勒的,我们没必要去一个个了解,我们先了解GPU的这个大概的历史,免得和人交谈说不出一和二,脱离菜鸟嘛!
GPU和CPU的区别可以参考下,讲的还算详细和明了。
http://www.cnblogs.com/viviman/archive/2012/11/26/2789113.html
可以这样的去理解,单核CPU多线程并行,是感官上的并行,世界上在CPU上还是串行指令在跑;而在GPU上,才叫真正的并行!需要介绍下CUDA1.0开发包是支持在CPU上模拟GPU开发的,其原理就是用多线程来模拟;而现在的版本就不支持了。最新的是5.0的,官网上是下不到之前的了,反正我是找到腿软了还没找到。1.0是古董了!如果你有,一定要给我开开眼见哦。
我们的CUDA编程,很明显是GPU和CPU一起来处理的嘛——异构编程!对的,我自己的理解是就一个工程而言:CPU处理串行计算业务,GPU处理并行计算业务,这里将的并行都是并行度可观的哦,不是说两个元素你也来GPU上计算,那样是不环保的——会浪费很多GPU的资源!
说道环保,我想多说一句,在GPU编程,很体现“绿色”理念。48个核,你如果写的好,48个核全在干活,而且干的是有意义的活,那么你是合理的利用资源,如果你只让一个核在干有意义的活,其他的都在空转,那么你很浪费电哦。
对于CPU和GPU分配自己的业务,稍微画个图失意一下,如图1:
图1 工程中GPU和CPU的分工
总的来说,GPU只是干计算并行度高的功能模块的活,一定不可以越权啦!
3 你好,CUDA!
3.1 开发环境配置
第一次和同事交谈,我说你和我说说C-U-D-A。他说:哭打……苦打…… ……。半天后,我说,你和我说说C-U-D-A,你说哭打是什么东西……它说哭打就是C-U-D-A。我操,顿时傻眼了,原来这东西行业里年哭打,对的,Cu - Da,连起来就是这么发音的,我的无知啊。我们还是用中文解释吧:CUDA的意思是统一计算设备架构。
CUDA的集成开发环境可以参考下:
http://www.cnblogs.com/viviman/archive/2012/11/05/2775100.html
在win7+vs2008+CUDA5.0的环境下体验,是一种新的尝试,我自己配的时候,很少或者就没有5.0的配置博客文档等等。因为5.0的那一场雪比2010年来的更晚一些……
5.0是和2.0、3.0、4.0都会有那么一点不同的,是集成了SDK和TOOL两个东西,之前是分开的,现在是合在一起的。其实是差不多的,但是,这一分一合,就会给人很不习惯。不过,我虽然愚钝,但是,试了几次之后还是摸索成功。当第一个helloworld输出后,心里是有那么一点小激动的,立马跑到小区门口,买了半斤羊肉吃了,因为为了配这一套环境,我中午饭都没吃!我这只哭打小菜鸟就是可怜啊!
3.2 特殊的"hello world"
搭建好环境,你肯定想看看我的helloworld程序,对的,但是我不想输出helloworld,我想干一件事情是:我在CPU上创建个变量,传到GPU中,然后,在GPU中赋值,然后传出来。这件事情如果成功了,是不是可以说明,通了+GPU工作了!网上一搜的CUDA的helloworld程序,都是在CPU上输出helloworld,那多不过瘾。
__global__ void hello(char *ch)
{
ch = {'h', 'e', 'l', 'l', 'o'};
}
int main()
{
……
hello<<<1,1>>>(dev_ch);
……
return 0;
}
I think you know my idea.
在你网上搜索到的程序的基础上,作这样的一个改变,相信自己动手的才是快乐的。
4 敲开编程的门
我习惯性的喜欢先看一门语言的关键字,CUDA的关键字很简单很少:
函数类型
__global__
用来修饰内核函数的,内核函数是什么呢,内核函数是跑在GPU上的函数;与之对应的是主机函数,用__host__修饰,也可以缺省,跑在CPU上。因此,CPU也叫主机,GPU也叫设备。通常定义这个内核函数,我喜欢在函数名前加个kernel作为修饰,让自己清楚点。
比如__global__ void kerneladd(float *a){}
__device__
也是用来修饰内核函数的,那和__global__有什么区别吗?对的。__global__修饰的内核函数只能被主机函数调用;__device__修饰的内核函数只能被内核函数调用,应该很好理解。
__host__
主机函数,供主机函数调用,可缺省哦,一般情况下,都是缺省的,知道这个东西就行。
存储类型
寄存器:在核函数内 int i即表示寄存器变量。
__global__:全局内存。在主机函数中开辟和释放。
__shared__:共享存储,每个block内的线程共享这个存储。
__constant__:常量存储,只读。定义在所有函数之外,作用范围整个文件。
__texture__:纹理存储,只读。内存不连续。
内建变量
dim3
threadId
blockId
gridId