OpenCL on GEGL: Results up to now

Hello everyone! I’m glad to show you the results up to now of my GSoC project about adding OpenCL support to the General Graphics Library.

What I’ve done

GEGL has two basic data types:

• GeglTile
• GeglBuffer

A GeglBuffer can be seen as a layer in a image editing tool, they can be translated, cut, duplicated, etc. A final image is a composition of buffers. A buffer is composed by many GeglTiles, which are rectangular regions of pixels with same size, so pixel data like color is stored in tiles. This architecture is very flexible and allows for example that tiles may be stored in the disk, in a network or compacted.

What I want in my project is to be able to process tiles using an OpenCL device, like GPUs or even a multi-core CPU, the solution I implemented is that each tile has two states, the host memory data and a pointer to a OpenCL memory buffer and each one has its revision number which are used for synchronization.

This synchronization is achieved through locks. For example, suppose gegl_buffer_get is called for a buffer which tiles are being processed in the GPU. This function asks for buffer data to be copied to a pointer, as such, each buffer’s tile is going to be locked for reading, this locking process will verify the revision numbers and move data from the GPU to the CPU accordingly. The picture below illustrate this architecture:

An Example of Use

I’ll show an example of use of gegl buffer iterators to implement a Brightness-Contrast filter using OpenCL.

First, we define the OpenCL kernel that will be executed for each tile:

    const char* kernel_source[] =
    {
    "sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |              \n",
    "                    CLK_ADDRESS_NONE            |              \n",
    "                    CLK_FILTER_NEAREST;                        \n",
    "__kernel void kernel_bc(__read_only  image2d_t in,             \n",
    "                        __write_only image2d_t out,            \n",
    "                         float brightness,                     \n",
    "                         float contrast)                       \n",
    "{                                                              \n",
    "  int2 gid = (int2)(get_global_id(0), get_global_id(1));       \n",
    "  float4 in_v  = read_imagef(in, sampler, gid);                \n",
    "  float4 out_v;                                                \n",
    "  out_v.xyz = (in_v.xyz - 0.5f) * contrast + brightness + 0.5f;\n",
    "  out_v.w   =  in_v.w;                                         \n",
    "  write_imagef(out, gid, out_v);                               \n",
    "}                                                              \n",
}

So, each tile is a OpenCL image2d_t type which can be read-only or write-only and be must be fetched through a sampler.

Now let’s see the iterator code:

    i = gegl_buffer_iterator_new (buffer_write, NULL, NULL, GEGL_BUFFER_CL_WRITE);
    index = gegl_buffer_iterator_add (i, buffer_read, NULL, NULL, GEGL_BUFFER_CL_READ);
    while (gegl_buffer_iterator_next (i))
      {
        GeglClTexture *in_tex  = i->cl_data[index];
        GeglClTexture *out_tex = i->cl_data[0];
        size_t global_worksize[2] = {i->roi[0].width, i->roi[0].height};

        CL_SAFE_CALL( errcode = gegl_clSetKernelArg(kernel, 0, sizeof(cl_mem),   (void*)&in_tex->data) );
        CL_SAFE_CALL( errcode = gegl_clSetKernelArg(kernel, 1, sizeof(cl_mem),   (void*)&out_tex->data) );
        CL_SAFE_CALL( errcode = gegl_clSetKernelArg(kernel, 2, sizeof(cl_float), (void*)&brightness) );
        CL_SAFE_CALL( errcode = gegl_clSetKernelArg(kernel, 3, sizeof(cl_float), (void*)&contrast) );

        CL_SAFE_CALL( errcode = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(), kernel, 2,
                                                      NULL, global_worksize, NULL,
                                                      0, NULL, NULL) );
        CL_SAFE_CALL( errcode = gegl_clFinish(gegl_cl_get_command_queue()) );
      }

The key point here is the GEGL_BUFFER_CL_WRITE and GEGL_BUFFER_CL_READ flags passed to the iterator. They mean that writing and reading will be done through whatever OpenCL device [GPU or CPU] we’re using. This code just executes the kernel defined above for each tile.

Before entering the iteration in buffer_read and buffer_write, all data from buffer_read is copied to the GPU [of course, only if it is the most recent]. At the end, the OpenCL revision numbers from buffer_write’s tiles are bumped.

If after all that we do this:

gegl_buffer_get (buffer_write, 1.0, NULL, NULL, buf_write, GEGL_AUTO_ROWSTRIDE);

This means we want to copy buffer_write’s data to a pointer in the host memory, so we have to synchronize host and GPU data versions before that. So all buffer functions will always return the most recent data version and, at the same time, memory transferences will be made only if necessary.

Here is a fluxogram  of what is happening in this code:

Full code

Performance Results

Running the Brightness-Contrast code with a 1 mega pixel image using a NVidia Tesla C2050 as OpenCL device and a Intel Xeon E5506 as comparison [just using one core, but the code uses SSE2].

Time of memory transferences to the GPU was considered in this benchmark [EDIT: this time considers transferring data back and forth between GPU and CPU].

• CPU Elapsed time: 526 milliseconds
• OpenCL Elapsed time: 483 milliseconds

Also, here is a chart from NVidia profiler showing how execution time was spent:

Almost 80% of total execution time has been spent in memory transferences to and from the GPU. This is a good result, because even with this overhead the results were reasonable. Consider that the typical use case of GEGL is doing many operations in sequence, so the ratio processing/memory transferences tends to be higher. In fact, the case present here is the worst-case.

Possible Improvements

There are a lot of things that can be done in order to increase current code speed:

Intercalate execution of tiles with memory transferences of others tiles

GPU hardware (at least modern NVidia GPUs) has separated  units for processing and memory transferences, we can use this to intercalate tiles processing and copying.

Tiles sharing the same OpenCL memory buffer

There is a lot of overhead in allocating a GPU texture for each tile. which is typically small [128×64]. I think the best way to tackle this problem is allocating a big chunk of memory and using offsets in this chunk when processing [it’s impossible to have pointers to GPU memory], the problem is that GEGL is supposed to abstract the user  this kind of stuff from the user. Another idea is to serialize execution by having a pool of textures which can be reused by tiles, this would be good also because GPU memory is smaller than Host memory in general, the direct mapping CPU<->GPU cannot stand in fact.

Multiple OpenCL Command Queues

Command Queues can be executed concurrently in the same device, the Fermi architecture from NVidia can run 16 kernels at the same time, for example. This can be used to solve the memory transference overhead also.

Next Steps

I have yet to finish the implementation of a operator interface for OpenCL and make some OpenCL operators in order to create a useful chain of processing only in the GPU.

As the time for a GSoC project is very limited, my mentor and I decided to let optimizations outside the project, but I intend to work on them as soon as I can 🙂

Conclusion

The use of locking in order to synchronize CPU and GPU data was the most challenging part of the implementation, but after extensive testing. I think it’s working now, though I took more time than I expected to make it run properly.

Moreover, results so far show that using OpenCL to speed up Gegl is feasible and very interesting, thought still there is some challenges to be tackled, the tiled structure of Gegl allows a lot of optimizations.

The Gegl OpenCL branch is here.