Test for OpenGL-OpenCL Interoperations

glclinterop.cpp

#if defined(__APPLE__)
#include <OpenGL/CGLCurrent.h>
#include <GLUT/glut.h>
#include <OpenCL/cl_gl_ext.h>
#else
#if defined(_WIN32) || defined(_WIN64)
#include <GL/glew.h>
#else
#include <GL/glxew.h>
#endif
#include <GL/glut.h>
#include <CL/cl_gl.h>
#endif

static const char cl_source[] =
  "__kernel void\n"
  "position(__global float *a)\n"
  "{\n"
    "size_t ix = get_global_id(0);\n"
    "size_t nx = get_global_size(0);\n"
    "size_t iy = get_global_id(1);\n"
    "float x = (float) ix * 2.0f - 1.0f;\n"
    "float y = (float) iy * 2.0f - 1.0f;\n"
    "vstore2((float2) (x, y), iy * nx + ix, a)\n;"
  "}\n"
  "__kernel void\n"
  "coord(__global float *a)\n"
  "{\n"
    "size_t ix = get_global_id(0);\n"
    "size_t nx = get_global_size(0);\n"
    "size_t iy = get_global_id(1);\n"
    "float x = (float) ix;\n"
    "float y = 1.0f - (float) iy;\n"
    "vstore2((float2) (x, y), iy * nx + ix, a);\n"
  "}\n"
  "__kernel void\n"
  "mandelbrot(__write_only image2d_t a)\n"
  "{\n"
    "size_t ix = get_global_id(0);\n"
    "size_t nx = get_global_size(0);\n"
    "size_t iy = get_global_id(1);\n"
    "size_t ny = get_global_size(1);\n"
    "float x = ((float) ix / (float) (nx - 1)) * 4.0f - 2.0f;\n"
    "float y = ((float) iy / (float) (ny - 1)) * 4.0f - 2.0f;\n"
    "float r = x;\n"
    "float i = y;\n"
    "int loop;\n"
    "for (loop = 0; loop < 255 && r * r + i * i <= 4.0f; loop++) {\n"
      "float newr = r * r - i * i + x;\n"
      "float newi = 2.0f * r * i + y;\n"
      "r = newr;\n"
      "i = newi;\n"
    "}\n"
    "float c = (float) loop / 255.0f;\n"
    "write_imagef(a, (int2) (ix, iy), (float4) (c, c, c, 1.0f));\n"
  "}\n";

static GLuint buffers[2], textures[1];

static void
display()
{
  glClear(GL_COLOR_BUFFER_BIT);
  glEnableClientState(GL_VERTEX_ARRAY);
  glEnableClientState(GL_TEXTURE_COORD_ARRAY);
  glBindBuffer(GL_ARRAY_BUFFER, buffers[0]);
  glVertexPointer(2, GL_FLOAT, 0, 0);
  glBindBuffer(GL_ARRAY_BUFFER, buffers[1]);
  glTexCoordPointer(2, GL_FLOAT, 0, 0);
  glBindTexture(GL_TEXTURE_2D, textures[0]);
  glDrawArrays(GL_TRIANGLE_STRIP, 0, 4);
  glFlush();
  glDisableClientState(GL_TEXTURE_COORD_ARRAY);
  glDisableClientState(GL_VERTEX_ARRAY);
}

int
main(int argc, char **argv)
{
  const size_t texwidth = 1024;
  const size_t texheight = 1024;

  glutInit(&argc, argv);
  glutCreateWindow(argv[0]);
#if !defined(__APPLE__)
  glewInit();
#endif

  glEnable(GL_TEXTURE_2D);
  glGenBuffers(2, buffers);
  glBindBuffer(GL_ARRAY_BUFFER, buffers[0]);
  glBufferData(GL_ARRAY_BUFFER, 8 * sizeof(GLfloat), 0, GL_STATIC_DRAW);
  glBindBuffer(GL_ARRAY_BUFFER, buffers[1]);
  glBufferData(GL_ARRAY_BUFFER, 8 * sizeof(GLfloat), 0, GL_STATIC_DRAW);
  glGenTextures(1, textures);
  glBindTexture(GL_TEXTURE_2D, textures[0]);
  glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, texwidth, texheight, 0,
               GL_RGBA, GL_UNSIGNED_BYTE, 0);
  glGenerateMipmap(GL_TEXTURE_2D);

  cl_device_id device = 0;
