Opencl API解释（二）

欢迎关注，转载引用请注明 http://blog.csdn.net/leonwei/article/details/8909897

 

这里将更深入的说明一些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间同步