OpenCL 从零开始学习OpenCL开发

分享一下我老师大神的人工智能教程吧。零基础!通俗易懂!风趣幽默!还带黄段子!希望你也加入到我们人工智能的队伍中来!http://www.captainbed.net

 多谢大家关注 转载本文请注明:http://blog.csdn.net/leonwei/article/details/8880012

 

本文将作为我《从零开始做OpenCL开发》系列文章的第一篇。

 

1 异构计算、GPGPU与OpenCL

  OpenCL是当前一个通用的由很多公司和组织共同发起的多CPU\GPU\其他芯片 异构计算(heterogeneous)的标准,它是跨平台的。旨在充分利用GPU强大的并行计算能力以及与CPU的协同工作,更高效的利用硬件高效的完成大规模的(尤其是并行度高的)计算。在过去利用GPU对图像渲染进行加速的技术非常成熟,但是我们知道GPU的芯片结构擅长大规模的并行计算(PC级的GPU可能就是CPU的上万倍),CPU则擅长逻辑控制,因此不只局限与图像渲染,人们希望将这种计算能力扩展到更多领域,所以这也被称为GPGPU(即通用处计算处理的GPU)。

    简单的说,我们的CPU并不适合计算,它是多指令单数据流(MISD)的体系结构,更加擅长的是做逻辑控制,而数据处理基本是单流水线的,所以我们的代码for(i=0;...;i++)这种在CPU上要重复迭代的跑很多遍,但是你的显卡GPU则不是这样,GPU是典型的单指令多数据(SIMD)的体系结构,它不擅长逻辑控制,但是确实天生的向量计算机器,对于for(i=0;...;i++)这样的代码有时只需要跑一遍,所以图形世界中那么多的顶点、片段才能快速的并行在显卡中渲染处理

OpenCL  从零开始学习OpenCL开发

 

GPU的晶体管可以到几十亿个,而CPU通常只有几个亿,

OpenCL  从零开始学习OpenCL开发

如上图是NVidia Femi100的结构,它有着大量的并行计算单元。

所以人们就想如何将更多的计算代码搬到GPU上,让他不知做rendering,而CPU只负责逻辑控制,这种一个CPU(控制单元)+几个GPU(有时可能再加几个CPU)(计算单元)的架构就是所谓的异构编程(heterogeneous),在这里面的GPU就是GPGPU。异构编程的前景和效率是非常振奋人心的,在很多领域,尤其是高并行度的计算中,效率提升的数量级不是几倍,而是百倍千倍。

   其实NVIDIA在很早就退出了利用其显卡的GPGPU计算 CUDA架构,当时的影响是很大的,将很多计算工作(科学计算、图像渲染、游戏)的问题提高了几个数量级的效率,记得那时NVIDIA来浙大介绍CUDA,演示了实时的ray tracing、大量刚体的互相碰撞等例子,还是激动了一下的,CUDA现在好像已经发展到了5.0,而且是NVDIA主力推的通用计算架构,但是CUDA最大的局限就是它只能使用N家自己的显卡,对于广大的A卡用户鞭长莫及。OpenCL则在之后应运而生,它由极大主流芯片商、操作系统、软件开发者、学术机构、中间件提供者等公司联合发起,它最初由Apple提出发起标准,随后Khronos Group成立工作组,协调这些公司共同维护这套通用的计算语言。Khronos Group听起来比较熟悉吧,图像绘制领域著名的软硬件接口API规范著名的OpenGL也是这个组织维护的,其实他们还维护了很多多媒体领域的规范,可能也是类似于Open***起名的(所以刚听到OpenCL的时候就在想它与OpenGl有啥关系),OpenCl没有一个特定的SDK,Khronos Group只是指定标准(你可以理解为他们定义头文件),而具体的implementation则是由不同参与公司来做,这样你会发现NVDIA将OpenCL做了实现后即成到它的CUDA SDK中,而AMD则将其实现后放在所谓是AMD APP (Accelerated Paral Processing)SDK中,而Intel也做了实现,所以目前的主流CPU和GPU都支持OpenCL架构,虽然不同公司做了不同的SDK,但是他们都遵照同样的OpenCL规范,也就是说原则上如果你用标准OpenCl头中定义的那些接口的话,使用NVIDIA的SDK编的程序可以跑在A家的显卡上的。但是不同的SDK会有针对他们芯片的特定扩展,这点类似于标砖OpenGL库和GL库扩展的关系。

  OpenGL的出现使得AMD在GPGPU领域终于迎头赶上的NVIDIA,但是NVIDIA虽为OpenCL的一员,但是他们似乎更加看重自己的独门武器CUDA,所以N家对OpenCL实现的扩展也要比AMD少,AMD由于同时做CPU和GPU,还有他们的APU,似乎对OpenCL更来劲一些。

 