#if defined(__APPLE__)
  cl_context_properties properties[] = {
    CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE,
    (cl_context_properties) CGLGetShareGroup(CGLGetCurrentContext()),
    0
  };
  cl_context context = clCreateContext(properties, 0, 0, 0, 0, 0);
  size_t apple_devices_size;
  clGetContextInfo(context, CL_CONTEXT_DEVICES,
                   0, 0, &apple_devices_size);
  cl_device_id *apple_devices =
    new cl_device_id[apple_devices_size / sizeof(cl_device_id)];
  clGetContextInfo(context, CL_CONTEXT_DEVICES,
                   apple_devices_size, apple_devices, 0);
  device = apple_devices[0];
  delete[] apple_devices;
#else
  cl_uint num_platforms;
  clGetPlatformIDs(0, 0, &num_platforms);
  cl_platform_id *platforms = new cl_platform_id[num_platforms];
  clGetPlatformIDs(num_platforms, platforms, 0);
#if 1
  clGetGLContextInfoKHR_fn clGetGLContextInfoKHR =
    (clGetGLContextInfoKHR_fn) clGetExtensionFunctionAddress("clGetGLContextInfoKHR");
#endif
  cl_context_properties properties[] = {
    CL_CONTEXT_PLATFORM, 0 /* properties[1] */,
#if defined(_WIN32) || defined(_WIN64)
    CL_GL_CONTEXT_KHR, (cl_context_properties) wglGetCurrentContext(),
    CL_WGL_HDC_KHR, (cl_context_properties) wglGetCurrentDC(),
#else
    CL_GL_CONTEXT_KHR, (cl_context_properties) glXGetCurrentContext(),
    CL_GLX_DISPLAY_KHR, (cl_context_properties) glXGetCurrentDisplay(),
#endif
    0
  };
  for (cl_uint i = 0; i < num_platforms; ++i) {
    properties[1] = (cl_context_properties) platforms[i];
    size_t param_value_size;
    clGetGLContextInfoKHR(properties, CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR,
                          0, 0, &param_value_size);
    if (param_value_size != 0) {
      clGetGLContextInfoKHR(properties, CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR,
                            sizeof(cl_device_id), &device, 0);
      break;
    }
  }
  delete[] platforms;
  if (device == 0) {
    return -1;
  }
  cl_context context = clCreateContext(properties, 1, &device, 0, 0, 0);
#endif
  cl_command_queue queue = clCreateCommandQueue(context, device, 0, 0);
  const char *source = cl_source;
  size_t sourcelen = sizeof(cl_source);
  cl_program program = clCreateProgramWithSource(context,
                                                 1, &source, &sourcelen, 0);
  clBuildProgram(program, 1, &device, 0, 0, 0);
  cl_kernel position = clCreateKernel(program, "position", 0);
  cl_kernel coord = clCreateKernel(program, "coord", 0);
  cl_kernel mandelbrot = clCreateKernel(program, "mandelbrot", 0);

  cl_mem memobjs[3];
  memobjs[0] = clCreateFromGLBuffer(context, CL_MEM_WRITE_ONLY, buffers[0], 0);
  memobjs[1] = clCreateFromGLBuffer(context, CL_MEM_WRITE_ONLY, buffers[1], 0);
  memobjs[2] = clCreateFromGLTexture2D(context, CL_MEM_WRITE_ONLY,
                                       GL_TEXTURE_2D, 0, textures[0], 0);
  clSetKernelArg(position, 0, sizeof(cl_mem), &memobjs[0]);
  clSetKernelArg(coord, 0, sizeof(cl_mem), &memobjs[1]);
  clSetKernelArg(mandelbrot, 0, sizeof(cl_mem), &memobjs[2]);
  glFinish();
  clEnqueueAcquireGLObjects(queue, 3, memobjs, 0, 0, 0);
  const size_t position_work_size[2] = { 2, 2 };
  clEnqueueNDRangeKernel(queue, position,
                         2, 0, position_work_size, 0, 0, 0, 0);
  const size_t coord_work_size[2] = { 2, 2 };
  clEnqueueNDRangeKernel(queue, coord,
                         2, 0, coord_work_size, 0, 0, 0, 0);
  const size_t mandelbrot_work_size[2] = { texwidth, texheight };
  clEnqueueNDRangeKernel(queue, mandelbrot,
                         2, 0, mandelbrot_work_size, 0, 0, 0, 0);
  clEnqueueReleaseGLObjects(queue, 3, memobjs, 0, 0, 0);
  clFinish(queue);

  glBindTexture(GL_TEXTURE_2D, textures[0]);
  glGenerateMipmap(GL_TEXTURE_2D);

  glutDisplayFunc(display);
  glutMainLoop();

  return 0;
}