I've been fiddling with OpenCL computations for different projects in the past. One such project is the implementation of a specific triangulation algorithm for a structured light measurement system. The other one is a parallelized implementation of a 3d binary thinning algorithm, which a student of mine is writing his bachelor thesis about (in time, i hope he will be willing to publish some small description about it online..). We are both using python and the excellent module pyOpenCL by Andreas Klöckner.
As fast as OpenCL might be, to visualize the result we generally have to copy the gpu-buffered data back to the main memory and let frameworks such as Enthoughts mayavi or matplotlib or others draw it. For the end result of a computation this might be a viable solution but if we wanted to visualize how the computation itself progresses, the visualization will ultimately get in the way.
Since the data we need resides in the memory of the graphic card (usually i would use the gpu even though there are cpu-based opencl libraries by now), i guessed the graphics card should be able to use them for the visualization without even bothering the main memory or the cpu for that matter. This article by Ian Johnson who is working on bringing OpenCL acceleration to blenders particle system - or rather create a new particle system implementation using OpenCL - got me started on the interaction between OpenGL and OpenCL, so here it goes. The complete source code will be cloneable via git in time. I would like to describe the fundmental functions i'm using here and will omit any extra code that is not needed per se.
This is how my module import section looks like:
-
from OpenGL import GL as gl, GLU as glu, GLUT as glut
-
from OpenGL.arrays import vbo
-
import numpy as np
-
import pyopencl as cl
-
import scipy
-
import os, sys
Starting off, we have to setup the glut-interface in which the OpenGL-buffer will be shown. This can be replaced by any GUI-framework, but i guess glut is very simple to setup.
-
glut.glutInit()
-
glut.glutInitDisplayMode(glut.GLUT_RGBA | glut.GLUT_DOUBLE | glut.GLUT_DEPTH)
-
glut.glutInitWindowSize(self.width, self.height)
-
glut.glutInitWindowPosition(0, 0)
-
self.win = glut.glutCreateWindow("Testing...")
This initializes glut, sets up the display mode to support RGBA with two buffers that can be swapped (GLUT_DOUBLE, it also works with only one buffer, but this seems more stable) and enables the depth field. After that we setup the window size, position and then create it with the title "Testing..."
.
After that we need to setup some callbacks for glut. The most important ones are
-
glut.glutDisplayFunc(self.draw)
-
glut.glutTimerFunc(10, self.timer, 10)
The first one sets up the function which draws the content of the glut window, the second one is called every 10 msecs and will be responsible for initiating the redraw.
Now we need to setup our scene. But before that we need some data, that we want to display. I chose a voxel field, gave it some color based on its position. Also, i need a numpy array with random values. Based on the random value of a voxel, i would like to let it disappear or reappear.
-
self.randarr = np.random.rand(size, size, size)
-
self.randarr = np.require(self.randarr, 'f')
-
arr = [(float(x)/(size-1)-.5,float(y)/(size-1)-.5,float(z)/(size-1)-.5,1.0) for x in xrange(size) for y in xrange(size) for z in xrange(size)]
-
self.arr = np.require(arr, 'f')
-
colarr = [(.5-float(x)/size+.2,float(y)/size-.5+.2,float(z)/size-.5+.2,1.0) for x in xrange(size) for y in xrange(size) for z in xrange(size)]
-
self.colarr = np.require(colarr, 'f')
-
self.arrvbo = vbo.VBO(data=self.arr, usage=gl.GL_DYNAMIC_DRAW, target=gl.GL_ARRAY_BUFFER)
-
self.arrvbo.bind()
-
self.colvbo = vbo.VBO(data=self.colarr, usage=gl.GL_DYNAMIC_DRAW, target=gl.GL_ARRAY_BUFFER)
-
self.colvbo.bind()
After creating the arrays in python they are transfered to OpenGL-memory via "Vertex Buffer Object"s (vbo). I'm still not sure, what the subsequent bind does but it is not optional anyway. It is important to mind the size of the vectors used. Here, i've used 4 dimensions for both the position (arr) and the color (colarr). The color is thereby an RGBA-color and it seems (though i'm not sure), that the fourth dimension of the position will be used as a scaling factor, so i set it to 1.0.
We can go directly to the drawing function:
-
def draw(self):
-
# execute the OpenCL computation
-
self.execute()
-
# clear the gl buffers
-
gl.glClear(gl.GL_COLOR_BUFFER_BIT | gl.GL_DEPTH_BUFFER_BIT)
-
# tell opengl to use the vertex-array and the color-array
-
gl.glEnableClientState(gl.GL_VERTEX_ARRAY)
-
gl.glEnableClientState(gl.GL_COLOR_ARRAY)
-
# tell opengl where to find the arrays
-
self.arrvbo.bind()
-
gl.glVertexPointer(4, gl.GL_FLOAT, 0, self.arrvbo)
-
self.colvbo.bind()
-
gl.glColorPointer(4, gl.GL_FLOAT, 0, self.colvbo)
-
# tell gl to draw the arrays (with offset and number of points)
-
gl.glDrawArrays(gl.GL_POINTS, 0, self.size**3)
-
# disable the state and unbind the buffers
-
gl.glDisableClientState(gl.GL_VERTEX_ARRAY)
-
gl.glDisableClientState(gl.GL_COLOR_ARRAY)
-
self.arrvbo.unbind()
-
self.colvbo.unbind()
-
# we use double buffer, so we have to swap the active buffer
-
# and get the new buffer to the foreground
-
glut.glutSwapBuffers()
To start the main loop of glut, we still have to call glut.glutMainLoop().
Now to the nice part: The interaction between OpenGL and OpenCL means, that we have to manipulate OpenGL data within OpenCL. This is a little bit of a illogic paradigm, since it should be OpenCL data that we want to visualize. Well, but once you've moved beyond this slight irregularity, it's easy-peasy:
First, we setup our OpenCL environment: the context and the queue.
-
plats = cl.get_platforms()
-
from pyopencl.tools import get_gl_sharing_context_properties
-
self.ctx = cl.Context(properties=[
-
(cl.context_properties.PLATFORM, plats[0])]
-
+ get_gl_sharing_context_properties(), devices=None)
-
self.queue = cl.CommandQueue(self.ctx)
We then load and build our program...
-
f = open(filename, 'r')
-
fstr = f.read()
-
self.program = cl.Program(self.ctx, fstr).build()
... and bind our OpenGL buffers to OpenCL pointers:
-
mf = cl.mem_flags
-
self.arrvbo.bind()
-
self.arr_cl = cl.GLBuffer(self.ctx, mf.READ_WRITE, int(self.arrvbo.buffers[0]))
-
self.colvbo.bind()
-
self.col_cl = cl.GLBuffer(self.ctx, mf.READ_WRITE, int(self.colvbo.buffers[0]))
-
self.randarr_cl = cl.Buffer(self.ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=self.randarr)
-
self.arrvbo.unbind()
-
self.colvbo.unbind()
-
self.queue.finish()
-
# set up the list of GL objects to share with opencl
-
self.gl_objects = [self.arr_cl, self.col_cl]
We need the list of gl_objects because they have to be acquired each time we want to use them in OpenCL and released afterwards. I'm not sure, but i think this way we ensure that OpenCL and OpenGL don't hamper with the data at the same time. With cl.Buffer we set up the randarr array within OpenCL memory. The OpenCL calculation is done in the aforementioned self.execute():
-
def execute(self):
-
self.randarr = np.random.rand(self.size, self.size, self.size)
-
self.randarr = np.require(self.randarr, 'f')
-
cl.enqueue_copy(self.queue, self.randarr_cl, self.randarr)
-
cl.enqueue_acquire_gl_objects(self.queue, self.gl_objects).wait()
-
global_size = (self.size**3,)
-
local_size = None
-
kernelargs = (self.arr_cl,
-
self.col_cl,
-
self.randarr_cl,
-
self.dt)
-
self.program.mykernel(self.queue, global_size, local_size, *(kernelargs))
-
cl.enqueue_release_gl_objects(self.queue, self.gl_objects).wait()
-
self.queue.finish()
The important bit is the function cl.enqueue_acquire_gl_objects and the corresponding cl.enqueue_release_gl_objects.
So now to the cl-code. It's actually fairly simple:
-
__kernel void mykernel(__global float4* pos, __global float4* color, __global float* rand, float dt)
-
{
-
unsigned int i = get_global_id(0);
-
if (rand[i] <.05) { color[i].w = 0; }
-
if (rand[i]> .999) { color[i].w = 1; }
-
}
Et voila: I've used some code snippets i found here but changed the image save procedure to use scipy.misc.imsave instead of the much slower matplotlib.
The result can be seen here:
Video not there? Download it: movie.ogv.ogv
Update: repository
The code can be found here: https://gitorious.org/openclgltest/openclgltest