static gboolean
cl_stretch_contrast (cl_mem in_tex,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *roi,
cl_float4 min,
cl_float4 diff)
{
cl_int cl_err = 0;
cl_err = gegl_clSetKernelArg(cl_data->kernel[2], 0, sizeof(cl_mem),
(void*)&in_tex);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[2], 1, sizeof(cl_mem),
(void*)&out_tex);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[2], 2, sizeof(cl_float4),
(void*)&min);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[2], 3, sizeof(cl_float4),
(void*)&diff);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[2], 1,
NULL, &global_worksize, NULL,
0, NULL, NULL);
CL_CHECK;
return FALSE;
error:
return TRUE;
}
/* OpenCL processing function */
static cl_int
cl_process (GeglOperation *op,
cl_mem in_tex,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *roi,
gint level)
{
/* Retrieve a pointer to GeglChantO structure which contains all the
* chanted properties
*/
GeglChantO *o = GEGL_CHANT_PROPERTIES (op);
gfloat in_range;
gfloat out_range;
gfloat in_offset;
gfloat out_offset;
gfloat scale;
cl_int cl_err = 0;
in_offset = o->in_low * 1.0;
out_offset = o->out_low * 1.0;
in_range = o->in_high-o->in_low;
out_range = o->out_high-o->out_low;
if (in_range == 0.0)
in_range = 0.00000001;
scale = out_range/in_range;
if (!cl_data)
{
const char *kernel_name[] = {"kernel_levels", NULL};
cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
}
if (!cl_data) return 1;
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&in_offset);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&out_offset);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float), (void*)&scale);
if (cl_err != CL_SUCCESS) return cl_err;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 1,
NULL, &global_worksize, NULL,
0, NULL, NULL);
if (cl_err != CL_SUCCESS) return cl_err;
return cl_err;
}
/* OpenCL processing function */
static gboolean
cl_process (GeglOperation *op,
cl_mem in_tex,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *roi,
int level)
{
/* Retrieve a pointer to GeglProperties structure which contains all the
* chanted properties
*/
GeglProperties *o = GEGL_PROPERTIES (op);
const gfloat *coeffs = o->user_data;
cl_int cl_err = 0;
if (! coeffs)
{
coeffs = o->user_data = preprocess (o);
}
if (!cl_data)
{
const char *kernel_name[] = {"gegl_color_temperature", NULL};
cl_data = gegl_cl_compile_and_build (color_temperature_cl_source, kernel_name);
}
if (!cl_data) return 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*)&out_tex);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&coeffs[0]);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&coeffs[1]);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float), (void*)&coeffs[2]);
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;
error:
return TRUE;
}
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);
gfloat color[4];
gegl_color_get_pixel (o->color, babl_format ("R'G'B'A float"), color);
if (!cl_data)
{
const char *kernel_name[] = {"cl_color_to_alpha",NULL};
cl_data = gegl_cl_compile_and_build (color_to_alpha_cl_source, kernel_name);
}
if (!cl_data)
return TRUE;
else
{
cl_int cl_err = 0;
cl_float4 f_color;
f_color.s[0] = color[0];
f_color.s[1] = color[1];
f_color.s[2] = color[2];
f_color.s[3] = color[3];
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float4),(void*)&f_color);
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;
error:
return TRUE;
}
/* OpenCL processing function */
static cl_int
cl_process (GeglOperation *op,
cl_mem in_tex,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *roi,
gint level)
{
/* Retrieve a pointer to GeglProperties structure which contains all the
* chanted properties
*/
GeglProperties *o = GEGL_PROPERTIES (op);
gfloat black_level = (gfloat) o->black_level;
gfloat diff;
gfloat exposure_negated = (gfloat) -o->exposure;
gfloat gain;
gfloat white;
cl_int cl_err = 0;
if (!cl_data)
{
const char *kernel_name[] = {"kernel_exposure", NULL};
cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
}
if (!cl_data) return 1;
white = exp2f (exposure_negated);
diff = MAX (white - black_level, 0.01);
gain = 1.0f / diff;
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&black_level);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&gain);
if (cl_err != CL_SUCCESS) return cl_err;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 1,
NULL, &global_worksize, NULL,
0, NULL, NULL);
if (cl_err != CL_SUCCESS) return cl_err;
return cl_err;
}
static cl_int
cl_edge_laplace (cl_mem in_tex,
cl_mem aux_tex,
cl_mem out_tex,
const GeglRectangle *src_rect,
const GeglRectangle *roi,
gint radius)
{
cl_int cl_err = 0;
size_t global_ws[2];
if (!cl_data)
{
const char *kernel_name[] = {"pre_edgelaplace", "knl_edgelaplace", NULL};
cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
}
if (!cl_data) return 1;
global_ws[0] = roi->width;
global_ws[1] = roi->height;
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&aux_tex);
if (cl_err != CL_SUCCESS) return cl_err;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 2,
NULL, global_ws, NULL,
0, NULL, NULL);
if (cl_err != CL_SUCCESS) return cl_err;
cl_err = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
if (CL_SUCCESS != cl_err) return cl_err;
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_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, global_ws, NULL,
0, NULL, NULL);
if (cl_err != CL_SUCCESS) return cl_err;
return cl_err;
}
/* OpenCL processing function */
static cl_int
cl_process (GeglOperation *op,
cl_mem in_tex,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *roi,
int level)
{
/* Retrieve a pointer to GeglChantO structure which contains all the
* chanted properties
*/
GeglChantO *o = GEGL_CHANT_PROPERTIES (op);
const gfloat *coeffs = o->chant_data;
cl_int cl_err = 0;
if (! coeffs)
{
coeffs = o->chant_data = preprocess (o);
}
if (!cl_data)
{
const char *kernel_name[] = {"kernel_temp", NULL};
cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
}
if (!cl_data) return 1;
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&coeffs[0]);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&coeffs[1]);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float), (void*)&coeffs[2]);
if (cl_err != CL_SUCCESS) return cl_err;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 1,
NULL, &global_worksize, NULL,
0, NULL, NULL);
if (cl_err != CL_SUCCESS) return cl_err;
return cl_err;
}
static gboolean
cl_process (GeglOperation *op,
cl_mem in_tex,
cl_mem aux_tex,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *roi,
gint level)
{
cl_int cl_err = 0;
int kernel;
gfloat value;
if (!cl_data)
{
const char *kernel_name[] = {"gegl_opacity_RaGaBaA_float", "gegl_opacity_RGBA_float", NULL};
cl_data = gegl_cl_compile_and_build (opacity_cl_source, kernel_name);
}
if (!cl_data) return TRUE;
value = GEGL_CHANT_PROPERTIES (op)->value;
kernel = (GEGL_CHANT_PROPERTIES (op)->chant_data == NULL)? 0 : 1;
cl_err = gegl_clSetKernelArg(cl_data->kernel[kernel], 0, sizeof(cl_mem), (void*)&in_tex);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[kernel], 1, sizeof(cl_mem), (aux_tex)? (void*)&aux_tex : NULL);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[kernel], 2, sizeof(cl_mem), (void*)&out_tex);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[kernel], 3, sizeof(cl_float), (void*)&value);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[kernel], 1,
NULL, &global_worksize, NULL,
0, NULL, NULL);
CL_CHECK;
return FALSE;
error:
return TRUE;
}
/* OpenCL processing function */
static cl_int
cl_process (GeglOperation *op,
cl_mem in_tex,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *roi,
gint level)
{
/* Retrieve a pointer to GeglChantO structure which contains all the
* chanted properties
*/
GeglChantO *o = GEGL_CHANT_PROPERTIES (op);
gfloat gain = powf(2.0, o->exposure);
gfloat offset = o->offset;
gfloat gamma = 1.0 / o->gamma;
cl_int cl_err = 0;
if (!cl_data)
{
const char *kernel_name[] = {"kernel_exposure", NULL};
cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
}
if (!cl_data) return 1;
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&gain);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&offset);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float), (void*)&gamma);
if (cl_err != CL_SUCCESS) return cl_err;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 1,
NULL, &global_worksize, NULL,
0, NULL, NULL);
if (cl_err != CL_SUCCESS) return cl_err;
return cl_err;
}
static gboolean
cl_bilateral_filter (cl_mem in_tex,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *roi,
gfloat radius,
gfloat preserve)
{
cl_int cl_err = 0;
size_t global_ws[2];
if (!cl_data)
{
const char *kernel_name[] = {"bilateral_filter", NULL};
cl_data = gegl_cl_compile_and_build (bilateral_filter_cl_source, kernel_name);
}
if (!cl_data) return TRUE;
global_ws[0] = roi->width;
global_ws[1] = roi->height;
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*)&out_tex);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&radius);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&preserve);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 2,
NULL, global_ws, NULL,
0, NULL, NULL);
CL_CHECK;
return FALSE;
error:
return TRUE;
}
static cl_int
cl_mono_mixer(cl_mem in_tex,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *roi,
gfloat red,
gfloat green,
gfloat blue)
{
cl_int cl_err = 0;
if (!cl_data)
{
const char *kernel_name[] = {"Mono_mixer_cl", NULL};
cl_data = gegl_cl_compile_and_build(kernel_source, kernel_name);
}
if (!cl_data) return 0;
{
cl_float4 color;
color.s[0] = red;
color.s[1] = green;
color.s[2] = blue;
color.s[3] = 1.0f;
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_float4), (void*)&color);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 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[0],
1, NULL,
&global_worksize, NULL,
0, NULL, NULL);
}
return cl_err;
}
static gboolean
cl_edge_sobel (cl_mem in_tex,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *roi,
gboolean horizontal,
gboolean vertical,
gboolean keep_signal,
gboolean has_alpha)
{
const size_t gbl_size[2] = {roi->width, roi->height};
cl_int n_horizontal = horizontal;
cl_int n_vertical = vertical;
cl_int n_keep_signal = keep_signal;
cl_int n_has_alpha = has_alpha;
cl_int cl_err = 0;
if (!cl_data)
{
const char *kernel_name[] = {"kernel_edgesobel", NULL};
cl_data = gegl_cl_compile_and_build (edge_sobel_cl_source, kernel_name);
}
if (!cl_data) return TRUE;
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*)&out_tex);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int), (void*)&n_horizontal);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_int), (void*)&n_vertical);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_int), (void*)&n_keep_signal);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_int), (void*)&n_has_alpha);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(),
cl_data->kernel[0], 2,
NULL, gbl_size, NULL,
0, NULL, NULL);
CL_CHECK;
return FALSE;
error:
return TRUE;
}
static gboolean
cl_snn_mean (cl_mem in_tex,
cl_mem out_tex,
const GeglRectangle *src_rect,
const GeglRectangle *roi,
gint radius,
gint pairs)
{
cl_int cl_err = 0;
size_t global_ws[2];
if (!cl_data)
{
const char *kernel_name[] = {"snn_mean", NULL};
cl_data = gegl_cl_compile_and_build (snn_mean_cl_source, kernel_name);
}
if (!cl_data) return TRUE;
global_ws[0] = roi->width;
global_ws[1] = roi->height;
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_int), (void*)&src_rect->width);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int), (void*)&src_rect->height);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_mem), (void*)&out_tex);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_int), (void*)&radius);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_int), (void*)&pairs);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 2,
NULL, global_ws, NULL,
0, NULL, NULL);
CL_CHECK;
return FALSE;
error:
return TRUE;
}
static gboolean
cl_process (GeglOperation *self,
cl_mem in_tex,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *roi,
gint level)
{
GeglChantO *o = GEGL_CHANT_PROPERTIES (self);
gint num_sampling_points;
gdouble *xs, *ys;
gfloat *ysf = NULL;
cl_mem cl_curve = NULL;
cl_ulong cl_max_constant_size;
cl_int cl_err = 0;
num_sampling_points = o->sampling_points;
if (!cl_data)
{
const char *kernel_name[] = {"cl_contrast_curve",NULL};
cl_data = gegl_cl_compile_and_build (contrast_curve_cl_source,
kernel_name);
}
if (!cl_data)
return TRUE;
if (num_sampling_points > 0)
{
xs = g_new (gdouble, num_sampling_points);
ys = g_new (gdouble, num_sampling_points);
gegl_curve_calc_values (o->curve, 0.0, 1.0, num_sampling_points, xs, ys);
g_free (xs);
/*We need to downscale the array to pass it to the GPU*/
ysf = g_new (gfloat, num_sampling_points);
copy_double_array_to_float_array (ys, ysf, num_sampling_points);
g_free (ys);
cl_err = gegl_clGetDeviceInfo (gegl_cl_get_device (),
CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE,
sizeof (cl_ulong),
&cl_max_constant_size,
NULL);
CL_CHECK;
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Max Constant Mem Size: %lu bytes",
(unsigned long) cl_max_constant_size);
if (sizeof (cl_float) * num_sampling_points < cl_max_constant_size)
{
cl_curve = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_USE_HOST_PTR | CL_MEM_READ_ONLY,
num_sampling_points * sizeof (cl_float),
ysf, &cl_err);
CL_CHECK;
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*) &out_tex);
CL_CHECK;
cl_err = gegl_clSetKernelArg (cl_data->kernel[0], 2, sizeof (cl_mem),
(void*) &cl_curve);
CL_CHECK;
cl_err = gegl_clSetKernelArg (cl_data->kernel[0], 3, sizeof (gint),
(void*) &num_sampling_points);
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;
cl_err = gegl_clReleaseMemObject (cl_curve);
CL_CHECK_ONLY (cl_err);
}
else
{
/*If the curve size doesn't fit constant memory is better to use CPU*/
GEGL_NOTE (GEGL_DEBUG_OPENCL,
"Not enough constant memory for the curve");
g_free (ysf);
return TRUE;
}
g_free (ysf);
return FALSE;
error:
if (ysf)
g_free (ysf);
if (cl_curve)
gegl_clReleaseMemObject (cl_curve);
return TRUE;
}
else /*If the curve doesn't have a lookup table is better to use CPU*/
{
GEGL_NOTE (GEGL_DEBUG_OPENCL,
"Curve not suitable to be computed in the GPU");
return TRUE;
}
}
static void
motion_blur_cl (GeglBuffer *src,
const GeglRectangle *src_rect,
GeglBuffer *dst,
const GeglRectangle *dst_rect,
const int num_steps,
const float offset_x,
const float offset_y)
{
const Babl * in_format = babl_format("RaGaBaA float");
const Babl *out_format = babl_format("RaGaBaA float");
/* AreaFilter general processing flow.
Loading data and making the necessary color space conversion. */
#include "gegl-cl-operation-area-filter-fw1.h"
///////////////////////////////////////////////////////////////////////////
/* Algorithm specific processing flow.
Build kernels, setting parameters, and running them. */
if (!cl_data)
{
const char *kernel_name[] = {
"motion_blur_CL", NULL
};
cl_data = gegl_cl_compile_and_build(kernel_source, kernel_name);
}
if (!cl_data) CL_ERROR;
cl_int cl_src_width = src_rect->width;
cl_int cl_src_height = src_rect->height;
cl_int cl_src_x = src_rect->x;
cl_int cl_src_y = src_rect->y;
cl_int cl_dst_x = dst_rect->x;
cl_int cl_dst_y = dst_rect->y;
cl_int cl_num_steps = num_steps;
cl_float cl_offset_x = offset_x;
cl_float cl_offset_y = offset_y;
CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&src_mem));
CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
cl_data->kernel[0], 1, sizeof(cl_int), (void*)&cl_src_width));
CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
cl_data->kernel[0], 2, sizeof(cl_int), (void*)&cl_src_height));
CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
cl_data->kernel[0], 3, sizeof(cl_int), (void*)&cl_src_x));
CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
cl_data->kernel[0], 4, sizeof(cl_int), (void*)&cl_src_y));
CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
cl_data->kernel[0], 5, sizeof(cl_mem), (void*)&dst_mem));
CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
cl_data->kernel[0], 6, sizeof(cl_int), (void*)&cl_dst_x));
CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
cl_data->kernel[0], 7, sizeof(cl_int), (void*)&cl_dst_y));
CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
cl_data->kernel[0], 8, sizeof(cl_int), (void*)&cl_num_steps));
CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
cl_data->kernel[0], 9, sizeof(cl_float), (void*)&cl_offset_x));
CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
cl_data->kernel[0], 10, sizeof(cl_float), (void*)&cl_offset_y));
CL_SAFE_CALL(errcode = gegl_clEnqueueNDRangeKernel(
gegl_cl_get_command_queue(), cl_data->kernel[0],
2, NULL,
gbl_size, NULL,
0, NULL, NULL));
errcode = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
if (CL_SUCCESS != errcode) CL_ERROR;
///////////////////////////////////////////////////////////////////////////
/* AreaFilter general processing flow.
Making the necessary color space conversion and Saving data. */
#include "gegl-cl-operation-area-filter-fw2.h"
}
static cl_int
cl_motion_blur (cl_mem in_tex,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *roi,
const GeglRectangle *src_rect,
gint num_steps,
gfloat offset_x,
gfloat offset_y)
{
cl_int cl_err = 0;
size_t global_ws[2];
if (!cl_data)
{
const char *kernel_name[] = {"motion_blur_CL", NULL};
cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
}
if (!cl_data) return 1;
global_ws[0] = roi->width;
global_ws[1] = roi->height;
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_int), (void*)&src_rect->width);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int), (void*)&src_rect->height);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_int), (void*)&src_rect->x);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_int), (void*)&src_rect->y);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_mem), (void*)&out_tex);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 6, sizeof(cl_int), (void*)&roi->x);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 7, sizeof(cl_int), (void*)&roi->y);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 8, sizeof(cl_int), (void*)&num_steps);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 9, sizeof(cl_float), (void*)&offset_x);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 10, sizeof(cl_float), (void*)&offset_y);
if (cl_err != CL_SUCCESS) return cl_err;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 2,
NULL, global_ws, NULL,
0, NULL, NULL);
if (cl_err != CL_SUCCESS) return cl_err;
return cl_err;
}
static gboolean
cl_process (GeglOperation *operation,
cl_mem in_tex,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *roi,
gint level)
{
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
gfloat scale;
gfloat radius0, radius1;
gint roi_x, roi_y,x;
gint midx, midy;
GeglRectangle *bounds = gegl_operation_source_get_bounding_box (operation, "input");
gfloat length = hypot (bounds->width, bounds->height)/2;
gfloat rdiff;
gfloat cost, sint;
gfloat color[4];
scale = bounds->width / (1.0 * bounds->height);
scale = scale * (o->proportion) + 1.0 * (1.0-o->proportion);
scale *= aspect_to_scale (o->squeeze);
length = (bounds->width/2.0);
if (scale > 1.0)
length /= scale;
gegl_color_get_pixel (o->color, babl_format ("RGBA float"), color);
for (x=0; x<3; x++) /* premultiply */
color[x] *= color[3];
radius0 = o->radius * (1.0-o->softness);
radius1 = o->radius;
rdiff = radius1-radius0;
if (fabs (rdiff) < 0.0001)
rdiff = 0.0001;
midx = bounds->x + bounds->width * o->x;
midy = bounds->y + bounds->height * o->y;
roi_x = roi->x;
roi_y = roi->y;
/* constant for all pixels */
cost = cos(-o->rotation * (G_PI*2/360.0));
sint = sin(-o->rotation * (G_PI*2/360.0));
if (!cl_data)
{
const char *kernel_name[] = {"vignette_cl",NULL};
cl_data = gegl_cl_compile_and_build (vignette_cl_source, kernel_name);
}
if (!cl_data) return TRUE;
{
const size_t gbl_size[2] = {roi->width, roi->height};
gint shape = (gint) o->shape;
gfloat gamma = o->gamma;
cl_int cl_err = 0;
cl_float4 f_color;
f_color.s[0] = color[0];
f_color.s[1] = color[1];
f_color.s[2] = color[2];
f_color.s[3] = color[3];
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*)&out_tex);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float4),(void*)&f_color);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&scale);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float), (void*)&cost);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_float), (void*)&sint);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 6, sizeof(cl_int), (void*)&roi_x);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 7, sizeof(cl_int), (void*)&roi_y);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 8, sizeof(cl_int), (void*)&midx);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 9, sizeof(cl_int), (void*)&midy);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 10, sizeof(cl_int), (void*)&shape);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 11, sizeof(cl_float), (void*)&gamma);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 12, sizeof(cl_float), (void*)&length);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 13, sizeof(cl_float), (void*)&radius0);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 14, sizeof(cl_float), (void*)&rdiff);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 2,
NULL, gbl_size, NULL,
0, NULL, NULL);
CL_CHECK;
}
return FALSE;
error:
return TRUE;
}
gboolean
gegl_operation_cl_set_kernel_args (GeglOperation *operation,
cl_kernel kernel,
gint *p,
cl_int *err)
{
GParamSpec **self;
GParamSpec **parent;
guint n_self;
guint n_parent;
gint prop_no;
self = g_object_class_list_properties (
G_OBJECT_CLASS (g_type_class_ref (G_OBJECT_CLASS_TYPE (GEGL_OPERATION_GET_CLASS(operation)))),
&n_self);
parent = g_object_class_list_properties (
G_OBJECT_CLASS (g_type_class_ref (GEGL_TYPE_OPERATION)),
&n_parent);
for (prop_no=0;prop_no<n_self;prop_no++)
{
gint parent_no;
gboolean found=FALSE;
for (parent_no=0;parent_no<n_parent;parent_no++)
if (self[prop_no]==parent[parent_no])
found=TRUE;
/* only print properties if we are an addition compared to
* GeglOperation
*/
/* Removing pads */
if (!strcmp(g_param_spec_get_name (self[prop_no]), "input") ||
!strcmp(g_param_spec_get_name (self[prop_no]), "output") ||
!strcmp(g_param_spec_get_name (self[prop_no]), "aux"))
continue;
if (!found)
{
if (g_type_is_a (G_PARAM_SPEC_VALUE_TYPE (self[prop_no]), G_TYPE_DOUBLE))
{
gdouble value;
cl_float v;
g_object_get (G_OBJECT (operation), g_param_spec_get_name (self[prop_no]), &value, NULL);
v = value;
*err = gegl_clSetKernelArg(kernel, (*p)++, sizeof(cl_float), (void*)&v);
}
else if (g_type_is_a (G_PARAM_SPEC_VALUE_TYPE (self[prop_no]), G_TYPE_FLOAT))
{
gfloat value;
cl_float v;
g_object_get (G_OBJECT (operation), g_param_spec_get_name (self[prop_no]), &value, NULL);
v = value;
*err = gegl_clSetKernelArg(kernel, (*p)++, sizeof(cl_float), (void*)&v);
}
else if (g_type_is_a (G_PARAM_SPEC_VALUE_TYPE (self[prop_no]), G_TYPE_INT))
{
gint value;
cl_int v;
g_object_get (G_OBJECT (operation), g_param_spec_get_name (self[prop_no]), &value, NULL);
v = value;
*err = gegl_clSetKernelArg(kernel, (*p)++, sizeof(cl_int), (void*)&v);
}
else if (g_type_is_a (G_PARAM_SPEC_VALUE_TYPE (self[prop_no]), G_TYPE_BOOLEAN))
{
gboolean value;
cl_bool v;
g_object_get (G_OBJECT (operation), g_param_spec_get_name (self[prop_no]), &value, NULL);
v = value;
*err = gegl_clSetKernelArg(kernel, (*p)++, sizeof(cl_bool), (void*)&v);
}
else
{
g_error ("Unsupported OpenCL kernel argument");
return FALSE;
}
}
}
if (self)
g_free (self);
if (parent)
g_free (parent);
return TRUE;
}
static gboolean
cl_buffer_get_min_max (cl_mem in_tex,
size_t global_worksize,
const GeglRectangle *roi,
gfloat min[4],
gfloat max[4])
{
cl_int cl_err = 0;
size_t local_ws, max_local_ws;
size_t work_groups;
size_t global_ws;
cl_mem cl_aux_min = NULL;
cl_mem cl_aux_max = NULL;
cl_mem cl_min_max = NULL;
cl_int n_pixels = (cl_int)global_worksize;
cl_float4 min_max_buf[2];
if (global_worksize < 1)
{
min[0] = min[1] = min[2] = min[3] = G_MAXFLOAT;
max[0] = max[1] = max[2] = max[3] = -G_MAXFLOAT;
return FALSE;
}
cl_err = gegl_clGetDeviceInfo (gegl_cl_get_device (),
CL_DEVICE_MAX_WORK_GROUP_SIZE,
sizeof (size_t), &max_local_ws, NULL);
CL_CHECK;
max_local_ws = MIN (max_local_ws,
MIN (cl_data->work_group_size[0],
cl_data->work_group_size[1]));
/* Needs to be a power of two */
local_ws = 256;
while (local_ws > max_local_ws)
local_ws /= 2;
work_groups = MIN ((global_worksize + local_ws - 1) / local_ws, local_ws);
global_ws = work_groups * local_ws;
cl_aux_min = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_READ_WRITE,
local_ws * sizeof(cl_float4),
NULL, &cl_err);
CL_CHECK;
cl_aux_max = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_READ_WRITE,
local_ws * sizeof(cl_float4),
NULL, &cl_err);
CL_CHECK;
cl_min_max = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_WRITE_ONLY,
2 * sizeof(cl_float4),
NULL, &cl_err);
CL_CHECK;
/* The full initialization is done in the two_stages_local_min_max_reduce
kernel */
#if 0
cl_err = gegl_clSetKernelArg(cl_data->kernel[3], 0, sizeof(cl_mem),
(void*)&cl_aux_min);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[3], 1, sizeof(cl_mem),
(void*)&cl_aux_max);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[3], 1,
NULL, &local_ws, &local_ws,
0, NULL, NULL);
CL_CHECK;
#endif
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*)&cl_aux_min);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_mem),
(void*)&cl_aux_max);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3,
sizeof(cl_float4) * local_ws, NULL);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 4,
sizeof(cl_float4) * local_ws, NULL);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_int),
(void*)&n_pixels);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 1,
NULL, &global_ws, &local_ws,
0, NULL, NULL);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem),
(void*)&cl_aux_min);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_mem),
(void*)&cl_aux_max);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 2, sizeof(cl_mem),
(void*)&cl_min_max);
CL_CHECK;
/* Only one work group */
cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
cl_data->kernel[1], 1,
NULL, &local_ws, &local_ws,
0, NULL, NULL);
CL_CHECK;
/* Read the memory buffer, probably better to keep it in GPU memory */
cl_err = gegl_clEnqueueReadBuffer (gegl_cl_get_command_queue (),
cl_min_max, CL_TRUE, 0,
2 * sizeof (cl_float4), &min_max_buf, 0,
NULL, NULL);
CL_CHECK;
min[0] = min_max_buf[0].x;
min[1] = min_max_buf[0].y;
min[2] = min_max_buf[0].z;
min[3] = min_max_buf[0].w;
max[0] = min_max_buf[1].x;
max[1] = min_max_buf[1].y;
max[2] = min_max_buf[1].z;
max[3] = min_max_buf[1].w;
cl_err = gegl_clReleaseMemObject (cl_aux_min);
CL_CHECK_ONLY (cl_err);
cl_err = gegl_clReleaseMemObject (cl_aux_max);
CL_CHECK_ONLY (cl_err);
cl_err = gegl_clReleaseMemObject (cl_min_max);
CL_CHECK_ONLY (cl_err);
return FALSE;
error:
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;
}
static gboolean
gegl_operation_point_filter_cl_process (GeglOperation *operation,
GeglBuffer *input,
GeglBuffer *output,
const GeglRectangle *result,
gint level)
{
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *out_format = gegl_operation_get_format (operation, "output");
GeglOperationClass *operation_class = GEGL_OPERATION_GET_CLASS (operation);
GeglOperationPointFilterClass *point_filter_class = GEGL_OPERATION_POINT_FILTER_GET_CLASS (operation);
GeglBufferClIterator *iter = NULL;
cl_int cl_err = 0;
gboolean err;
/* non-texturizable format! */
if (!gegl_cl_color_babl (in_format, NULL) ||
!gegl_cl_color_babl (out_format, NULL))
{
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Non-texturizable format!");
return FALSE;
}
GEGL_NOTE (GEGL_DEBUG_OPENCL, "GEGL_OPERATION_POINT_FILTER: %s", operation_class->name);
/* Process */
iter = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE);
gegl_buffer_cl_iterator_add (iter, input, result, in_format, GEGL_CL_BUFFER_READ, GEGL_ABYSS_NONE);
while (gegl_buffer_cl_iterator_next (iter, &err))
{
if (err)
return FALSE;
if (point_filter_class->cl_process)
{
err = point_filter_class->cl_process (operation, iter->tex[1], iter->tex[0],
iter->size[0], &iter->roi[0], level);
if (err)
{
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", operation_class->name);
gegl_buffer_cl_iterator_stop (iter);
return FALSE;
}
}
else if (operation_class->cl_data)
{
gint p = 0;
GeglClRunData *cl_data = operation_class->cl_data;
cl_err = gegl_clSetKernelArg (cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&iter->tex[1]);
CL_CHECK;
cl_err = gegl_clSetKernelArg (cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&iter->tex[0]);
CL_CHECK;
gegl_operation_cl_set_kernel_args (operation, cl_data->kernel[0], &p, &cl_err);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
cl_data->kernel[0], 1,
NULL, &iter->size[0], NULL,
0, NULL, NULL);
CL_CHECK;
}
else
{
g_warning ("OpenCL support enabled, but no way to execute");
gegl_buffer_cl_iterator_stop (iter);
return FALSE;
}
}
return TRUE;
error:
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", gegl_cl_errstring (cl_err));
if (iter)
gegl_buffer_cl_iterator_stop (iter);
return FALSE;
}
static gboolean
cl_box_max (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_max_hor", "kernel_max_ver", NULL};
cl_data = gegl_cl_compile_and_build (box_max_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;
}
