[gegl] Added OpenCL support to noise-hurl
- From: Téo Mazars <teom src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] Added OpenCL support to noise-hurl
- Date: Thu, 31 Oct 2013 11:01:20 +0000 (UTC)
commit cdd84b4e07c63436fdacd4ddd667533c8edda551
Author: Carlos Zubieta <czubieta dev gmail com>
Date: Wed Sep 11 05:51:18 2013 -0500
Added OpenCL support to noise-hurl
opencl/noise-hurl.cl | 56 ++++++++++++++++++
opencl/noise-hurl.cl.h | 58 ++++++++++++++++++
operations/common/noise-hurl.c | 125 +++++++++++++++++++++++++++++++++++----
3 files changed, 226 insertions(+), 13 deletions(-)
---
diff --git a/opencl/noise-hurl.cl b/opencl/noise-hurl.cl
new file mode 100644
index 0000000..6cc0412
--- /dev/null
+++ b/opencl/noise-hurl.cl
@@ -0,0 +1,56 @@
+/* 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_hurl(__global float4 *src,
+ __global const int *random_data,
+ __global const long *random_primes,
+ int x_offset,
+ int y_offset,
+ int roi_width,
+ int whole_region_width,
+ int seed,
+ float pct_random,
+ int offset)
+{
+ int gid = get_global_id(0);
+ int gidy = gid / roi_width;
+ int gidx = gid - gidy*roi_width;
+
+ int x = gidx + x_offset;
+ int y = gidy + y_offset;
+ int n = 4 * (x + whole_region_width * y + offset);
+
+ float4 src_v = src[gid];
+
+ float pc = gegl_cl_random_float_range (random_data, random_primes,
+ seed, x, y, 0, n, 0, 100);
+ float red_noise = gegl_cl_random_float (random_data, random_primes,
+ seed, x, y, 0, n+1);
+ float green_noise = gegl_cl_random_float (random_data, random_primes,
+ seed, x, y, 0, n+2);
+ float blue_noise = gegl_cl_random_float (random_data, random_primes,
+ seed, x, y, 0, n+3);
+
+ if(pc <= pct_random)
+ {
+ src_v.x = red_noise;
+ src_v.y = green_noise;
+ src_v.z = blue_noise;
+ }
+ src[gid] = src_v;
+}
diff --git a/opencl/noise-hurl.cl.h b/opencl/noise-hurl.cl.h
new file mode 100644
index 0000000..9cc9f78
--- /dev/null
+++ b/opencl/noise-hurl.cl.h
@@ -0,0 +1,58 @@
+static const char* noise_hurl_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_noise_hurl(__global float4 *src, \n"
+" __global const int *random_data, \n"
+" __global const long *random_primes, \n"
+" int x_offset, \n"
+" int y_offset, \n"
+" int roi_width, \n"
+" int whole_region_width, \n"
+" int seed, \n"
+" float pct_random, \n"
+" int offset) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" int gidy = gid / roi_width; \n"
+" int gidx = gid - gidy*roi_width; \n"
+" \n"
+" int x = gidx + x_offset; \n"
+" int y = gidy + y_offset; \n"
+" int n = 4 * (x + whole_region_width * y + offset); \n"
+" \n"
+" float4 src_v = src[gid]; \n"
+" \n"
+" float pc = gegl_cl_random_float_range (random_data, random_primes, \n"
+" seed, x, y, 0, n, 0, 100); \n"
+" float red_noise = gegl_cl_random_float (random_data, random_primes, \n"
+" seed, x, y, 0, n+1); \n"
+" float green_noise = gegl_cl_random_float (random_data, random_primes, \n"
+" seed, x, y, 0, n+2); \n"
+" float blue_noise = gegl_cl_random_float (random_data, random_primes, \n"
+" seed, x, y, 0, n+3); \n"
+" \n"
+" if(pc <= pct_random) \n"
+" { \n"
+" src_v.x = red_noise; \n"
+" src_v.y = green_noise; \n"
+" src_v.z = blue_noise; \n"
+" } \n"
+" src[gid] = src_v; \n"
+"} \n"
+;
diff --git a/operations/common/noise-hurl.c b/operations/common/noise-hurl.c
index b047ca0..f2a88a3 100644
--- a/operations/common/noise-hurl.c
+++ b/operations/common/noise-hurl.c
@@ -61,12 +61,14 @@ process (GeglOperation *operation,
const GeglRectangle *roi,
gint level)
{
- GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
- gfloat *GEGL_ALIGNED in_pixel;
- gfloat *GEGL_ALIGNED out_pixel;
- gfloat *out_pix;
- gint i, cnt;
- gint x, y, n;
+ GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
+ gfloat *GEGL_ALIGNED in_pixel;
+ gfloat *GEGL_ALIGNED out_pixel;
+ gfloat *out_pix;
+ GeglRectangle *whole_region;
+ gint total_size;
+ gint i, cnt;
+ gint offset;
in_pixel = in_buf;
out_pixel = out_buf;
@@ -79,32 +81,37 @@ process (GeglOperation *operation,
in_pixel += 1;
}
- n = 0;
+ whole_region = gegl_operation_source_get_bounding_box (operation, "input");
+ total_size = whole_region->width*whole_region->height;
+ offset = 0;
for (cnt = 0; cnt < o->repeat; cnt++)
{
- x = roi->x;
- y = roi->y;
+ gint x = roi->x;
+ gint y = roi->y;
out_pix = out_pixel;
for (i = 0; i < n_pixels; i++)
{
gfloat red, green, blue, alpha;
+ gint idx, n;
red = out_pix[0];
green = out_pix[1];
blue = out_pix[2];
alpha = out_pix[3];
+ idx = x + whole_region->width * y;
+ n = 4 * (idx + offset);
+
if (gegl_random_float_range (o->seed, x, y, 0, n, 0.0, 100.0) <=
o->pct_random)
{
- red = gegl_random_float_range (o->seed, x, y, 0, n+1, 0.0, 1.0);
- green = gegl_random_float_range (o->seed, x, y, 0, n+2, 0.0, 1.0);
- blue = gegl_random_float_range (o->seed, x, y, 0, n+3, 0.0, 1.0);
+ red = gegl_random_float (o->seed, x, y, 0, n+1);
+ green = gegl_random_float (o->seed, x, y, 0, n+2);
+ blue = gegl_random_float (o->seed, x, y, 0, n+3);
}
- n += 4;
out_pix[0] = red;
out_pix[1] = green;
@@ -120,11 +127,101 @@ process (GeglOperation *operation,
y++;
}
}
+ offset += total_size;
}
return TRUE;
}
+#include "opencl/gegl-cl.h"
+#include "opencl/noise-hurl.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);
+ GeglRectangle *wr = gegl_operation_source_get_bounding_box (operation,
+ "input");
+ cl_int cl_err = 0;
+ cl_mem cl_random_data = NULL;
+ cl_mem cl_random_primes = NULL;
+ cl_float pct_random = o->pct_random;
+ cl_int seed = o->seed;
+ cl_int x_offset = roi->x;
+ cl_int y_offset = roi->y;
+ cl_int roi_width = roi->width;
+ cl_int wr_width = wr->width;
+ int total_size = wr->width*wr->height;
+ cl_int offset;
+ int it;
+
+ GEGL_CL_BUILD(noise_hurl, "cl_noise_hurl")
+
+ {
+ 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_err = gegl_clEnqueueCopyBuffer(gegl_cl_get_command_queue(),
+ in , out , 0 , 0 ,
+ global_worksize * sizeof(cl_float4),
+ 0, NULL, NULL);
+ CL_CHECK;
+
+ GEGL_CL_ARG_START(cl_data->kernel[0])
+ GEGL_CL_ARG(cl_mem, out)
+ GEGL_CL_ARG(cl_mem, cl_random_data)
+ GEGL_CL_ARG(cl_mem, cl_random_primes)
+ GEGL_CL_ARG(cl_int, x_offset)
+ GEGL_CL_ARG(cl_int, y_offset)
+ GEGL_CL_ARG(cl_int, roi_width)
+ GEGL_CL_ARG(cl_int, wr_width)
+ GEGL_CL_ARG(cl_int, seed)
+ GEGL_CL_ARG(cl_float, pct_random)
+ GEGL_CL_ARG_END
+
+ offset = 0;
+
+ for(it = 0; it < o->repeat; ++it)
+ {
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 9, sizeof(cl_int),
+ (void*)&offset);
+ 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;
+
+ offset += total_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 void
gegl_chant_class_init (GeglChantClass *klass)
{
@@ -135,7 +232,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:noise-hurl",
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
