[gegl/soc-2013-opecl-ops: 4/6] Added partial OpenCL support to contrast-curve

[Carlos Zubieta <czubieta src gnome org>]

commit c159ac65b889164f0b21b23d56dcc62538f8f927
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           |   20 ++++++
 opencl/contrast-curve.cl.h         |   22 +++++++
 operations/common/contrast-curve.c |  122 ++++++++++++++++++++++++++++++++++++
 3 files changed, 164 insertions(+), 0 deletions(-)
---
diff --git a/opencl/contrast-curve.cl b/opencl/contrast-curve.cl
new file mode 100644
index 0000000..48d6982
--- /dev/null
+++ b/opencl/contrast-curve.cl
@@ -0,0 +1,20 @@
+__kernel void cl_contrast_curve(__global const float2 *in,
+                                __global       float2 *out,
+                                __constant     float  *curve,
+                                constant 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..2ad559d
--- /dev/null
+++ b/opencl/contrast-curve.cl.h
@@ -0,0 +1,22 @@
+static const char* cl_contrast_curve_source =
+"__kernel void cl_contrast_curve(__global const float2 *in,                    \n"
+"                                __global       float2 *out,                   \n"
+"                                __constant     float  *curve,                 \n"
+"                                constant 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..c1f2e9c 100644
--- a/operations/common/contrast-curve.c
+++ b/operations/common/contrast-curve.c
@@ -41,6 +41,126 @@ 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 (cl_contrast_curve_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;
+          }
+        }
+      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 +236,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",