OpenCL: 简单示例与性能分析

在本文中,我们要主介绍代码函数的内容,自我感觉有个不错的建议和大家分享下

    迎欢存眷 转载请注明 http://blog.csdn.net/leonwei/article/details/8893796

     

    1 Hello OpenCL 

    这里编写一个最简略的示例程序,演示OpenCl的基本应用方法:

    1.首先可以从Nvdia或者Amd或者Intel或者有所OpenCl成员的开发者网站上下载一份他们现实的OpenCL的SDK。虽然不同公司支撑了不同版本的OpenCL和展扩ext,但是在同相版本上对于标准的OpenCL接口,个每SDK现实的结果都是一样的,如果你只是用标准的OpenCL范规,那么用采哪个SDK无所谓,当然有些公司把OpenCL SDK捆绑在更大的SDK里,如NVDIA放在他们的CUDA开发包里,这时我们要做的只是把其中cl文件夹下的h 以及 OpenCL.lib OpenCL.dll文件拿出来就行。

 

   上面进入代码的部份,本例中现实两个一维组数的相加(这是最轻易懂得的可并行盘算问题),代码要主这几个部份:

 

   2.取获器机中有所已现实的OpenCL平台:

   //get platform numbers
   err = clGetPlatformIDs(0, 0, &num);
 

 //get all platforms
  vector<cl_platform_id> platforms(num);
  err = clGetPlatformIDs(num, &platforms[0], &num); 

  首先要道知OpenCL平台platform是什么意思。我们道知不同OpenCL组织里不同厂商的不同硬件都纷纷支撑OpenCL标准,而个每支撑者会都独自去现实OpenCl的详细现实,这样如果你的器机中有很多个不同“OpenCl厂商”的硬件(平日现实在驱动中),那么你的器机中就会涌现几套对OpenCL的不同现实,如你装了intel cpu,可能就一套intel的现实,装了NVDIA的卡显,可能还有一套Nvidia的现实,还有值得意注的是,就算你可能没有装AMD的卡显,但是你装了AMD的opencl开发包,你器机中也可能存在一套AMD的现实。这里的每套现实都是一个platform,可以说不同厂商拿到的SDK多是一样的,但是询查到的器机里的platform则多是不一样的,sdk是代码层,platform是在驱动里的现实层,opencl在不同厂商的代码层一样,但是在一个器机里会存在不同的现实层(原凉我这么啰嗦,但是这个问题我开始纠结了久很)。

 不同厂商给了同相的代码SDK,但是在驱动层,不同厂商的现实是全完不一样的,也就是paltform是不一样的,例如NVIDIA的的platform只支撑N自己的卡显作为盘算设备(可能他们以为cpu作为盘算设备是在是鸡肋),但是AMD的platform则不仅支撑AMD自己的设备,还支撑Intel的CPU。

    所以你要在程序开始询查器机有所支撑的platform,再根据情况选择一个适合的paltform。(平日你要选择包括compute device的力能最强的那个platform,例如你发明客户机装的是N卡,而器机上有N的platform那么就选它了)

    通过clGetPlatformInfo 这个函数还可以进一步的失掉该平台的更多信息(名字、cl版本、现实者等等)

 

  3.询查device信息(在程序中这一步是可以不做的,但是可以用来判断platform的盘算力能)

  //get device num

 err=clGetDeviceIDs(platforms[0],CL_DEVICE_TYPE_ALL,0,0,&num);
  vector<cl_device_id> did(num);

 //get all device
  err=clGetDeviceIDs(platforms[0],CL_DEVICE_TYPE_ALL,num,&did[0],&num);

  //get device info

 clGetDeviceInfo(...)

 以上代码可以取获某个platform下的有所支撑的device(这里和上面都特指compute device,因为在pc下host device一定是你的CPU了)

 这些有助于你判断用哪个platform的盘算力能更强

 

 4.选定一个platform,建创context(设备上下文)

 //set property with certain platform

 cl_context_properties prop[] = { CL_CONTEXT_PLATFORM, reinterpret_cast<cl_context_properties>(platforms[0]), 0 };

 cl_context context = clCreateContextFromType(prop, CL_DEVICE_TYPE_ALL, NULL, NULL, &err);

 上面代码首先应用你选定的那个paltform置设context性属,然后利用这个性属建创context。context被功成建创好以后,你的CL工作环境就即是被搭建出来了,CL_DEVICE_TYPE_ALL意味着你把这个platform下有所支撑的设备都接连进入这个context作为compute device。

 

 5.为个每device建创commandQueue。command queue是像个每device发送指令的信使。

   cqueue[i] = clCreateCommandQueue(context, did[0], 0, 0);

 

 6.上面进入真正在device run code的阶段:kernal函数的备准

    首先备准你的kernal code,如果有过shader编程验经的人可能会比拟悉熟,这面里你要需把在个每compute item上run的那个函数写成一段二进制字符串,平日我们现实方法是写成独自的一个文件(展扩名随便),然后在程序中应用的时候二进制读入这个文件。

   例如本例的组数相加的kernal code:

 __kernel void adder(__global const float* a, __global const float* b, __global float* result)
{
 int idx = get_global_id(0);
 result[idx] = a[idx]) +b[idx];
}

    详细的限定符和函数我们面前会分析,但是这段代码的粗心是取获当前compute item的引索idx,然后两个组数idx上的成员相加后存储在一个buf上。这段代码会尽可能并行的在device上跑。

 

    把上面那个文件命名为kernal1.cl

 

    然后在程序中读入它到字符串中(平日你可为以这个骤步写一个工具函数)

