[gegl/soc-2013-opecl-ops] Added OpenCL support to noise-slur
- From: Carlos Zubieta <czubieta src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl/soc-2013-opecl-ops] Added OpenCL support to noise-slur
- Date: Wed, 11 Sep 2013 12:24:49 +0000 (UTC)
commit 58b2180a8dbedf3d1f90780da1f5762a10207ab0
Author: Carlos Zubieta <czubieta dev gmail com>
Date: Wed Sep 11 05:52:44 2013 -0500
Added OpenCL support to noise-slur
opencl/noise-slur.cl | 77 +++++++++++++++
opencl/noise-slur.cl.h | 81 ++++++++++++++++
operations/common/noise-slur.c | 208 +++++++++++++++++++++++++++++++++++++++-
3 files changed, 361 insertions(+), 5 deletions(-)
---
diff --git a/opencl/noise-slur.cl b/opencl/noise-slur.cl
new file mode 100644
index 0000000..6b5db77
--- /dev/null
+++ b/opencl/noise-slur.cl
@@ -0,0 +1,77 @@
+/* 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_noise_slur(__global const float4 *src,
+ __global float4 *dst,
+ __global const int *random_data,
+ __global const long *random_primes,
+ int x_offset,
+ int y_offset,
+ int src_width,
+ int whole_region_width,
+ int radius,
+ int seed,
+ float pct_random,
+ int offset)
+{
+ int gidx = get_global_id(0);
+ int gidy = get_global_id(1);
+
+ int x = gidx + x_offset;
+ int y = gidy + y_offset;
+ int n = 2 * (x + whole_region_width * y + offset);
+
+ int src_idx = (gidx + radius) + src_width * (gidy + radius);
+ int dst_idx = gidx + get_global_size(0) * gidy;
+
+ float4 src_v[4] = {src[src_idx],
+ src[src_idx - src_width - 1],
+ src[src_idx - src_width],
+ src[src_idx - src_width + 1]};
+
+ float pct = gegl_cl_random_float_range (random_data, random_primes,
+ seed, x, y, 0, n, 0.0, 100.0);
+ int k = gegl_cl_random_int_range (random_data, random_primes,
+ seed, x, y, 0, n+1, 0, 10);
+
+ float4 tmp = src_v[0];
+ if (pct <= pct_random)
+ {
+ if(k == 0)
+ tmp = src_v[1];
+ else if(k == 9)
+ tmp = src_v[3];
+ else
+ tmp = src_v[2];
+ }
+ dst[dst_idx] = tmp;
+}
+
+__kernel void copy_out_to_aux(__global const float4 *src,
+ __global float4 *dst,
+ int src_width,
+ int radius)
+{
+ int gidx = get_global_id(0);
+ int gidy = get_global_id(1);
+ int dst_idx = (gidx + radius) + src_width * (gidy + radius);
+ int src_idx = gidx + get_global_size(0) * gidy;
+
+ float4 tmp = src[src_idx];
+ dst[dst_idx] = tmp;
+}
diff --git a/opencl/noise-slur.cl.h b/opencl/noise-slur.cl.h
new file mode 100644
index 0000000..2cb200d
--- /dev/null
+++ b/opencl/noise-slur.cl.h
@@ -0,0 +1,81 @@
+/* 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>
+ */
+
+static const char* noise_slur_cl_source =
+"__kernel void cl_noise_slur(__global const float4 *src, \n"
+" __global float4 *dst, \n"
+" __global const int *random_data, \n"
+" __global const long *random_primes, \n"
+" int x_offset, \n"
+" int y_offset, \n"
+" int src_width, \n"
+" int whole_region_width, \n"
+" int radius, \n"
+" int seed, \n"
+" float pct_random, \n"
+" int offset) \n"
+"{ \n"
+" int gidx = get_global_id(0); \n"
+" int gidy = get_global_id(1); \n"
+" \n"
+" int x = gidx + x_offset; \n"
+" int y = gidy + y_offset; \n"
+" int n = 2 * (x + whole_region_width * y + offset); \n"
+" \n"
+" int src_idx = (gidx + radius) + src_width * (gidy + radius); \n"
+" int dst_idx = gidx + get_global_size(0) * gidy; \n"
+" \n"
+" float4 src_v[4] = {src[src_idx], \n"
+" src[src_idx - src_width - 1], \n"
+" src[src_idx - src_width], \n"
+" src[src_idx - src_width + 1]}; \n"
+" \n"
+" float pct = gegl_cl_random_float_range (random_data, random_primes, \n"
+" seed, x, y, 0, n, 0.0, 100.0); \n"
+" int k = gegl_cl_random_int_range (random_data, random_primes, \n"
+" seed, x, y, 0, n+1, 0, 10); \n"
+" \n"
+" float4 tmp = src_v[0]; \n"
+" if (pct <= pct_random) \n"
+" { \n"
+" if(k == 0) \n"
+" tmp = src_v[1]; \n"
+" else if(k == 9) \n"
+" tmp = src_v[3]; \n"
+" else \n"
+" tmp = src_v[2]; \n"
+" } \n"
+" dst[dst_idx] = tmp; \n"
+"} \n"
+" \n"
+" \n"
+"__kernel void copy_out_to_aux(__global const float4 *src, \n"
+" __global float4 *dst, \n"
+" int src_width, \n"
+" int radius) \n"
+"{ \n"
+" int gidx = get_global_id(0); \n"
+" int gidy = get_global_id(1); \n"
+" int dst_idx = (gidx + radius) + src_width * (gidy + radius); \n"
+" int src_idx = gidx + get_global_size(0) * gidy; \n"
+" \n"
+" float4 tmp = src[src_idx]; \n"
+" dst[dst_idx] = tmp; \n"
+"} \n"
+;
+
diff --git a/operations/common/noise-slur.c b/operations/common/noise-slur.c
index 9d59d12..588bdda 100644
--- a/operations/common/noise-slur.c
+++ b/operations/common/noise-slur.c
@@ -63,6 +63,181 @@ prepare (GeglOperation *operation)
gegl_operation_set_format (operation, "output", babl_format ("RGBA float"));
}
+#include "opencl/gegl-cl.h"
+#include "buffer/gegl-buffer-cl-iterator.h"
+#include "gegl/gegl-config.h"
+#include "opencl/noise-slur.cl.h"
+
+GEGL_CL_STATIC
+
+static gboolean
+cl_noise_slur (cl_mem in_tex,
+ cl_mem aux_tex,
+ cl_mem out_tex,
+ const GeglRectangle *src_roi,
+ const GeglRectangle *roi,
+ const GeglRectangle *wr,
+ gint radius,
+ gint seed,
+ gfloat pct_random,
+ gint repeat)
+{
+ cl_int cl_err = 0;
+ cl_mem cl_random_data = NULL;
+ cl_mem cl_random_primes = NULL;
+ const size_t gbl_size[2] = {roi->width, roi->height};
+ const size_t src_size = src_roi->width*src_roi->height;
+
+ GEGL_CL_BUILD(noise_slur, "cl_noise_slur", "copy_out_to_aux")
+
+ cl_random_data = gegl_cl_load_random_data(&cl_err);
+ CL_CHECK;
+ cl_random_primes = gegl_cl_load_random_primes(&cl_err);
+ CL_CHECK;
+
+ {
+ cl_int cl_roi_x = roi->x;
+ cl_int cl_roi_y = roi->y;
+ cl_int cl_src_width = src_roi->width;
+ cl_int cl_wr_width = wr->width;
+ cl_int cl_radius = radius;
+ cl_int cl_seed = seed;
+ cl_float cl_pct_random = pct_random;
+ cl_int offset;
+ int wr_size = wr->width*wr->height;
+ int it;
+
+ cl_err = gegl_clEnqueueCopyBuffer(gegl_cl_get_command_queue(),
+ in_tex , aux_tex , 0 , 0 ,
+ src_size * sizeof(cl_float4),
+ 0, NULL, NULL);
+ CL_CHECK;
+
+ /*Arguments for the first kernel*/
+ GEGL_CL_ARG_START(cl_data->kernel[0])
+ GEGL_CL_ARG(cl_mem, aux_tex)
+ GEGL_CL_ARG(cl_mem, out_tex)
+ GEGL_CL_ARG(cl_mem, cl_random_data)
+ GEGL_CL_ARG(cl_mem, cl_random_primes)
+ GEGL_CL_ARG(cl_int, cl_roi_x)
+ GEGL_CL_ARG(cl_int, cl_roi_y)
+ GEGL_CL_ARG(cl_int, cl_src_width)
+ GEGL_CL_ARG(cl_int, cl_wr_width)
+ GEGL_CL_ARG(cl_int, cl_radius)
+ GEGL_CL_ARG(cl_int, cl_seed)
+ GEGL_CL_ARG(cl_float, cl_pct_random)
+ GEGL_CL_ARG_END
+
+ /*Arguments for the second kernel*/
+ GEGL_CL_ARG_START(cl_data->kernel[1])
+ GEGL_CL_ARG(cl_mem, out_tex)
+ GEGL_CL_ARG(cl_mem, aux_tex)
+ GEGL_CL_ARG(cl_int, cl_src_width)
+ GEGL_CL_ARG(cl_int, cl_radius)
+ GEGL_CL_ARG_END
+
+ offset = 0;
+
+ for(it = 0; it < repeat; ++it)
+ {
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 11, sizeof(cl_int),
+ (void*)&offset);
+ CL_CHECK;
+ cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+ cl_data->kernel[0], 2,
+ NULL, gbl_size, NULL,
+ 0, NULL, NULL);
+ CL_CHECK;
+ if(it < repeat -1) /*No need for the last copy*/
+ {
+ cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+ cl_data->kernel[1], 2,
+ NULL, gbl_size, NULL,
+ 0, NULL, NULL);
+ CL_CHECK;
+ }
+
+ offset += wr_size;
+ }
+
+ cl_err = gegl_clFinish(gegl_cl_get_command_queue ());
+ CL_CHECK;
+
+ GEGL_CL_RELEASE(cl_random_data)
+ GEGL_CL_RELEASE(cl_random_primes)
+ }
+
+ return FALSE;
+
+error:
+ if(cl_random_data)
+ GEGL_CL_RELEASE(cl_random_data)
+ if(cl_random_primes)
+ GEGL_CL_RELEASE(cl_random_primes)
+ return TRUE;
+}
+
+static gboolean
+cl_process (GeglOperation *operation,
+ GeglBuffer *input,
+ GeglBuffer *output,
+ const GeglRectangle *result)
+{
+ const Babl *in_format = gegl_operation_get_format (operation, "input");
+ const Babl *out_format = gegl_operation_get_format (operation, "output");
+ GeglRectangle *wr = gegl_operation_source_get_bounding_box (operation,
+ "input");
+
+ gint err;
+
+ GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation);
+ GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
+
+ GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output,
+ result,
+ out_format,
+ GEGL_CL_BUFFER_WRITE);
+
+ gint read = gegl_buffer_cl_iterator_add_2 (i,
+ input,
+ result,
+ in_format,
+ GEGL_CL_BUFFER_READ,
+ op_area->left,
+ op_area->right,
+ op_area->top,
+ op_area->bottom,
+ GEGL_ABYSS_NONE);
+
+ gint aux = gegl_buffer_cl_iterator_add_2 (i,
+ NULL,
+ result,
+ in_format,
+ GEGL_CL_BUFFER_AUX,
+ op_area->left,
+ op_area->right,
+ op_area->top,
+ op_area->bottom,
+ GEGL_ABYSS_NONE);
+
+ GEGL_CL_BUFFER_ITERATE_START(i, err)
+ {
+ err = cl_noise_slur (i->tex[read],
+ i->tex[aux],
+ i->tex[0],
+ &i->roi[read],
+ &i->roi[0],
+ wr,
+ 1,
+ o->seed,
+ o->pct_random,
+ o->repeat);
+ }
+ GEGL_CL_BUFFER_ITERATE_END(err)
+
+ return TRUE;
+}
+
static gboolean
process (GeglOperation *operation,
GeglBuffer *input,
@@ -80,9 +255,20 @@ process (GeglOperation *operation,
gint n_pixels = result->width * result->height;
gint width = result->width;
GeglRectangle src_rect;
+ GeglRectangle whole_region;
gint total_pixels;
+ gint whole_region_size;
+ gint offset;
gint i;
+ if (gegl_cl_is_accelerated ())
+ {
+ if (cl_process (operation, input, output, result))
+ return TRUE;
+ else
+ gegl_cl_disable();
+ }
+
tmp = gegl_buffer_new (result, babl_format ("RGBA float"));
src_rect.x = result->x - op_area->left;
@@ -97,13 +283,17 @@ process (GeglOperation *operation,
gegl_buffer_copy (input, NULL, tmp, NULL);
+ whole_region = *(gegl_operation_source_get_bounding_box (operation,
+ "input"));
+ whole_region_size = whole_region.width*whole_region.height;
+ offset = 0;
+
for (i = 0; i < o->repeat; i++)
{
- gint x, y, n;
+ gint x, y;
x = result->x;
y = result->y;
- n = 0;
n_pixels = result->width * result->height;
@@ -116,12 +306,16 @@ process (GeglOperation *operation,
while (n_pixels--)
{
- gint b;
+ gint b, n, idx;
+
+ /* n is independent from the roi, but from the whole image*/
+ idx = x + whole_region.width * y;
+ n = 2 * (idx + offset);
- if (gegl_random_float_range (o->seed, x, y, 0, n++, 0.0, 100.0) <=
+ if (gegl_random_float_range (o->seed, x, y, 0, n, 0.0, 100.0) <=
o->pct_random)
{
- gint k = gegl_random_int_range (o->seed, x, y, 0, n++, 0, 10);
+ gint k = gegl_random_int_range (o->seed, x, y, 0, n+1, 0, 10);
for (b = 0; b < 4; b++)
{
@@ -162,6 +356,8 @@ process (GeglOperation *operation,
}
}
+ offset += whole_region_size;
+
gegl_buffer_set (tmp, result, 0,
babl_format ("RGBA float"), dst_buf,
GEGL_AUTO_ROWSTRIDE);
@@ -189,6 +385,8 @@ gegl_chant_class_init (GeglChantClass *klass)
operation_class->prepare = prepare;
filter_class->process = process;
+ operation_class->opencl_support = TRUE;
+
gegl_operation_class_set_keys (operation_class,
"name", "gegl:noise-slur",
"categories", "noise",
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