2.关于在GPU上写代码的那些事儿

  OpenCL也是通过在GPU上写代码来加速,只不过他把CPU、GPU、其他什么芯片给统一封装了起来,更高了一层,对开发者也更友好。说到这里突然很想赘述一些在GPU上写代码的那些历史。。

  其实最开始显卡是不存在的,最早的图形处理是放在CPU上,后来发现可以再主板上放一个单独的芯片来加速图形绘制,那时还叫图像处理单元,直到NVIDIA把这东西做强做大,并且第一给它改了个NB的称呼,叫做GPU,也叫图像处理器,后来GPU就以比CPU高几倍的速度增长性能。

   开始的时候GPU不能编程,也叫固定管线的,就是把数据按照固定的通路走完

   和CPU同样作为计算处理器,顺理成章就出来了可编程的GPU,但是那时候想在GPU上编程可不是容易的事,你只能使用GPU汇编来写GPU程序,GPU汇编?听起来就是很高级的玩意儿,所以那时使用GPU绘制很多特殊效果的技能只掌握在少数图形工程师身上,这种方式叫可编程管线。

    很快这种桎桍被打破,GPU上的高级编程语言诞生,在当时更先进的一些显卡上(记忆中应该是3代显卡开始吧),像C一样的高级语言可以使程序员更加容易的往GPU写代码,这些语言代表有nvidia和微软一起创作的CG,微软的HLSL,openGl的GLSL等等,现在它们也通常被称为高级着色语言(Shading Language),这些shader目前已经被广泛应用于我们的各种游戏中。

   在使用shading language的过程中,一些科研人员发现很多非图形计算的问题(如数学、物理领域的并行计算)可以伪装成图形问题利用Shading Language实现在GPU上计算,而这结果是在CPU上跑速度的N倍,人们又有了新的想法,想着利用GPU这种性能去解决所有大量并行计算的问题(不只图形领域),这也叫做通用处理的GPU(GPGPU),很多人尝试这样做了,一段时间很多论文在写怎样怎样利用GPU算了哪个东东。。。但是这种工作都是伪装成图形处理的形式做的,还没有一种天然的语言来让我们在GPU上做通用计算。这时又是NVIDIA带来了革新,09年前后推出的GUDA架构,可以让开发者在他们的显卡上用高级语言编写通用计算程序,一时CUDA热了起来,直到现在N卡都印着大大的CUDA logo,不过它的局限就是硬件的限制。

  OpenCL则突破了硬件的壁垒,试图在所有支持的硬件上搭建起通用计算的协同平台,不管你是cpu还是gpu通通一视同仁,都能进行计算,可以说OpenCL的意义在于模糊了主板上那两种重要处理器的界限,并使在GPU上跑代码变得更容易。

 

  3 OpenCL架构

  3.1 硬件层:

  上面说的都是关于通用计算以及OpenCL是什么,下面就提纲挈领的把OpenCL的架构总结一下:

  以下是OpenCL硬件层的抽象

  OpenCL  从零开始学习OpenCL开发

 

 它是一个Host(控制处理单元,通常由一个CPU担任)和一堆Computer Device(计算处理单元,通常由一些GPU、CPU其他支持的芯片担任),其中Compute Device切分成很多Processing Element(这是独立参与单数据计算的最小单元,这个不同硬件实现都不一样,如GPU可能就是其中一个Processor,而CPU可能是一个Core,我猜的。。因为这个实现对开发者是隐藏的),其中很多个Processing Element可以组成组为一个Computer Unit,一个Unit内的element之间可以方便的共享memory,也只有一个Unit内的element可以实现同步等操作。