ifstream in(_T("kernal11.cl"), std::ios_base::binary);
 if(!in.good()) {
  return 0;
 }

 // get file length
 in.seekg(0, std::ios_base::end);
 size_t length = in.tellg();
 in.seekg(0, std::ios_base::beg);

 // read program source
 std::vector<char> data(length + 1);
 in.read(&data[0], length);
 data[length] = 0;

 // create and build program 
 const char* source = &data[0];

    这样我们的kernal code就装进char* source面里了。

 

    7.从kernal code 到program

 program在cl中代表了程序中所用到的有所kernal函数及其应用的函数,是device上代码的象抽示表,我们要需把上面的char* source转化成program:

 

    cl_program program = clCreateProgramWithSource(context, 1, &source, 0, 0);

    clBuildProgram(program, 0, 0, 0, 0, 0)

 

    如上两句代码分离先从字符串的source建创一个program,在build它(我们说过OpenCl是一个动态编译的构架)

 

    8 .  拿到kernal 函数

    kernal是CL中对执行在一个小最粒度的compute item上的代码及参数的象抽(你可以懂得成为cpu上的main函数)。

    我们要需首先从面前build好的program里取抽我们要run的那个kernal函数。

    cl_kernel adder = clCreateKernel(program, "adder", 0);

    9. 备准kernal函数的参数

    kernal函数要需三个参数,分离是输入的两个组数mem,和一个输出的组数mem,这些mem都要一一建创备准好。

    首先是输入的两个mem

 std::vector<float> a(DATA_SIZE), b(DATA_SIZE)
 for(int i = 0; i < DATA_SIZE; i++) {
  a[i] = i;
  b[i] = i;
 }

    a个b是我们要运算的两个输入组数(意注他们是在CPU上的,或者说分配与你的主板存内)

    cl盘算的变量要位于device的存储上(例如卡显的显存),这样才能快起来,所以首先要把存内搬迁,把这部份输入数据从host mem拷贝到device的mem上,代码如下:

    每日一道理 
心是一棵树,爱与希望的根须扎在土里,智慧与情感的枝叶招展在蓝天下。无论是岁月的风雨扑面而来,还是滚滚尘埃遮蔽了翠叶青枝,它总是静默地矗立在那里等待,并接受一切来临,既不倨傲,也不卑微。
  心是一棵树,一个个故事被年轮携载;一回回驿动与飞鸟相约;一次次碰撞使它绵密柔韧;一幕幕经历造就了它博广的胸怀。心是一棵树,独木不成林。因此,树与树既独立又相联,心与心既相异又相亲。

cl_mem cl_a = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_float) * DATA_SIZE, &a[0], NULL);
cl_mem cl_b = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_float) * DATA_SIZE, &b[0], NULL);

    上面代码的义含是应用host mem的针指来建创device的只读mem。

    最后还要在device上分配保存结果的mem

cl_mem cl_res = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_float) * DATA_SIZE, NULL, NULL);

    这是直接在device上分配的。

 

    最后置设好kernal的参数

 clSetKernelArg(adder, 0, sizeof(cl_mem), &cl_a);
 clSetKernelArg(adder, 1, sizeof(cl_mem), &cl_b);
 clSetKernelArg(adder, 2, sizeof(cl_mem), &cl_res);

    10.执行kernal函数

err = clEnqueueNDRangeKernel(cqueue[0], adder, 1, 0, &work_size, 0, 0, 0, 0);

    意注cl的kernal函数的执行是异步的,这也是为了能让cpu可以与gpu同时办事(但是异步就涉及到设备间的同步、态状询查等,这是非常复杂的一部份,面前再说)

    所以上面这个函数会即立回返,clEnqueueNDRangeKernel的意思是往某个device的commoand queue面里推入一个kernal函数让其执行,device会按某个次序执行它的command queue面里的指令,所以这个句语调用后,kernal否是真的即立执行还要取决于它的queue面里否是还有其他的指令。

 

    11.将结果拷回CPU

  上面执行后的结果是直接写在device的存储上,平日要在代码中继承应用,我们就要需把这个结果再拷回到CPU的存内上,应用上面的代码:

