Added OpenCL support to red-eyed-removal

3 files changed, 128 insertions, 0 deletions

diff --git a/opencl/red-eye-removal.cl b/opencl/red-eye-removal.cl
new file mode 100644
index 00000000..3de51d54
--- /dev/null
+++ b/opencl/red-eye-removal.cl
@@ -0,0 +1,42 @@
+/* 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>
+ */
+
+#define RED_FACTOR    0.5133333
+#define GREEN_FACTOR  1
+#define BLUE_FACTOR   0.1933333
+
+__kernel void cl_red_eye_removal(__global const float4 *in,
+                                 __global       float4 *out,
+                                                float threshold)
+{
+  int gid     = get_global_id(0);
+  float4 in_v = in[gid];
+  float adjusted_red       = in_v.x * RED_FACTOR;
+  float adjusted_green     = in_v.y * GREEN_FACTOR;
+  float adjusted_blue      = in_v.z * BLUE_FACTOR;
+  float adjusted_threshold = (threshold - 0.4) * 2;
+  float tmp;
+
+  if (adjusted_red >= adjusted_green - adjusted_threshold &&
+      adjusted_red >= adjusted_blue  - adjusted_threshold)
+    {
+      tmp = (adjusted_green + adjusted_blue) / (2.0 * RED_FACTOR);
+      in_v.x = clamp(tmp, 0.0f, 1.0f);
+    }
+  out[gid]  = in_v;
+}
diff --git a/opencl/red-eye-removal.cl.h b/opencl/red-eye-removal.cl.h
new file mode 100644
index 00000000..0c517b51
--- /dev/null
+++ b/opencl/red-eye-removal.cl.h
@@ -0,0 +1,44 @@
+static const char* red_eye_removal_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"
+"#define RED_FACTOR    0.5133333                                               \n"
+"#define GREEN_FACTOR  1                                                       \n"
+"#define BLUE_FACTOR   0.1933333                                               \n"
+"                                                                              \n"
+"__kernel void cl_red_eye_removal(__global const float4 *in,                   \n"
+"                                 __global       float4 *out,                  \n"
+"                                                float threshold)              \n"
+"{                                                                             \n"
+"  int gid     = get_global_id(0);                                             \n"
+"  float4 in_v = in[gid];                                                      \n"
+"  float adjusted_red       = in_v.x * RED_FACTOR;                             \n"
+"  float adjusted_green     = in_v.y * GREEN_FACTOR;                           \n"
+"  float adjusted_blue      = in_v.z * BLUE_FACTOR;                            \n"
+"  float adjusted_threshold = (threshold - 0.4) * 2;                           \n"
+"  float tmp;                                                                  \n"
+"                                                                              \n"
+"  if (adjusted_red >= adjusted_green - adjusted_threshold &&                  \n"
+"      adjusted_red >= adjusted_blue  - adjusted_threshold)                    \n"
+"    {                                                                         \n"
+"      tmp = (adjusted_green + adjusted_blue) / (2.0 * RED_FACTOR);            \n"
+"      in_v.x = clamp(tmp, 0.0f, 1.0f);                                        \n"
+"    }                                                                         \n"
+"  out[gid]  = in_v;                                                           \n"
+"}                                                                             \n"
+;
diff --git a/operations/common/red-eye-removal.c b/operations/common/red-eye-removal.c
index fb7f0d27..f8d3dd46 100644
--- a/operations/common/red-eye-removal.c
+++ b/operations/common/red-eye-removal.c
@@ -101,6 +101,46 @@ process (GeglOperation       *operation,
   return TRUE;
 }
 
+#include "opencl/gegl-cl.h"
+#include "opencl/red-eye-removal.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   threshold    = o->threshold;
+
+  GEGL_CL_BUILD(red_eye_removal, "cl_red_eye_removal")
+
+  {
+  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, threshold)
+  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)
 {
@@ -111,7 +151,9 @@ gegl_chant_class_init (GeglChantClass *klass)
   point_filter_class = GEGL_OPERATION_POINT_FILTER_CLASS (klass);
 
   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:red-eye-removal",