3.2 内存架构

OpenCL  从零开始学习OpenCL开发

其中Host有自己的内存,而在compute Device上则比较复杂,首先有个常量内存,是所有人能用的,通常也是访问最快的但是最稀少的,然后每个element有自己的memory,这是private的,一个组内的element有他们共用的一个local memery。仔细分析,这是一个高效优雅的内存组织方式。数据可以沿着Host-》gloabal-》local-》private的通道流动(这其中可能跨越了很多个硬件)。

3.3软件层面的组成

 这些在SDK中都有对应的数据类型

 setup相关:

  Device:对应一个硬件(标准中特别说明多core的CPU是一个整个Device)

 

  Context:环境上下文,一个Context包含几个device(单个Cpu或GPU),一个Context就是这些device的一个联系纽带,只有在一个Context上的那些Device才能彼此交流工作,你的机器上可以同时存在很多Context。你可以用一个CPu创建context,也可以用一个CPU和一个GPU创建一个。

 

Command queue:这是个给每个Device提交的指令序列

 

内存相关:

Buffers:这个好理解,一块内存

Images:毕竟并行计算大多数的应用前景在图形图像上,所以原生带有几个类型,表示各种维度的图像。

 

gpu代码执行相关:

 Program:这是所有代码的集合,可能包含Kernel是和其他库,OpenCl是一个动态编译的语言,代码编译后生成一个中间文件(可实现为虚拟机代码或者汇编代码,看不同实现),在使用时连接进入程序读入处理器。

 Kernel:这是在element跑的核函数及其参数组和,如果把计算设备看做好多人同时为你做一个事情,那么Kernel就是他们每个人做的那个事情,这个事情每个人都是同样的做,但是参数可能是不同的,这就是所谓的单指令多数据体系。

 WorkI tem:这就是代表硬件上的一个Processing Element,最基本的计算单元。

 

同步相关:

Events:在这样一个分布式计算的环境中,不同单元之间的同步是一个大问题,event是用来同步的

 

他们的关系如下图

OpenCL  从零开始学习OpenCL开发

 

   上面就是OpenCL的入门介绍,其实说实话在10年左右就跟踪过GPGPU相关的东西,那时很多相关技术还存在于实验室,后来的CUDA出现后,也激动过,学习过一阵,不过CUDA过度依赖于特定硬件,产业应用前景并不好,只能做做工程试验,你总不能让用户装个游戏的同时,让他顺便换个高配的N卡吧。所以一度也对这个领域不太感兴趣,最近看到OpenCL的出现,发现可能这个架构还是有很好的应用前景的,也是众多厂商目前合力力推的一个东西。想想一下一个迭代10000次的for循环一遍过,还是很激动的一件事。

  在游戏领域,OpenCL已经有了很多成功的实践,好像EA的F1就已经应用了OpenCL,还有一些做海洋的lib应用OpenCL(海面水波的FFT运算在过去是非常慢的),另外还有的库干脆利用OpenCL去直接修改现有的C代码,加速for循环等,甚至还有OpenCl版本的C++ STL,叫thrust,所以我觉得OpenCL可能会真正的给我们带来些什么~

 

以下是一些关于OpenCL比较重要的资源:

http://www.khronos.org/opencl/ 组织的主页

https://developer.nvidia.com/opencl N家的主页

http://developer.amd.com/resources/heterogeneous-computing/opencl-zone/ A家的主页

http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/ 标准的reference

http://developer.amd.com/wordpress/media/2012/10/opencl-1.2.pdf 必看 最新的1.2版本标准

http://www.khronos.org/assets/uploads/developers/library/overview/opencl-overview.pdf 必看,入门的review

http://www.kimicat.com/opencl-1/opencl-jiao-xue-yi  一个教学网站

-------------------------------------------------------------------------------------------------------------------------------------

 

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上是无意义的。)


--------------------------------------------------------------------------------------------------------------

 

