Compute Shaders in OpenGL

[quirin]

The functionality of CUDA/OpenCL is exposed to Direct3D developers through "Computer Shaders" aka DirectCompute.

While at first sight, these APIs seem equivalent, there are differences that result in both advantages and disadvantages.

From a graphics programmers' perspective, DirectCompute has one important advantage, right now: interoperability works a lot better and is faster, because there is no context switch between the compute context and the graphics context. There is only one context, if I am informed correctly.

Is anyone aware of attempts to have compute shaders in GLSL?
Or even better: is there a way of running OpenCL in the same context as OpenGL?

[Alfonse Reinheart]

What would it mean to have them running "in the same context"? An OpenGL context isn't the same thing as an OpenCL context. They're two different concepts.

The issue with interop is that you basically have to finish all your GL stuff (with glFinish) before you can start your CL stuff. And vice-versa. That has nothing to do with contexts, and everything to do with the specification design. This hard synchronization is due to neither spec really mentioning how they're supposed to interact.

[Chris Lux]

There are extensions to GL and CL to handle the interop, but sadly they are currently not supported (at least by Nvidia):

CL event to OpenGL sync
http://www.opengl.org/registry/specs/ARB/cl_event.txt

OpenGL sync to CL event
http://www.khronos.org/registry/cl/s.../gl_event.html

[Zyx_2000]

Some of the things one might want to do with OpenCL can be done efficiently with transform feedback.

[marshats]

Are you saying clEnqueueAcquireGLObjects/clEnqueueReleaseGLObjects are inefficient? I have seen no problems on NVIDIA hardware with interop between openGL and openCL ... this is a snippet from OpenCL Book Chapter 10 GLInterop.cpp. I ran it on my Linux, GeForce GTX 465, OpenGL version 4.1.0, OpenCL 1.1 setup.

    // Note, we should ensure GL is completed with any commands that might affect this VBO
    // before we issue OpenCL commands
    glFinish();
 
    // map OpenGL buffer object for writing from OpenCL
    clEnqueueAcquireGLObjects(cqCommandQue, 1, &vbo_cl, 0,0,0);
 
    // Set work size and execute the kernel
    clEnqueueNDRangeKernel(cqCommandQue, ckKernel, 2, NULL, szGlobalWorkSize, NULL, 0,0,0 );
 
    // unmap openGL buffer object
    clEnqueueReleaseGLObjects(cqCommandQue, 1, &vbo_cl, 0,0,0);
 
    // Note, we should ensure OpenCL is finished with any commands that might affect the VBO
    clFinish(cqCommandQue);

Also, my NVIDIA card reports back "cl_khr_gl_sharing" is a supported CL_DEVICE_EXTENSIONS (see link) .

[quirin]

Some of the things one might want to do with OpenCL can be done efficiently with transform feedback.

You are right. But I never got quite the bandwidth that I got with CUDA

[quirin]

What would it mean to have them running "in the same context"? An OpenGL context isn't the same thing as an OpenCL context. They're two different concepts.

Exactly. I am aware of that. It seems to me that in D3D11 there is just one context, i.e. one concept, instead of two. That is why I asked in the first place.

My problem is that, before every draw call I run a CUDA kernel taking about 0.1 ms. I have a cost of about 0.15ms when switching between CUDA and OpenGL. So I fear, my algorithm is context-switch bound.

[quirin]

Also, my NVIDIA card reports back "cl_khr_gl_sharing" is a supported CL_DEVICE_EXTENSIONS.

Sweet :-) Looks like the way to go...

Only one little thing: Could you please be so kind and report timings how fast you can switch between the GL and CL world on your card. That would be really nice to know. Thanks in advance!

[marshats]

On further reading (OpenCL Book by Munshi et al Chapter 10 pg 348-350 "Syncronization between OpenGL and OpenCL"), there is a code snippet to remove glFinish/clFinish. It requires two extensions GL_ARB_sync and cl_khr_event. My pseudo code changes to above post are as follows (note untested)

   cl_event release_event;
 
    GLsync sync = glFenceSync(GLSYNC_GPU_COMMANDS_COMPLETE, 0);
    gl_event = clCreateEventFromGLSyncKHR(context,sync,NULL);
 
    // map OpenGL buffer object for writing from OpenCL
    clEnqueueAcquireGLObjects(cqCommandQue, 1, &vbo_cl, 0,&gl_event,NULL);
 
    // Set work size and execute the kernel
    clEnqueueNDRangeKernel(cqCommandQue, ckKernel, 2, NULL, szGlobalWorkSize, NULL, 0,0,0 );
 
    // unmap openGL buffer object
    clEnqueueReleaseGLObjects(cqCommandQue, 1, &vbo_cl, 0,0,&release_event);
 
    GLsync cl_sync = glCreateSyncFromCLeventARB(context,release_event,0);
    glWaitSync(cl_sync, 0, GL_TIMEOUT_IGNORED);

Sadly my GPU has only one of the two required extensions required to try this; GL_ARB_sync is available but cl_khr_gl_event is not.

[Chris Lux]

Sadly my GPU has only one of the two required extensions required to try this; GL_ARB_sync is available but cl_khr_gl_event is not.

you are missing more than one extension. you need the cl_khr_gl_event extension to OpenCL and the ARB_cl_event extension to OpenGL. (see my post on top).