[gegl] Added OpenCL support to posterize

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

commit 8111876e4b57eafa6e9b0cbdff1fdc37697667ce
Author: Carlos Zubieta <czubieta dev gmail com>
Date:   Wed Sep 11 07:59:49 2013 -0500

    Added OpenCL support to posterize

 opencl/posterize.cl           |   28 ++++++++++++++++++++++++++
 opencl/posterize.cl.h         |   30 +++++++++++++++++++++++++++
 operations/common/posterize.c |   44 ++++++++++++++++++++++++++++++++++++++++-
 3 files changed, 101 insertions(+), 1 deletions(-)
---
diff --git a/opencl/posterize.cl b/opencl/posterize.cl
new file mode 100644
index 0000000..5ba9e8b
--- /dev/null
+++ b/opencl/posterize.cl
@@ -0,0 +1,28 @@
+/* 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_posterize(__global const float4 *in,
+                           __global       float4 *out,
+                                          float  levels)
+{
+  int gid     = get_global_id(0);
+  float4 in_v = in[gid];
+
+  in_v.xyz  = trunc(in_v.xyz * levels + (float3)(0.5f)) / levels;
+  out[gid]  = in_v;
+}
diff --git a/opencl/posterize.cl.h b/opencl/posterize.cl.h
new file mode 100644
index 0000000..e7717c5
--- /dev/null
+++ b/opencl/posterize.cl.h
@@ -0,0 +1,30 @@
+static const char* posterize_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_posterize(__global const float4 *in,                         \n"
+"                           __global       float4 *out,                        \n"
+"                                          float  levels)                      \n"
+"{                                                                             \n"
+"  int gid     = get_global_id(0);                                             \n"
+"  float4 in_v = in[gid];                                                      \n"
+"                                                                              \n"
+"  in_v.xyz  = trunc(in_v.xyz * levels + (float3)(0.5f)) / levels;             \n"
+"  out[gid]  = in_v;                                                           \n"
+"}                                                                             \n"
+;
diff --git a/operations/common/posterize.c b/operations/common/posterize.c
index 0fcca17..f5a1592 100644
--- a/operations/common/posterize.c
+++ b/operations/common/posterize.c
@@ -64,6 +64,46 @@ static gboolean process (GeglOperation       *operation,
   return TRUE;
 }
 
+#include "opencl/gegl-cl.h"
+#include "opencl/posterize.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_float    levels = o->levels;
+
+  GEGL_CL_BUILD(posterize, "cl_posterize")
+
+  {
+  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_float, levels)
+  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)
 {
@@ -73,7 +113,9 @@ gegl_chant_class_init (GeglChantClass *klass)
   operation_class    = GEGL_OPERATION_CLASS (klass);
   point_filter_class = GEGL_OPERATION_POINT_FILTER_CLASS (klass);
 
-  point_filter_class->process = process;
+  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:posterize",