[gegl] Added partial OpenCL support to contrast-curve

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

commit 6aea9f790ef65ede7733e07113de9d21c2b78726
Author: Carlos Zubieta <czubieta dev gmail com>
Date:   Tue Jul 23 18:13:14 2013 -0500

    Added partial OpenCL support to contrast-curve

 opencl/contrast-curve.cl           |   38 +++++++++++
 opencl/contrast-curve.cl.h         |   39 +++++++++++
 operations/common/contrast-curve.c |  125 ++++++++++++++++++++++++++++++++++++
 3 files changed, 202 insertions(+), 0 deletions(-)
---
diff --git a/opencl/contrast-curve.cl b/opencl/contrast-curve.cl
new file mode 100644
index 0000000..4abf13e
--- /dev/null
+++ b/opencl/contrast-curve.cl
@@ -0,0 +1,38 @@
+/* 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_contrast_curve(__global const float2 *in,
+                                __global       float2 *out,
+                                __global       float  *curve,
+                                               int     num_sampling_points)
+{
+  int gid     = get_global_id(0);
+  float2 in_v = in[gid];
+  int x       = in_v.x * num_sampling_points;
+  float y;
+
+  if(x < 0)
+    y = curve[0];
+  else if (x < num_sampling_points)
+    y = curve[x];
+  else
+    y = curve[num_sampling_points - 1];
+
+  out[gid] = (float2) (y, in_v.y);
+}
+
diff --git a/opencl/contrast-curve.cl.h b/opencl/contrast-curve.cl.h
new file mode 100644
index 0000000..7912cd7
--- /dev/null
+++ b/opencl/contrast-curve.cl.h
@@ -0,0 +1,39 @@
+static const char* contrast_curve_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_contrast_curve(__global const float2 *in,                    \n"
+"                                __global       float2 *out,                   \n"
+"                                __global       float  *curve,                 \n"
+"                                               int     num_sampling_points)   \n"
+"{                                                                             \n"
+"  int gid     = get_global_id(0);                                             \n"
+"  float2 in_v = in[gid];                                                      \n"
+"  int x       = in_v.x * num_sampling_points;                                 \n"
+"  float y;                                                                    \n"
+"                                                                              \n"
+"  if(x < 0)                                                                   \n"
+"    y = curve[0];                                                             \n"
+"  else if (x < num_sampling_points)                                           \n"
+"    y = curve[x];                                                             \n"
+"  else                                                                        \n"
+"    y = curve[num_sampling_points - 1];                                       \n"
+"                                                                              \n"
+"  out[gid] = (float2) (y, in_v.y);                                            \n"
+"}                                                                             \n"
+;
diff --git a/operations/common/contrast-curve.c b/operations/common/contrast-curve.c
index 9fbb2b9..d015478 100644
--- a/operations/common/contrast-curve.c
+++ b/operations/common/contrast-curve.c
@@ -41,6 +41,129 @@ static void prepare (GeglOperation *operation)
   gegl_operation_set_format (operation, "output", format);
 }
 
+#include "opencl/gegl-cl.h"
+#include "gegl/gegl-debug.h"
+
+#include "opencl/contrast-curve.cl.h"
+
+/* TODO : Replace gegl_curve_calc_values and gegl_curve_calc_value in cl_process
+          with something more suitable for the cl version*/
+
+static void
+copy_double_array_to_float_array (gdouble *in,
+                                  gfloat  *out,
+                                  gint     size)
+{
+  gint i;
+  for (i = 0; i < size; ++i)
+    out[i] = (gfloat) in[i];
+}
+
+static GeglClRunData * cl_data = NULL;
+
+static gboolean
+cl_process (GeglOperation       *self,
+            cl_mem               in_tex,
+            cl_mem               out_tex,
+            size_t               global_worksize,
+            const GeglRectangle *roi,
+            gint                 level)
+{
+  GeglChantO *o = GEGL_CHANT_PROPERTIES (self);
+  gint       num_sampling_points;
+  gdouble    *xs, *ys;
+  gfloat     *ysf;
+  cl_mem     cl_curve;
+  cl_ulong   cl_max_constant_size;
+  cl_int     cl_err = 0;
+
+  num_sampling_points = o->sampling_points;
+
+  if (num_sampling_points > 0)
+    {
+      xs = g_new (gdouble, num_sampling_points);
+      ys = g_new (gdouble, num_sampling_points);
+
+      gegl_curve_calc_values (o->curve, 0.0, 1.0, num_sampling_points, xs, ys);
+      g_free(xs);
+
+      /*We need to downscale the array to pass it to the GPU*/
+      ysf = g_new (gfloat, num_sampling_points);
+      copy_double_array_to_float_array (ys, ysf, num_sampling_points);
+      g_free (ys);
+
+      cl_err = gegl_clGetDeviceInfo (gegl_cl_get_device (),
+                                     CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE,
+                                     sizeof (cl_ulong),
+                                     &cl_max_constant_size,
+                                     NULL);
+      CL_CHECK;
+
+      GEGL_NOTE (GEGL_DEBUG_OPENCL, "Max Constant Mem Size: %lu bytes",
+                 (unsigned long) cl_max_constant_size);
+
+      if (sizeof (cl_float) * num_sampling_points < cl_max_constant_size)
+        {
+          if (!cl_data)
+            {
+              const char *kernel_name[] = {"cl_contrast_curve",NULL};
+              cl_data = gegl_cl_compile_and_build (contrast_curve_cl_source,
+                                                   kernel_name);
+            }
+          if (!cl_data)
+            return TRUE;
+
+          cl_curve = gegl_clCreateBuffer (gegl_cl_get_context (),
+                                          CL_MEM_USE_HOST_PTR | CL_MEM_READ_ONLY,
+                                          num_sampling_points * sizeof (cl_float),
+                                          ysf, &cl_err);
+          CL_CHECK;
+          cl_err = gegl_clSetKernelArg (cl_data->kernel[0], 0, sizeof (cl_mem),
+                                        (void*) &in_tex);
+          CL_CHECK;
+          cl_err = gegl_clSetKernelArg (cl_data->kernel[0], 1, sizeof (cl_mem),
+                                        (void*) &out_tex);
+          CL_CHECK;
+          cl_err = gegl_clSetKernelArg (cl_data->kernel[0], 2, sizeof (cl_mem),
+                                        (void*) &cl_curve);
+          CL_CHECK;
+          cl_err = gegl_clSetKernelArg (cl_data->kernel[0], 3, sizeof (gint),
+                                        (void*) &num_sampling_points);
+          CL_CHECK;
+          cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
+                                                cl_data->kernel[0], 1,
+                                                NULL, &global_worksize, NULL,
+                                                0, NULL, NULL);
+          CL_CHECK;
+
+          cl_err = gegl_clFinish (gegl_cl_get_command_queue ());
+          CL_CHECK;
+
+          cl_err= gegl_clReleaseMemObject (cl_curve);
+          CL_CHECK;
+        }
+      else
+        {
+          /*If the curve size doesn't fit constant memory is better to use CPU*/
+          GEGL_NOTE (GEGL_DEBUG_OPENCL,
+                     "Not enough constant memory for the curve");
+          g_free (ysf);
+          return TRUE;
+        }
+
+      g_free(ysf);
+      return FALSE;
+error:
+      return TRUE;
+    }
+  else  /*If the curve doesn't have a lookup table is better to use CPU*/
+    {
+      GEGL_NOTE (GEGL_DEBUG_OPENCL,
+                 "Curve not suitable to be computed in the GPU");
+      return TRUE;
+    }
+}
+
 static gboolean
 process (GeglOperation       *op,
          void                *in_buf,
@@ -116,7 +239,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->prepare = prepare;
+  operation_class->opencl_support = TRUE;
 
   gegl_operation_class_set_keys (operation_class,
     "name"       , "gegl:contrast-curve",