opencl/opengl interop using clCreateFromGLTexture fails to draw to texture (texture black)

Question

The setup is a little complicated so I will try my best to detail it.

First I am attempting to use openCL/openGL interop. The code works when the interop cl::ImageGL is not used so the basics are there. The project is using 3 main openGL context (took a lot to get the openCL context to open in that). I am using QGLWidgets for the context. There is a hidden QGLWidget created first and the other 2 share the hidden's context. Each of the 2 QGLWidgets are running in their own threads. The hidden QGLWidget is transfer to a thread where the openCL context is created.

QGLFormat qglFormat;
qglFormat.setVersion(3, 3); 
qglFormat.setProfile(QGLFormat::CoreProfile);

m_hiddenGl=new GLHiddenWidget(qglFormat);
m_hiddenGl->setVisible(false);

m_view1=new GLWidget(qglFormat, m_hiddenGl);
m_view2=new GLWidget(qglFormat, m_hiddenGl);
 ...
 QThread *processThread=m_process.qThread();

m_hiddenGl->doneCurrent();
m_hiddenGl->context()->moveToThread(processThread);

GLWidget is a custom class that launches its own thread and moves the context, GLHiddenWidget again custom class but basically just overrides all functions needed to keep makeCurrent from being called by the main thread.

Inside the process thread at start is the following

m_hiddenGl->makeCurrent();
hdc=wglGetCurrentDC();
glHandle=wglGetCurrentContext();

cl_context_properties clContextProps[]={ 
    CL_CONTEXT_PLATFORM, (cl_context_properties)m_openCLPlatform(),
    CL_WGL_HDC_KHR, (intptr_t) hdc,
    CL_GL_CONTEXT_KHR, (intptr_t) glHandle, 0
};

m_openCLContext=cl::Context(m_openCLDevice, clContextProps, NULL, NULL, &error);

This all is a go. From this point several kernels are executed sequentially on incoming data which is images. All of the kernels succeed (no errors) however the one kernel using a openGL texture to write out data fails to write anything. When using an openCL cl::Image2d it works fine (produces the correct output) even if the openCL context is created as interop.

The openGL texture is created after all of the openGL contexts are created and after the openCL context is created (also in the same thread as the openCL context). At the beginning of the processThread the texture is generated by the hidden QGLWidget with glGenTextures. Then all of the kernels are created along with other openCL buffers and images. Right before the kernel is executed the following is done.

if(initBuffer) //runs only if buffer size is changed, allways runs first time
{
    m_hiddenGl->makeCurrent();

    glBindTexture(GL_TEXTURE_2D, m_texture);

    //I have attempted to put data in, result is always black from kernel
    glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, width, height, 0, GL_RGBA, GL_UNSIGNED_BYTE, NULL); 
    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);
    glBindTexture(GL_TEXTURE_2D, 0);

    glFinish();

    //m_flags is CL_MEM_READ_WRITE
    m_imageGL=cl::ImageGL(m_openCLContext, m_flags, GL_TEXTURE_2D, 0, m_texture, &error);
}

and a simplified kernel.

__kernel void kernel1(__read_only image2d_t src1, __read_only image2d_t src2, __write_only image2d_t dst)
{
    int2 coord=(int2)(get_global_id(0), get_global_id(1));
    uint4 value=255;

    write_imageui(dst, coord, convert_uint4(value));
}

Even if I do not display the texture the image is still black. The image is read back from openCL and saved to the hard disk, when using cl::Image2d or cl::ImageGL. With cl::Image2d it is correct with cl::ImageGL it is black.

Answer 1

Well as noted above in my comment I found my problem was related to the definition of the texture for openGL. Since I used write_imageui in the openCL kernel the openGL texture had to be defined as follows:

glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8UI, width, height, 0, GL_RGBA_INTEGER, GL_UNSIGNED_BYTE, NULL);

I have made some time since then and created a example program that uses both openGL/openCL interop and threaded QGLWidgets. You can read my comments on it here http://www.krazer.com/?p=109 and/or you can get the source from github.com

Answer 2

I've had the same problem and using write_imagef() (don't forget to divide RGBA values by 255.0!), instead of write_imageui() in the OpenCL kernel solved it, without changing the pixel format of the OpenGL texture (you may not have the opportunity to do that, especially if it's an existing texture, created somewhere else in a large app).

Answer 3

You'd have to post more code for us to find the error.

In any case, I've written an example program which does something very similar. It uses OpenCL to calculate a mandelbrot fractal then draws it with OpenGL inside a QGLWidget. The source is here: https://github.com/kylelutz/compute/blob/master/example/mandelbrot.cpp.

Hope that helps.