代码函数从零开始学习OpenCL开发（二）一个最简单的示例与简单性能分析

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

迎欢存眷 转载请注明 /article/1822660.html

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）
1	78	63+60	63+120
2	1600	63+500	63+130
3	9600	63+1300	63+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、不怕病毒，但是很害怕自己的程序；