Skip to content

Instantly share code, notes, and snippets.

@nattoheaven

nattoheaven/glclinterop.cpp

Created February 7, 2015 04:47
Show Gist options
  • Save nattoheaven/f15eeba34e2e78bd902f to your computer and use it in GitHub Desktop.
Save nattoheaven/f15eeba34e2e78bd902f to your computer and use it in GitHub Desktop.
Download ZIP
Test for OpenGL-OpenCL Interoperations
Raw
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;
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment