[gegl] gegl:over with OpenCL

[Ãyvind KolÃs <ok src gnome org>]

commit db83482455d106f4215994ffcc5de970b3d65cd0
Author: Victor Oliveira <victormatheus gmail com>
Date:   Thu Feb 16 16:32:36 2012 -0200

    gegl:over with OpenCL

 operations/common/over.c |   59 ++++++++++++++++++++++++++++++++++++++++++++++
 1 files changed, 59 insertions(+), 0 deletions(-)
---
diff --git a/operations/common/over.c b/operations/common/over.c
index 026253e..092af65 100644
--- a/operations/common/over.c
+++ b/operations/common/over.c
@@ -70,6 +70,63 @@ process (GeglOperation        *op,
   return TRUE;
 }
 
+#include "opencl/gegl-cl.h"
+
+static const char* kernel_source =
+"__kernel void kernel_over(__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"
+"  out_v.xyz = aux_v.xyz + in_v.xyz * (1.0f - aux_v.w);         \n"
+"  out_v.w   = aux_v.w + in_v.w - aux_v.w * in_v.w;             \n"
+"  out[gid]  = out_v;                                           \n"
+"}                                                              \n";
+
+static gegl_cl_run_data *cl_data = NULL;
+
+static cl_int
+cl_process (GeglOperation       *op,
+            cl_mem              in_tex,
+            cl_mem              aux_tex,
+            cl_mem              out_tex,
+            size_t              global_worksize,
+            const GeglRectangle *roi)
+{
+  /* Retrieve a pointer to GeglChantO structure which contains all the
+   * chanted properties
+   */
+
+  cl_int cl_err = 0;
+
+  if (aux_tex == NULL)
+    return 0;
+    
+  if (!cl_data)
+    {
+      const char *kernel_name[] = {"kernel_over", NULL};
+      cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
+    }
+
+  if (!cl_data) return 1;
+
+  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_mem),   (void*)&aux_tex);
+  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);
+  if (cl_err != CL_SUCCESS) return cl_err;
+
+  return cl_err;
+}
+
 /* Fast paths */
 static gboolean operation_process (GeglOperation        *operation,
                                    GeglOperationContext *context,
@@ -128,9 +185,11 @@ gegl_chant_class_init (GeglChantClass *klass)
   operation_class->process = operation_process;
 
   point_composer_class->process = process;
+  point_composer_class->cl_process = cl_process;
 
   operation_class->compat_name = "gegl:over";
   operation_class->name        = "svg:src-over";
+  operation_class->opencl_support = TRUE;
   operation_class->description =
         _("Porter Duff operation over (d = cA + cB * (1 - aA))");
   operation_class->categories  = "compositors:porter-duff";