极客笔记为了要确保共享对象的数据完整性，OpenCL扩展提供了CL与GL对象访问的同步机制。例如，当我们要从一个OpenGL纹理对象来创建一个OpenCL图像类型的存储器对象时，倘若OpenGL的纹理数据没有生成完，那么OpenCL那边所获得的图像数据将是不完整的。我们尽管可以通过在OpenGL上下文中调用glFinish这种同步API，但是效率非常低！因为此API与clFinish（）语义差不多，要将当前OpenGL所有命令都执行完后才能让当前主机端的线程继续执行下去，否则当前线程将会被阻塞。这会导致很多不必要的计算资源的浪费。所以，OpenCL扩展API中提供了clEnqueueAcquireGLObjects这个函数，以获得相应的一个OpenCL事件对象。在OpenCL上下文中可以直接用此事件对象来做相应的OpenCL命令执行的同步就能更高效地利用计算资源了。下面介绍一下这个函数：
cl_int clEnqueueAcquireGLObjects(cl_command_queue command_queue,
cl_uint num_objects.
const cl_mem ＊mem_objects,
cl_uint num_events_in_wait_list,
const cl_event ＊event_wait_list,
cl_event ＊event)
- 参数num_objects指定了参数mem_objects所指向的一维数组中包含的所要获得的从GL对象所创建出的OpenCL存储器对象的个数。
- 参数mem_objects就指向一个包含所需要获得的OpenCL存储器对象的数组首地址。
后三个参数与其他clEnqueue系API的一样。
当这个函数调用成功后会返回CL_SUCCESS，如果是其他值说明调用失败。另外，我们往往需要最后一个参数event所生成的事件对象来追踪mem_objects中的所有OpenCL存储器对象的数据是否都在OpenGL上下文中生成完毕。所以，我们需要利用该事件对象对其他排入当前命令队列的命令进行同步。
在调用了clEnqueueAcquireGLObjects之后，必须在OpenGL上下文中再次使用OpenCL存储器对象所绑定的OpenGL对象之前调用clEnqueueReleaseGLObjects函数来释放所获得的OpenCL对象。该函数原型为：
cl_int clEnqueueReleaseGLObjects (cl_command_queue command_queue,
cl_uint num_objects,
const cl_mem ＊mem_objects,
cl_uint num_events_in_wait_list,
const cl_event ＊event_wait_list,
cl_event ＊event)
该函数参数与clEnqueueAcquireGLObjects的一模一样。我们在调用此函数时，所传递的num_objects和mem_objects参数也要同调用clEnqueueAcquireGLObjects时完全一致。
下面就用上面的核心代码稍作修改来给出使用clEnqueueAcquireGLObjects和clEnqueue-ReleaseGLObjects的例子。
- (void)generateTextureWithCLMemFromRBO
{
/＊ 做OpenCL初始化 ＊/
cl_platform_id oclPlatform = NULL;
cl_device_id oclDevice = NULL;
//要被创建的OpenCL上下文对象
cl_context context = NULL;
cl_command_queue commandQueue = NULL;
cl_program oclProgram = NULL;
cl_kernel kernel = NULL;
//从OpenGL的RBO创建用于输出的OpenCL存储器对象
cl_mem dstMem = NULL;
//从OpenGL的RBO创建用于输入的OpenCL存储器对象
cl_mem srcMem = NULL;
//访问图像对象的采样器
cl_sampler sampler = NULL;
#ifdef _APPLE_
CGLContextObj cgl_context = CGLGetCurrentContext();
CGLShareGroupObj sharegroup = CGLGetShareGroup(cgl_context);
gcl_gl_set_sharegroup(sharegroup);
#endif
do
{
//获得当前OpenCL平台
cl_int status = clGetPlatformIDs(1, &oclPlatform, NULL);
if(status != CL_SUCCESS)
{
NSLog(@"OpenCL platform get failed!");
break;
}
//获得当前GPU设备。严格地来说,
//此GPU设备也应该是OpenGL所使用的设备
status = clGetDeviceIDs(oclPlatform, CL_DEVICE_TYPE_GPU, 1,
&oclDevice, NULL);
if(status != CL_SUCCESS)
{
NSLog(@"OpenCL GPU cannot be found!");
break;
}
//设置用于创建OpenCL上下文的属性列表
cl_context_properties properties[] =
{
#ifdef WIN32
CL_GL_CONTEXT_KHR ,
(cl_context_properties)wglGetCurrentContext(),
CL_WGL_HDC_KHR ,
(cl_context_properties)wglGetCurrentDC(),
#endif
#ifdef _linux_
CL_GL_CONTEXT_KHR ,
(cl_context_properties)glXGetCurrentContext(),
CL_GLX_DISPLAY_KHR ,
(cl_context_properties)glXGetCurrentDisplay(),
#endif
#ifdef _APPLE_
CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE,
(cl_context_properties)sharegroup,
#endif
0
};
//创建OpenCL上下文
context = clCreateContext(properties, 1, &oclDevice,
NULL, NULL, NULL);
//创建命令队列
commandQueue = clCreateCommandQueue(context, oclDevice,
0, NULL);
//编译内核程序
NSString ＊kernelPath = [[NSBundle mainBundle]
pathForResource:
@"compute" ofType:@"ocl"];
const char ＊aSource = [[NSString stringWithContentsOfFile:
kernelPath
encoding:NSUTF8StringEncoding error:nil] UTF8String];
size_t kernelLength = strlen(aSource);
oclProgram = clCreateProgramWithSource(context, 1, &aSource,
&kernelLength, NULL);
if(oclProgram == NULL)
{
NSLog(@"OpenCL program create failed!");
break;
}
//构建程序
status = clBuildProgram(oclProgram, 1, &oclDevice,
NULL, NULL, NULL);
if(status != CL_SUCCESS)
{
NSLog(@"OpenCL kernel build failed!");
break;
}
//创建与GL纹理对象共享的CL图像存储器对象,用于纹理输出
dstMem = clCreateFromGLTexture(context, CL_MEM_WRITE_ONLY,
GL_TEXTURE_2D, 0, mTexName,
NULL);
//从GL渲染缓存创建与CL共享的存储器对象,作为输入源
srcMem = clCreateFromGLRenderbuffer(context,
CL_MEM_READ_ONLY,
mRBO, NULL);
//创建采样器对象
sampler = clCreateSampler(context, CL_FALSE,
CL_ADDRESS_CLAMP_TO_EDGE,
CL_FILTER_LINEAR, NULL);
//创建内核对象
kernel = clCreateKernel(oclProgram, "ImageProcessing",
NULL);
//设置内核参数
status = clSetKernelArg(kernel, 0, sizeof(dstMem), &dstMem);
status |= clSetKernelArg(kernel, 1, sizeof(srcMem),
&srcMem);
status |= clSetKernelArg(kernel, 2, sizeof(cl_sampler),
(void ＊)&sampler);
if(status != CL_SUCCESS)
{
NSLog(@"Kernel parameters pass failed!");
break;
}
//用于跟踪获得GL对象的事件
cl_event acquireEvt = NULL;
//获得从GL对象所创建的CL对象,等待GL绘制完成
clEnqueueAcquireGLObjects(commandQueue, 2, (cl_mem[])
{
dstMem, srcMem
}, 0, NULL, &acquireEvt);
//这里我们总共使用mImageWidth ＊ mImageHeight个工作项,
//每个工作项来处理一个像素
size_t global_work_size[] = { mRenderWidth, mRenderHeight };
//然后设置一个工作组中的工作项个数。
//要注意,x维度与y维度两个数相乘
//不能大于工作组中最多可容纳的工作项的个数
size_t local_work_size[] = { 16, 16 };
//运行内核程序
status |= clEnqueueNDRangeKernel(commandQueue, kernel, 2,
NULL, global_work_size,
local_work_size, 1, &acquireEvt, NULL);
//这里直接用clFinish进行同步,确保顶点坐标以及相应的颜色值全都设置好
clFinish(commandQueue);
//调用过clEnqueueAcquireGLObjects之后必须调用此函数,
//并且在后续使用相关OpenGL对象之前完成调用
clEnqueueReleaseGLObjects(commandQueue, 2, (cl_mem[]){
dstMem, srcMem
}, 0, NULL, NULL);
clReleaseEvent(acquireEvt);
if(status != CL_SUCCESS)
{
NSLog(@"OpenCL kernel run error!");
}
}
while(NO);
//释放OpenCL各种对象
if(dstMem != NULL)
clReleaseMemObject(dstMem);
if(srcMem != NULL)
clReleaseMemObject(srcMem);
if(sampler != NULL)
clReleaseSampler(sampler);
if(kernel != NULL)
clReleaseKernel(kernel);
if(oclProgram != NULL)
clReleaseProgram(oclProgram);
if(commandQueue != NULL)
clReleaseCommandQueue(commandQueue);
if(context != NULL)
clReleaseContext(context);
}
- (void)drawRect:(NSRect)dirtyRect
{
//绑定设置好的帧缓存对象与渲染缓存对象
glBindFramebuffer(GL_FRAMEBUFFER, mFBO);
glBindRenderbuffer(GL_RENDERBUFFER, mRBO);
//准备在绑定的渲染缓存上绘制彩色的圆
[self doOpenCLComputing];
glClear(GL_COLOR_BUFFER_BIT);
glDrawArrays(GL_TRIANGLE_FAN, 0, 362);
//这里将使用acquire接口进行同步,而不直接使用glFinish()
//glFinish();
//准备将渲染缓存中的图像变为纹理,然后交给OpenCL做进一步的处理
[self prepareForDraw];
[self generateTextureWithCLMemFromRBO];
//使用默认帧缓存与渲染缓存,准备绘制
glBindFramebuffer(GL_FRAMEBUFFER, 0);
glBindRenderbuffer(GL_RENDERBUFFER, 0);
glClear(GL_COLOR_BUFFER_BIT);
glDrawArrays(GL_TRIANGLE_STRIP, 0, 4);
glFlush();
[[self openGLContext] flushBuffer];
}
在drawRect方法中，我们看到之前的glFinish（）的调用被注释掉了，取而代之的就是在generateTextureWithCLMemFromRBO方法中使用clEnqueueAcquireGLObjects进行同步。通过该函数所返回的acquireEvt事件对象作为clEnqueueNDRangeKernel的同步事件。在完成dstMem与srcMem对象的数据生成之前，OpenCL内核命令的执行会一直被阻塞。在内核命令执行完之后，我们即可调用clEnqueueReleaseGLObjects来释放与OpenGL对象相关联的OpenCL存储器对象。紧接着generateTextureWithCLMemFromRBO方法下面，我们就开始做OpenGL端的图形绘制。这时，OpenCL端的工作全部完成，且OpenCL上下文及各种存储器对象也可以全部销毁。
