近期计划
- 了解gpu硬件组成和软件运行时执行过程
- 尝试优化arm opencl benchmark 版本的opencl开销,降低计算耗时
- opencl 最多只是一套标准的api,要想深入了解还需要对GPU的工作方式深入了解.
OpenCl API
clFinish( )
阻塞，直到该命令队列中所有命令发送给设备,并且执行完成.不建议使用,可用wait event代替该操作.
clCreateCommandQueue( )
CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE 乱序执行该命令队列
CL_QUEUE_PROFILING_ENABLE 允许profile ,不建议release版使用
命令队列可以压入对context创建的memory, program和kernel 对象的操作命令,同一队列的命令可以指定为顺序执行,多个命令队列可以让程序在不用同步的前提下执行多个有序的命令.如果多个命令队列操作的对象之间有共享的地方就需要用同步机制保证.
乱序执行的时候设置clEnqueueNDRangeKernel 的event_wait_list 参数,指定某个命令等待特定命令执行完毕之后再执行,可以实现乱序执行的同步机制.
在队列中可以wait event(clEnqueueMarkerWithWaitList),等待事件可以保证被等待事件指定的命令执行完毕之后再执行队列中的后续事件.也可以插入barrier(clEnqueueBarrierWithWaitList)命令,barrier命令保证之前的命令都被执行完毕之后再执行后续的命令.在乱序执行的时候可由clEnqueueNDRangeKernel,clEnqueueTask ,clEnqueueNativeKernel返回的event通过wait event或者barrier命令加入到命令队列中保证等待上述命令执行结束再执行后续命令.
clCreateBuffer( )
先了解opencl的内存模型和内存机制之后再看相关的内存API会比较容易接受一点,内存操作很大程度上决定程序的整体效率.在嵌入式领域,Mali GPU是没有独显的,和CPU共用一块内存,在这种情况下可以省去不必要的内存拷贝操作.
请参阅ARM Optimizing OpenCL for Mali GPUs
CL_MEM_USE_HOST_PTR 显存,含内存拷贝
在device中开辟空间用host指针内存区域内的数据来初始化,device执行的时候直接使用device上的内存.在Mali GPU上,这个操作会将数据拷贝到到另一块区域,map的时候再拷贝回来,浪费时间,在Mali GPU上不建议使用该选项分配内存对象.但是对于有单独显存的设备,此种分配内存对象的方式还有有意义的.
CL_MEM_ALLOC_HOST_PTR 零拷贝内存
从host指针可以访问的内存区域申请内存.这块内存属于零拷贝内存.在Mali系列的GPU上,这将有利于系统性能提升.但对于有单独显存的设备则不适合.
CL_MEM_COPY_HOST_PTR 会将host内存数据拷贝到driver分配的内存中.和CL_MEM_ALLOC_HOST_PTR一起使用,用来初始化这块内存.
clEnqueueNDRangeKerne ()
- global_work_size
每一维上work-item的数量 - local_work_size
每一维上一个workgroup中的workitem的数量.除非要用到共享内存,否则这个值传入NULL即可,由opencl决定最好的workgroup size.
编译
- 代码是在opencl运行时候编译的,规避不同平台之间的差异,build program之后,创建的kernel可以复用,减小运行时开销
在嵌入式设备上动态编译估计开销很大,还需研究一下是否可以静态编译
vector(向量化)
Arm 本身有一套SIMD指令--NEON,在目前看来,Arm正在积极扩展NEON指令功能,力求在Arm平台上尽可以的提升多媒体处理能力.Arm官方给出的两个优化选项(向量化和展开循环)尝试了一番,并不知道怎么用,日后还需研究一下.
opencl本身就是用于数据并发,而在访存和计算方面,向量化无疑会巨大的提升opencl的整体表现.
vloadn/vstoren n: 2,4,8,16.最多单次读写16个操作数,并且对这些向量计算只需要一次即可完成.复合Mali对内存的访问要求.
测试发现使用向量化的opencl性能比NEON提升十倍多(MT6753平台,矩阵乘).效果十分显著,opencl+vectorization 简直就是开挂,NEON的处境就比较尴尬了.
