OpenGL 3.3 / OpenCL 1.1 texture sharing

I'll present a sample that create an OpenGL texture, update it with an OpenCL kernel, and display it on a screen space quad.

The complete source code is located on sourceforge : virtrev

A CMakeList (CMake) is provided to create the project type you want.

Requirements

• An OpenCL (1.0 or 1.1) SDK supporting the gl_sharing_khr extension
• Freeglut libraries is used for windows management
• An OpenGL 3.3 SDK
• On linux the latest glext.h should be sufficient
• On windows you will need glew 1.5.5

Code
To be more concise, almost all message printing and error management has been removed. But remember that error management is essential to know exactly what's happening, and information printing could be interesting.

Also here I use OpenGL, Glut, and OpenCL API directly, it's easier for teaching. But in a real project I guess I should be better to use higher lever libraries.

Main

int main(int argc, char * argv[])
{
  glutInit(&argc, argv);
  glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE | GLUT_ALPHA | GLUT_DEPTH | GLUT_STENCIL);

  glutCreateWindow("clgl_sharing");
  glutSetOption( GLUT_ACTION_ON_WINDOW_CLOSE, GLUT_ACTION_GLUTMAINLOOP_RETURNS );
   
  glutDisplayFunc( display );
  glutReshapeFunc( reshape );
  glutIdleFunc( idle );

  init_gl();
  glFinish(); //Wait until all opengl commands are processed
  init_cl();

  glutMainLoop();

  destroy_gl();
  destroy_cl();
}

Global data

size_t const TEXTURE_WIDTH = 1024;
size_t const TEXTURE_HEIGHT = 1024;

OpenGL data

namespace gl
{
  GLuint positionBuffer = 0;
  GLuint texcoordBuffer = 0;
  GLuint texture = 0;
  GLuint vertexShader = 0;
  GLuint fragmentShader = 0;
  GLuint program = 0;
  GLuint vertexArray = 0;
  GLuint sampler = 0;
  
  float const positionData[] = { -1,-1,0, -1,1,0, 1,1,0, 1,-1,0 };
  float const texcoordData[] = { 0,0,   0,1,   1,1,   1,0   };
  
  char const * const vertexShaderSrc = "\
    #version 330 core \n\
    \n\
    in Vertex{ \n\
      vec3 position; \n\
      vec2 texcoord; \n\
    } vertex; \n\
    \n\
    out vec2 texcoord; \n\
    \n\
    void main() \n\
    { \n\
      texcoord = vertex.texcoord; \n\
      gl_Position = vec4(vertex.position, 1); \n\
    } \n\
  ";

  char const * const fragmentShaderSrc = "\
    #version 330 core \n\
    \n\
    uniform sampler2D texture; \n\
    \n\
    in vec2 texcoord; \n\
    \n\
    out vec4 fragColor; \n\
    \n\
    void main() \n\
    { \n\
      fragColor = texture2D( texture, texcoord ); \n\
    } \n\
  ";
  
  GLuint positionAttrib;
  GLuint texcoordAttrib;
  GLuint samplerLocation;
}

OpenCL data

namespace cl
{
  cl_device_type const DEVICE_TYPE = CL_DEVICE_TYPE_GPU;
  cl_platform_id platform = 0;
  cl_device_id device = 0;
  cl_context context = 0;
  cl_mem texture = 0;
  cl_command_queue command_queue = 0;
  cl_program program = 0;
  cl_kernel kernel = 0;
  
  char const * programSrc = "\
   float m1_1_to_0_1( float val ) \n\
  { \n\
    return (val + 1.f ) / 2.f; \n\
  } \n\
  \n\
  __kernel void writeTexture( __write_only image2d_t image, float t ) \n\
  { \n\
    int2 coordi = (int2)( get_global_id(0), get_global_id(1) ); \n\
    float2 coordf = convert_float2(coordi) / (float2)( get_global_size(0), get_global_size(1) ); \n\
    float4 color = (float4)( \n\
      m1_1_to_0_1 ( cos( (coordf.x+0.5f) * 3.14f * 10.f + t ) * sin( coordf.y * 3.14f * 10.f + t ) * tan( coordf.y * coordf.x * 3.14f * 10.f + t ) ), \n\
     1.f -  m1_1_to_0_1 ( cos( (coordf.x+0.5f) * 3.14f * 10.f + t ) * sin( coordf.y * 3.14f * 10.f + t ) * tan( coordf.y * coordf.x * 3.14f * 10.f + t ) ), \n\
      0.f, \n\
      1.f \n\
    ); \n\
    \n\
    write_imagef( \n\
      image, \n\
      coordi, \n\
      color \n\
    );\n\
  } \n\
  ";
  