这里将更深入的说明一些OpenCL API的功能

1. 创建buffer

涉及到内存与显存的操作总是复杂麻烦的,这个函数也一样。。。

 

cl_memclCreateBuffer ( cl_context context,
  cl_mem_flags flags,
  size_t size,
  void *host_ptr,
  cl_int *errcode_ret)

 

 

函数将创建(或分配)一片buffer,并返回。这里创建的mem可以是globla也可以是local或private,具体要看kernal中怎样声明限定符。cl会根据执行情况自动管理global到更进一层如private的copy。这里的buffer概念是用于kernal函数计算的(或者说是用于device访问的,什么是device?host是C++写的那段控制程序,一定运行在CPU,device就是执行kernal计算的,运行在所有有计算能力的处理器上,有时你的CPU同时扮演host与device,有时用GPU做device),这里模糊了host与device的内存,也就是说根据flag的不同,可以是在host上的,也可以是在device上的,反正只有这里分配的内存可以用于kernal函数的执行。

 

主要的参数在 flags,这些参数可以|

1  CL_MEM_READ_WRITE:在device上开辟一段kernal可读可写的内存,这是默认

2  CL_MEM_WRITE_ONLY:在device上开辟一段kernal只可以写的内存

3  CL_MEM_READ_ONLY:在device上开辟一段kernal只可以读的内存

 

4  CL_MEM_USE_HOST_PTR:直接使用host上一段已经分配的mem供device使用,注意:这里虽然是用了host上已经存在的内存,但是这个内存的值不一定会和经过kernal函数计算后的实际的值,即使用clEnqueueReadBuffer函数拷贝回的内存和原本的内存是不一样的,或者可以认为opencl虽然借用了这块内存作为cl_mem,但是并不保证同步的,不过初始的值是一样的,(可以使用mapmem等方式来同步)

5  CL_MEM_ALLOC_HOST_PTR:在host上新开辟一段内存供device使用

6  CL_MEM_COPY_HOST_PTR:在device上开辟一段内存供device使用,并赋值为host上一段已经存在的mem

 

7  CL_MEM_HOST_WRITE_ONLY:这块内存是host只可写的

8  CL_MEM_HOST_READ_ONLY:这块内存是host只可读的

9  CL_MEM_HOST_NO_ACCESS:这块内存是host可读可写的

 

谈谈这些flag,这些flag看起来行为比较复杂和乱,因为Opencl是一个跨硬件平台的框架,所以要照顾到方方面面,更统一就要更抽象。

首先456的区别,他们都是跟host上内存有关,区别是,4是直接使用已有的,5是新开辟,6是在device上开内存,但是初值与host相同(45都是在host上开内存)

然后看看123 和789,123是针对kernal函数的访问说的,而789是针对host的访问说的,kernal函数是device的访问,而除了kernal函数的访问基本都是host的访问(如enqueueRead/write这些操作)

通常使用host上的内存计算的效率是没有使用device上的效率高的,而创建只读内存比创建可写内存又更加高效(我们都知道GPU上分很多种内存区块,最快的是constant区域,那里通常用于创建只读device内存)

通常用各种方式开内存你的程序都work,但这里就要考验不同情况下优化的功力了

size参数:要开的内存的大小

host_ptr参数:只有在4.6两种情况用到,其他都为NULL

 当然这些内存都要使用clReleaseMemObject释放

 

内存的call_back:

有些方式 ,如CL_MEM_USE_HOST_PTR,cl_mem使用的存储空间实际就在host mem上,所以我们要小心处理这块主存,比如你删了它,但是cl_mem还在用呢,就会出现问题,而clReleaseMemObject并不一定会马上删除这个Cl_mem,它只是一个引用计数的消减,这里需要一个回调,告诉我们什么时候这块主存可以被放心的清理掉,就是clSetMemObjectDestructorCallback

CL的规范中特别说明最好不要在这个callback里面加入耗时的系统和cl API。

 

2.内存操作

1 从Cl_mem读回host mem(就算Cl_mem是直接使用host mem实现的,想读它的内容,还是要这样读回来,可以看做cl_mem是更高一层封装)

  clEnqueueReadBuffer

