[gegl] color-exchange: add opencl support
- From: Thomas Manni <tmanni src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] color-exchange: add opencl support
- Date: Tue, 7 Jul 2015 09:32:06 +0000 (UTC)
commit d0bfae0970c6a3b8934600ec09e5ef121cc117e0
Author: Thomas Manni <thomas manni free fr>
Date: Mon Jul 6 23:50:36 2015 +0200
color-exchange: add opencl support
opencl/color-exchange.cl | 44 +++++++++++++++++++++
opencl/color-exchange.cl.h | 46 ++++++++++++++++++++++
operations/common/color-exchange.c | 73 ++++++++++++++++++++++++++++++++---
3 files changed, 156 insertions(+), 7 deletions(-)
---
diff --git a/opencl/color-exchange.cl b/opencl/color-exchange.cl
new file mode 100644
index 0000000..2a6be6c
--- /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 0000000..99cac4e
--- /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 750764b..23d385b 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",
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
