[gegl/soc-2013-opecl-ops] Added OpenCL support to alien-map

[Carlos Zubieta <czubieta src gnome org>]

commit 4b7909ad06ba5977e0f8f962997d97b2de3b5237
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           |   34 ++++++++++++++++++++++
 opencl/alien-map.cl.h         |   36 +++++++++++++++++++++++
 operations/common/alien-map.c |   62 +++++++++++++++++++++++++++++++++++++++-
 3 files changed, 130 insertions(+), 2 deletions(-)
---
diff --git a/opencl/alien-map.cl b/opencl/alien-map.cl
new file mode 100644
index 0000000..f4de4b8
--- /dev/null
+++ b/opencl/alien-map.cl
@@ -0,0 +1,34 @@
+/* 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 tmp  = 0.5*(1.0 
+                     + sin((2 * in_v.xyz - 1.0) * 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..45283e7
--- /dev/null
+++ b/opencl/alien-map.cl.h
@@ -0,0 +1,36 @@
+/* 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>
+ */
+
+static const char* alien_map_cl_source =
+"__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 tmp  = 0.5*(1.0                                                      \n"
+"                     + sin((2 * in_v.xyz - 1.0) * 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",