2 使用host_mem的值写cl_mem

  clEnqueueWriteBuffer

3 在Cl_mem和host mem之间做映射

 clEnqueueMapBuffer

这个函数比较特殊,回顾上面一节在创建buf时有一种方法CL_MEM_USE_HOST_PTR,是直接让device使用host上已有的一块的mem(p1)做buf,但是这个产生的CL_mem(p2)经过计算后值会改变,p2改变后通常p1不会被改变,因为虽然用的一块物理空间,但是cl_mem是高层封装,和host上的mem还是不一样的,要想使p1同步到p2的最新值,就要调用这句map

MAP与CopyBack的性能对比

后来我想了想,这和使用clEnqueueReadBuffer从p2read到p1有什么区别呢?map的方法按道理更快,因为p1p2毕竟一块物理地址吗,map是不是就做个转换,而read则多一遍copy的操作。而且应该在CPU做device时map速度更快,但是事实是这样的吗?本着刨根问题的精神,我真的做了一下实验,

我的实验结果是这样的,如果使用CPU做host,GPU做device,那么CopyBack反而更快,但是如果使用CPU做host,CPU也做device,那么MAP更快(不跨越硬件),而且总体上CPU+GPU的方式更快。

这个实验结果彻底颠覆了我最初的一些想法,实验数据说明1.不考虑硬件差异,MAP确实比CopyBack更快,跟我理解一样,从CPU做device的两组数据就可看出。2.至少在我的这个实验中,主存与显存间的数据copy比主存到主存自己的数据copy更快,所以在CPU+GPU的架构中,由于CopyBack方式采用的是主存显存拷贝,而MAP值涉及主存上的操作,所以CopyBack更快。不过这里我仍存在疑虑,我的分析很可能不对,或存在其他因素没考虑,关于这点,要再继续查查关于pinned memory和内存显存传递数据的一些知识。

所以在这种异构计算领域,性能和你的硬件架构、性能、组合有着非常重要的关联,所以最好的方法就是实际做实验对比。

4  在Cl_mem直接做copy

 clEnqueueCopyBuffer

 

这些函数都跟执行kernal一样是投入到device的command queue里的,但是他们又都带有一个参数blocking_read,可以指定函数是否在执行完毕后返回。

 

3.Program

3.1.compile build link

有两种从文本创建program的方式

  • 直接build:clBuildProgram
  • 先complie好,根据情况动态的link,即把上面的过程拆分为两个步骤

   clCompileProgram   clLinkProgram

但是1.2的方式不保险,这是CL1.2中加入的,而目前不是所有的platform都支持到1.2,NVIDIA好像就才1.1

opencl实际上会根据不同的硬件把通样一份代码编译成不同的机器语言,如CPu汇编或GPU汇编

 

4.Kernal的执行

这里是精华

1.设置kernal的参数

  clSetKernelArg

2.执行kernal

  clEnqueueNDRangeKernel

 

先给一段kernal代码,方便下边参数的解释,另外这里需要一些空间想象能力~

    kernal代码

 __kernel void adder(__global const float* a, __global const float* b, __global float* result)
{
 int idx = get_global_id(0);//得到当前单元格的0维度上的序号
 result[idx] = a[idx] + b[idx];
}

参数说明:

command_queue :执行那个device的命令序列

kernel:待执行的kernal obj

work_dim:我们知道CL的执行是放在一个个独立的compute unit中进行的,你可以想像这些unit是排成一条线的,或是一个二维方阵,甚至是一个立体魔方,或着更高维,这里参数就描述这个执行的维度,从1到CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS之间

global_work_size :每个维度的unit的数量,这样总共拥有的计算单元的数量将是global_work_size[0]*global_work_size[1]...

global_work_offset :这里就是规定上面代码里每个维度上第一个get_global_id()0得到的id,默认为0,例如计算一个一维的长度为255work_size的工作,cl会自动虚拟出255个计算单元,每个单元各自计算0-254位置的数相加,而如果你把他设为3,那么cl会从3开始算,也就是说3-254位置的unit会计算出结果,而0 -2这些unit根本不去参与计算。

