[gegl] Added OpenCL support to alien-map

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

commit a2feb9419ad61657adf9f050ebb4236a9d04b979
Author: Carlos Zubieta <czubieta dev gmail com>
Date:   Wed Sep 11 06:10:18 2013 -0500

    Added OpenCL support to alien-map

 opencl/alien-map.cl           |   35 +++++++++++++++++++++++
 opencl/alien-map.cl.h         |   37 ++++++++++++++++++++++++
 operations/common/alien-map.c |   62 +++++++++++++++++++++++++++++++++++++++-
 3 files changed, 132 insertions(+), 2 deletions(-)
---
diff --git a/opencl/alien-map.cl b/opencl/alien-map.cl
new file mode 100644
index 0000000..cf3dabe
--- /dev/null
+++ b/opencl/alien-map.cl
@@ -0,0 +1,35 @@
+/* This file is an image processing operation for GEGL
+ *
+ * GEGL is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 3 of the License, or (at your option) any later version.
+ *
+ * GEGL is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with GEGL; if not, see <http://www.gnu.org/licenses/>.
+ *
+ * Copyright 2013 Carlos Zubieta <czubieta dev gmail com>
+ */
+
+__kernel void cl_alien_map(__global const float4 *in,
+                           __global       float4 *out,
+                                          float3 freq,
+                                          float3 phaseshift,
+                                          int3   keep)
+{
+  int gid     = get_global_id(0);
+  float4 in_v = in[gid];
+  float3 unit = (float3) (1.0f, 1.0f, 1.0f);
+  float3 tmp  = 0.5f * (unit
+                        + sin((2.0f * in_v.xyz - unit) * freq.xyz + phaseshift.xyz));
+  float4 out_v;
+
+  out_v.xyz = keep.xyz ? in_v.xyz : tmp;
+  out_v.w   = in_v.w;
+  out[gid]  = out_v;
+}
diff --git a/opencl/alien-map.cl.h b/opencl/alien-map.cl.h
new file mode 100644
index 0000000..4d7bfef
--- /dev/null
+++ b/opencl/alien-map.cl.h
@@ -0,0 +1,37 @@
+static const char* alien_map_cl_source =
+"/* This file is an image processing operation for GEGL                        \n"
+" *                                                                            \n"
+" * GEGL is free software; you can redistribute it and/or                      \n"
+" * modify it under the terms of the GNU Lesser General Public                 \n"
+" * License as published by the Free Software Foundation; either               \n"
+" * version 3 of the License, or (at your option) any later version.           \n"
+" *                                                                            \n"
+" * GEGL is distributed in the hope that it will be useful,                    \n"
+" * but WITHOUT ANY WARRANTY; without even the implied warranty of             \n"
+" * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU          \n"
+" * Lesser General Public License for more details.                            \n"
+" *                                                                            \n"
+" * You should have received a copy of the GNU Lesser General Public           \n"
+" * License along with GEGL; if not, see <http://www.gnu.org/licenses/>.       \n"
+" *                                                                            \n"
+" * Copyright 2013 Carlos Zubieta <czubieta dev gmail com>                     \n"
+" */                                                                           \n"
+"                                                                              \n"
+"__kernel void cl_alien_map(__global const float4 *in,                         \n"
+"                           __global       float4 *out,                        \n"
+"                                          float3 freq,                        \n"
+"                                          float3 phaseshift,                  \n"
+"                                          int3   keep)                        \n"
+"{                                                                             \n"
+"  int gid     = get_global_id(0);                                             \n"
+"  float4 in_v = in[gid];                                                      \n"
+"  float3 unit = (float3) (1.0f, 1.0f, 1.0f);                                  \n"
+"  float3 tmp  = 0.5f * (unit                                                  \n"
+"                        + sin((2.0f * in_v.xyz - unit) * freq.xyz + phaseshift.xyz));\n"
+"  float4 out_v;                                                               \n"
+"                                                                              \n"
+"  out_v.xyz = keep.xyz ? in_v.xyz : tmp;                                      \n"
+"  out_v.w   = in_v.w;                                                         \n"
+"  out[gid]  = out_v;                                                          \n"
+"}                                                                             \n"
+;
diff --git a/operations/common/alien-map.c b/operations/common/alien-map.c
index 791945c..dcc0df7 100644
--- a/operations/common/alien-map.c
+++ b/operations/common/alien-map.c
@@ -140,6 +140,62 @@ process (GeglOperation       *op,
   return TRUE;
 }
 
+#include "opencl/gegl-cl.h"
+#include "opencl/alien-map.cl.h"
+
+GEGL_CL_STATIC
+
+static gboolean
+cl_process (GeglOperation       *operation,
+            cl_mem              in,
+            cl_mem              out,
+            size_t              global_worksize,
+            const GeglRectangle *roi,
+            gint                level)
+{
+  GeglChantO *o   = GEGL_CHANT_PROPERTIES (operation);
+  cl_float3   freq;
+  cl_float3   phaseshift;
+  cl_int3     keep;
+
+  GEGL_CL_BUILD(alien_map, "cl_alien_map")
+
+  freq.s[0] = o->cpn_1_frequency * G_PI;
+  freq.s[1] = o->cpn_2_frequency * G_PI;
+  freq.s[2] = o->cpn_3_frequency * G_PI;
+
+  phaseshift.s[0] = G_PI * o->cpn_1_phaseshift / 180.0;
+  phaseshift.s[1] = G_PI * o->cpn_2_phaseshift / 180.0;
+  phaseshift.s[2] = G_PI * o->cpn_3_phaseshift / 180.0;
+
+  keep.s[0] = (cl_int)o->cpn_1_keep;
+  keep.s[1] = (cl_int)o->cpn_2_keep;
+  keep.s[2] = (cl_int)o->cpn_3_keep;
+
+  {
+  cl_int cl_err = 0;
+
+  GEGL_CL_ARG_START(cl_data->kernel[0])
+  GEGL_CL_ARG(cl_mem,    in)
+  GEGL_CL_ARG(cl_mem,    out)
+  GEGL_CL_ARG(cl_float3, freq)
+  GEGL_CL_ARG(cl_float3, phaseshift)
+  GEGL_CL_ARG(cl_int3,   keep)
+  GEGL_CL_ARG_END
+
+  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+                                       cl_data->kernel[0], 1,
+                                       NULL, &global_worksize, NULL,
+                                       0, NULL, NULL);
+  CL_CHECK;
+  }
+
+  return  FALSE;
+
+error:
+  return TRUE;
+}
+
 static void
 gegl_chant_class_init (GeglChantClass *klass)
 {
@@ -149,8 +205,10 @@ gegl_chant_class_init (GeglChantClass *klass)
   operation_class    = GEGL_OPERATION_CLASS (klass);
   point_filter_class = GEGL_OPERATION_POINT_FILTER_CLASS (klass);
 
-  operation_class->prepare    = prepare;
-  point_filter_class->process = process;
+  operation_class->prepare        = prepare;
+  operation_class->opencl_support = TRUE;
+  point_filter_class->process     = process;
+  point_filter_class->cl_process  = cl_process;
 
   gegl_operation_class_set_keys (operation_class,
     "name",        "gegl:alien-map",