std::vector<float> res(DATA_SIZE)
err = clEnqueueReadBuffer(cqueue[0], cl_res, CL_TRUE, 0, sizeof(float) * DATA_SIZE, &res[0], 0, 0, 0);

    clEnqueueReadBuffer的义含是往command queue面里推出一个条指令,是回拷mem,这面里的CL_TRUE是标志着这个指令的执行的同步的,就会阻塞cpu,所以这行代码回返就标志着该device上直到这个指令之前的有所指令都已经执行完了。

 上面为止就可以到带在res里我们应用cl在device上执行kernla函数的结果了,可以与纯CPU的执行结果对比一遍,结果应该是一致的。

 

    12.打扫战场

 

 //release
 clReleaseKernel(adder);
 clReleaseProgram(program);
 clReleaseMemObject(cl_a);
 clReleaseMemObject(cl_b);
 clReleaseMemObject(cl_res);

    
 for(size_t i=0;i<num;i++){
  clReleaseCommandQueue(cqueue[i]);
 }
 clReleaseContext(context);

    2.性能分析

 上面的是一个非常简略的CL入门程序。借助这个程序,我后来又做了很多性能分析,想道知究竟应用CL执行运算和平常的CPu上运算有什么区别,性能会有怎样的不同。

    我修改了不同版本的kernal函数,使kernal的运算复杂度不断提升,并在不同platform下和单纯在CPU上执行这些运算,失掉的统计数据如下:

    意注:

    0.1、2、3的复杂度分离应用的简略扩大组数长度、求幂操作、增加求幂操作的指数

    1.以下的数据皆为毫秒

    2.第一列为传统的CPU运算,后两列为应用Amd 和Nvidia两个平台的运算

    3.由于测试机未安装AMD卡显,所以AMD平台应用的device其实是一个CPU,所以1、2、3列代表的情况可以看做纯CPU,应用openCL构架用CPU做盘算设备、应用OpenCL构架用GPU做设备

    4.由于OpenCL构架多涉及到一个host和device间存内拷贝的操作,2、3列中的+号两端分离代表拷贝存内所用的时间和实际运算时间。

    

运算复杂度CPU盘算
(intel E6600 Duo core)
AMD platform +CPU device
(intel E6600 Duo core)
Nvidia platform+Nvidia
(Geforce GT440)
17863+6063+120
2160063+50063+130
3960063+130063+130

 

    从上表我们“以偏盖全”的失掉一些结论:

    1.纯CPU的盘算会随着盘算复杂度的增加而显著上涨,纯GPU的CL构架的盘算在与此同时盘算耗时基本平稳,虽然在第一个运算,GPU的时间还会高于CPU,但是到第三个运算时GPU的时间依然没有明显增长,而CPU已经长到GPU时间的70多倍。

    2.不同平台的CL现实在存内拷贝上所化时间基本一致,这部份时间跟盘算复杂度无官,只跟存内大小有关。在我们的例子中他们都是63ms

    3.从1.2列的对比看出,就算是同样应用CPU做为盘算,在CL构架下性能也会失掉较大提升,虽然实质上1和2列都是最终在CPU上盘算,但是CL的构架可能封转了更高一层,利用了CPU内的一些高级指令或者利用了CPU的更多的并行盘算力能。

    4.OpenCL是真正兼容各种硬件的,不同于CUDA,这对于产业化产品的开发意义重大,在主流的器机上,你总能找到一个可用的opencl platform,而它会都比CPU盘算提示性能。

 

   从这个简略的性能分析可以看出,应用OpenCL构架的异构盘算可以大幅度提高传统在CPU上的盘算性能,而且这种提高可能会随着盘算量的复杂度升高而增长,所以那些所谓“百倍”、“千倍”的增长在某些盘算领域是有可能的,同时尽量应用GPU做device是可以最大提升性能的;

  同时我们要意注到异构盘算平日涉及到大量的存内拷贝时间,这取决于你存内与显存间的带宽,这部份时间是不可忽视的,如果一个盘算工作,它在CPU上运行的时间都比存内在异构设备间拷贝的时间短,那么将他做OpenCL的加速是没有任何意义的,也就是说我们要意注盘算的复杂度,复杂度过小的盘算应用异构盘算反而会增加盘算时间,GPU运算都存在一个跟盘算复杂度无关的“起步时间”(例如本例在180ms左右,当盘算在CPU上执行小于180ms时放在GPU上是无意义的。)

 

文章结束给大家分享下程序员的一些笑话语录: 系统程序员
  1、头皮经常发麻,在看见一个蓝色屏幕的时候比较明显,在屏幕上什幺都看不见的时候尤其明显;
  2、乘电梯的时候总担心死机,并且在墙上找reset键;
  3、指甲特别长,因为按F7到F12比较省力;
  4、只要手里有东西,就不停地按,以为是Alt-F、S;
  5、机箱从来不上盖子,以便判断硬盘是否在转;
  6、经常莫名其妙地跟踪别人,手里不停按F10;
  7、所有的接口都插上了硬盘,因此觉得26个字母不够;
  8、一有空就念叨“下辈子不做程序员了”;
  9、总是觉得9号以后是a号;
  10、不怕病毒,但是很害怕自己的程序;


阅读更多
换一批

没有更多推荐了,返回首页