从零开始学习OpenCL开发(三)深入API
这里将更深入的说明一些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做数组相加的效率,
这里其实看不太出什么,直觉对这个应用实例是组越少越快,但是其中也不是严格的线性关系,无论在CPU还是GPU上这个关系都是近似的,所以在实际开发中,我们选择什么维度?选择什么样的组大小?我的答案是:多做实验吧,或者要偷懒的话就置0吧,交给CL为你做(实时上Cl中很多函数都有这个NULL的自适应选项。。)
关于维度、偏移、worksize这里有个原版的图,说明的更加形象
后面几个参数就跟同步有关了
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上同步
在多个device间同步