老所工作室

OpenCL，并行计算模型的统一

十一月 15th, 2008 by Soloman |

| 2,035 views

2007年的GPU领域，nVidia无疑是耀眼的，这主要归功于其统一计算架构CUDA的推出与普及。在那些支持CUDA的GPU设备上（比如GeForce 8系列和Tesla），其强大的计算能力被释放到了通用计算上，而不仅仅局限于图形处理。我曾经做过一个基于CUDA的程序，用粒子跟踪算法（Particle Tracking）来从海量的试验数据中分析粒子的布朗运动路径，虽然这个算法不是非常适合用并行计算来计算，但在Tesla设备上，其效率较之CPU提高了近40倍。

虽然我们的日常生活表面上看上去，用不到什么海量的计算，但在科学、气象、金融甚至是军事领域，高速的计算仍然是一个巨大的需求。于是乎，Intel、ATI、IBM都纷纷研制自己的通用并行计算平台。一时间，场面比较混乱。比如我们如果利用CUDA开发了一个高速视频转换的软件，可以大大提高将普通视频文件转换成适合手持设备观看的视频。这是一个非常有用的应用。但是，我们的程序只能运行在支持CUDA的nVidia设备上，如果你使用的是ATI显卡，那么这个程序将一无是处。

天下大势，分久必合。这次关于通用计算平台的统一，发起于苹果公司。苹果电脑一向以拥有高端的显卡著称，自然而然地，他们想将这些计算能力释放出来。这个新的标准就是OpenCL，英文全称是 Open Computing Language。他们将这个标准提交到了The Khronos Group，这是一个业界的标准化组织，目前管理着OpenGL和OpenAL。这两个标准主要应用于图形和声音计算，显然，OpenCL将是这两个标准的升级，甚至替代？

OpenCL将统一管理一台电脑上的所有计算资源，比如你的多核CPU，以及多个GPU。OpenCL将这些资源统一看待，算作计算单元，并配置各种级别的内存：private, local 和 global。每个计算单元可以容纳多个工作单元，类似CUDA里的thread概念。

在OpenCL里，每个工作单元执行一段kernel代码，比如下面这段计算FFT的代码：

// This kernel computes FFT of length 1024. The 1024 length FFT is decomposed into
// calls to a radix 16 function, another radix 16 function and then a radix 4 function
__kernel void fft1D_1024 (__global float2 *in, __global float2 *out,
                          __local float *sMemx, __local float *sMemy) {
  int tid = get_local_id(0);
  int blockIdx = get_group_id(0) * 1024 + tid;
  float2 data[16];
  // starting index of data to/from global memory
  in = in + blockIdx;  out = out + blockIdx;
  globalLoads(data, in, 64); // coalesced global reads
  fftRadix16Pass(data);      // in-place radix-16 pass
  twiddleFactorMul(data, tid, 1024, 0);
  // local shuffle using local memory
  localShuffle(data, sMemx, sMemy, tid, (((tid & 15) * 65) + (tid >> 4)));
  fftRadix16Pass(data);               // in-place radix-16 pass
  twiddleFactorMul(data, tid, 64, 4); // twiddle factor multiplication
  localShuffle(data, sMemx, sMemy, tid, (((tid >> 4) * 64) + (tid & 15)));
  // four radix-4 function calls
  fftRadix4Pass(data); fftRadix4Pass(data + 4);
  fftRadix4Pass(data + 8); fftRadix4Pass(data + 12);
  // coalesced global writes
  globalStores(data, out, 64);
}

这非常类似CUDA里的kernel代码，而在上层对这些kernel代码，OpenCL也提供了统一的调度：

// create a compute context with GPU device
context = clCreateContextFromType(CL_DEVICE_TYPE_GPU);

// create a work-queue
queue = clCreateWorkQueue(context, NULL, NULL, 0);

// allocate the buffer memory objects
memobjs[0] = clCreateBuffer(context,
    CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
    sizeof(float)*2*num_entries, srcA);
memobjs[1] = clCreateBuffer(context,
    CL_MEM_READ_WRITE,
    sizeof(float)*2*num_entries, NULL);

// create the compute program
program = clCreateProgramFromSource(context,
    1, &fft1D_1024_kernel_src, NULL);

// build the compute program executable
clBuildProgramExecutable(program, false, NULL, NULL);

// create the compute kernel
kernel = clCreateKernel(program, “fft1D_1024”);

