[gegl] noise-hsv: Add CL implementation and tests
- From: Øyvind Kolås <ok src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] noise-hsv: Add CL implementation and tests
- Date: Mon, 22 May 2017 21:18:07 +0000 (UTC)
commit be32baa1d7f6524159748c7f492cf93fb4b6f822
Author: Anton Gorenko <anton streamcomputing eu>
Date: Thu May 11 11:26:00 2017 +0600
noise-hsv: Add CL implementation and tests
(pippin: removed test and associated (1MB!) test image, before checking in to
GEGL repo, avoiding increasing the initial clone for a full source history of
GEGL by 1mb - such testing is now covered by the reference compositions that
only store the hash of the resulting raster instead of a full raster.)
opencl/noise-hsv.cl | 120 ++++++++++++++++++++++++++++++++++++++++
opencl/noise-hsv.cl.h | 122 +++++++++++++++++++++++++++++++++++++++++
operations/common/noise-hsv.c | 80 +++++++++++++++++++++++++++
3 files changed, 322 insertions(+), 0 deletions(-)
---
diff --git a/opencl/noise-hsv.cl b/opencl/noise-hsv.cl
new file mode 100644
index 0000000..3cff79a
--- /dev/null
+++ b/opencl/noise-hsv.cl
@@ -0,0 +1,120 @@
+/* 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/>.
+ *
+ */
+
+float
+randomize_value (__global const int *random_data,
+ const GeglRandom rand,
+ float now,
+ float min_val,
+ float max_val,
+ bool wraps_around,
+ float rand_max,
+ int holdness,
+ int x,
+ int y,
+ int n)
+{
+ const float steps = max_val - min_val;
+ float rand_val = gegl_cl_random_float (random_data, rand, x, y, 0, n++);
+
+ for (int i = 1; i < holdness; i++)
+ {
+ float tmp = gegl_cl_random_float (random_data, rand, x, y, 0, n++);
+ rand_val = min (rand_val, tmp);
+ }
+
+ const int flag = (gegl_cl_random_float (random_data, rand, x, y, 0, n) < 0.5)
+ ? -1 : 1;
+ float new_val = now + flag * fmod (rand_max * rand_val, steps);
+
+ if (new_val < min_val)
+ {
+ if (wraps_around)
+ new_val += steps;
+ else
+ new_val = min_val;
+ }
+
+ if (max_val < new_val)
+ {
+ if (wraps_around)
+ new_val -= steps;
+ else
+ new_val = max_val;
+ }
+
+ return new_val;
+}
+
+__kernel void cl_noise_hsv(__global const float4 *in,
+ __global float4 *out,
+ __global const int *random_data,
+ const GeglRandom rand,
+ int x_offset,
+ int y_offset,
+ int roi_width,
+ int whole_region_width,
+ int holdness,
+ float hue_distance,
+ float saturation_distance,
+ float value_distance)
+{
+ const int gid = get_global_id(0);
+ const int gidy = gid / roi_width;
+ const int gidx = gid - gidy * roi_width;
+
+ const int x = gidx + x_offset;
+ const int y = gidy + y_offset;
+
+ int n = (3 * holdness + 4) * (x + whole_region_width * y);
+
+ const float4 in_v = in[gid];
+
+ float hue = in_v.s0;
+ float saturation = in_v.s1;
+ float value = in_v.s2;
+ float alpha = in_v.s3;
+
+ /* there is no need for scattering hue of desaturated pixels here */
+ if ((hue_distance > 0) && (saturation > 0))
+ hue = randomize_value (random_data, rand,
+ hue, 0.0, 1.0, true,
+ hue_distance, holdness,
+ x, y, n);
+
+ n += holdness + 1;
+ if (saturation_distance > 0)
+ {
+ /* desaturated pixels get random hue before increasing saturation */
+ if (saturation == 0)
+ hue = gegl_cl_random_float_range (random_data, rand,
+ x, y, 0, n, 0.0, 1.0);
+ saturation = randomize_value (random_data, rand,
+ saturation, 0.0, 1.0, false,
+ saturation_distance, holdness,
+ x, y, n + 1);
+ }
+
+ n += holdness + 2;
+ if (value_distance > 0)
+ value = randomize_value (random_data, rand,
+ value, 0.0, 1.0, false,
+ value_distance, holdness,
+ x, y, n);
+
+ out[gid] = (float4)(hue, saturation, value, alpha);
+}
diff --git a/opencl/noise-hsv.cl.h b/opencl/noise-hsv.cl.h
new file mode 100644
index 0000000..528be47
--- /dev/null
+++ b/opencl/noise-hsv.cl.h
@@ -0,0 +1,122 @@
+static const char* noise_hsv_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"
+" */ \n"
+" \n"
+"float \n"
+"randomize_value (__global const int *random_data, \n"
+" const GeglRandom rand, \n"
+" float now, \n"
+" float min_val, \n"
+" float max_val, \n"
+" bool wraps_around, \n"
+" float rand_max, \n"
+" int holdness, \n"
+" int x, \n"
+" int y, \n"
+" int n) \n"
+"{ \n"
+" const float steps = max_val - min_val; \n"
+" float rand_val = gegl_cl_random_float (random_data, rand, x, y, 0, n++); \n"
+" \n"
+" for (int i = 1; i < holdness; i++) \n"
+" { \n"
+" float tmp = gegl_cl_random_float (random_data, rand, x, y, 0, n++); \n"
+" rand_val = min (rand_val, tmp); \n"
+" } \n"
+" \n"
+" const int flag = (gegl_cl_random_float (random_data, rand, x, y, 0, n) < 0.5)\n"
+" ? -1 : 1; \n"
+" float new_val = now + flag * fmod (rand_max * rand_val, steps); \n"
+" \n"
+" if (new_val < min_val) \n"
+" { \n"
+" if (wraps_around) \n"
+" new_val += steps; \n"
+" else \n"
+" new_val = min_val; \n"
+" } \n"
+" \n"
+" if (max_val < new_val) \n"
+" { \n"
+" if (wraps_around) \n"
+" new_val -= steps; \n"
+" else \n"
+" new_val = max_val; \n"
+" } \n"
+" \n"
+" return new_val; \n"
+"} \n"
+" \n"
+"__kernel void cl_noise_hsv(__global const float4 *in, \n"
+" __global float4 *out, \n"
+" __global const int *random_data, \n"
+" const GeglRandom rand, \n"
+" int x_offset, \n"
+" int y_offset, \n"
+" int roi_width, \n"
+" int whole_region_width, \n"
+" int holdness, \n"
+" float hue_distance, \n"
+" float saturation_distance, \n"
+" float value_distance) \n"
+"{ \n"
+" const int gid = get_global_id(0); \n"
+" const int gidy = gid / roi_width; \n"
+" const int gidx = gid - gidy * roi_width; \n"
+" \n"
+" const int x = gidx + x_offset; \n"
+" const int y = gidy + y_offset; \n"
+" \n"
+" int n = (3 * holdness + 4) * (x + whole_region_width * y); \n"
+" \n"
+" const float4 in_v = in[gid]; \n"
+" \n"
+" float hue = in_v.s0; \n"
+" float saturation = in_v.s1; \n"
+" float value = in_v.s2; \n"
+" float alpha = in_v.s3; \n"
+" \n"
+" /* there is no need for scattering hue of desaturated pixels here */ \n"
+" if ((hue_distance > 0) && (saturation > 0)) \n"
+" hue = randomize_value (random_data, rand, \n"
+" hue, 0.0, 1.0, true, \n"
+" hue_distance, holdness, \n"
+" x, y, n); \n"
+" \n"
+" n += holdness + 1; \n"
+" if (saturation_distance > 0) \n"
+" { \n"
+" /* desaturated pixels get random hue before increasing saturation */ \n"
+" if (saturation == 0) \n"
+" hue = gegl_cl_random_float_range (random_data, rand, \n"
+" x, y, 0, n, 0.0, 1.0); \n"
+" saturation = randomize_value (random_data, rand, \n"
+" saturation, 0.0, 1.0, false, \n"
+" saturation_distance, holdness, \n"
+" x, y, n + 1); \n"
+" } \n"
+" \n"
+" n += holdness + 2; \n"
+" if (value_distance > 0) \n"
+" value = randomize_value (random_data, rand, \n"
+" value, 0.0, 1.0, false, \n"
+" value_distance, holdness, \n"
+" x, y, n); \n"
+" \n"
+" out[gid] = (float4)(hue, saturation, value, alpha); \n"
+"} \n"
+;
diff --git a/operations/common/noise-hsv.c b/operations/common/noise-hsv.c
index fc5455d..d6fe9c2 100644
--- a/operations/common/noise-hsv.c
+++ b/operations/common/noise-hsv.c
@@ -178,6 +178,84 @@ process (GeglOperation *operation,
return TRUE;
}
+#include "opencl/gegl-cl.h"
+#include "opencl/noise-hsv.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);
+ GeglRectangle *wr = gegl_operation_source_get_bounding_box (operation,
+ "input");
+ cl_int cl_err = 0;
+ cl_mem cl_random_data = NULL;
+ cl_int x_offset = roi->x;
+ cl_int y_offset = roi->y;
+ cl_int roi_width = roi->width;
+ cl_int wr_width = wr->width;
+ cl_ushort4 rand;
+ cl_int holdness;
+ cl_float hue_distance;
+ cl_float saturation_distance;
+ cl_float value_distance;
+
+ gegl_cl_random_get_ushort4 (o->rand, &rand);
+
+ if (!cl_data)
+ {
+ const char *kernel_name[] = { "cl_noise_hsv", NULL };
+ cl_data = gegl_cl_compile_and_build (noise_hsv_cl_source, kernel_name);
+ }
+
+ if (!cl_data)
+ return TRUE;
+
+ cl_random_data = gegl_cl_load_random_data (&cl_err);
+ CL_CHECK;
+
+ holdness = o->holdness;
+ hue_distance = o->hue_distance / 360.0;
+ saturation_distance = o->saturation_distance;
+ value_distance = o->value_distance;
+
+ gegl_cl_set_kernel_args (cl_data->kernel[0],
+ sizeof(cl_mem), &in,
+ sizeof(cl_mem), &out,
+ sizeof(cl_mem), &cl_random_data,
+ sizeof(cl_ushort4), &rand,
+ sizeof(cl_int), &x_offset,
+ sizeof(cl_int), &y_offset,
+ sizeof(cl_int), &roi_width,
+ sizeof(cl_int), &wr_width,
+ sizeof(cl_int), &holdness,
+ sizeof(cl_float), &hue_distance,
+ sizeof(cl_float), &saturation_distance,
+ sizeof(cl_float), &value_distance,
+ 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;
+
+ cl_err = gegl_clFinish (gegl_cl_get_command_queue ());
+ CL_CHECK;
+
+ return FALSE;
+
+error:
+ return TRUE;
+}
+
static void
gegl_op_class_init (GeglOpClass *klass)
{
@@ -188,7 +266,9 @@ gegl_op_class_init (GeglOpClass *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-hsv",
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
