[gegl] Add OpenCL support for gegl:mono-mixer

[Victor Matheus de Araujo Oliveira <vmaolive src gnome org>]

commit a2c3d88e3b970afb85584f2470376394a3cffec7
Author: Zhang Peixuan <zhangpeixuan cn gmail com>
Date:   Tue Mar 6 17:35:38 2012 +0800

    Add OpenCL support for gegl:mono-mixer

 operations/common/mono-mixer.c |   94 ++++++++++++++++++++++++++++++++++++++++
 1 files changed, 94 insertions(+), 0 deletions(-)
---
diff --git a/operations/common/mono-mixer.c b/operations/common/mono-mixer.c
index fe5f306..75e29ec 100644
--- a/operations/common/mono-mixer.c
+++ b/operations/common/mono-mixer.c
@@ -44,6 +44,94 @@ static void prepare (GeglOperation *operation)
   gegl_operation_set_format (operation, "output", babl_format ("YA float"));
 }
 
+#include "opencl/gegl-cl.h"
+#include "buffer/gegl-buffer-cl-iterator.h"
+
+static const char* kernel_source =
+"__kernel void Mono_mixer_cl(__global const float4 *src_buf,              \n"
+"                            float4                 color,                \n"
+"                            __global float2       *dst_buf)              \n"
+"{                                                                        \n"
+"  int gid = get_global_id(0);                                            \n"
+"  float4 tmp = src_buf[gid] * color;                                     \n"
+"  dst_buf[gid].x = tmp.x + tmp.y + tmp.z;                                \n"
+"  dst_buf[gid].y = tmp.w;                                                \n"
+"}                                                                        \n";
+
+static gegl_cl_run_data * cl_data = NULL;
+
+static cl_int
+cl_mono_mixer(cl_mem                in_tex,
+              cl_mem                out_tex,
+              size_t                global_worksize,
+              const GeglRectangle  *roi,
+              gfloat                red,
+              gfloat                green,
+              gfloat                blue)
+{
+  cl_int cl_err = 0;
+  if (!cl_data)
+  {
+    const char *kernel_name[] = {"Mono_mixer_cl", NULL};
+    cl_data = gegl_cl_compile_and_build(kernel_source, kernel_name);
+  }
+  if (!cl_data) return 0;
+
+  {
+  cl_float4 color;
+  color.s[0] = red;
+  color.s[1] = green;
+  color.s[2] = blue;
+  color.s[3] = 1.0f;
+
+  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem),    (void*)&in_tex);
+  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_float4), (void*)&color);
+  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_mem),    (void*)&out_tex);
+  if (cl_err != CL_SUCCESS) return cl_err;
+
+  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(), cl_data->kernel[0],
+                                       1, NULL,
+                                       &global_worksize, NULL,
+                                       0, NULL, NULL);
+  }
+
+  return cl_err;
+}
+
+static gboolean
+cl_process (GeglOperation       *operation,
+            GeglBuffer          *input,
+            GeglBuffer          *output,
+            const GeglRectangle *result)
+{
+  const Babl *in_format  = gegl_operation_get_format (operation, "input");
+  const Babl *out_format = gegl_operation_get_format (operation, "output");
+  gint err;
+  gint j;
+  cl_int cl_err;
+
+  GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
+
+  GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output,   result, out_format, GEGL_CL_BUFFER_WRITE, GEGL_ABYSS_NONE);
+  gint read = gegl_buffer_cl_iterator_add (i, input, result, in_format,  GEGL_CL_BUFFER_READ, GEGL_ABYSS_NONE);
+  while (gegl_buffer_cl_iterator_next (i, &err))
+  {
+    if (err) return FALSE;
+    for (j=0; j < i->n; j++)
+    {
+      cl_err = cl_mono_mixer(i->tex[read][j], i->tex[0][j], i->size[0][j], &i->roi[0][j], o->red ,o->green , o->blue);
+
+      if (cl_err != CL_SUCCESS)
+      {
+        g_warning("[OpenCL] Error in gegl:mono-mixer: %s", gegl_cl_errstring(cl_err));
+        return FALSE;
+      }
+    }
+  }
+  return TRUE;
+}
+
+
 static gboolean
 process (GeglOperation       *operation,
          GeglBuffer          *input,
@@ -58,6 +146,10 @@ process (GeglOperation       *operation,
   gfloat     *in_buf;
   gfloat     *out_buf;
 
+  if (gegl_cl_is_accelerated ())
+    if (cl_process (operation, input, output, result))
+      return TRUE;
+
  if ((result->width > 0) && (result->height > 0))
  {
      gint num_pixels = result->width * result->height;
@@ -102,6 +194,8 @@ gegl_chant_class_init (GeglChantClass *klass)
   filter_class->process = process;
   operation_class->prepare = prepare;
 
+  operation_class->opencl_support = TRUE;
+
   gegl_operation_class_set_keys (operation_class,
     "name"       , "gegl:mono-mixer",
     "categories" , "color",