// create N-D range object with work-item dimensions
global_work_size[0] = n;
local_work_size[0] = 64;
range = clCreateNDRangeContainer(
    context, 0, 1, global_work_size, local_work_size);

// set the args values
clSetKernelArg(kernel, 0, (void *)&memobjs[0], sizeof(cl_mem), NULL);
clSetKernelArg(kernel, 1, (void *)&memobjs[1], sizeof(cl_mem), NULL);
clSetKernelArg(kernel, 2, NULL, sizeof(float)*(local_work_size[0]+1)*16, NULL);
clSetKernelArg(kernel, 3, NULL, sizeof(float)*(local_work_size[0]+1)*16, NULL);

 // execute kernel
clExecuteKernel(queue, kernel, NULL, range, NULL, 0, NULL);

这些概念和CUDA非常相似，事实上，苹果公司并不打算自己制造轮子，今年年初，CUDA就已经提供了mac操作系统的支持。

“Apple knows a lot about CUDA,” Huang said, implying the company might be ready to formally embrace Nvidia’s technology to make it easier to exploit graphics chips inside Macs. Apple’s implementation “won’t be called CUDA, but it will be called something else,” Huang said in an interview here at Nvidia’s headquarters on Wednesday.

苹果公司还表示，将在未来的操作系统Snow Leopard中集成OpenCL的支持。我们知道，并行计算的算法和普通顺序计算的算法是不一样的，在设计时，必须考虑到数据的切分，使各个计算单元可以相对独立地计算。所以，要普及并行计算，一个统一的标准是非常必要的，一来各个硬件厂家不会像无头苍蝇一样去各自研发自己的系统和规范，二来程序员们也不用针对不同的硬件编制相同的算法。而将这种平台集成进操作系统，使得我们可以方便地扩充一个系统的计算能力。想想，在未来，我们如果觉得一个机器的计算能力不够了，我们唯一所需要做的，就是去购买一个新的计算设备（类似现在的Tesla卡），将其插入我们的计算机，即可大大提高我们的计算速度。

下图是Tesla系统，可以通过增加Tesla卡来增强计算能力，但这种系统基于CUDA，所以只能使用nVidia的设备。这种专一性非常不利于并行计算的普及，OpenCL有望打破这种情形：

目前OpenCL还未正式发布，预计将在明年发布，并且已经得到不少大厂商的支持：

届时，我也需要将我的CUDA程序更新为OpenCL了：）顺便说句题外话，估计到时候，咱们的1024位RSA秘钥也会存在威胁了吧，毕竟，暴力破解所需的时间将会大大缩减，呵呵。

写了这一堆，总结下OpenCL和CUDA的不同：

OpenCL提供统一的标准，让不同厂家的设备来支持，使得我们可以统一看待计算资源。这个计算资源不但包括GPU，也包括CPU。
OpenCL更关注于多个计算资源的管理，不但在计算单元内部可以实现并行，而且对于一台宿主机上的所有计算设备的并行化进行管理。

cuda, OpenCL, 并行计算

This entry was posted on 星期六, 十一月 15th, 2008 at 16:17 and is filed under 编程与开发平台. You can follow any responses to this entry through the RSS 2.0 feed. You can leave a response, or trackback from your own site.

本文有64条评论

sofish 说: [回复]
十一月 15th, 2008 at 16:42

必得。
Soloman 说: [回复]
十一月 15th, 2008 at 16:43

@sofish 恭喜，还“必得”，哈哈
sofish 说: [回复]
十一月 15th, 2008 at 16:43

好吧，抢到沙发之后就是看内容了。但是，我刚才略看了一下。除了那些LOGO认识几个，其他的都与我没过交集。囧。
Soloman 说: [回复]
十一月 15th, 2008 at 16:45

@sofish 没关系，抢沙发和内容无关
Shawn 说: [回复]
十一月 15th, 2008 at 16:53

太深奥了。。。老所你加入松鼠会吧。。。
Soloman 说: [回复]
十一月 15th, 2008 at 17:03

@Shawn 松鼠会是啥？
Soloman 说: [回复]
十一月 15th, 2008 at 17:24

@Shawn 呵呵，搜索了下，原来是“科学松鼠会”，不过我的这篇不是科幻，是近在咫尺的技术，下一代apple操作系统就将集成OpenCL了。
Shawn 说: [回复]
十一月 15th, 2008 at 17:26

