color-exchange: add opencl support

3 files changed, 156 insertions, 7 deletions

diff --git a/opencl/color-exchange.cl b/opencl/color-exchange.cl
new file mode 100644
index 00000000..2a6be6c3
--- /dev/null
+++ b/opencl/color-exchange.cl
@@ -0,0 +1,44 @@
+/* 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 2015 Thomas Manni <thomas.manni@free.fr>
+ */
+
+__kernel void cl_color_exchange(__global const float4 *in,
+                                __global       float4 *out,
+                                               float3 color_diff,
+                                               float3 min,
+                                               float3 max)
+{
+  int gid     = get_global_id(0);
+  float4 in_v = in[gid];
+  float4 out_v;
+
+  if(in_v.x > min.x && in_v.x < max.x &&
+     in_v.y > min.y && in_v.y < max.y &&
+     in_v.z > min.z && in_v.z < max.z)
+    {
+      out_v.x = clamp(in_v.x + color_diff.x, 0.0f, 1.0f);
+      out_v.y = clamp(in_v.y + color_diff.y, 0.0f, 1.0f);
+      out_v.z = clamp(in_v.z + color_diff.z, 0.0f, 1.0f);
+    }
+  else
+    {
+      out_v.xyz = in_v.xyz;
+    }
+
+  out_v.w  = in_v.w;
+  out[gid] = out_v;
+}
diff --git a/opencl/color-exchange.cl.h b/opencl/color-exchange.cl.h
new file mode 100644
index 00000000..99cac4e9
--- /dev/null
+++ b/opencl/color-exchange.cl.h
@@ -0,0 +1,46 @@
+static const char* color_exchange_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 2015 Thomas Manni <thomas.manni@free.fr>                         \n"
+" */                                                                           \n"
+"                                                                              \n"
+"__kernel void cl_color_exchange(__global const float4 *in,                    \n"
+"                                __global       float4 *out,                   \n"
+"                                               float3 color_diff,             \n"
+"                                               float3 min,                    \n"
+"                                               float3 max)                    \n"
+"{                                                                             \n"
+"  int gid     = get_global_id(0);                                             \n"
+"  float4 in_v = in[gid];                                                      \n"
+"  float4 out_v;                                                               \n"
+"                                                                              \n"
+"  if(in_v.x > min.x && in_v.x < max.x &&                                      \n"
+"     in_v.y > min.y && in_v.y < max.y &&                                      \n"
+"     in_v.z > min.z && in_v.z < max.z)                                        \n"
+"    {                                                                         \n"
+"      out_v.x = clamp(in_v.x + color_diff.x, 0.0f, 1.0f);                     \n"
+"      out_v.y = clamp(in_v.y + color_diff.y, 0.0f, 1.0f);                     \n"
+"      out_v.z = clamp(in_v.z + color_diff.z, 0.0f, 1.0f);                     \n"
+"    }                                                                         \n"
+"  else                                                                        \n"
+"    {                                                                         \n"
+"      out_v.xyz = in_v.xyz;                                                   \n"
+"    }                                                                         \n"
+"                                                                              \n"
+"  out_v.w  = in_v.w;                                                          \n"
+"  out[gid] = out_v;                                                           \n"
+"}                                                                             \n"
+;
diff --git a/operations/common/color-exchange.c b/operations/common/color-exchange.c
index 750764be..23d385b4 100644
--- a/operations/common/color-exchange.c
+++ b/operations/common/color-exchange.c
@@ -161,23 +161,82 @@ process (GeglOperation       *operation,
   return TRUE;
 }
 
+#include "opencl/gegl-cl.h"
+#include "opencl/color-exchange.cl.h"
+
+static GeglClRunData *cl_data = NULL;
+
+static gboolean
+cl_process (GeglOperation       *operation,
+            cl_mem              in,
+            cl_mem              out,
+            size_t              global_worksize,
+            const GeglRectangle *roi,
+            gint                level)
+{
+  GeglProperties *o = GEGL_PROPERTIES (operation);
+  CeParamsType   *params = (CeParamsType*) o->user_data;
+  cl_float3   color_diff;
+  cl_float3   min;
+  cl_float3   max;
+  cl_int      cl_err = 0;
+  gint        i;
+
+  if (!cl_data)
+    {
+      const char *kernel_name[] = {"cl_color_exchange",
+                                   NULL};
+      cl_data = gegl_cl_compile_and_build (color_exchange_cl_source, kernel_name);
+    }
+
+  if (!cl_data)
+    return TRUE;
+
+  for (i = 0; i < 3; i++)
+    {
+      color_diff.s[i] = params->color_diff[i];
+      min.s[i] = params->min[i];
+      max.s[i] = params->max[i];
+    }
+
+  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
+                                    sizeof(cl_mem),    &in,
+                                    sizeof(cl_mem),    &out,
+                                    sizeof(cl_float3), &color_diff,
+                                    sizeof(cl_float3), &min,
+                                    sizeof(cl_float3), &max,
+                                    NULL);
+  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;
+
+  return  FALSE;
+
+error:
+  return TRUE;
+}
+
 static void
 gegl_op_class_init (GeglOpClass *klass)
 {
   GObjectClass                  *object_class;
   GeglOperationClass            *operation_class;
-  GeglOperationPointFilterClass *filter_class;
+  GeglOperationPointFilterClass *point_filter_class;
 
-  object_class    = G_OBJECT_CLASS (klass);
-  operation_class = GEGL_OPERATION_CLASS (klass);
-  filter_class    = GEGL_OPERATION_POINT_FILTER_CLASS (klass);
+  object_class       = G_OBJECT_CLASS (klass);
+  operation_class    = GEGL_OPERATION_CLASS (klass);
+  point_filter_class = GEGL_OPERATION_POINT_FILTER_CLASS (klass);
 
   object_class->finalize = finalize;
 
-  operation_class->prepare = prepare;
-  operation_class->opencl_support = FALSE;
+  operation_class->prepare     = prepare;
 
-  filter_class->process    = process;
+  point_filter_class->process    = process;
+  point_filter_class->cl_process = cl_process;
 
   gegl_operation_class_set_keys (operation_class,
     "name",        "gegl:color-exchange",