  char const * build_options = "";
  
  size_t const GLOBAL_WORK_SIZE[2] = {TEXTURE_WIDTH, TEXTURE_HEIGHT};
}

OpenGL initialization

void init_gl()
{
// Buffers
  glGenBuffers(1, &gl::positionBuffer);
  glGenBuffers(1, &gl::texcoordBuffer);
  
  glBindBuffer( GL_ARRAY_BUFFER, gl::positionBuffer );
  glBufferData( GL_ARRAY_BUFFER, sizeof(gl::positionData), gl::positionData, GL_STATIC_DRAW );
  
  glBindBuffer( GL_ARRAY_BUFFER, gl::texcoordBuffer );
  glBufferData( GL_ARRAY_BUFFER, sizeof(gl::texcoordData), gl::texcoordData, GL_STATIC_DRAW );
 
  glBindBuffer( GL_ARRAY_BUFFER, 0 );
  
// Texture
  glActiveTexture( GL_TEXTURE0 );

  glGenTextures(1, &gl::texture);
  glBindTexture( GL_TEXTURE_2D, gl::texture );
  
  GLfloat * textureData = new GLfloat[TEXTURE_WIDTH * TEXTURE_HEIGHT];
  glTexImage2D( GL_TEXTURE_2D, 0, GL_RGBA, TEXTURE_WIDTH, TEXTURE_HEIGHT, 0, GL_RED, GL_FLOAT, textureData );
  delete textureData;
  glGenerateMipmap( GL_TEXTURE_2D );
  
  glBindTexture( GL_TEXTURE_2D, 0 );
  
// Sampler
  glGenSamplers(1, &gl::sampler );
  glSamplerParameterf( gl::sampler, GL_TEXTURE_MIN_FILTER, GL_NEAREST );
  glSamplerParameterf( gl::sampler, GL_TEXTURE_MAG_FILTER, GL_NEAREST );  
  
// Shaders
  gl::vertexShader = create_shader( GL_VERTEX_SHADER, gl::vertexShaderSrc );
  gl::fragmentShader = create_shader( GL_FRAGMENT_SHADER, gl::fragmentShaderSrc );
  
  gl::program = glCreateProgram();
  
  glAttachShader( gl::program, gl::vertexShader );
  glAttachShader( gl::program, gl::fragmentShader );
  
  glLinkProgram( gl::program );
  
  GLint status;
  glGetProgramiv( gl::program, GL_LINK_STATUS, &status ); 
  
  if( status == GL_FALSE)
  {
    GLint length;
    glGetProgramiv( gl::program, GL_INFO_LOG_LENGTH, &length );
    
    GLchar * infoLog = new GLchar[length];
    
    glGetProgramInfoLog( gl::program, length, NULL, infoLog);

    std::cerr << "shader link log : " << infoLog << std::endl;
  }
  
//Attribute
  gl::positionAttrib = glGetAttribLocation( gl::program, "Vertex.position" );
  gl::texcoordAttrib = glGetAttribLocation( gl::program, "Vertex.texcoord" );
  
// Uniform
  gl::samplerLocation = glGetUniformLocation( gl::program, "texture" );
  
  glUseProgram( gl::program );
  
  glUniform1i( gl::samplerLocation, 0 );
  
  glUseProgram( 0 );
  
// Vertex array
  glGenVertexArrays( 1, &gl::vertexArray );
   
  glBindVertexArray( gl::vertexArray );
  
  glBindBuffer( GL_ARRAY_BUFFER, gl::positionBuffer );
  glVertexAttribPointer( gl::positionAttrib, 3, GL_FLOAT, false, 0, 0);
  
  glBindBuffer( GL_ARRAY_BUFFER, gl::texcoordBuffer );
  glVertexAttribPointer( gl::texcoordAttrib, 2, GL_FLOAT, false, 0, 0);
  
  glEnableVertexAttribArray( gl::positionAttrib );
  glEnableVertexAttribArray( gl::texcoordAttrib );
  
  glBindBuffer( GL_ARRAY_BUFFER, 0 );
  
  glBindVertexArray( 0 );
  
  glFinish();
}

