[gegl/opencl-ops: 12/14] Add OpenCL support for noise-reduction
- From: Ãyvind KolÃs <ok src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl/opencl-ops: 12/14] Add OpenCL support for noise-reduction
- Date: Tue, 20 Mar 2012 13:50:47 +0000 (UTC)
commit 800286b9650130b8979752a9de0fccdebc7aea4c
Author: Victor Oliveira <victormatheus gmail com>
Date: Sun Mar 18 01:32:12 2012 -0300
Add OpenCL support for noise-reduction
operations/common/noise-reduction.c | 222 +++++++++++++++++++++++++++++++++++
1 files changed, 222 insertions(+), 0 deletions(-)
---
diff --git a/operations/common/noise-reduction.c b/operations/common/noise-reduction.c
index 5ef4a52..c65bec1 100644
--- a/operations/common/noise-reduction.c
+++ b/operations/common/noise-reduction.c
@@ -15,6 +15,8 @@
*
* Ali Alsam, Hans Jakob Rivertz, Ãyvind KolÃs (c) 2011
*/
+#include "config.h"
+#include <glib/gi18n-lib.h>
#ifdef GEGL_CHANT_PROPERTIES
@@ -144,6 +146,219 @@ static void prepare (GeglOperation *operation)
gegl_operation_set_format (operation, "output", babl_format ("R'G'B'A float"));
}
+#include "opencl/gegl-cl.h"
+#include "buffer/gegl-buffer-cl-iterator.h"
+
+static const char* kernel_source =
+"#define NEIGHBOURS 8 \n"
+"#define AXES (NEIGHBOURS/2) \n"
+" \n"
+"#define POW2(a) ((a)*(a)) \n"
+" \n"
+"#define GEN_METRIC(before, center, after) POW2((center) * 2 - (before) - (after)) \n"
+" \n"
+"#define BAIL_CONDITION(new,original) ((new) < (original)) \n"
+" \n"
+"#define SYMMETRY(a) (NEIGHBOURS - (a) - 1) \n"
+" \n"
+"#define O(u,v) (((u)+((v) * (src_stride)))) \n"
+" \n"
+"__kernel void noise_reduction_cl (__global float4 *src_buf, \n"
+" int src_stride, \n"
+" __global float4 *dst_buf, \n"
+" int dst_stride) \n"
+"{ \n"
+" int gidx = get_global_id(0); \n"
+" int gidy = get_global_id(1); \n"
+" \n"
+" __global float4 *center_pix = src_buf + (gidy + 1) * src_stride + gidx + 1; \n"
+" int dst_offset = dst_stride * gidy + gidx; \n"
+" \n"
+" int offsets[NEIGHBOURS] = { \n"
+" O(-1, -1), O( 0, -1), O( 1, -1), \n"
+" O(-1, 0), O( 1, 0), \n"
+" O(-1, 1), O( 0, 1), O( 1, 1) \n"
+" }; \n"
+" \n"
+" float4 sum; \n"
+" int4 count; \n"
+" float4 cur; \n"
+" float4 metric_reference[AXES]; \n"
+" \n"
+" for (int axis = 0; axis < AXES; axis++) \n"
+" { \n"
+" float4 before_pix = *(center_pix + offsets[axis]); \n"
+" float4 after_pix = *(center_pix + offsets[SYMMETRY(axis)]); \n"
+" metric_reference[axis] = GEN_METRIC (before_pix, *center_pix, after_pix); \n"
+" } \n"
+" \n"
+" cur = sum = *center_pix; \n"
+" count = 1; \n"
+" \n"
+" for (int direction = 0; direction < NEIGHBOURS; direction++) \n"
+" { \n"
+" float4 pix = *(center_pix + offsets[direction]); \n"
+" float4 value = (pix + cur) * (0.5f); \n"
+" int axis; \n"
+" int4 mask = {1, 1, 1, 0}; \n"
+" \n"
+" for (axis = 0; axis < AXES; axis++) \n"
+" { \n"
+" float4 before_pix = *(center_pix + offsets[axis]); \n"
+" float4 after_pix = *(center_pix + offsets[SYMMETRY(axis)]); \n"
+" \n"
+" float4 metric_new = GEN_METRIC (before_pix, \n"
+" value, \n"
+" after_pix); \n"
+" mask = BAIL_CONDITION (metric_new, metric_reference[axis]) & mask; \n"
+" } \n"
+" sum += mask >0 ? value : 0; \n"
+" count += mask >0 ? 1 : 0; \n"
+" } \n"
+" dst_buf[dst_offset] = (sum/convert_float4(count)); \n"
+" dst_buf[dst_offset].w = cur.w; \n"
+"} \n"
+"__kernel void transfer(__global float4 * in, \n"
+" int in_width, \n"
+" __global float4 * out) \n"
+"{ \n"
+" int gidx = get_global_id(0); \n"
+" int gidy = get_global_id(1); \n"
+" int width = get_global_size(0); \n"
+" out[gidy * width + gidx] = in[gidy * in_width + gidx]; \n"
+"} \n";
+
+static gegl_cl_run_data *cl_data = NULL;
+
+static cl_int
+cl_noise_reduction (cl_mem in_tex,
+ cl_mem aux_tex,
+ cl_mem out_tex,
+ size_t global_worksize,
+ const GeglRectangle *src_roi,
+ const GeglRectangle *roi,
+ const int iterations)
+{
+ int i = 0;
+ size_t gbl_size_tmp[2];
+
+ cl_int n_src_stride = roi->width + iterations * 2;
+ cl_int cl_err = 0;
+
+ cl_mem temp_tex;
+
+ gint stride = 16; /*R'G'B'A float*/
+
+ if (!cl_data)
+ {
+ const char *kernel_name[] ={"noise_reduction_cl","transfer", NULL};
+ cl_data = gegl_cl_compile_and_build(kernel_source, kernel_name);
+ }
+ if (!cl_data) return 0;
+
+ temp_tex = gegl_clCreateBuffer (gegl_cl_get_context(),
+ CL_MEM_READ_WRITE,
+ src_roi->width * src_roi->height * stride,
+ NULL, &cl_err);
+ if (cl_err != CL_SUCCESS) return cl_err;
+
+
+ cl_err = gegl_clEnqueueCopyBuffer(gegl_cl_get_command_queue(),
+ in_tex , temp_tex , 0 , 0 ,
+ src_roi->width * src_roi->height * stride,
+ NULL, NULL, NULL);
+
+ cl_err = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
+ if (CL_SUCCESS != cl_err) return cl_err;
+
+ cl_mem tmptex = temp_tex;
+ for (i = 0;i<iterations;i++)
+ {
+ if (i > 0)
+ {
+ cl_mem temp = aux_tex;
+ aux_tex = temp_tex;
+ temp_tex = temp;
+ }
+ gbl_size_tmp[0] = roi->width + 2 * (iterations - 1 -i);
+ gbl_size_tmp[1] = roi->height + 2 * (iterations - 1 -i);
+
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&temp_tex);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_int), (void*)&n_src_stride);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_mem), (void*)&aux_tex);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_int), (void*)&n_src_stride);
+ if (cl_err != CL_SUCCESS) return cl_err;
+
+ cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(), cl_data->kernel[0],
+ 2, NULL, gbl_size_tmp, NULL,
+ 0, NULL, NULL);
+ cl_err = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
+ if (CL_SUCCESS != cl_err) return cl_err;
+ }
+
+ gbl_size_tmp[0] = roi->width ;
+ gbl_size_tmp[1] = roi->height;
+
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem), (void*)&aux_tex);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_int), (void*)&n_src_stride);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 2, sizeof(cl_mem), (void*)&out_tex);
+ if (cl_err != CL_SUCCESS) return cl_err;
+
+ cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(), cl_data->kernel[1],
+ 2, NULL, gbl_size_tmp, NULL,
+ 0, NULL, NULL);
+
+ cl_err = gegl_clFinish(gegl_cl_get_command_queue());
+ if (CL_SUCCESS != cl_err) return cl_err;
+
+ if (tmptex) gegl_clReleaseMemObject (tmptex);
+
+ return cl_err;
+}
+
+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;
+ gint j;
+ cl_int cl_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);
+ 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);
+
+ while (gegl_buffer_cl_iterator_next (i, &err))
+ {
+ if (err) return FALSE;
+ for (j=0; j < i->n; j++)
+ {
+ cl_err = cl_noise_reduction(i->tex[read][j],
+ i->tex[aux][j],
+ i->tex[0][j],
+ i->size[0][j],
+ &i->roi[read][j],
+ &i->roi[0][j],
+ o->iterations);
+ if (cl_err != CL_SUCCESS)
+ {
+ g_warning("[OpenCL] Error in noise-reduction");
+ return FALSE;
+ }
+ }
+ }
+ return TRUE;
+}
+
#define INPLACE 1
static gboolean
@@ -153,6 +368,11 @@ process (GeglOperation *operation,
const GeglRectangle *result)
{
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
+
+ if (cl_state.is_accelerated)
+ if(cl_process(operation, input, output, result))
+ return TRUE;
+
int iteration;
int stride;
float *src_buf;
@@ -239,6 +459,8 @@ gegl_chant_class_init (GeglChantClass *klass)
filter_class->process = process;
operation_class->prepare = prepare;
+ operation_class->opencl_support = TRUE;
+
operation_class->get_bounding_box = get_bounding_box;
operation_class->name = "gegl:noise-reduction";
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
