glitter Example: OpenCL

Summary

This program will open a GLUT window and render an animated particle simulation. The particle simulation is computed in OpenCL, and the results are written directly to an OpenGL array buffer.

Front matter

Module docstring

The module docstring is used as a description of this example in the generated documentation:

"""Basic OpenGL/OpenCL interoperability example.

@author: Stephan Wenger
@date: 2012-02-29
"""

Imports

We use numpy for creating the particle data:

import numpy

PyOpenCL (documentation, PyPi, download) is a convenient Python wrapper for OpenCL:

import pyopencl as cl

For creating an OpenCL context that shares data with OpenGL, we need get_gl_sharing_context_properties:

from pyopencl.tools import get_gl_sharing_context_properties

We can usually import classes and functions contained in glitter submodules directly from glitter:

from glitter import ArrayBuffer, VertexArray, get_default_program

Modules with external dependencies other than numpy, such as platform dependent parts like methods for the generation of an OpenGL context, however, have to be imported from their respective submodules:

from glitter.contexts.glut import GlutWindow, main_loop

OpenCL

OpenCL kernel

The OpenCL kernel (adapted from a CMSoft tutorial) computes a single iteration of a simple Eulerian particle simulation. Kernel code can be loaded from a file or defined inline as a Python string:

kernel_code = """
__kernel void animate(__global float4* positions,
                      __global float4* colors,
                      __global float4* velocities,
                      __global float4* initial_positions,
                      __global float4* initial_velocities,
                      float dt) {
    unsigned int i = get_global_id(0);
    float4 p = positions[i];
    float4 v = velocities[i];
    float life = velocities[i].w;

    life -= dt;
    if (life <= 0.0) {
        p = initial_positions[i];
        v = initial_velocities[i];
        life = 1.0;    
    }

    v.z -= 9.8 * dt;
    p.xyz += v.xyz * dt;
    v.w = life;

    positions[i] = p;
    velocities[i] = v;
    colors[i].w = life;
}
"""

OpenCL interaction

The CLCode class will comprise all OpenCL interaction:

class CLCode(object):

Initialization

The constructor receives OpenGL buffers for positions and colors as well as a numpy array of velocities and a timestep for the simulation:

    def __init__(self, gl_positions, gl_colors, velocities, dt=0.001):

First, we have to initialize the OpenCL context. That means we have to get a list of available platforms and select one:

        platform = cl.get_platforms()[0]

Then, we can create a context. Passing get_gl_sharing_context_properties() as a property ensures that we share state with the active OpenGL context:

        self.ctx = cl.Context(properties=[(cl.context_properties.PLATFORM, platform)] +
                get_gl_sharing_context_properties(), devices=None)

A command queue is necessary for serializing OpenCL commands:

        self.queue = cl.CommandQueue(self.ctx)

Finally, we can compile the kernel:

        self.program = cl.Program(self.ctx, kernel_code).build()

The constructor parameters are stored for later use:

        self.gl_positions = gl_positions
        self.gl_colors = gl_colors
        self.velocities = velocities

The dt value will later be passed to an OpenCL kernel as a 32-bit float. We therefore wrap it in a numpy float32 object:

        self.dt = numpy.float32(dt)

Next, we generate OpenCL buffers. The positions and colors are contained in OpenGL buffers, which we wrap in PyOpenCL's GLBuffer<(/code> class:


        self.cl_positions = cl.GLBuffer(self.ctx, cl.mem_flags.READ_WRITE, self.gl_positions._id)
        self.cl_colors = cl.GLBuffer(self.ctx, cl.mem_flags.READ_WRITE, self.gl_colors._id)


 Note how we had to extract the _ids from the
  ArrayBuffer objects. In pure glitter code, you
  should never (have to) access this value; however for interaction
  with other OpenGL-related libraries, this cannot always be avoided.

 The velocities are given as a numpy array, which is simply uploaded
  into a new OpenCL Buffer object along with the initial
  values of the positions and colors:

        self.cl_velocities = cl.Buffer(self.ctx, cl.mem_flags.READ_ONLY |
                cl.mem_flags.COPY_HOST_PTR, hostbuf=velocities)
        self.cl_initial_positions = cl.Buffer(self.ctx, cl.mem_flags.READ_ONLY |
                cl.mem_flags.COPY_HOST_PTR, hostbuf=self.gl_positions.data)
        self.cl_initial_velocities = cl.Buffer(self.ctx, cl.mem_flags.READ_ONLY |
                cl.mem_flags.COPY_HOST_PTR, hostbuf=self.velocities)


 
Execution
  The execute function executes the OpenCL kernel several
  times in a row:

    def execute(self, sub_intervals):


 First, we have to make sure that OpenGL is done using the buffer objects:

        cl.enqueue_acquire_gl_objects(self.queue, [self.cl_positions, self.cl_colors])


 Now, we can safely call the kernel. Its arguments are buffer objects:

        args = (self.cl_positions, self.cl_colors, self.cl_velocities,
                self.cl_initial_positions, self.cl_initial_velocities, self.dt)


 The kernel will be executed several times with a small step size.
  This increases the accuracy with respect to a single step with a
  large step size. However, it is not necessary to display all the
  intermediate results.

        for i in xrange(0, sub_intervals):


 In each step, the animate kernel function is called.
  Its arguments are the queue object that schedules its execution,
  the global and local block sizes, and any arguments that will be
  passed to the actual kernel.

            self.program.animate(self.queue, [len(self.gl_positions)], None, *args)


 Finally, we allow OpenGL to access the buffers again:

        cl.enqueue_release_gl_objects(self.queue, [self.cl_positions, self.cl_colors])


 