@Soloman 松鼠会不止关注前沿技术，这类即将上马的技术也是关注的重点，你可以投稿去。
leehow 说: [回复]
十一月 15th, 2008 at 17:44

终于盼到你写了，虽然我看不懂…
Soloman 说: [回复]
十一月 15th, 2008 at 17:46

@leehow 但是你的精神很可嘉：）你啥时更新啊？
iColor 说: [回复]
十一月 15th, 2008 at 21:26

文中如果没有”苹果”二字,我想我是太不下去的,,,,
很期待 Snow Leopard 又是一个小小革命…..
iColor 说: [回复]
十一月 15th, 2008 at 21:27

@Soloman Leehow更新也很慢很慢地….
Soloman 说: [回复]
十一月 15th, 2008 at 21:30

@iColor 唉，苹果粉丝都有点“病态”…..呵呵
leehow 说: [回复]
十一月 15th, 2008 at 21:41

@Soloman 你不是要我写shawn的嘛，我在酝酿啊酝酿。
Soloman 说: [回复]
十一月 15th, 2008 at 21:43

@leehow 好，好，好话题啊~~
老时说: [回复]
十一月 15th, 2008 at 21:47

说实话，没看懂。唉，落伍了，要被射灰淘汰了。
Soloman 说: [回复]
十一月 15th, 2008 at 21:50

@老时其实你博客的话题也很潮流的。
iColor 说: [回复]
十一月 15th, 2008 at 21:54

@老时没落伍,看不懂就对了,这里可能只有老所一人明白…
iColor 说: [回复]
十一月 15th, 2008 at 21:55

@leehow 我也期待…
iColor 说: [回复]
十一月 15th, 2008 at 21:57

@Soloman 病态美…-_-!!!
P.S. 你这的回复按钮的提示没改啊,,哈哈哈,,还是默认的…我的好坏也改了一下…
山上渔夫说: [回复]
十一月 15th, 2008 at 21:57

你的回复的最后几个都不显示，这是特色呀
其他人的博客没有这种现象，不知道其他朋友有没有反映这个问题：）
iColor 说: [回复]
十一月 15th, 2008 at 21:59

@山上渔夫没有这问题啊,我看着正常啊
Soloman 说: [回复]
十一月 15th, 2008 at 21:59

@山上渔夫？？怎么会，大家都好好的啊，你是什么浏览器？
Soloman 说: [回复]
十一月 15th, 2008 at 22:00

@山上渔夫我测试了IE7, FF都没问题啊，手边暂时没有IE6
山上渔夫说: [回复]
十一月 15th, 2008 at 22:23

问一个很傻冒的问题，怎么看是ie几？
我用ff测试了一下也没问题
不过用ie和傲游都不行
可能就是ie6吧
Soloman 说: [回复]
十一月 15th, 2008 at 22:29

@山上渔夫我用ie6测试了下，确实，不过不是一直没有，是拉到某个程度就没了，往上拉点又有了，晕。网页方面俺不是很懂啊，这ie6的bug还不知道怎么弄哦。不光回复，我右上角的美女RSS图ie6里也显示不了，只要是用了float的好像都有问题。看ie几，看“帮助”-〉关于吧，不过遨游的内核也是IE的。有个js代码可以sniffer出浏览器，http://www.webreference.com/tools/browser/javascript.html ,不过要自己写个简单的页面，然后用浏览器打开就可以清楚地看到浏览器类型和浏览器内核类型。
Soloman 说: [回复]
十一月 15th, 2008 at 22:46

@山上渔夫

我刚才找了个在线浏览器检测的地址：

http://www.webreference.com/tools/browser/_browserhawk.html
山上渔夫说: [回复]
十一月 15th, 2008 at 22:47

呵呵，水落石出啦
睡觉啦
晚安
Jor 说: [回复]
十一月 15th, 2008 at 23:05

……你是做什么的？
basil 说: [回复]
十一月 15th, 2008 at 23:05

@老所，
你可以用ietester，来看看不同版本的显示，
另外可以看一下CSS mastery这本书，里面关于float有很好的解释。
Soloman 说: [回复]
十一月 15th, 2008 at 23:06

@Jor 我目前就是做这个并行计算的：）
Soloman 说: [回复]
十一月 15th, 2008 at 23:07

@basil 谢谢，改天去当当买本这书。不过得有时间啊：）
Jor 说: [回复]
十一月 15th, 2008 at 23:22

