[gegl] Added OpenCL support to red-eyed-removal
- From: Téo Mazars <teom src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] Added OpenCL support to red-eyed-removal
- Date: Thu, 31 Oct 2013 11:01:30 +0000 (UTC)
commit 2c5dfa494545435a49ddf5e68d74babfd6a3934e
Author: Carlos Zubieta <czubieta dev gmail com>
Date: Wed Sep 11 06:51:26 2013 -0500
Added OpenCL support to red-eyed-removal
opencl/red-eye-removal.cl | 42 +++++++++++++++++++++++++++++++++
opencl/red-eye-removal.cl.h | 44 +++++++++++++++++++++++++++++++++++
operations/common/red-eye-removal.c | 42 +++++++++++++++++++++++++++++++++
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 0000000..3de51d5
--- /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 0000000..0c517b5
--- /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 fb7f0d2..f8d3dd4 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",
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]