local_work_size :前面介绍过CL的的unit是可以组合成组的(同组内可以互相通信)这个参数就决定了Cl的每个组的各维度的大小,NULL时CL会自动给你找个合适的,这里贴下我试着用不同大小的group做数组相加的效率,

OpenCL  从零开始学习OpenCL开发

这里其实看不太出什么,直觉对这个应用实例是组越少越快,但是其中也不是严格的线性关系,无论在CPU还是GPU上这个关系都是近似的,所以在实际开发中,我们选择什么维度?选择什么样的组大小?我的答案是:多做实验吧,或者要偷懒的话就置0吧,交给CL为你做(实时上Cl中很多函数都有这个NULL的自适应选项。。)

关于维度、偏移、worksize这里有个原版的图,说明的更加形象

OpenCL  从零开始学习OpenCL开发

 

后面几个参数就跟同步有关了

event_wait_list和num_events_in_wait_list:说明这个command的执行要等这些event执行了之后

event:将返回这个command相关联的event

很明显,通过这几个参数的event可以控制command之间的执行顺序。

 

5.指令执行顺序和同步

command的执行默认都是异步的,这才有利于并行度提高效率,在并行的问题中我们有时经常要做些同步的事情,或者等待某个异步的操作完成,这里有两种方法:

  • 使用EnqueueRead/write这些操作可以指定他们为同步的(即执行完毕才在host上返回)
  • 使用event来跟踪,像clEnqueueNDRangeKernel这样的操作都会关联一个event

event:

  • clEnqueue这样的操作都会关联返回一个event
  • 用户可以自己创建一个自定义的event clCreateUserEvent,要使用clReleaseEvent释放

关于event的操作:

        正是通过event同步不同的command:

  •  设置event状态:     

       设置用户自定义event的状态,clSetUserEventStatus 状态只可以被设定一次,只可以为CL_COMPLETE或者一个负值,CL_COMPLETE代表这个event完成了,等待它的那些command得以执行,而负值表示引起错误,所有等待他的那些command都被取消执行。其实event的状态还有CL_RUNNING  CL_SUBMITTED   CL_QUEUED,只是不能在这里设置。

  • 等待event

        clWaitForEvents;可以在host中等待某些event的结束,如clEnqueueNDRangeKernel这样的异步操作,你可以等待他的event结束,就标志着它执行完了

  • 查询event信息:clGetEventInfo clGetEventProfilingInfo
  • 设置回调:clSetEventCallback

不同device上的event:

      clEnqueueNDRangeKernel这样的操作等待的只能是处于相同queue里面的event(也就是同一个device上的),而同步不同queue上的event则只能用显示的方法,如clWaitForEvents等。

marker:

      marker是这样一个object,它可以看做是一个投入queue的空指令,专门用于同步,它可以向其他comman一样设定需要等待的event,操作有clEnqueueMarkerWithWaitList

barrier:

   barrier和marker十分类似,但是从名字上就可以看出最大的不同点是:marker在等待到它的依赖event之后会自动执行完毕,让后续指令执行,而barrier会阻塞在这里,直到他关联的event被显示的设置成完成状态

marker和barrier的实现在1.1和1.2版本上存在着较大的不同

同步是CL的大问题,关于同步,原版overview上也有一个非常生动的图,贴在这里吧:

在同一个device上同步

OpenCL  从零开始学习OpenCL开发

在多个device间同步

 

OpenCL  从零开始学习OpenCL开发

--------------------------------------------------------------------------------------------------------------------------------

 

      这里介绍关于OpenCL中program函数的写法,program函数通常是文本形式的,然后使用clCreateProgramWithSource这样的接口load进来。在Shader编程中也经常使用这种形式书写GPU上运行的代码,所以为了表述清楚和理解方便,这里姑且把这些program函数的源码文本称为OpenCL的shader吧

 下面都是写在shader中的一些语法

 

1 shader是类C的语言,派生自C99标准(99年ansi C接受的标准,也是C最新的标准)

不支持:

   头文件、函数指针、递归、变长数组(这个VS也不支持)