@Soloman 看成了OpenGL…
Soloman 说: [回复]
十一月 15th, 2008 at 23:25

@Jor 不过确实和OpenGL有关的，都是同一个组织维护的业界标准。
Soloman 说: [回复]
十一月 15th, 2008 at 23:26

@Jor OpenGL主要关注与图形计算，还有个OpenAL关注与音频计算，OpenCL则是通用计算，用GPU来做通用数学计算。
herb 说: [回复]
十一月 16th, 2008 at 15:13

GPU比CPU的运算能力强是很久以前的事情了…..软件方式和硬件方式就是这个区别。
Soloman 说: [回复]
十一月 16th, 2008 at 15:37

@herb 各有所长, CPU是通用处理器,GPU则是专用浮点计算处理器,而且支持并行处理.而很多好的GPU如果只用来做图形计算,有点浪费了,OpenCL就是要统一看待CPU和GPU,当然,只是在数学计算方面.
Soloman 说: [回复]
十一月 16th, 2008 at 15:39

@herb 现在已经不再是光提高一块芯片的计算能力的年代了,而是要利用多个芯片来提高能力的时代了,未来的几年里,并行计算的平台和应用将会大大增加.
OpenHero 说: [回复]
十一月 16th, 2008 at 21:44

@herb OpenCL值得关注，看过OpenCL的内部文档，和Nvidia的CUDA有多地方都差不多，其实CUDA也可以把kernel放到不同的平台上运行，在Nvidia的内部是有这样的platform的，只是还没开放出来。并行计算在接下来的几年，会有很大的发展，就像C++的新标准也会支持并行一样，值得关注：）
Soloman 说: [回复]
十一月 16th, 2008 at 21:53

@OpenHero 哈哈,你来啦,欢迎欢迎,你是高手哦!!!
Soloman 说: [回复]
十一月 16th, 2008 at 21:54

@OpenHero 我最近做的项目就是类似的工作,不过现在OpenCL要出来了,我预计明年要把代码转像OpenCL
OpenHero 说: [回复]
十一月 16th, 2008 at 22:06

@Soloman 哈哈，礼尚往来，多交流才会有收获:) 其实如果不是代码像SSE这样太被core给限制住，按照spmd的思路来写代码，在未来的平台上，也是会可以让程序很快的移植过去的，就像intel的tbb的库，很方便的写的多线程的程序，可以很快的移植到CUDA这样的平台上来，应该写程序时候的思路是按照spmd的思路走，如果扣得太深，按照simd来写一个小部分，就要把这个部分封装起来。我测试过用intel的tbb的库写的多线程的程序，拿到cuda上，也可以很快的移植过去，而且效果还不错，不过这个必须是要对两边的平台都要深入的了解。其实CUDA这样的平台也是可以被MPI这样的lib库支持的，不过现在还没有mpi的库支持GPU，我敢大胆的预测，MPI的接下来的版本，很快就可以gpu上的编程，呵呵，也可能会对OpenCL有冲击，OpenCL现在看到的对手只是CUDA，其实隐藏在背后的对手其实还有几个啦，呵呵，有了竞争才会有发展：）
Soloman 说: [回复]
十一月 16th, 2008 at 22:12

@OpenHero MPI我还没用过，MPI支持GPU么？我想了下，现在这些处理器，好像也就CPU和GPU可以统一起来做计算，NP这些，应该是不适合做计算的。
Soloman 说: [回复]
十一月 16th, 2008 at 22:13

@OpenHero 我觉得光靠移植肯定不是个办法，还是应该统一，这个是趋势。
OpenHero 说: [回复]
十一月 16th, 2008 at 22:21

@Soloman 恩，其实自己也可以做一个类似的编译平台，让一个kernel两边都支持，但是有的时候，得设计到硬件的相关的部分，还得自己来优化，例如cpu的SSE这部分，gpu是没有的，CPU的L1这些缓存的预取的方式，gpu也是没有的；NV的gpu的架构的shared的mem cpu上也没有；虽然两者可以同一起来，但是有些时候，编译器就得做很多工作，这样比较智能的编译器可能还需要一段路要走~
Soloman 说: [回复]
十一月 16th, 2008 at 22:26

@OpenHero 要统一，就得牺牲一些细节。细节却是对提升效率有帮助，但是从长远看，统一后，应用的开发会加速，这个利应该是大于弊的。而且，有了标准，今后的硬件在设计时也会尽量注意这些的。

