# OpenCL-OpenGL Interop

OpenCL is a cross-platform interface to compute devices.

OpenGL is a cross-platform interface to graphic devices.

Interop means sharing data directly on the device (GPU), without having to transfer via the host (CPU). This provides significant efficiency savings.

So far this page focuses on writing OpenGL textures from OpenCL kernels.

# 1 Device Support

To check if the OpenCL platform supports OpenGL sharing, check the platform extensions string:

// TODO: code example to follow

To check if the OpenCL device supports OpenGL sharing, check the device extensions string:

cl_int err
size_t size = 0;
err = clGetDeviceInfo(device, CL_DEVICE_EXTENSIONS, 0, 0, &size);
char *extensions = malloc(size + 1);
err = clGetDeviceInfo(device, CL_DEVICE_EXTENSIONS, size, extensions, 0);
extensions[size] = 0;
const char *extension = "cl_khr_gl_sharing";
int len = strlen(extension);
char *match = strstr(extensions, extension);
bool have_gl_sharing = !!match && (match[len] == 0 || match[len] == ' ');

I have support with rusticl but not Clover, using Mesa 25.0.7-2 on AMD RX 580.

# 2 Context Initialisation

OpenGL context etc must be provided at OpenCL context creation time, so you need to initialize OpenGL first.

# 2.1 SDL

Using OpenCL C API and SDL2, with the SDL GL context current:

#include <SDL_syswm.h>

#ifdef SDL_VIDEO_DRIVER_X11
#include <GL/glx.h>
#endif

#ifdef SDL_VIDEO_DRIVER_WINDOWS
#include <wingdi.h>
#endif

cl_context_properties properties[8];
memset(properties, 0, sizeof(properties));
properties[0] = CL_CONTEXT_PLATFORM;
properties[1] = platform;

SDL_SysWMinfo info;
SDL_VERSION(&info.version);
SDL_GetWindowWMInfo(window, &info);
switch (info.subsystem)
{

#ifdef SDL_VIDEO_DRIVER_X11
  case SDL_SYSWM_X11:
    properties[2] = CL_GLX_DISPLAY_KHR;
    properties[3] = (intptr_t) info.info.x11.display;
    properties[4] = CL_GL_CONTEXT_KHR;
    properties[5] = (intptr_t) glXGetCurrentContext();
    break;
#endif

#ifdef SDL_VIDEO_DRIVER_WINDOWS
  case SDL_SYSWM_WINDOWS:
    properties[2] = CL_WGL_HDC_KHR;
    properties[3] = (intptr_t) info.info.win.hdc;
    properties[4] = CL_GL_CONTEXT_KHR;
    properties[5] = (intptr_t) wglGetCurrentContext();
    break;
#endif

// TODO: Wayland, Apple, ...
}

cl_context *context = clCreateContext(properties, num_devices, devices, 0, 0, &err);

# 2.2 GLFW

I found an OpenCL-OpenGL interop example using GLFW3.

Key things to modify from above SDL snippet:

#ifdef USE_X11
#define GLFW_EXPOSE_NATIVE_X11
#define GLFW_EXPOSE_NATIVE_GLX
#endif

#ifdef USE_WINDOWS
#define GLFW_EXPOSE_NATIVE_WIN32
#define GLFW_EXPOSE_NATIVE_WGL
#endif

#include <GLFW/glfw3.h>
#include <GLFW/glfw3native.h>

#ifdef USE_X11
properties[3] = (intptr_t) glfwGetX11Display();
properties[5] = (intptr_t) glfwGetGLXContext(window);
#endif

#ifdef USE_WINDOWS
properties[3] = (intptr_t) GetDC(glfwGetWin32Window(window));
properties[5] = (intptr_t) glfwGetWGLContext(window);
#endif

// TODO: Wayland, Apple, ...

# 3 Textures / Images

For writing textures from OpenCL:

# 3.1 Kernel Image Functions

In OpenCL kernel source:

__kernel foo(write_only image2d_t out)
{
  write_imagef(out, (int2)(x, y), (float4)(...));
  write_imagei(out, (int2)(x, y), (int4)(...));
  write_imageui(out, (int2)(x, y), (uint4)(...));
}

The correct write function to use depends on the OpenGL texture internal format. Needs four channels even if texture has fewer.

# 3.2 Image Creation

Create OpenGL texture as normal. After each glTexImage2D() (which invalidates any corresponding OpenCL image), for writing from OpenCL:

glFinish();
int mipmap_level = 0;
cl_mem image = clCreateFromGLTexture(context, CL_MEM_WRITE_ONLY, GL_TEXTURE_2D, mipmap_level, texture, &err);

Connect to kernel with:

clSetKernelArg(kernel, 0, sizeof(image), &image);

Need to clReleaseMemObject() when done.

# 3.3 OpenCL Image Write

See Synchronisation below.

glFinish(); // see Synchronisation below
cl_event acquired_event, kernel_event;
clEnqueueAcquireGLObjects(command_queue, 1, &image, 0, 0, &acquired_event);
clEnqueueNDRangeKernel(command_queue, kernel, dim, offset, global_size, local_size, 1, &acquired_event, &kernel_event);
clEnqueueReleaseGLObjects(command_queue, 1, &image, 1, &kernel_event, 0);
clFinish(command_queue); // see Synchronisation below
clReleaseEvent(acquired_event);
clReleaseEvent(kernel_event);

# 3.4 OpenGL Texture Read

After the image has been written from OpenCL, the texture needs to be rebound in OpenGL:

glBindTexture(GL_TEXTURE_2D, texture);

# 4 Synchronisation

glFinish() and clFinish() are expensive. Synchronisation through GLsync and cl_event is preferable but not yet supported by Mesa 25.0.7-2.

# 4.1 `cl_khr_gl_event`

OpenGL 3.2+ has:

GLsync glFenceSync(GL_SYNC_GPU_COMMANDS_COMPLETE, 0);

OpenCL extension: cl_khr_gl_event provides:

cl_event clCreateEventFromGLsyncKHR(cl_context context, GLsync sync, cl_int *err);

These events can only be used by clEnqueueAcquireGLObjects().

Need to glDeleteSync() and clReleaseEvent() after use.

Unfortunately does not seem to be widely supported (I don’t have it in rusticl or Clover from Mesa 25.0.7-2 on AMD RX 580).

# 4.2 `GL_ARB_cl_event`

OpenGL extension GL_ARB_cl_event provides:

GLsync glCreateSyncFromCLeventARB(cl_context context, cl_event event, 0);

Only events from clEnqueueReleaseGLObjects() can be used.

Then OpenGL 3.2+ has:

void glWaitSync(GLsync sync, 0, GL_TIMEOUT_IGNORED); // server (GPU) waits
void glClientWaitSync(GLsync sync, GLbitfield flags, GLuint64 timeout_nanoseconds); // client (CPU) waits

There is a potential issue with client wait and flushing.