I will start my blog series with a subseries on technology that I know and come to love best - OpenCL from the Khronos Group. OpenCL aims to provide functional portability across computer systems comprised of general-purpose processors (e.g. CPUs) and special-purpose accelerators (e.g. GPUs), to which I will now and then refer to as accelerated systems.


OpenCL is often criticised for presenting a steep learning curve for software engineers but I don't think that's fair: the verbosity of the OpenCL API is really a consequence of its generality. Once you have written a couple of OpenCL programs, you realise that most of the code running on the host processor is actually boilerplate.


I will sidestep this verbosity of OpenCL altogether by using PyOpenCL - a neat Python module written by Andreas Klöckner. (If you are reading this Andreas, keep up the good work!)


Just install PyOpenCL and NumPy - and you are ready to roll!

In [1]:

import pyopencl as cl

import numpy as np



Trivial example

Suppose we want to create an array containing integers from 0 to 15.

In [2]:
N = 16


That's trivial using NumPy:

In [3]:

np_range = np.arange(N, dtype=np.int32)



array([ 0,  1,  2,  3,  4,  5,  6,  7,  8,  9, 10, 11, 12, 13, 14, 15], dtype=int32)


but our PyOpenCL example will fill in a similar array using OpenCL:

In [4]:

cl_range = np.zeros(N, dtype=np.int32)



array([0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0], dtype=int32)



Gimme some context!

Creating a context could hardly be easier:

In [5]:
In [6]:



Building and running

An OpenCL C program equivalent to np.arange(N) follows:

In [7]:

source = '''

kernel void arange(global int * buffer)


    const size_t gid = get_global_id(0); 

    buffer[gid] = convert_int(gid);




The kernel will be launched as N work-items over a one-dimensional range [0, N-1]. Each work-item will get its unique index gid in the range (that is, an integer between 0 and N-1 inclusive) and write it into argument buffer at offset gid.


Let's build the program:

In [8]:
program = cl.Program(context, source).build()


allocate a memory buffer:

In [9]:
memory_flags = cl.mem_flags.WRITE_ONLY | cl.mem_flags.ALLOC_HOST_PTR
memory = cl.Buffer(context, flags=memory_flags, size=cl_range.nbytes)


launch the kernel:

In [10]:
kernel = program.arange(queue, [N], None, memory)


and copy the results from the buffer to cl_range:

In [11]:

cl.enqueue_copy(queue, cl_range, memory, wait_for=[kernel])



array([ 0,  1,  2,  3,  4,  5,  6,  7,  8,  9, 10, 11, 12, 13, 14, 15], dtype=int32)


Finally, let's confirm that arrays np_range and cl_range match element-wise:

In [12]:
np.all(np_range == cl_range)




Concluding remarks

I am very interested in getting feedback. I am aiming primarily at software engineers already familiar with OpenCL so I can jump straight to more advanced topics. But if you feel you like a little more explanation, just ask me: I will be happy to explain, or point you to some excellent learning resources out there, or indeed consider filling in details in my later posts...


Also, I prepared this post using the wonderful IPython Notebook environment ("Introducing PyOpenCL" as IPython notebook). I haven't seen much in the way of using PyOpenCL and IPython together, so would be very grateful for any links.


In my next post, I start discussing how to optimise OpenCL kernels for the ARM® Mali-T600 and Mali-T700 GPU series.