当前位置：极客笔记 > OpenCL 教程 > OpenCL 与OpenGL互操作 > OpenCL 共享OpenGL渲染缓存

OpenCL 共享OpenGL渲染缓存

OpenCL除了能共享OpenGL的顶点缓存对象（VBO）与纹理对象之外，还能共享OpenGL的渲染缓存对象（Render Buffer Object，RBO）。在某些OpenGL环境能够直接将与渲染缓存所绑定的帧缓存显示到屏幕上，而在其他环境下可能需要采取其他方式来显示帧缓存中的数据。因此，对于前者，我们可以直接采用OpenCL对最后的渲染缓存做最终的后处理，然后再输出到屏幕上。而对于后者情况，我们可以将渲染缓存中的内容映射为一个纹理，然后再进行绘制到能够显示到屏幕上的上下文中。
下面介绍从OpenGL的渲染缓存创建CL存储器对象的方法：

cl_mem clCreateFromGLRenderbuffer(cl_context context,
                                        cl_mem_flags flags,
                                        GLuint renderbuffer,
                                        cl_int ＊errcode_ret)

这里，flags只能是CL_MEM_READ_ONLY、CL_MEM_WRITE_ONLY或CL_MEM_READ_WRITE三者之一。用于指明所创建的存储器是只读、只写还是可读、可写的。不过由于我们这里所创建的存储器对象用于图像对象，因此基本只能用只读或只写这两者之一。
参数renderbuffer就是在OpenGL上下文中所创建的渲染缓存对象。在调用此函数前，渲染缓存对象必须已经创建好，并绑定到指定的帧缓存上。
这个函数所返回的存储器对象就是一个图像类型的存储器对象。
下面这个demo示例将作为一个综合性的示例给出。这个示例的实现是先利用我们第一个示例中的实现来绘制一个彩色的圆，背景也是采用灰色。然后，将它绘制到指定的渲染缓存，再将该渲染缓存作为一个OpenCL内核程序的输入图像存储器对象。随即，我们再创建一个纹理对象作为OpemCL输出图像存储器对象。最后，我们在OpenCL中将整帧图像变为黑白色，然后输出到纹理，再将该纹理显示到显示器上。因此，这个示例涵盖了所有CL与GL共享存储器对象的使用方式。由于代码比较多，我们就给出主机端新增的，比较核心部分的代码。由于GLSL代码与OpenCL内核代码与前两个例子都一样，因此不再重复给出了。

-(void)prepareForDraw
{
    glGenVertexArrays(1, &mVAOForDraw);
    glBindVertexArray(mVAOForDraw);
    //设置顶点VBO
    glGenBuffers(1, &mVBOVertForDraw);
    glBindBuffer(GL_ARRAY_BUFFER, mVBOVertForDraw);
    //将顶点坐标数据拷贝到mVBOVertices对象的缓存中
    glBufferData(GL_ARRAY_BUFFER, sizeof(sVertexCoords),
                    sVertexCoords, GL_STATIC_DRAW);
    //将顶点VBO绑定到属性0
    glEnableVertexAttribArray(0);
    glVertexAttribPointer(0, 4, GL_FLOAT, GL_FALSE, 0,
                              (const GLvoid ＊)0);
    //设置纹理坐标VBO
    glGenBuffers(1, &mVBOTextureCoords);
    glBindBuffer(GL_ARRAY_BUFFER, mVBOTextureCoords);
    //将sTextureCoords中的数据拷贝到mVBOTextureCoords对象的缓存中
    glBufferData(GL_ARRAY_BUFFER, sizeof(sTextureCoords),
                    sTextureCoords, GL_STATIC_DRAW);
    //将纹理坐标VBO绑定到属性1
    glEnableVertexAttribArray(1);
glVertexAttribPointer(1, 2, GL_FLOAT, GL_FALSE, 0,
                              (const GLvoid ＊)0);
    //解绑
    glBindBuffer(GL_ARRAY_BUFFER, 0);
    //设置纹理
    glPixelStorei(GL_UNPACK_ALIGNMENT, 8);
    glActiveTexture(GL_TEXTURE0);
    glGenTextures(1, &mTexName);
    glBindTexture(GL_TEXTURE_2D, mTexName);
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_REPEAT);
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_REPEAT);
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER,
                      GL_LINEAR);
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER,
                      GL_LINEAR);
    //这里只分配纹理大小及属性,而不对它进行图像数据的拷贝
    //纹理的实际图像数据将在OpenCL内核程序中给出
    glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, mRenderWidth,
                  mRenderHeight, 0, GL_BGRA,
                  GL_UNSIGNED_INT_8_8_8_8_REV, NULL);
    if(mProgram != 0)
    {
        glDeleteProgram(mProgram);
        mProgram = 0;
    }
    //加载着色器并构建OpenGL程序
    if(![self loadShaders:@"draw"])
        return;
    glUseProgram(mProgram);
    //将采样器对象映射到GL_TEXTURE0,使得它对0号纹理单元进行采样
    glUniform1i(mSamplerLocation, 0);
    glViewport(0, 0, mRenderWidth, mRenderHeight);
    glClearColor(1.0f, 1.0f, 1.0f, 1.0f);
}
- (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;
    //从GL的RBO创建用于输出的CL存储器对象
    cl_mem dstMem = NULL;
    //从GL的RBO创建用于输入的CL存储器对象
    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;
        }
        //这里我们总共使用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, 0,
                                                NULL, NULL);
        //这里直接用clFinish进行同步,确保顶点坐标以及相应的颜色值全都设置好
        clFinish(commandQueue);
        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);
    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);
    [[self openGLContext] flushBuffer];
}

Camera课程

Camera API2 教程

API2 教程
Camera 性能教程

Perf 教程
Camera Native Framework 教程

Native FW 教程
Camera QCOM ISP 教程

ISP 教程
Camera QCOM Sensor 教程

Sensor 教程

Python教程

Python 教程

Python 教程
NumPy 教程

NumPy 教程
Django 教程

Django 教程

Java教程

Java 教程

Java 教程

Web教程

JavaScript 教程

JavaScript 教程
jQuery 教程

jQuery 教程
CSS 教程

CSS 教程
PHP 教程

PHP 教程
Laravel 教程

Laravel 教程
TypeScript 教程

TypeScript 教程
WordPress 教程

WordPress 教程

数据库教程

SQL 教程

SQL 教程
SQLite 教程

SQLite 教程
MySQL 教程

MySQL 教程
PL/SQL 教程

PL/SQL 教程
MongoDB 教程

MongoDB 教程

图形图像教程

OpenCV 教程

OpenCV 教程

办公软件教程

Excel 教程

Excel 教程

Linux教程

Linux命令全集

Linux命令全集
Linux内核API

Linux内核API
LVGL 教程

LVGL 教程

计算机教程

操作系统 教程

操作系统 教程
计算机网络 教程

计算机网络 教程
C语言 教程

C语言 教程
C++ 教程

C++ 教程
Swift 教程

Swift 教程

大数据教程

Hadoop 教程

Hadoop 教程
Spark 教程

Spark 教程
Scala 教程

Scala 教程

开发工具教程

Git 教程

Git 教程
Jenkins 教程

Jenkins 教程
ChatGPT 教程

ChatGPT 教程
IntelliJ 教程

IntelliJ 教程

OpenCL 与OpenGL互操作