改章节朋友在上海逛街的时候突然想到的...之前就有想写几篇关于内存执行的笔记，所以回家到以后就奋笔疾书的写出来发表了
欢送关注，转载引用请注明 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，并返回。这里创建的都是global的mem。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间同步
http://www.cnblogs.com/jisi5789/archive/2013/05/18/3085738.html