OpenCL platform/device filtering
We only want one platform/device supporting gl_sharing_khr.

bool use_this_cl_platform(cl_platform_id const & platform)
{
  size_t extensions_size;
  char * extensions;
  
  clGetPlatformInfo( platform, CL_PLATFORM_EXTENSIONS, 0, NULL, &extensions_size);
  extensions = new char[extensions_size];
  clGetPlatformInfo( platform, CL_PLATFORM_EXTENSIONS, extensions_size, extensions, NULL);
  
  std::string ext = extensions;
  bool use_this_platform = ext.find("cl_khr_gl_sharing") != std::string::npos;
  
  delete[] extensions;
  
  return use_this_platform;
}

bool use_this_cl_device(cl_device_id const & device)
{
  size_t exensions_size;
  char * extensions;

  clGetDeviceInfo( device, CL_DEVICE_EXTENSIONS, 0, NULL, &exensions_size);
  extensions = new char[exensions_size];
  clGetDeviceInfo( device, CL_DEVICE_EXTENSIONS, exensions_size, extensions, NULL);
  
  std::string ext = extensions;
  bool use_this_device = ext.find("cl_khr_gl_sharing") != std::string::npos;
  
  delete[] extensions;
  
  return use_this_device;
}

OpenCL initialization

void init_cl()
{
//Platforms
  cl_uint num_platforms;
  cl_platform_id * platforms;
    
  clGetPlatformIDs( 0, NULL, &num_platforms);
  
  platforms = new cl_platform_id[num_platforms];
   
  clGetPlatformIDs( num_platforms, platforms, NULL);
  
  for(size_t p=0; p<num_platforms; ++p)
  {
    bool use_this_platform = false;
    if( cl::platform == 0 && use_this_cl_platform( platforms[p] ) )
    {
      cl::platform = platforms[p];
      use_this_platform = true;
    }
    
//Devices
    cl_uint num_devices;
    cl_device_id *devices;
    clGetDeviceIDs( platforms[p], cl::DEVICE_TYPE, 0, NULL, &num_devices);
    devices = new cl_device_id[num_devices];
    clGetDeviceIDs( platforms[p], cl::DEVICE_TYPE, num_devices, devices, NULL);
    
    for(size_t d=0; d<num_devices; ++d)
    {    
      bool use_this_device = false;
      if( use_this_platform && cl::device == 0 && use_this_cl_device( devices[d] ) )
      {
         cl::device = devices[d];
         use_this_device = true;
      }
    }
    
    delete[] devices;
  }
  
  delete[] platforms;

  if( cl::platform == 0 )
  {
    std::cerr << "no cl platform found with  cl_khr_gl_sharing extension" << std::endl;
    exit( 1 );
  }
  
  if( cl::device == 0 )
  {
    std::cerr << "no cl device found with  cl_khr_gl_sharing extension" << std::endl;
    exit( 1 );
  }
  
//Context
#ifdef BUILD_UNIX
  Display * display = XOpenDisplay(NULL);
#endif

  cl_context_properties properties[] =
  {
    CL_CONTEXT_PLATFORM, reinterpret_cast<cl_context_properties>(cl::platform),
#ifdef BUILD_UNIX
    CL_GL_CONTEXT_KHR, reinterpret_cast<cl_context_properties>(glXGetCurrentContext()),
    CL_GLX_DISPLAY_KHR, reinterpret_cast<cl_context_properties>(display),
#endif
#ifdef BUILD_WINDOWS
    CL_GL_CONTEXT_KHR, reinterpret_cast<cl_context_properties>( wglGetCurrentContext() ),
    CL_WGL_HDC_KHR, reinterpret_cast<cl_context_properties>( wglGetCurrentDC() ),
#endif
    0
  };
  
  cl::context = clCreateContext( properties, 1, &cl::device, NULL, NULL, NULL);

#ifdef BUILD_UNIX
  XCloseDisplay(display);
#endif

//Texture 2D
  cl::texture = clCreateFromGLTexture2D ( cl::context, CL_MEM_WRITE_ONLY, GL_TEXTURE_2D, 0, gl::texture, NULL);
  
//Command queue
  cl::command_queue = clCreateCommandQueue( cl::context, cl::device, 0, NULL); 
   
//Program
  cl::program = clCreateProgramWithSource  ( cl::context, 1, &cl::programSrc, NULL, NULL);
    
  clBuildProgram ( cl::program, 1, &cl::device, cl::build_options, NULL, NULL );

  cl_build_status buildStatus;
  
  clGetProgramBuildInfo ( cl::program, cl::device, CL_PROGRAM_BUILD_STATUS, sizeof( cl_build_status ), &buildStatus, NULL);
  
  if( buildStatus != CL_BUILD_SUCCESS )
  {
    std::cerr << "Build status : " << buildStatus << std::endl;
    char * log = NULL;
    size_t logSize;
    clGetProgramBuildInfo ( cl::program, cl::device, CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize);
    log = new char[logSize];
    clGetProgramBuildInfo ( cl::program, cl::device, CL_PROGRAM_BUILD_LOG, logSize, log, NULL);
    
    std::cerr << "Build log : " << log << std::endl;
    
    delete[] log;
  }
   
//Kernel
  cl::kernel = clCreateKernel( cl::program, "writeTexture", NULL);

  clSetKernelArg( cl::kernel, 0, sizeof( cl_mem), &cl::texture );
}