我目前做的项目，是基于monte carlo method的算法，为了验证这个结果，所以必须写两套代码，一套用cpu执行，另一套用gpu执行，造成开发效率的下降。所以，我盼望统一，即使要放弃一些硬件的特性。
OpenHero 说: [回复]
十一月 16th, 2008 at 22:33

@Soloman 这个还有点难，不过统一确实有好处，对开发实际项目来说，确实很有效率的提升。蒙特卡洛的这样的需要在cpu和gpu上的都执行代码，可以直接用CUDA的kernel的代码编译到cpu上，呵呵，不过这个只有NV内部有这样的转换编译工具，还没公开……不知道会不会公开
OpenHero 说: [回复]
十一月 16th, 2008 at 22:34

@Soloman 那就得看谁牺牲谁的功能，呵呵，或者以后多一些各个平台支持库，把内部的实现的部分隐藏，透明起来
Soloman 说: [回复]
十一月 16th, 2008 at 22:35

@OpenHero 你说这个工具应该是MCUDA，没怎么发布，也不知道好不好用。不过我们不是为了某一个具体算法而写代码，而是要搭一个框架来实现一些统计学计算，所以好的平台是有用处的。所以我们在上层是用的C++来包装的，可以方便地替换算法中的一些对象或者模型。
OpenHero 说: [回复]
十一月 16th, 2008 at 22:38

@Soloman 恩，做实际项目，就得考虑这些，呵呵，过去在工作的时候，也会这样的构架系统；不过现在又读书了，得在一个上面做到极限，所以就得很扣底层，有点累~~呵呵
Soloman 说: [回复]
十一月 16th, 2008 at 22:39

@OpenHero 好，很高兴在这见到你，以后常互相串门哈：）你好像是CSDN里CUDA方面的高手，不过CSDN这个东西真是的，网站做的真烂，太不美观不说，用Firefox好多功能都有问题，Feed也不给全文，唉，我真是晕死了。以前在CSDN写了几篇，后来就懒得去那写了，不过你不一样，你在CSDN有25万流量了：）
OpenHero 说: [回复]
十一月 16th, 2008 at 22:50

@Soloman 呵呵，好啊，加我msn吧~~QQ满了~~陌生人都有200多了~= 我的msn可以在我的csdn的资料里面看到，这里就不写了，呵呵~我也想自己搭建一个blog，不过一直没时间，然后csdn的访问已经有20多w索性就在哪里写了，呵呵，以后有时间了，自己搭建一个blog~~呵呵，多交流
OpenHero 说: [回复]
十一月 16th, 2008 at 22:51

多了，可以关注一下语言本身支持并行的消息，这方面 erlang做得不错，然后c++的新标准，也会在新标准里面支持并行，值得期待
i.robot 说: [回复]
十一月 16th, 2008 at 23:45

看到一半就不行了，头晕。。。
Soloman 说: [回复]
十一月 16th, 2008 at 23:52

@i.robot 呵呵，晚安，早点休息。
醉倚西风说: [回复]
十一月 17th, 2008 at 15:23

老所加你连接了
Soloman 说: [回复]
十一月 17th, 2008 at 15:26

@醉倚西风哈哈，我在你博客里写的是让你和“老时”友情链接，你看成我了吧，哈哈，因为我觉得你现在开始倾向老时的style了。不过加我也好，但是我目前还没有地方显示友情链接，我还在考虑在那显示，等做好了我再加你，谢谢哈。
小墨说: [回复]
十一月 19th, 2008 at 10:43

我还以为写的OpenGL，一看不是
感觉应用到普通生活还需时日吧
NetPuter 说: [回复]
十一月 19th, 2008 at 13:07

对硬件..哎..
算了..跳走…
Soloman 说: [回复]
十一月 20th, 2008 at 02:52

@小墨计算机发展的速度，有的时候是超出我们的想象的
Soloman 说: [回复]
十一月 20th, 2008 at 02:53

@NetPuter 哪里走！：）
aa 说: [回复]
十二月 2nd, 2008 at 03:22

哈哈~~
日后可以在家里搞核爆运算咯
好戏说: [回复]
一月 29th, 2009 at 14:12

请问这个opencl与cuda影响我现在买显卡吗？
Soloman 说: [回复]
三月 4th, 2009 at 02:49

@好戏现在cuda 只在nvidia卡中有，而且是要8xxx以上的。OpenCL 还没成熟，不是很清楚。