Applications with PyOpenCL (PIL & PyOpenGL)

I think I wrote it just recently, but it's the 15th day of GPGPU Advent Calendar. You can finish it already, right?

Since PyOpenCL makes it a little easier to program GPGPU, I will show you how to easily create an application using GPGPU using other Python libraries. There are two examples, one using OpenCL's Image object and Python's image processing library PIL, and the other using OpenCL's OpenGL Interoperability and PyOpenGL.

PyOpenCL + PIL

Use pyopencl.Image to work with Image objects in PyOpenCL. That's right. The ImageFormat object gives the representation format of the image. This time, RGBA is 8 bits per pixel, 32 bits each, so it looks like cl.ImageFormat (cl.channel_order.RGBA, cl.channel_type.UNSIGNED_INT8). If you create a numpy ndarray that handles host memory like numpy.empty ((height, width, 4), dtype = numpy.uint8), you can receive pixel information directly from the Image object. If you use PIL, you can easily save the image information that you have in ndarray to a file.

The entire source code is below. By the way, when writing kernel functions in Python docstring with PyOpenCL, use Syntax highlighting for Vim. , OpenCL C language part is also colored, which is convenient.

`polarchecker.py`


# vim: filetype=pyopencl.python
import numpy as np
from PIL import Image
import pyopencl as cl
 
SRC = '''//CL//
__constant const int CHECKER_SIZE = 10;
 
float2 polar(const int2);
int disc(const int2);
 
 
float2 polar(const int2 pos)
{
    const float x = pos.x;
    const float y = pos.y;
    return (float2)(sqrt(hypot(x, y)), atan2(x, y));
}
 
 
int disc(const int2 pos)
{
    const float2 ppos = polar(pos);
    return ((int)(ppos.x) + (int)(ppos.y * CHECKER_SIZE / M_PI_F)) % 2;
}
 
 
__kernel void gen(write_only image2d_t dest)
{
    const int2 pos = (int2)(get_global_id(0), get_global_id(1));
    uint4 pix;
    if (disc(pos)) {
        pix = (uint4)(0, 0, 255, 255);
    } else {
        pix = (uint4)(255, 255, 255, 255);
    }
    write_imageui(dest, pos, pix);
}
'''
 
ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx)
prg = cl.Program(ctx, SRC).build()
 
width = 640
height = 480
fmt = cl.ImageFormat(cl.channel_order.RGBA, cl.channel_type.UNSIGNED_INT8)
buf = cl.Image(ctx, cl.mem_flags.WRITE_ONLY, fmt, shape=(width, height))
 
prg.gen(queue, (width, height), None, buf)
 
result = np.empty((height, width, 4), dtype=np.uint8)
cl.enqueue_copy(queue, result, buf, origin=(0, 0), region=(width, height))
 
image = Image.fromarray(result)
image.save('checker.png')

I think that the kernel function part is written in a way unique to the OpenCL C language. If you look at this area such as OpenCL API 1.2 Reference Card, you may discover various new things. Hmm.

When executed, an image like this will be generated.

結果画像

It's good because the kernel function is processed in pixel units and the visibility of the program is improved.

PyOpenCL + PyOpenGL

If you take PyOpenCL and PyOpenGL together seriously, there are many steps and it is a little difficult to explain, so Sample included with PyOpenCL Will be explained. This program only outputs a sine wave graph with OpenGL, and calculates the vertex position using OpenCL at the time of initialization.

The main part of the program is as follows.

if __name__ == '__main__':
    import sys
    glutInit(sys.argv)
    if len(sys.argv) > 1:
        n_vertices = int(sys.argv[1])
    glutInitWindowSize(800, 160)
    glutInitWindowPosition(0, 0)
    glutCreateWindow('OpenCL/OpenGL Interop Tutorial: Sin Generator')
    glutDisplayFunc(display)
    glutReshapeFunc(reshape)
    initialize()
    glutMainLoop()

Only the contents of initialize are important in OpenCL / OpenGL cooperation, so let's look at them in order.

First, create an OpenCL Context with OpenGL integration enabled.

    platform = cl.get_platforms()[0]

    from pyopencl.tools import get_gl_sharing_context_properties
    import sys
    if sys.platform == "darwin":
        ctx = cl.Context(properties=get_gl_sharing_context_properties(),
                devices=[])
    else:
        # Some OSs prefer clCreateContextFromType, some prefer
        # clCreateContext. Try both.
        try:
            ctx = cl.Context(properties=[
                (cl.context_properties.PLATFORM, platform)]
                + get_gl_sharing_context_properties())
        except:
            ctx = cl.Context(properties=[
                (cl.context_properties.PLATFORM, platform)]
                + get_gl_sharing_context_properties(),
                devices = [platform.get_devices()[0]])

I'm writing something that seems confusing, but in most cases I think that only the part inside the try clause is sufficient. My environment enters the first if condition on Mac, but on the contrary it fails with an error so I commented it out, but I feel like it works normally.

Next, initialize the OpenGL Buffer. The point to pay attention to is getting vbo with glGenBuffers. The value obtained here is necessary to generate the GLBuffer object of PyOpenCL.

    glClearColor(1, 1, 1, 1)
    glColor(0, 0, 1)
    vbo = glGenBuffers(1)
    glBindBuffer(GL_ARRAY_BUFFER, vbo)
    rawGlBufferData(GL_ARRAY_BUFFER, n_vertices * 2 * 4, None, GL_STATIC_DRAW)
    glEnableClientState(GL_VERTEX_ARRAY)
    glVertexPointer(2, GL_FLOAT, 0, None)

The following line is the creation of the GLBuffer object.

    coords_dev = cl.GLBuffer(ctx, cl.mem_flags.READ_WRITE, int(vbo))

Finally, operate the OpenGL Buffer from OpenCL. When operating the OpenGL Buffer on the OpenCL side, sandwich it between enqueue_acquire_gl_objects and enqueue_release_gl_objects.

    cl.enqueue_acquire_gl_objects(queue, [coords_dev])
    prog.generate_sin(queue, (n_vertices,), None, coords_dev)
    cl.enqueue_release_gl_objects(queue, [coords_dev])

That's all the basics. When you run the entire program, the following screen will be displayed. 結果画面

It will be much easier than writing in C language. However, some knowledge of OpenGL may be required.

Demos that move a little more are this area and [this area](http://enja.org/2011/03/22/adventures-in-pyopencl-part- It's in 2-particles-with-pyopengl /).