GPU resources destruction
It's really important to destroy data before exiting, otherwise the easyier way to free space on the GPU is to reboot the system.

void destroy_gl()
{
  glDeleteBuffers( 1, &gl::positionBuffer );
  glDeleteBuffers( 1, &gl::texcoordBuffer );
  glDeleteTextures( 1, &gl::texture );

  glDeleteShader( gl::vertexShader );
  glDeleteShader( gl::fragmentShader );
  glDeleteProgram( gl::program );
  
  glDeleteVertexArrays( 1, &gl::vertexArray );
 
  glDeleteSamplers( 1, &gl::sampler );
}

void destroy_cl()
{
  clReleaseMemObject( cl::texture );
  clReleaseKernel( cl::kernel );
  clReleaseProgram( cl::program );
  clReleaseCommandQueue( cl::command_queue );
  clReleaseContext( cl::context );
}

Display

void display()
{
  static float t = 0; // a "time" variable
  t += 0.01f;

//OPENCL
//Acquire shared objects
  clEnqueueAcquireGLObjects ( cl::command_queue, 1, &cl::texture, 0, NULL, NULL );

  clSetKernelArg( cl::kernel, 1, sizeof( cl_float ), &t );

//Run OpenCL
  clEnqueueNDRangeKernel(
    cl::command_queue,
    cl::kernel,
    2,
    NULL, cl::GLOBAL_WORK_SIZE, NULL,
    0, NULL, NULL
  );
 
//Release shared Objects
  clEnqueueReleaseGLObjects ( cl::command_queue, 1, &cl::texture, 0, NULL, NULL );
  
//Wait until all OpenCL commands are processed
  clFinish( cl::command_queue );

//OPENGL
  glClear( GL_COLOR_BUFFER_BIT );
  glUseProgram( gl::program );
  
  glActiveTexture( GL_TEXTURE0 + 0 );
  glBindTexture( GL_TEXTURE_2D, gl::texture );
  glBindSampler( 0, gl::sampler );
  
  glBindVertexArray( gl::vertexArray );
  glDrawArrays(GL_TRIANGLE_FAN, 0, 4);
  glBindVertexArray( 0 );
  
  glBindTexture( GL_TEXTURE_2D, 0 );
  
  glutSwapBuffers(); //After this all OpenGL commands are processed
}

Other functions

void idle()
{
  glutPostRedisplay();
}
void reshape(int width, int height)
{
  glViewport(0, 0, width, height);
}

Testing

• on kubuntu 10.4 64bits with a nVidia 9800gt (256.35 drivers).
• on windows vista 64 bits with nVidia gtx275