额外加入的类型:

   vector 类型 char2 ushaort4 int8  这些最后都会变成长度对齐的

   图像类型 image2d_t image3d_t  sampler_t ...

   event类型 event_t(关联于API中CL_event)

 

 

2.work item和 work Group相关函数

OpenCL  从零开始学习OpenCL开发

 

 

 3.vector 操作

vector的前一半为lo,后一半为hi

int4 v=(int4) 7 =(int4)(7,7,7,7)

 v=(in4)(1,2,3,4)

int2 v2=v.lo ->(1,2)

v2=v.hi ->(3,4)

v2.v.odd ->(2,4)

对vector做四则运算、abs等于对每个元素分别计算

 

 4.寻址空间描述符,写在变量的最前面,用于变量所处于的地址空间

__global

__local

__private

__constant

这四个分别对应了CL架构中的存储区域(设备全局、work group、compute unit 、设备constant)

 

  • 前面的__也可去掉
  • 目前global一定是constant的,也就是声明global时必须赋值 (global就等于 global constant)
  • 不同地址空间上的指针转化是没有定义的

 5.类型转换

  5.1convert类型转换;这是按照变量语意的类型转换

写成convert_destType<_sat><_roundingMode>形式,

如float4 f4=(float4)(1.0f,2.0f,3.0f,4.0f)

int4 i4=convert_int4_sat_rte(f4)

 

destType:目标类型

_sat:超出范围自动归结为最大或最小显示的数

_roundingMode:

   _rte:表示成最接近的偶数

   _rtz:朝0接近

   _rtp:朝正无穷大

  _rtn:朝负无穷大

这里面的规则比较复杂,详见http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/convert_T.html 

 

5.2 as转换:这是根据bit值重新解释的类型转换

写成as_desttype

其中转换前后的类型的vetctor size是要一样的,desttype是目标类型,这个转换会保持bit值不变,在此基础上根据desttype重新解释数值

as转换和convert转换有着本质的区别!

如float4 f4=(float4)(1.0f,2.0f,3.0f,4.0f)

int4 i4=as_int4(f4)

 

6.内建函数:

6.1各种各样超多的数学函数
:详见http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/的Built_in Function一节

贴个简表

OpenCL  从零开始学习OpenCL开发

6.2Work_group函数:

主要用于一个group内的computer item间的交互

            一个goup内的所有item必须全部执行完这个barrier函数之后才能继续进行后续的事情,也可看做这是所有item的一个同步点,不管谁快谁慢,必须到这个点停一下,大家都到了这个点之后,再继续进行。

        这里的参数分两种情况:

          CLK_LOCAL_MEM_FENCE和CLK_GLOBAL_MEM_FENCE

      这个参数我现在也没搞得很懂,大意是加入一个mem fence保证这时loca mem或者globalmem 的同步正常,关于mem fence 的概念还要再看看opencl的描述

       

  • 异步的内存copy和prefetch函数

         async_work_group_copy:他会完成global与local之间的异步的内存拷贝,这种拷贝可能会使用DMA 引擎的(DMA的数据传输不使用传统的硬件中断,会很快),这个函数是异步的,所以会返回一个事件event_t用于同步

        使用wait_group_events函数来等待上面的event返回,用于同步

        async_work_group_strided_copy:文档上说它用于gather数据从src到dest,但是文档中gather的意思不能让人很好的理解,仔细的分析一下,这个函数同  async_work_group_copy的差异在于stride,他也是完成异步的拷贝,但是它可以从src抽取一部分域出来到dst中。例如在图形学中我们经常用一个大数组表示颜色、法向、纹理坐标等等,他们是连在一起的,如{color1,ccolor2,color3,tex0,tex1,color1,color2,color3,text0,tex1,....},这时我们需要抽取其中的color信息出来,那就要用到这个stride copy。

  

再分享一下我老师大神的人工智能教程吧。零基础!通俗易懂!风趣幽默!还带黄段子!希望你也加入到我们人工智能的队伍中来!http://www.captainbed.net

上一篇:linux运维21


下一篇:DESCRIBEFIELD