[gegl] gegl:checkerboard: Add OpenCL version for float4 formats

[Daniel Sabo <daniels src gnome org>]

commit be8583c3b373865d8c5b4023d8178195c81e0ca5
Author: Daniel Sabo <DanielSabo gmail com>
Date:   Sun Oct 20 18:28:00 2013 -0700

    gegl:checkerboard: Add OpenCL version for float4 formats
    
    This only supports float4 formats for now because OpenCL doesn't
    necessarily allow byte addressing. Other formats will fall back
    to the CPU version.

 operations/common/checkerboard.c |  189 ++++++++++++++++++++++++++++++++++++--
 1 files changed, 179 insertions(+), 10 deletions(-)
---
diff --git a/operations/common/checkerboard.c b/operations/common/checkerboard.c
index 918b83f..8b9459d 100644
--- a/operations/common/checkerboard.c
+++ b/operations/common/checkerboard.c
@@ -46,8 +46,8 @@ gegl_chant_color    (color2, _("Other color"),
                      "white",
                      _("The other cell color (defaults to 'white')"))
 
-gegl_chant_format (format, _("Babl Format"),
-                   _("The babl format of the output"))
+gegl_chant_format   (format, _("Babl Format"),
+                     _("The babl format of the output"))
 
 #else
 
@@ -55,7 +55,8 @@ gegl_chant_format (format, _("Babl Format"),
 #define GEGL_CHANT_C_FILE "checkerboard.c"
 
 #include "gegl-chant.h"
-#include <gegl-utils.h>
+#include <gegl-buffer-cl-iterator.h>
+#include <gegl-debug.h>
 
 static void
 prepare (GeglOperation *operation)
@@ -74,17 +75,135 @@ get_bounding_box (GeglOperation *operation)
   return gegl_rectangle_infinite_plane ();
 }
 
+static const char* checkerboard_cl_source =
+"static inline int tile_index (int coordinate, int stride)   \n"
+"{                                                           \n"
+"  if (coordinate >= 0)                                      \n"
+"    return coordinate / stride;                             \n"
+"  else                                                      \n"
+"    return ((coordinate + 1) / stride) - 1;                 \n"
+"}                                                           \n"
+"                                                            \n"
+"__kernel void kernel_checkerboard (__global float4 *out,    \n"
+"                                   float4 color1,           \n"
+"                                   float4 color2,           \n"
+"                                   int square_width,        \n"
+"                                   int square_height,       \n"
+"                                   int x_offset,            \n"
+"                                   int y_offset,            \n"
+"                                   int roi_x,               \n"
+"                                   int roi_y,               \n"
+"                                   int roi_width)           \n"
+"{                                                           \n"
+"    int gidx = 0;                                           \n"
+"    int gidy = get_global_id(0);                            \n"
+"    float4 cur_color;                                       \n"
+"    bool in_color1;                                         \n"
+"                                                            \n"
+"    int x = roi_x + gidx - x_offset;                        \n"
+"    int y = roi_y + gidy - y_offset;                        \n"
+"                                                            \n"
+"    int tilex = tile_index (x, square_width);               \n"
+"    int tiley = tile_index (y, square_height);              \n"
+"                                                            \n"
+"    if ((tilex + tiley) % 2 == 0)                           \n"
+"      {                                                     \n"
+"        cur_color = color1;                                 \n"
+"        in_color1 = true;                                   \n"
+"      }                                                     \n"
+"    else                                                    \n"
+"      {                                                     \n"
+"        cur_color = color2;                                 \n"
+"        in_color1 = false;                                  \n"
+"      }                                                     \n"
+"                                                            \n"
+"    int stripe_end = (tilex + 1) * square_width;            \n"
+"    int stripe_width = stripe_end - x;                      \n"
+"    int gidx_max = roi_width;                               \n"
+"                                                            \n"
+"    while (gidx < gidx_max)                                 \n"
+"      {                                                     \n"
+"        out[gidx++ + gidy * roi_width] = cur_color;         \n"
+"        stripe_width--;                                     \n"
+"                                                            \n"
+"        if (stripe_width == 0)                              \n"
+"          {                                                 \n"
+"            stripe_width = square_width;                    \n"
+"                                                            \n"
+"            if (in_color1)                                  \n"
+"              cur_color = color2;                           \n"
+"            else                                            \n"
+"              cur_color = color1;                           \n"
+"            in_color1 = !in_color1;                         \n"
+"          }                                                 \n"
+"      }                                                     \n"
+"}                                                           \n";
+
 #define TILE_INDEX(coordinate,stride) \
   (((coordinate) >= 0)?\
       (coordinate) / (stride):\
       ((((coordinate) + 1) /(stride)) - 1))
 
+
+static GeglClRunData *cl_data = NULL;
+
 static gboolean
