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!
import pyopencl as cl
import numpy as np
Suppose we want to create an array containing integers from 0 to 15.
That's trivial using NumPy:
np_range = np.arange(N, dtype=np.int32)
np_range
Out[3]:
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:
cl_range = np.zeros(N, dtype=np.int32)
cl_range
array([0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0], dtype=int32)
Creating a context could hardly be easier:
Ditto creating a command queue:
An OpenCL C program equivalent to np.arange(N) follows:
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:
allocate a memory buffer:
launch the kernel:
and copy the results from the buffer to cl_range:
cl.enqueue_copy(queue, cl_range, memory, wait_for=[kernel])
Finally, let's confirm that arrays np_range and cl_range match element-wise:
True
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.
FYI there are also Mali based odroid boards available
The ARM based Samsung Chromebook with a Mali T604 is a good choice. I suspect the upcoming Chromebook 2 will also be a great choice but given the hardware isn't available yet it's too early. The Arndale board is also a good choice.
I like the odroid boards but like the odroid-xu for instance doesn't include OpenCL support for Linux unless you sign an NDA with Imagination.
Is there any affordable ARM Mali GPU/ OpenCL enabled that can be used for development?