[gegl] Added OpenCL support to box-min
- From: Téo Mazars <teom src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] Added OpenCL support to box-min
- Date: Thu, 31 Oct 2013 11:00:39 +0000 (UTC)
commit 13b09cab60952d79e1c99b27e91f3f0304bd6647
Author: Carlos Zubieta <czubieta dev gmail com>
Date: Tue Jul 23 18:47:59 2013 -0500
Added OpenCL support to box-min
opencl/box-min.cl | 49 ++++++++++++++
opencl/box-min.cl.h | 51 ++++++++++++++
operations/workshop/box-min.c | 144 ++++++++++++++++++++++++++++++++++++++++-
3 files changed, 242 insertions(+), 2 deletions(-)
---
diff --git a/opencl/box-min.cl b/opencl/box-min.cl
new file mode 100644
index 0000000..f0978de
--- /dev/null
+++ b/opencl/box-min.cl
@@ -0,0 +1,49 @@
+__kernel void kernel_min_hor (__global const float4 *in,
+ __global float4 *aux,
+ int width, int radius)
+{
+ const int in_index = get_global_id(0) * (width + 2 * radius)
+ + (radius + get_global_id (1));
+
+ const int aux_index = get_global_id(0) * width + get_global_id (1);
+ int i;
+ float4 min;
+ float4 in_v;
+
+ min = (float4)(1000000000.0f);
+
+ if (get_global_id(1) < width)
+ {
+ for (i=-radius; i <= radius; i++)
+ {
+ in_v = in[in_index + i];
+ min = min > in_v ? in_v : min;
+ }
+ aux[aux_index] = min;
+ }
+}
+
+__kernel void kernel_max_ver (__global const float4 *aux,
+ __global float4 *out,
+ int width, int radius)
+{
+
+ const int out_index = get_global_id(0) * width + get_global_id (1);
+ int aux_index = out_index;
+ int i;
+ float4 min;
+ float4 aux_v;
+
+ min = (float4)(1000000000.0f);
+
+ if(get_global_id(1) < width)
+ {
+ for (i=-radius; i <= radius; i++)
+ {
+ aux_v = aux[aux_index];
+ min = min > aux_v ? aux_v : min;
+ aux_index += width;
+ }
+ out[out_index] = min;
+ }
+}
diff --git a/opencl/box-min.cl.h b/opencl/box-min.cl.h
new file mode 100644
index 0000000..1233258
--- /dev/null
+++ b/opencl/box-min.cl.h
@@ -0,0 +1,51 @@
+static const char* box_min_cl_source =
+"__kernel void kernel_min_hor (__global const float4 *in, \n"
+" __global float4 *aux, \n"
+" int width, int radius) \n"
+"{ \n"
+" const int in_index = get_global_id(0) * (width + 2 * radius) \n"
+" + (radius + get_global_id (1)); \n"
+" \n"
+" const int aux_index = get_global_id(0) * width + get_global_id (1); \n"
+" int i; \n"
+" float4 min; \n"
+" float4 in_v; \n"
+" \n"
+" min = (float4)(1000000000.0f); \n"
+" \n"
+" if (get_global_id(1) < width) \n"
+" { \n"
+" for (i=-radius; i <= radius; i++) \n"
+" { \n"
+" in_v = in[in_index + i]; \n"
+" min = min > in_v ? in_v : min; \n"
+" } \n"
+" aux[aux_index] = min; \n"
+" } \n"
+"} \n"
+" \n"
+"__kernel void kernel_min_ver (__global const float4 *aux, \n"
+" __global float4 *out, \n"
+" int width, int radius) \n"
+"{ \n"
+" \n"
+" const int out_index = get_global_id(0) * width + get_global_id (1); \n"
+" int aux_index = out_index; \n"
+" int i; \n"
+" float4 min; \n"
+" float4 aux_v; \n"
+" \n"
+" min = (float4)(1000000000.0f); \n"
+" \n"
+" if(get_global_id(1) < width) \n"
+" { \n"
+" for (i=-radius; i <= radius; i++) \n"
+" { \n"
+" aux_v = aux[aux_index]; \n"
+" min = min > aux_v ? aux_v : min; \n"
+" aux_index += width; \n"
+" } \n"
+" out[out_index] = min; \n"
+" } \n"
+"} \n"
+;
diff --git a/operations/workshop/box-min.c b/operations/workshop/box-min.c
index 529a411..16aec0d 100644
--- a/operations/workshop/box-min.c
+++ b/operations/workshop/box-min.c
@@ -162,6 +162,136 @@ static void 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 "opencl/box-min.cl.h"
+
+static GeglClRunData *cl_data = NULL;
+
+static gboolean
+cl_box_min (cl_mem in_tex,
+ cl_mem aux_tex,
+ cl_mem out_tex,
+ size_t global_worksize,
+ const GeglRectangle *roi,
+ gint radius)
+{
+ cl_int cl_err = 0;
+ size_t global_ws_hor[2], global_ws_ver[2];
+ size_t local_ws_hor[2], local_ws_ver[2];
+
+ if (!cl_data)
+ {
+ const char *kernel_name[] = {"kernel_min_hor", "kernel_min_ver", NULL};
+ cl_data = gegl_cl_compile_and_build (box_min_cl_source, kernel_name);
+ }
+ if (!cl_data) return TRUE;
+
+ local_ws_hor[0] = 1;
+ local_ws_hor[1] = 256;
+ global_ws_hor[0] = roi->height + 2 * radius;
+ global_ws_hor[1] = ((roi->width + local_ws_hor[1] -1)/local_ws_hor[1]) * local_ws_hor[1];
+
+ local_ws_ver[0] = 1;
+ local_ws_ver[1] = 256;
+ global_ws_ver[0] = roi->height;
+ global_ws_ver[1] = ((roi->width + local_ws_ver[1] -1)/local_ws_ver[1]) * local_ws_ver[1];
+
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex);
+ CL_CHECK;
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&aux_tex);
+ CL_CHECK;
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int), (void*)&roi->width);
+ CL_CHECK;
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_int), (void*)&radius);
+ CL_CHECK;
+
+ cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+ cl_data->kernel[0], 2,
+ NULL, global_ws_hor, local_ws_hor,
+ 0, NULL, NULL);
+ CL_CHECK;
+
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem), (void*)&aux_tex);
+ CL_CHECK;
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_mem), (void*)&out_tex);
+ CL_CHECK;
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 2, sizeof(cl_int), (void*)&roi->width);
+ CL_CHECK;
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 3, sizeof(cl_int), (void*)&radius);
+ CL_CHECK;
+
+ cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+ cl_data->kernel[1], 2,
+ NULL, global_ws_ver, local_ws_ver,
+ 0, NULL, NULL);
+ CL_CHECK;
+
+ return FALSE;
+
+error:
+ 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");
+
+ 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,
+ 0,
+ 0,
+ op_area->top,
+ op_area->bottom,
+ GEGL_ABYSS_NONE);
+
+ while (gegl_buffer_cl_iterator_next (i, &err))
+ {
+ if (err) return FALSE;
+
+ err = cl_box_min(i->tex[read],
+ i->tex[aux],
+ i->tex[0],
+ i->size[0],
+ &i->roi[0],
+ ceil (o->radius));
+
+ if (err) return FALSE;
+ }
+
+ return TRUE;
+}
+
static gboolean
process (GeglOperation *operation,
GeglBuffer *input,
@@ -172,8 +302,16 @@ process (GeglOperation *operation,
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
GeglRectangle input_rect = gegl_operation_get_required_for_output (operation, "input", result);
- hor_min ( input, &input_rect, output, result, o->radius);
- ver_min (output, result, output, result, o->radius);
+ if (gegl_cl_is_accelerated ())
+ {
+ if (cl_process (operation, input, output, result))
+ return TRUE;
+ else
+ gegl_cl_disable();
+ }
+
+ hor_mim ( input, &input_rect, output, result, o->radius);
+ ver_mim (output, result, output, result, o->radius);
return TRUE;
}
@@ -191,6 +329,8 @@ gegl_chant_class_init (GeglChantClass *klass)
filter_class->process = process;
operation_class->prepare = prepare;
+ operation_class->opencl_support = TRUE;
+
gegl_operation_class_set_keys (operation_class,
"name" , "gegl:box-min",
"categories" , "misc",
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
