从零开始学习OpenCL开发(二)一个最简单的示例与简单性能分析
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上是无意义的。)