-process (GeglOperation       *operation,
-         void                *out_buf,
-         glong                n_pixels,
-         const GeglRectangle *roi,
-         gint                 level)
+checkerboard_cl_process (GeglOperation       *operation,
+                         cl_mem               out_tex,
+                         size_t               global_worksize,
+                         const GeglRectangle *roi,
+                         gint                 level)
+{
+  GeglChantO   *o           = GEGL_CHANT_PROPERTIES (operation);
+  const Babl   *out_format  = gegl_operation_get_format (operation, "output");
+  const size_t  gbl_size[1] = {roi->height};
+  cl_int        cl_err      = 0;
+  float         color1[4];
+  float         color2[4];
+
+  if (!cl_data)
+  {
+    const char *kernel_name[] = {"kernel_checkerboard", NULL};
+    cl_data = gegl_cl_compile_and_build (checkerboard_cl_source, kernel_name);
+
+    if (!cl_data)
+      return TRUE;
+  }
+
+  gegl_color_get_pixel (o->color1, out_format, color1);
+  gegl_color_get_pixel (o->color2, out_format, color2);
+
+  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
+                                    sizeof(cl_mem), &out_tex,
+                                    sizeof(color1), &color1,
+                                    sizeof(color2), &color2,
+                                    sizeof(cl_int), &o->x,
+                                    sizeof(cl_int), &o->y,
+                                    sizeof(cl_int), &o->x_offset,
+                                    sizeof(cl_int), &o->y_offset,
+                                    sizeof(cl_int), &roi->x,
+                                    sizeof(cl_int), &roi->y,
+                                    sizeof(cl_int), &roi->width,
+                                    NULL);
+  CL_CHECK;
+
+  cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
+                                        cl_data->kernel[0], 1,
+                                        NULL, gbl_size, NULL,
+                                        0, NULL, NULL);
+  CL_CHECK;
+
+  return FALSE;
+error:
+  return TRUE;
+}
+
+static gboolean
+checkerboard_process (GeglOperation       *operation,
+                      void                *out_buf,
+                      glong                n_pixels,
+                      const GeglRectangle *roi,
+                      gint                 level)
 {
   GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
   const Babl *out_format = gegl_operation_get_format (operation, "output");
@@ -137,19 +256,69 @@ process (GeglOperation       *operation,
         }
     }
 
-  return  TRUE;
+  return TRUE;
+}
+
+
+static gboolean
+operation_source_process (GeglOperation       *operation,
+                          GeglBuffer          *output,
+                          const GeglRectangle *result,
+                          gint                 level)
+{
+  const Babl *out_format = gegl_operation_get_format (operation, "output");
+
+  if ((result->width > 0) && (result->height > 0))
+    {
+      GeglBufferIterator *iter;
+      if (gegl_operation_use_opencl (operation) &&
+          babl_format_get_n_components (out_format) == 4 &&
+          babl_format_get_type (out_format, 0) == babl_type ("float"))
+        {
+          GeglBufferClIterator *cl_iter;
+          gboolean err;
+
+          GEGL_NOTE (GEGL_DEBUG_OPENCL, "GEGL_OPERATION_POINT_RENDER: %s", GEGL_OPERATION_GET_CLASS 
(operation)->name);
+
+          cl_iter = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE);
+
+          while (gegl_buffer_cl_iterator_next (cl_iter, &err) && !err)
+            {
+              err = checkerboard_cl_process (operation, cl_iter->tex[0], cl_iter->size[0], &cl_iter->roi[0], 
level);
+
+              if (err)
+                {
+                  gegl_buffer_cl_iterator_stop (cl_iter);
+                  break;
+                }
+            }
+
+          if (err)
+            GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", GEGL_OPERATION_GET_CLASS (operation)->name);
+          else
+            return TRUE;
+        }
+
+      iter = gegl_buffer_iterator_new (output, result, level, out_format, GEGL_BUFFER_WRITE, 
GEGL_ABYSS_NONE);
+
+      while (gegl_buffer_iterator_next (iter))
+          checkerboard_process (operation, iter->data[0], iter->length, &iter->roi[0], level);
+    }
+  return TRUE;
 }
 
 static void
 gegl_chant_class_init (GeglChantClass *klass)
 {
   GeglOperationClass            *operation_class;
+  GeglOperationSourceClass      *source_class;
   GeglOperationPointRenderClass *point_render_class;
 
   operation_class = GEGL_OPERATION_CLASS (klass);
+  source_class = GEGL_OPERATION_SOURCE_CLASS (klass);
   point_render_class = GEGL_OPERATION_POINT_RENDER_CLASS (klass);
 
-  point_render_class->process = process;
+  source_class->process = operation_source_process;
   operation_class->get_bounding_box = get_bounding_box;
   operation_class->prepare = prepare;