[gegl] Added OpenCL support to weighted-blend

[Téo Mazars <teom src gnome org>]

commit 911d312b7369b13447becd6a417585dbd6dfeed0
Author: Carlos Zubieta <czubieta dev gmail com>
Date:   Tue Jul 23 18:09:56 2013 -0500

    Added OpenCL support to weighted-blend

 opencl/weighted-blend.cl           |   29 ++++++++++++++
 opencl/weighted-blend.cl.h         |   31 +++++++++++++++
 operations/common/weighted-blend.c |   73 +++++++++++++++++++++++++++++++++--
 3 files changed, 128 insertions(+), 5 deletions(-)
---
diff --git a/opencl/weighted-blend.cl b/opencl/weighted-blend.cl
new file mode 100644
index 0000000..68cec63
--- /dev/null
+++ b/opencl/weighted-blend.cl
@@ -0,0 +1,29 @@
+__kernel void cl_copy_weigthed_blend(__global const float4 *in,
+                                     __global       float4 *out)
+{
+  int gid = get_global_id(0);
+  float4 in_v = in[gid];
+  out[gid] = in_v;
+}
+
+__kernel void cl_weighted_blend(__global const float4 *in,
+                                __global const float4 *aux,
+                                __global       float4 *out)
+{
+  int gid = get_global_id(0);
+  float4 in_v = in[gid];
+  float4 aux_v = aux[gid];
+  float4 out_v;
+  float in_weight;
+  float aux_weight;
+  float total_alpha = in_v.w + aux_v.w;
+
+  total_alpha = total_alpha == 0 ? 1 : total_alpha;
+
+  in_weight = in_v.w / total_alpha;
+  aux_weight = 1.0f - in_weight;
+
+  out_v.xyz = in_weight * in_v.xyz + aux_weight * aux_v.xyz;
+  out_v.w = total_alpha;
+  out[gid] = out_v;
+}
diff --git a/opencl/weighted-blend.cl.h b/opencl/weighted-blend.cl.h
new file mode 100644
index 0000000..d516482
--- /dev/null
+++ b/opencl/weighted-blend.cl.h
@@ -0,0 +1,31 @@
+static const char* weighted_blend_cl_source =
+"__kernel void cl_copy_weigthed_blend(__global const float4 *in,               \n"
+"                                     __global       float4 *out)              \n"
+"{                                                                             \n"
+"  int gid = get_global_id(0);                                                 \n"
+"  float4 in_v = in[gid];                                                      \n"
+"  out[gid] = in_v;                                                            \n"
+"}                                                                             \n"
+"                                                                              \n"
+"__kernel void cl_weighted_blend(__global const float4 *in,                    \n"
+"                                __global const float4 *aux,                   \n"
+"                                __global       float4 *out)                   \n"
+"{                                                                             \n"
+"  int gid = get_global_id(0);                                                 \n"
+"  float4 in_v = in[gid];                                                      \n"
+"  float4 aux_v = aux[gid];                                                    \n"
+"  float4 out_v;                                                               \n"
+"  float in_weight;                                                            \n"
+"  float aux_weight;                                                           \n"
+"  float total_alpha = in_v.w + aux_v.w;                                       \n"
+"                                                                              \n"
+"  total_alpha = total_alpha == 0 ? 1 : total_alpha;                           \n"
+"                                                                              \n"
+"  in_weight = in_v.w / total_alpha;                                           \n"
+"  aux_weight = 1.0f - in_weight;                                              \n"
+"                                                                              \n"
+"  out_v.xyz = in_weight * in_v.xyz + aux_weight * aux_v.xyz;                  \n"
+"  out_v.w = total_alpha;                                                      \n"
+"  out[gid] = out_v;                                                           \n"
+"}                                                                             \n"
+;
diff --git a/operations/common/weighted-blend.c b/operations/common/weighted-blend.c
index 960b468..707429f 100644
--- a/operations/common/weighted-blend.c
+++ b/operations/common/weighted-blend.c
@@ -22,8 +22,6 @@
 
 #ifdef GEGL_CHANT_PROPERTIES
 
-gegl_chant_double (value, _("Value"), -G_MAXDOUBLE, G_MAXDOUBLE, 0.0, _("global value used if aux doesn't 
contain data"))
-
 #else
 
 #define GEGL_CHANT_TYPE_POINT_COMPOSER
@@ -40,6 +38,68 @@ static void prepare (GeglOperation *operation)
   gegl_operation_set_format (operation, "output", format);
 }
 
+#include "opencl/gegl-cl.h"
+#include "opencl/weighted-blend.cl.h"
+
+static GeglClRunData *cl_data = NULL;
+
+static gboolean
+cl_process (GeglOperation       *self,
+            cl_mem               in_tex,
+            cl_mem               aux_tex,
+            cl_mem               out_tex,
+            size_t               global_worksize,
+            const GeglRectangle *roi,
+            gint                 level)
+{
+  gint cl_err = 0;
+
+  if (!cl_data)
+    {
+      const char *kernel_name[] = {"cl_copy_weigthed_blend",
+                                   "cl_weighted_blend",
+                                   NULL};
+      cl_data = gegl_cl_compile_and_build (weighted_blend_cl_source,
+                                           kernel_name);
+    }
+  if (!cl_data) return TRUE;
+
+
+  if (!aux_tex)
+    {
+      cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex);
+      CL_CHECK;
+      cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex);
+      CL_CHECK;
+
+      cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+                                           cl_data->kernel[0], 1,
+                                           NULL, &global_worksize, NULL,
+                                           0, NULL, NULL);
+      CL_CHECK;
+    }
+  else
+    {
+      cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem), (void*)&in_tex);
+      CL_CHECK;
+      cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_mem), (void*)&aux_tex);
+      CL_CHECK;
+      cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 2, sizeof(cl_mem), (void*)&out_tex);
+      CL_CHECK;
+
+      cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+                                           cl_data->kernel[1], 1,
+                                           NULL, &global_worksize, NULL,
+                                           0, NULL, NULL);
+      CL_CHECK;
+    }
+
+  return FALSE;
+
+error:
+  return TRUE;
+}
+
 static gboolean
 process (GeglOperation       *op,
          void                *in_buf,
@@ -58,7 +118,6 @@ process (GeglOperation       *op,
     {
       /* there is no auxilary buffer.
        * output the input buffer.
-       * gfloat value = GEGL_CHANT_PROPERTIES (op)->value;
        */
       for (i = 0; i < n_pixels; i++)
         {
@@ -67,6 +126,8 @@ process (GeglOperation       *op,
             {
               out[j] = in[j];
             }
+          in  += 4;
+          out += 4;
         }
     }
   else
@@ -114,8 +175,10 @@ gegl_chant_class_init (GeglChantClass *klass)
   operation_class      = GEGL_OPERATION_CLASS (klass);
   point_composer_class = GEGL_OPERATION_POINT_COMPOSER_CLASS (klass);
 
-  point_composer_class->process = process;
-  operation_class->prepare      = prepare;
+  point_composer_class->process    = process;
+  point_composer_class->cl_process = cl_process;
+  operation_class->prepare         = prepare;
+  operation_class->opencl_support  = TRUE;
 
   gegl_operation_class_set_keys (operation_class,
     "name"       , "gegl:weighted-blend",