[gegl] operations: Clean up coding style and remove macros
- From: Daniel Sabo <daniels src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] operations: Clean up coding style and remove macros
- Date: Thu, 31 Oct 2013 21:55:34 +0000 (UTC)
commit 0b6c741f67bb8674fdd3a43bce3c32ffc42f7927
Author: Daniel Sabo <DanielSabo gmail com>
Date: Thu Oct 31 13:36:56 2013 -0700
operations: Clean up coding style and remove macros
operations/common/alien-map.c | 42 +++++++------
operations/common/dot.c | 113 +++++++++++++++++++---------------
operations/common/noise-hurl.c | 75 +++++++++++++----------
operations/common/posterize.c | 31 ++++++---
operations/common/red-eye-removal.c | 31 ++++++---
operations/common/stretch-contrast.c | 102 ++++++++++++++++++-------------
6 files changed, 231 insertions(+), 163 deletions(-)
---
diff --git a/operations/common/alien-map.c b/operations/common/alien-map.c
index dcc0df7..7efe29d 100644
--- a/operations/common/alien-map.c
+++ b/operations/common/alien-map.c
@@ -143,7 +143,7 @@ process (GeglOperation *op,
#include "opencl/gegl-cl.h"
#include "opencl/alien-map.cl.h"
-GEGL_CL_STATIC
+static GeglClRunData *cl_data = NULL;
static gboolean
cl_process (GeglOperation *operation,
@@ -157,8 +157,17 @@ cl_process (GeglOperation *operation,
cl_float3 freq;
cl_float3 phaseshift;
cl_int3 keep;
+ cl_int cl_err = 0;
- GEGL_CL_BUILD(alien_map, "cl_alien_map")
+ if (!cl_data)
+ {
+ const char *kernel_name[] = {"cl_alien_map",
+ NULL};
+ cl_data = gegl_cl_compile_and_build (alien_map_cl_source, kernel_name);
+ }
+
+ if (!cl_data)
+ return TRUE;
freq.s[0] = o->cpn_1_frequency * G_PI;
freq.s[1] = o->cpn_2_frequency * G_PI;
@@ -172,23 +181,20 @@ cl_process (GeglOperation *operation,
keep.s[1] = (cl_int)o->cpn_2_keep;
keep.s[2] = (cl_int)o->cpn_3_keep;
- {
- cl_int cl_err = 0;
-
- GEGL_CL_ARG_START(cl_data->kernel[0])
- GEGL_CL_ARG(cl_mem, in)
- GEGL_CL_ARG(cl_mem, out)
- GEGL_CL_ARG(cl_float3, freq)
- GEGL_CL_ARG(cl_float3, phaseshift)
- GEGL_CL_ARG(cl_int3, keep)
- GEGL_CL_ARG_END
-
- cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
- cl_data->kernel[0], 1,
- NULL, &global_worksize, NULL,
- 0, NULL, NULL);
+ gegl_cl_set_kernel_args (cl_data->kernel[0],
+ sizeof(cl_mem), &in,
+ sizeof(cl_mem), &out,
+ sizeof(cl_float3), &freq,
+ sizeof(cl_float3), &phaseshift,
+ sizeof(cl_int3), &keep,
+ 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;
- }
return FALSE;
diff --git a/operations/common/dot.c b/operations/common/dot.c
index 8a15e40..d6070f0 100644
--- a/operations/common/dot.c
+++ b/operations/common/dot.c
@@ -67,7 +67,7 @@ prepare (GeglOperation *operation)
#include "opencl/dot.cl.h"
#include <stdio.h>
-GEGL_CL_STATIC
+static GeglClRunData *cl_data = NULL;
static gboolean
cl_dot (cl_mem in,
@@ -79,7 +79,14 @@ cl_dot (cl_mem in,
cl_int cl_err = 0;
cl_mem block_colors = NULL;
- GEGL_CL_BUILD(dot, "cl_calc_block_colors", "cl_dot")
+ if (!cl_data)
+ {
+ const char *kernel_name[] ={"cl_calc_block_colors", "cl_dot" , NULL};
+ cl_data = gegl_cl_compile_and_build(dot_cl_source, kernel_name);
+ }
+
+ if (!cl_data)
+ return TRUE;
{
cl_int cl_size = size;
@@ -94,64 +101,67 @@ cl_dot (cl_mem in,
cl_int block_count_y = CELL_Y(roi->y + roi->height - 1, size) - cy0 + 1;
size_t glbl_s[2] = {block_count_x, block_count_y};
- block_colors = gegl_clCreateBuffer(gegl_cl_get_context(),
- CL_MEM_READ_WRITE,
- glbl_s[0] * glbl_s[1] * sizeof(cl_float4),
- NULL, &cl_err);
+ block_colors = gegl_clCreateBuffer (gegl_cl_get_context(),
+ CL_MEM_READ_WRITE,
+ glbl_s[0] * glbl_s[1] * sizeof(cl_float4),
+ NULL, &cl_err);
CL_CHECK;
- GEGL_CL_ARG_START(cl_data->kernel[0])
- GEGL_CL_ARG(cl_mem, in)
- GEGL_CL_ARG(cl_mem, block_colors)
- GEGL_CL_ARG(cl_int, cx0)
- GEGL_CL_ARG(cl_int, cy0)
- GEGL_CL_ARG(cl_int, cl_size)
- GEGL_CL_ARG(cl_float, weight)
- GEGL_CL_ARG(cl_int, roi_x)
- GEGL_CL_ARG(cl_int, roi_y)
- GEGL_CL_ARG(cl_int, line_width)
- GEGL_CL_ARG_END
+ gegl_cl_set_kernel_args (cl_data->kernel[0],
+ sizeof(cl_mem), &in,
+ sizeof(cl_mem), &block_colors,
+ sizeof(cl_int), &cx0,
+ sizeof(cl_int), &cy0,
+ sizeof(cl_int), &cl_size,
+ sizeof(cl_float), &weight,
+ sizeof(cl_int), &roi_x,
+ sizeof(cl_int), &roi_y,
+ sizeof(cl_int), &line_width,
+ NULL);
+ CL_CHECK;
/* calculate the average color of all the blocks */
- cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
- cl_data->kernel[0], 2,
- NULL, glbl_s, NULL,
- 0, NULL, NULL);
+ cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
+ cl_data->kernel[0], 2,
+ NULL, glbl_s, NULL,
+ 0, NULL, NULL);
CL_CHECK;
glbl_s[0] = roi->width;
glbl_s[1] = roi->height;
- GEGL_CL_ARG_START(cl_data->kernel[1])
- GEGL_CL_ARG(cl_mem, block_colors)
- GEGL_CL_ARG(cl_mem, out)
- GEGL_CL_ARG(cl_int, cx0)
- GEGL_CL_ARG(cl_int, cy0)
- GEGL_CL_ARG(cl_int, cl_size)
- GEGL_CL_ARG(cl_float, radius2)
- GEGL_CL_ARG(cl_int, roi_x)
- GEGL_CL_ARG(cl_int, roi_y)
- GEGL_CL_ARG(cl_int, block_count_x)
- GEGL_CL_ARG_END
+ gegl_cl_set_kernel_args (cl_data->kernel[1],
+ sizeof(cl_mem), &block_colors,
+ sizeof(cl_mem), &out,
+ sizeof(cl_int), &cx0,
+ sizeof(cl_int), &cy0,
+ sizeof(cl_int), &cl_size,
+ sizeof(cl_float), &radius2,
+ sizeof(cl_int), &roi_x,
+ sizeof(cl_int), &roi_y,
+ sizeof(cl_int), &block_count_x,
+ NULL);
+ CL_CHECK;
/* set each pixel to the average color of the block it belongs to */
- cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
- cl_data->kernel[1], 2,
- NULL, glbl_s, NULL,
- 0, NULL, NULL);
+ cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
+ cl_data->kernel[1], 2,
+ NULL, glbl_s, NULL,
+ 0, NULL, NULL);
CL_CHECK;
- cl_err = gegl_clFinish(gegl_cl_get_command_queue ());
+ cl_err = gegl_clFinish (gegl_cl_get_command_queue ());
CL_CHECK;
- GEGL_CL_RELEASE(block_colors)
+ cl_err = gegl_clReleaseMemObject (block_colors);
+ CL_CHECK;
}
return FALSE;
error:
- if(block_colors)
- GEGL_CL_RELEASE(block_colors)
+ if (block_colors)
+ gegl_clReleaseMemObject (block_colors);
return TRUE;
}
@@ -164,7 +174,7 @@ cl_process (GeglOperation *operation,
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *out_format = gegl_operation_get_format (operation, "output");
- gint err;
+ gint err = 0;
GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation);
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
@@ -185,17 +195,22 @@ cl_process (GeglOperation *operation,
op_area->bottom,
GEGL_ABYSS_NONE);
- GEGL_CL_BUFFER_ITERATE_START(i, err)
+ while (gegl_buffer_cl_iterator_next (i, &err) && !err)
{
- err = cl_dot(i->tex[read],
- i->tex[0],
- &i->roi[0],
- o->size,
- o->ratio);
+ err = cl_dot (i->tex[read],
+ i->tex[0],
+ &i->roi[0],
+ o->size,
+ o->ratio);
+
+ if (err)
+ {
+ gegl_buffer_cl_iterator_stop (i);
+ break;
+ }
}
- GEGL_CL_BUFFER_ITERATE_END(err)
- return TRUE;
+ return !err;
}
static void
diff --git a/operations/common/noise-hurl.c b/operations/common/noise-hurl.c
index f83a57f..5a50475 100644
--- a/operations/common/noise-hurl.c
+++ b/operations/common/noise-hurl.c
@@ -111,7 +111,7 @@ process (GeglOperation *operation,
#include "opencl/gegl-cl.h"
#include "opencl/noise-hurl.cl.h"
-GEGL_CL_STATIC
+static GeglClRunData *cl_data = NULL;
static gboolean
cl_process (GeglOperation *operation,
@@ -137,62 +137,73 @@ cl_process (GeglOperation *operation,
cl_int offset;
int it;
- GEGL_CL_BUILD(noise_hurl, "cl_noise_hurl")
+ if (!cl_data)
+ {
+ const char *kernel_name[] ={"cl_noise_hurl", NULL};
+ cl_data = gegl_cl_compile_and_build(noise_hurl_cl_source, kernel_name);
+ }
+
+ if (!cl_data)
+ return TRUE;
{
- cl_random_data = gegl_cl_load_random_data(&cl_err);
+ cl_random_data = gegl_cl_load_random_data (&cl_err);
CL_CHECK;
- cl_random_primes = gegl_cl_load_random_primes(&cl_err);
+ 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_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
+ gegl_cl_set_kernel_args (cl_data->kernel[0],
+ sizeof(cl_mem), &out,
+ sizeof(cl_mem), &cl_random_data,
+ sizeof(cl_mem), &cl_random_primes,
+ sizeof(cl_int), &x_offset,
+ sizeof(cl_int), &y_offset,
+ sizeof(cl_int), &roi_width,
+ sizeof(cl_int), &wr_width,
+ sizeof(cl_int), &seed,
+ sizeof(cl_float), &pct_random,
+ NULL);
+ CL_CHECK;
offset = 0;
for(it = 0; it < o->repeat; ++it)
{
- cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 9, sizeof(cl_int),
- (void*)&offset);
+ 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_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_err = gegl_clFinish (gegl_cl_get_command_queue ());
CL_CHECK;
- GEGL_CL_RELEASE(cl_random_data)
- GEGL_CL_RELEASE(cl_random_primes)
+ cl_err = gegl_clReleaseMemObject (cl_random_data);
+ CL_CHECK;
+
+ cl_err = gegl_clReleaseMemObject (cl_random_primes);
+ CL_CHECK;
}
return FALSE;
error:
- if(cl_random_data)
- GEGL_CL_RELEASE(cl_random_data)
- if(cl_random_primes)
- GEGL_CL_RELEASE(cl_random_primes)
+ if (cl_random_data)
+ gegl_clReleaseMemObject (cl_random_data);
+ if (cl_random_primes)
+ gegl_clReleaseMemObject (cl_random_primes);
return TRUE;
}
diff --git a/operations/common/posterize.c b/operations/common/posterize.c
index f5a1592..c4b6915 100644
--- a/operations/common/posterize.c
+++ b/operations/common/posterize.c
@@ -67,7 +67,7 @@ static gboolean process (GeglOperation *operation,
#include "opencl/gegl-cl.h"
#include "opencl/posterize.cl.h"
-GEGL_CL_STATIC
+static GeglClRunData *cl_data = NULL;
static gboolean
cl_process (GeglOperation *operation,
@@ -80,21 +80,30 @@ cl_process (GeglOperation *operation,
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
cl_float levels = o->levels;
- GEGL_CL_BUILD(posterize, "cl_posterize")
+ if (!cl_data)
+ {
+ const char *kernel_name[] = {"cl_posterize",
+ NULL};
+ cl_data = gegl_cl_compile_and_build (posterize_cl_source, kernel_name);
+ }
+
+ if (!cl_data)
+ return 1;
{
cl_int cl_err = 0;
- GEGL_CL_ARG_START(cl_data->kernel[0])
- GEGL_CL_ARG(cl_mem, in)
- GEGL_CL_ARG(cl_mem, out)
- GEGL_CL_ARG(cl_float, levels)
- GEGL_CL_ARG_END
+ gegl_cl_set_kernel_args (cl_data->kernel[0],
+ sizeof(cl_mem), &in,
+ sizeof(cl_mem), &out,
+ sizeof(cl_float), &levels,
+ 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_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
+ cl_data->kernel[0], 1,
+ NULL, &global_worksize, NULL,
+ 0, NULL, NULL);
CL_CHECK;
}
diff --git a/operations/common/red-eye-removal.c b/operations/common/red-eye-removal.c
index f8d3dd4..309b31d 100644
--- a/operations/common/red-eye-removal.c
+++ b/operations/common/red-eye-removal.c
@@ -104,7 +104,7 @@ process (GeglOperation *operation,
#include "opencl/gegl-cl.h"
#include "opencl/red-eye-removal.cl.h"
-GEGL_CL_STATIC
+static GeglClRunData *cl_data = NULL;
static gboolean
cl_process (GeglOperation *operation,
@@ -117,21 +117,30 @@ cl_process (GeglOperation *operation,
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
cl_float threshold = o->threshold;
- GEGL_CL_BUILD(red_eye_removal, "cl_red_eye_removal")
+
+ if (!cl_data)
+ {
+ const char *kernel_name[] = {"cl_red_eye_removal", NULL};
+ cl_data = gegl_cl_compile_and_build(red_eye_removal_cl_source, kernel_name);
+ }
+
+ if (!cl_data)
+ return TRUE;
{
cl_int cl_err = 0;
- GEGL_CL_ARG_START(cl_data->kernel[0])
- GEGL_CL_ARG(cl_mem, in)
- GEGL_CL_ARG(cl_mem, out)
- GEGL_CL_ARG(cl_float, threshold)
- GEGL_CL_ARG_END
+ gegl_cl_set_kernel_args (cl_data->kernel[0],
+ sizeof(cl_mem), &in,
+ sizeof(cl_mem), &out,
+ sizeof(cl_float), &threshold,
+ 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_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
+ cl_data->kernel[0], 1,
+ NULL, &global_worksize, NULL,
+ 0, NULL, NULL);
CL_CHECK;
}
diff --git a/operations/common/stretch-contrast.c b/operations/common/stretch-contrast.c
index caa0400..9094293 100644
--- a/operations/common/stretch-contrast.c
+++ b/operations/common/stretch-contrast.c
@@ -90,16 +90,24 @@ get_cached_region (GeglOperation *operation,
#include "buffer/gegl-buffer-cl-iterator.h"
#include "opencl/stretch-contrast.cl.h"
-GEGL_CL_STATIC
+static GeglClRunData *cl_data = NULL;
static gboolean
-cl_build_kernels( void )
+cl_build_kernels (void)
{
- GEGL_CL_BUILD( stretch_contrast,
- "two_stages_local_min_max_reduce",
- "global_min_max_reduce",
- "cl_stretch_contrast",
- "init_stretch")
+ if (!cl_data)
+ {
+ const char *kernel_name[] = {"two_stages_local_min_max_reduce",
+ "global_min_max_reduce",
+ "cl_stretch_contrast",
+ "init_stretch",
+ NULL};
+ cl_data = gegl_cl_compile_and_build (stretch_contrast_cl_source, kernel_name);
+ }
+
+ if (!cl_data)
+ return TRUE;
+
return FALSE;
}
@@ -141,20 +149,20 @@ cl_buffer_get_min_max (cl_mem in_tex,
global_ws = work_groups * local_ws;
- cl_aux_min = gegl_clCreateBuffer(gegl_cl_get_context(),
- CL_MEM_READ_WRITE,
- local_ws * sizeof(cl_float),
- NULL, &cl_err);
+ cl_aux_min = gegl_clCreateBuffer (gegl_cl_get_context (),
+ CL_MEM_READ_WRITE,
+ local_ws * sizeof(cl_float),
+ NULL, &cl_err);
CL_CHECK;
- cl_aux_max = gegl_clCreateBuffer(gegl_cl_get_context(),
- CL_MEM_READ_WRITE,
- local_ws * sizeof(cl_float),
- NULL, &cl_err);
+ cl_aux_max = gegl_clCreateBuffer (gegl_cl_get_context (),
+ CL_MEM_READ_WRITE,
+ local_ws * sizeof(cl_float),
+ NULL, &cl_err);
CL_CHECK;
- cl_min_max = gegl_clCreateBuffer(gegl_cl_get_context(),
- CL_MEM_WRITE_ONLY,
- 2 * sizeof(cl_float),
- NULL, &cl_err);
+ cl_min_max = gegl_clCreateBuffer (gegl_cl_get_context (),
+ CL_MEM_WRITE_ONLY,
+ 2 * sizeof(cl_float),
+ NULL, &cl_err);
CL_CHECK;
/* The full initialization is done in the two_stages_local_min_max_reduce
@@ -217,7 +225,7 @@ cl_buffer_get_min_max (cl_mem in_tex,
CL_CHECK;
/* Read the memory buffer, probably better to keep it in GPU memory */
- cl_err = gegl_clEnqueueReadBuffer (gegl_cl_get_command_queue(),
+ cl_err = gegl_clEnqueueReadBuffer (gegl_cl_get_command_queue (),
cl_min_max, CL_TRUE, 0,
2 * sizeof (cl_float), &min_max_buf, 0,
NULL, NULL);
@@ -226,19 +234,25 @@ cl_buffer_get_min_max (cl_mem in_tex,
*min = min_max_buf[0];
*max = min_max_buf[1];
- GEGL_CL_RELEASE(cl_aux_min)
- GEGL_CL_RELEASE(cl_aux_max)
- GEGL_CL_RELEASE(cl_min_max)
+
+ cl_err = gegl_clReleaseMemObject (cl_aux_min);
+ CL_CHECK;
+
+ cl_err = gegl_clReleaseMemObject (cl_aux_max);
+ CL_CHECK;
+
+ cl_err = gegl_clReleaseMemObject (cl_min_max);
+ CL_CHECK;
return FALSE;
error:
- if(cl_aux_min)
- GEGL_CL_RELEASE(cl_aux_min)
- if(cl_aux_max)
- GEGL_CL_RELEASE(cl_aux_max)
- if(cl_min_max)
- GEGL_CL_RELEASE(cl_min_max)
+ if (cl_aux_min)
+ gegl_clReleaseMemObject (cl_aux_min);
+ if (cl_aux_max)
+ gegl_clReleaseMemObject (cl_aux_max);
+ if (cl_min_max)
+ gegl_clReleaseMemObject (cl_min_max);
return TRUE;
}
@@ -293,10 +307,11 @@ cl_process (GeglOperation *operation,
gfloat min = 1.0f;
gfloat max = 0.0f;
gfloat i_min, i_max, diff;
- gint err, read;
+ cl_int err = 0;
+ gint read;
GeglBufferClIterator *i;
- if(cl_build_kernels())
+ if (cl_build_kernels ())
return FALSE;
i = gegl_buffer_cl_iterator_new (input,
@@ -339,20 +354,23 @@ cl_process (GeglOperation *operation,
0,
GEGL_ABYSS_NONE);
- while (gegl_buffer_cl_iterator_next (i, &err))
+ while (gegl_buffer_cl_iterator_next (i, &err) && !err)
{
- if (err) return FALSE;
-
- err = cl_stretch_contrast(i->tex[read],
- i->tex[0],
- i->size[0],
- &i->roi[0],
- min,
- diff);
- if (err) return FALSE;
+ err = cl_stretch_contrast (i->tex[read],
+ i->tex[0],
+ i->size[0],
+ &i->roi[0],
+ min,
+ diff);
+
+ if (err)
+ {
+ gegl_buffer_cl_iterator_stop (i);
+ break;
+ }
}
- return TRUE;
+ return !err;
}
static gboolean
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
