OpenCL 访问共享对象的OpenCL与OpenGL之间的同步

为了要确保共享对象的数据完整性,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上下文及各种存储器对象也可以全部销毁。

赞(1)
未经允许不得转载:极客笔记 » OpenCL 访问共享对象的OpenCL与OpenGL之间的同步
分享到: 更多 (0)

评论 抢沙发

  • 昵称 (必填)
  • 邮箱 (必填)
  • 网址