[gegl] Add OpenCL support for gegl:levels

[Victor Matheus de Araujo Oliveira <vmaolive src gnome org>]

commit 25fe8d705b9047f646ff2e2283c1c53b94969f1e
Author: Victor Oliveira <victormatheus gmail com>
Date:   Wed Mar 7 09:31:41 2012 -0300

    Add OpenCL support for gegl:levels

 operations/common/levels.c |   79 ++++++++++++++++++++++++++++++++++++++++++++
 1 files changed, 79 insertions(+), 0 deletions(-)
---
diff --git a/operations/common/levels.c b/operations/common/levels.c
index 1274973..7edb318 100644
--- a/operations/common/levels.c
+++ b/operations/common/levels.c
@@ -85,6 +85,82 @@ process (GeglOperation       *op,
   return TRUE;
 }
 
+#include "opencl/gegl-cl.h"
+
+static const char* kernel_source =
+"__kernel void kernel_levels(__global const float4     *in,      \n"
+"                            __global       float4     *out,     \n"
+"                            float in_offset,                    \n"
+"                            float out_offset,                   \n"
+"                            float scale)                        \n"
+"{                                                               \n"
+"  int gid = get_global_id(0);                                   \n"
+"  float4 in_v  = in[gid];                                       \n"
+"  float4 out_v;                                                 \n"
+"  out_v.xyz = (in_v.xyz - in_offset) * scale + out_offset;      \n"
+"  out_v.w   =  in_v.w;                                          \n"
+"  out[gid]  =  out_v;                                           \n"
+"}                                                               \n";
+
+static gegl_cl_run_data *cl_data = NULL;
+
+/* OpenCL processing function */
+static cl_int
+cl_process (GeglOperation       *op,
+            cl_mem               in_tex,
+            cl_mem               out_tex,
+            size_t               global_worksize,
+            const GeglRectangle *roi,
+            gint                 level)
+{
+  /* Retrieve a pointer to GeglChantO structure which contains all the
+   * chanted properties
+   */
+
+  GeglChantO *o = GEGL_CHANT_PROPERTIES (op);
+
+  gfloat      in_range;
+  gfloat      out_range;
+  gfloat      in_offset;
+  gfloat      out_offset;
+  gfloat      scale;
+
+  cl_int cl_err = 0;
+
+  in_offset = o->in_low * 1.0;
+  out_offset = o->out_low * 1.0;
+  in_range = o->in_high-o->in_low;
+  out_range = o->out_high-o->out_low;
+
+  if (in_range == 0.0)
+    in_range = 0.00000001;
+
+  scale = out_range/in_range;
+
+  if (!cl_data)
+    {
+      const char *kernel_name[] = {"kernel_levels", 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*)&out_tex);
+  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&in_offset);
+  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&out_offset);
+  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float), (void*)&scale);
+  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;
+}
+
+
 
 static void
 gegl_chant_class_init (GeglChantClass *klass)
@@ -96,6 +172,9 @@ gegl_chant_class_init (GeglChantClass *klass)
   point_filter_class = GEGL_OPERATION_POINT_FILTER_CLASS (klass);
 
   point_filter_class->process = process;
+  point_filter_class->cl_process = cl_process;
+
+  operation_class->opencl_support = TRUE;
 
   gegl_operation_class_set_keys (operation_class,
     "name"       , "gegl:levels",