summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
-rw-r--r--opencl/posterize.cl28
-rw-r--r--opencl/posterize.cl.h30
-rw-r--r--operations/common/posterize.c44
3 files changed, 101 insertions, 1 deletions
diff --git a/opencl/posterize.cl b/opencl/posterize.cl
new file mode 100644
index 00000000..5ba9e8b2
--- /dev/null
+++ b/opencl/posterize.cl
@@ -0,0 +1,28 @@
+/* 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>
+ */
+
+__kernel void cl_posterize(__global const float4 *in,
+ __global float4 *out,
+ float levels)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+
+ in_v.xyz = trunc(in_v.xyz * levels + (float3)(0.5f)) / levels;
+ out[gid] = in_v;
+}
diff --git a/opencl/posterize.cl.h b/opencl/posterize.cl.h
new file mode 100644
index 00000000..e7717c59
--- /dev/null
+++ b/opencl/posterize.cl.h
@@ -0,0 +1,30 @@
+static const char* posterize_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"
+"__kernel void cl_posterize(__global const float4 *in, \n"
+" __global float4 *out, \n"
+" float levels) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" \n"
+" in_v.xyz = trunc(in_v.xyz * levels + (float3)(0.5f)) / levels; \n"
+" out[gid] = in_v; \n"
+"} \n"
+;
diff --git a/operations/common/posterize.c b/operations/common/posterize.c
index 0fcca178..f5a15924 100644
--- a/operations/common/posterize.c
+++ b/operations/common/posterize.c
@@ -64,6 +64,46 @@ static gboolean process (GeglOperation *operation,
return TRUE;
}
+#include "opencl/gegl-cl.h"
+#include "opencl/posterize.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 levels = o->levels;
+
+ GEGL_CL_BUILD(posterize, "cl_posterize")
+
+ {
+ 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, levels)
+ 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)
{
@@ -73,7 +113,9 @@ gegl_chant_class_init (GeglChantClass *klass)
operation_class = GEGL_OPERATION_CLASS (klass);
point_filter_class = GEGL_OPERATION_POINT_FILTER_CLASS (klass);
- point_filter_class->process = process;
+ 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:posterize",