Main class

 We wrap all the OpenGL interaction in a class. The class will contain an
  __init__() method to set up all OpenGL objects, any required
  callback methods, as well as a run() method to trigger execution
  of the GLUT main loop.

class OpenCLExample(object):


 
Initialization

 When a OpenCLExample instance is created, we need to
  initialize a few OpenGL objects.

    def __init__(self):


 First, we create a window; this also creates an OpenGL context.

        self.window = GlutWindow(double=True, alpha=True, depth=True)


 Then, we set the GLUT display and keyboard callback functions which
  will be defined later.

        self.window.display_callback = self.display
        self.window.keyboard_callback = self.keyboard


 Here, we generate numpy arrays to hold the positions, colors, and
  velocities of the particles:

        num = 200000
        positions = numpy.empty((num, 4), dtype=numpy.float32)
        colors = numpy.empty((num, 4), dtype=numpy.float32)
        velocities = numpy.empty((num, 4), dtype=numpy.float32)


 So far, the array contents are undefined. We have to initialize them with meaningful values:

        positions[:, 0] = numpy.sin(numpy.arange(0, num) * 2 * numpy.pi / num) * (numpy.random.random_sample((num,)) / 3 + 0.2)
        positions[:, 1] = numpy.cos(numpy.arange(0, num) * 2 * numpy.pi / num) * (numpy.random.random_sample((num,)) / 3 + 0.2)
        positions[:, 2:] = 0, 1
        colors[:] = 0, 1, 0, 1
        velocities[:, :2] = 2 * positions[:, :2]
        velocities[:, 2] = 3
        velocities[:, 3] = numpy.random.random_sample((num,))


 Instead of simply generating a vertex array from the position and color
  data, we first generate array buffers for them:

        gl_positions = ArrayBuffer(data=positions, usage="DYNAMIC_DRAW")
        gl_colors = ArrayBuffer(data=colors, usage="DYNAMIC_DRAW")


 These array buffers will later also be used by OpenCL. We do not need to
  wrap velocities in this way, as it will only be used by
  OpenCL and can be wrapped in an OpenCL buffer directly.

 We now create a vertex array that will pass the position and color data
  to the shader. The vertex array constructor accepts
  ArrayBuffer instances:

        self.vao = VertexArray(gl_positions, gl_colors)


 In the OpenGL core profile, there is no such thing as a "standard pipeline"
  any more. We use the minimalistic defaultpipeline from the
  glitter.convenience module to create a shader program instead:

        self.shader = get_default_program()
        


 Here, we create the CLCode object that manages OpenCL
  interaction. It is passed the OpenGL buffer objects as well as a numpy
  array of velocities.

        self.clcode = CLCode(gl_positions, gl_colors, velocities)


 
Callback functions

 
Display function

 Here we define the display function. It will be called by GLUT whenever the
  screen has to be redrawn.

    def display(self):


 First we clear the default framebuffer:

        self.window.clear()
        


 To draw the vertex array, we use:

        self.vao.draw()


 After all rendering commands have been issued, we swap the back buffer to
  the front, making the rendered image visible all at once:

        self.window.swap_buffers()


 
Timer function

 The animation is controlled by a GLUT timer. The timer callback animates the
  particle system, schedules the next timer event, and causes a screen redraw:

    def timer(self):


 We first tell an instance of the CLCode class to execute the
  OpenCL kernel:

        self.clcode.execute(10)


 The following line schedules the next timer event to execute after one millisecond.

        self.window.add_timer(1, self.timer)


 Finally, we tell GLUT to redraw the screen.

        self.window.post_redisplay()


 
Keyboard function

 To further illustrate the concept of GLUT callbacks, here's a keyboard
  handler that will simply make the program exit when any key is pressed:

    def keyboard(self, key, x, y):
        raise SystemExit


 
Running

 We will call the run() method later to run the OpenGL code.

    def run(self):


 To start the animation, we call the timer once; all subsequent timer
  calls will be scheduled by the timer function itself.

        self.timer()


 The default program is bound by using a with statement. At
  the same time, we can pass in additional uniform variables, such as the
  modelview matrix:

        with self.shader(modelview_matrix=((1, 0, 0, 0), (0, 0, 1, 0), (0, 1, 0, 0), (0, 0, 0, 2))):


 With the shader bound, we enter the GLUT main loop.

            main_loop()


 When the main loop exits, control is handed back to the script,
  unless SystemExit has been raised by the keyboard
  handler.

 
Main section

 Finally, if this program is being run from the command line, we instanciate
  the main class and run it.

if __name__ == "__main__":
    OpenCLExample().run()