static gboolean
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_in[2];
size_t global_ws_aux[2];
if (!cl_data)
{
const char *kernel_name[] = {"pre_edgelaplace", "knl_edgelaplace", NULL};
cl_data = gegl_cl_compile_and_build (edge_laplace_cl_source, kernel_name);
}
if (!cl_data) return TRUE;
global_ws_in[0] = roi->width + LAPLACE_RADIUS;
global_ws_in[1] = roi->height + LAPLACE_RADIUS;
global_ws_aux[0] = roi->width;
global_ws_aux[1] = roi->height;
cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
sizeof (cl_mem), &in_tex,
sizeof (cl_mem), &aux_tex,
NULL);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
cl_data->kernel[0], 2,
NULL, global_ws_in, NULL,
0, NULL, NULL);
CL_CHECK;
cl_err = gegl_cl_set_kernel_args (cl_data->kernel[1],
sizeof (cl_mem), &aux_tex,
sizeof (cl_mem), &out_tex,
NULL);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
cl_data->kernel[1], 2,
NULL, global_ws_aux, NULL,
0, NULL, NULL);
CL_CHECK;
return FALSE;
error:
return TRUE;
}
static gboolean
cl_gaussian_blur (cl_mem in_tex,
cl_mem out_tex,
const GeglRectangle *roi,
cl_mem cl_cmatrix,
gint clen,
GeglOrientation orientation)
{
cl_int cl_err = 0;
size_t global_ws[2];
gint kernel_num;
if (!cl_data)
{
const char *kernel_name[] = {"fir_ver_blur", "fir_hor_blur", NULL};
cl_data = gegl_cl_compile_and_build (gblur_1d_cl_source, kernel_name);
}
if (!cl_data)
return TRUE;
if (orientation == GEGL_ORIENTATION_VERTICAL)
kernel_num = 0;
else
kernel_num = 1;
global_ws[0] = roi->width;
global_ws[1] = roi->height;
cl_err = gegl_cl_set_kernel_args (cl_data->kernel[kernel_num],
sizeof(cl_mem), (void*)&in_tex,
sizeof(cl_mem), (void*)&out_tex,
sizeof(cl_mem), (void*)&cl_cmatrix,
sizeof(cl_int), (void*)&clen,
NULL);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
cl_data->kernel[kernel_num], 2,
NULL, global_ws, NULL,
0, NULL, NULL);
CL_CHECK;
cl_err = gegl_clFinish (gegl_cl_get_command_queue ());
CL_CHECK;
return FALSE;
error:
return TRUE;
}
static gboolean
cl_process (GeglOperation *operation,
cl_mem in_buf,
cl_mem out_buf,
const size_t n_pixels,
const GeglRectangle *roi,
gint level)
{
GeglOperationClass *operation_class = GEGL_OPERATION_GET_CLASS (operation);
GeglClRunData *cl_data = operation_class->cl_data;
GeglProperties *o = GEGL_PROPERTIES (operation);
const size_t gbl_size[2] = {roi->width, roi->height};
const size_t gbl_off[2] = {roi->x, roi->y};
cl_int cl_err = 0;
cl_mem filter_pat = NULL;
if (!cl_data)
goto error;
filter_pat = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
pattern_width[o->pattern] *
pattern_height[o->pattern] * sizeof(cl_int),
(void*)pattern[o->pattern],
&cl_err);
CL_CHECK;
cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
sizeof(cl_mem), &in_buf,
sizeof(cl_mem), &out_buf,
sizeof(cl_mem), &filter_pat,
sizeof(cl_int), &pattern_width[o->pattern],
sizeof(cl_int), &pattern_height[o->pattern],
sizeof(cl_int), &o->additive,
sizeof(cl_int), &o->rotated,
NULL);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
cl_data->kernel[0], 2,
gbl_off, gbl_size, NULL,
0, NULL, NULL);
CL_CHECK;
cl_err = gegl_clFinish (gegl_cl_get_command_queue ());
CL_CHECK;
cl_err = gegl_clReleaseMemObject (filter_pat);
CL_CHECK;
return FALSE;
error:
if (filter_pat)
gegl_clReleaseMemObject (filter_pat);
return TRUE;
}
static gboolean
cl_process (GeglOperation *operation,
cl_mem in_tex,
cl_mem aux_tex,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *roi,
gint level)
{
GeglOperationClass *operation_class = GEGL_OPERATION_GET_CLASS (operation);
cl_int cl_err = 0;
/* The kernel will have been compiled by our parent class */
if (!operation_class->cl_data)
return TRUE;
cl_err = gegl_cl_set_kernel_args (operation_class->cl_data->kernel[0],
sizeof(cl_mem), &in_tex,
sizeof(cl_mem), &aux_tex,
sizeof(cl_mem), &out_tex,
NULL);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
operation_class->cl_data->kernel[0], 1,
NULL, &global_worksize, NULL,
0, NULL, NULL);
CL_CHECK;
return FALSE;
error:
return TRUE;
}
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;
}
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;
}
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);
cl_float3 freq;
cl_float3 phaseshift;
cl_int3 keep;
cl_int cl_err = 0;
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;
freq.s[2] = o->cpn_3_frequency * G_PI;
phaseshift.s[0] = G_PI * o->cpn_1_phaseshift / 180.0;
phaseshift.s[1] = G_PI * o->cpn_2_phaseshift / 180.0;
phaseshift.s[2] = G_PI * o->cpn_3_phaseshift / 180.0;
keep.s[0] = (cl_int)o->cpn_1_keep;
keep.s[1] = (cl_int)o->cpn_2_keep;
keep.s[2] = (cl_int)o->cpn_3_keep;
cl_err = 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;
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;
}
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);
CeParamsType *params = (CeParamsType*) o->user_data;
cl_float3 color_diff;
cl_float3 min;
cl_float3 max;
cl_int cl_err = 0;
gint i;
if (!cl_data)
{
const char *kernel_name[] = {"cl_color_exchange",
NULL};
cl_data = gegl_cl_compile_and_build (color_exchange_cl_source, kernel_name);
}
if (!cl_data)
return TRUE;
for (i = 0; i < 3; i++)
{
color_diff.s[i] = params->color_diff[i];
min.s[i] = params->min[i];
max.s[i] = params->max[i];
}
cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
sizeof(cl_mem), &in,
sizeof(cl_mem), &out,
sizeof(cl_float3), &color_diff,
sizeof(cl_float3), &min,
sizeof(cl_float3), &max,
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;
error:
return TRUE;
}
/* 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_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
checkerboard_cl_process (GeglOperation *operation,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *roi,
gint level)
{
GeglProperties *o = GEGL_PROPERTIES (operation);
const Babl *out_format = gegl_operation_get_format (operation, "output");
const size_t gbl_size[2] = {roi->width, roi->height};
const size_t gbl_offs[2] = {roi->x, roi->y};
cl_int cl_err = 0;
float color1[4];
float color2[4];
if (!cl_data)
{
const char *kernel_name[] = {"kernel_checkerboard", NULL};
cl_data = gegl_cl_compile_and_build (checkerboard_cl_source, kernel_name);
if (!cl_data)
return TRUE;
}
gegl_color_get_pixel (o->color1, out_format, color1);
gegl_color_get_pixel (o->color2, out_format, color2);
cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
sizeof(cl_mem), &out_tex,
sizeof(color1), &color1,
sizeof(color2), &color2,
sizeof(cl_int), &o->x,
sizeof(cl_int), &o->y,
sizeof(cl_int), &o->x_offset,
sizeof(cl_int), &o->y_offset,
NULL);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
cl_data->kernel[0], 2,
gbl_offs, gbl_size, 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 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;
}
/* 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 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 *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;
}
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);
cl_float threshold = o->threshold;
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_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_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_pixelize (cl_mem in_tex,
cl_mem aux_tex,
cl_mem out_tex,
const GeglRectangle *src_rect,
const GeglRectangle *roi,
gint xsize,
gint ysize,
gfloat xratio,
gfloat yratio,
gfloat bg_color[4],
gint norm,
GeglRectangle *image_extent)
{
cl_int cl_err = 0;
const size_t gbl_size[2]= {roi->width, roi->height};
gint cx0 = block_index (roi->x, xsize);
gint cy0 = block_index (roi->y, ysize);
gint block_count_x = block_index (roi->x + roi->width + xsize - 1, xsize) - cx0;
gint block_count_y = block_index (roi->y + roi->height + ysize - 1, ysize) - cy0;
cl_int4 bbox = {{ image_extent->x, image_extent->y,
image_extent->x + image_extent->width,
image_extent->y + image_extent->height }};
cl_int line_width = roi->width + 2 * xsize;
size_t gbl_size_tmp[2] = {block_count_x, block_count_y};
if (!cl_data)
{
const char *kernel_name[] = {"calc_block_color", "kernel_pixelize", NULL};
cl_data = gegl_cl_compile_and_build (pixelize_cl_source, kernel_name);
}
if (!cl_data) return 1;
cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
sizeof(cl_mem), (void*)&in_tex,
sizeof(cl_mem), (void*)&aux_tex,
sizeof(cl_int), (void*)&xsize,
sizeof(cl_int), (void*)&ysize,
sizeof(cl_int), (void*)&roi->x,
sizeof(cl_int), (void*)&roi->y,
sizeof(cl_int4), &bbox,
sizeof(cl_int), (void*)&line_width,
sizeof(cl_int), (void*)&block_count_x,
NULL);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
cl_data->kernel[0], 2,
NULL, gbl_size_tmp, NULL,
0, NULL, NULL);
CL_CHECK;
cl_err = gegl_cl_set_kernel_args (cl_data->kernel[1],
sizeof(cl_mem), (void*)&aux_tex,
sizeof(cl_mem), (void*)&out_tex,
sizeof(cl_int), (void*)&xsize,
sizeof(cl_int), (void*)&ysize,
sizeof(cl_float), (void*)&xratio,
sizeof(cl_float), (void*)&yratio,
sizeof(cl_int), (void*)&roi->x,
sizeof(cl_int), (void*)&roi->y,
sizeof(cl_float4),(void*)bg_color,
sizeof(cl_int), (void*)&norm,
sizeof(cl_int), (void*)&block_count_x,
NULL);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
cl_data->kernel[1], 2,
NULL, gbl_size, NULL,
0, NULL, NULL);
CL_CHECK;
return FALSE;
error:
return TRUE;
}
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;
}
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 inline gboolean
_gegl_buffer_cl_cache_flush2 (GeglTileHandlerCache *cache,
const GeglRectangle *roi)
{
size_t size;
GList *elem;
GeglRectangle tmp;
cl_int cl_err = 0;
gpointer data;
gboolean need_cl = FALSE;
for (elem=cache_entries; elem; elem=elem->next)
{
CacheEntry *entry = elem->data;
if (entry->valid && entry->tile_storage->cache == cache
&& (!roi || gegl_rectangle_intersect (&tmp, roi, &entry->roi)))
{
entry->valid = FALSE;
entry->used ++;
gegl_cl_color_babl (entry->buffer->soft_format, &size);
data = g_malloc(entry->roi.width * entry->roi.height * size);
cl_err = gegl_clEnqueueReadBuffer(gegl_cl_get_command_queue(),
entry->tex, CL_TRUE, 0, entry->roi.width * entry->roi.height * size, data,
0, NULL, NULL);
/* tile-ize */
gegl_buffer_set (entry->buffer, &entry->roi, 0, entry->buffer->soft_format, data, GEGL_AUTO_ROWSTRIDE);
entry->used --;
need_cl = TRUE;
g_free(data);
CL_CHECK;
}
}
if (need_cl)
{
cl_err = gegl_clFinish (gegl_cl_get_command_queue ());
CL_CHECK;
g_mutex_lock (&cache_mutex);
while (cache_entry_find_invalid (&data))
{
CacheEntry *entry = data;
#if 1
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Removing from cl-cache: %p %s {%d %d %d %d}", entry->buffer, babl_get_name(entry->buffer->soft_format),
entry->roi.x, entry->roi.y, entry->roi.width, entry->roi.height);
#endif
gegl_clReleaseMemObject(entry->tex);
memset (entry, 0x0, sizeof (CacheEntry));
g_slice_free (CacheEntry, data);
cache_entries = g_list_remove (cache_entries, data);
}
g_mutex_unlock (&cache_mutex);
}
return TRUE;
error:
g_mutex_lock (&cache_mutex);
while (cache_entry_find_invalid (&data))
{
g_slice_free (CacheEntry, data);
cache_entries = g_list_remove (cache_entries, data);
}
g_mutex_unlock (&cache_mutex);
/* XXX : result is corrupted */
return FALSE;
}
gboolean
gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
{
GeglBufferClIterators *i = (gpointer)iterator;
gboolean result = FALSE;
gint no, j;
cl_int cl_err = 0;
if (i->is_finished)
g_error ("%s called on finished buffer iterator", G_STRFUNC);
if (i->iteration_no == 0)
{
for (no=0; no<i->iterators;no++)
{
if (i->buffer[no])
{
gint j;
gboolean found = FALSE;
for (j=0; j<no; j++)
if (i->buffer[no]==i->buffer[j])
{
found = TRUE;
break;
}
if (!found)
gegl_buffer_lock (i->buffer[no]);
if (i->flags[no] == GEGL_CL_BUFFER_WRITE
|| (i->flags[no] == GEGL_CL_BUFFER_READ
&& (i->area[no][0] > 0 || i->area[no][1] > 0 || i->area[no][2] > 0 || i->area[no][3] > 0)))
{
gegl_buffer_cl_cache_flush (i->buffer[no], &i->rect[no]);
}
}
}
}
else
{
/* complete pending write work */
for (no=0; no<i->iterators;no++)
{
if (i->flags[no] == GEGL_CL_BUFFER_WRITE)
{
/* Wait Processing */
cl_err = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
if (cl_err != CL_SUCCESS) CL_ERROR;
/* color conversion in the GPU (output) */
if (i->conv[no] == GEGL_CL_COLOR_CONVERT)
for (j=0; j < i->n; j++)
{
cl_err = gegl_cl_color_conv (i->tex_op[no][j], i->tex_buf[no][j], i->size[no][j],
i->format[no], i->buffer[no]->soft_format);
if (cl_err == FALSE) CL_ERROR;
}
/* Wait Processing */
cl_err = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
if (cl_err != CL_SUCCESS) CL_ERROR;
/* GPU -> CPU */
for (j=0; j < i->n; j++)
{
gpointer data;
/* tile-ize */
if (i->conv[no] == GEGL_CL_COLOR_NOT_SUPPORTED)
{
data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_op[no][j], CL_TRUE,
CL_MAP_READ,
0, i->size[no][j] * i->op_cl_format_size [no],
0, NULL, NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* color conversion using BABL */
gegl_buffer_set (i->buffer[no], &i->roi[no][j], 0, i->format[no], data, GEGL_AUTO_ROWSTRIDE);
cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_op[no][j], data,
0, NULL, NULL);
if (cl_err != CL_SUCCESS) CL_ERROR;
}
else
#ifdef OPENCL_USE_CACHE
{
gegl_buffer_cl_cache_new (i->buffer[no], &i->roi[no][j], i->tex_buf[no][j]);
/* don't release this texture */
i->tex_buf[no][j] = NULL;
}
#else
{
data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_buf[no][j], CL_TRUE,
CL_MAP_READ,
0, i->size[no][j] * i->buf_cl_format_size [no],
0, NULL, NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* color conversion using BABL */
gegl_buffer_set (i->buffer[no], &i->roi[no][j], i->format[no], data, GEGL_AUTO_ROWSTRIDE);
cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_buf[no][j], data,
0, NULL, NULL);
if (cl_err != CL_SUCCESS) CL_ERROR;
}
#endif
}
}
}
/* Run! */
cl_err = gegl_clFinish(gegl_cl_get_command_queue());
if (cl_err != CL_SUCCESS) CL_ERROR;
for (no=0; no < i->iterators; no++)
for (j=0; j < i->n; j++)
{
if (i->tex_buf_from_cache [no][j])
{
gboolean ok = gegl_buffer_cl_cache_release (i->tex_buf[no][j]);
g_assert (ok);
}
if (i->tex_buf[no][j] && !i->tex_buf_from_cache [no][j])
gegl_clReleaseMemObject (i->tex_buf[no][j]);
if (i->tex_op [no][j])
gegl_clReleaseMemObject (i->tex_op [no][j]);
i->tex [no][j] = NULL;
i->tex_buf[no][j] = NULL;
i->tex_op [no][j] = NULL;
}
}
g_assert (i->iterators > 0);
result = (i->roi_no >= i->rois)? FALSE : TRUE;
i->n = MIN(GEGL_CL_NTEX, i->rois - i->roi_no);
/* then we iterate all */
for (no=0; no<i->iterators;no++)
{
for (j = 0; j < i->n; j++)
{
GeglRectangle r = {i->rect[no].x + i->roi_all[i->roi_no+j].x - i->area[no][0],
i->rect[no].y + i->roi_all[i->roi_no+j].y - i->area[no][2],
i->roi_all[i->roi_no+j].width + i->area[no][0] + i->area[no][1],
i->roi_all[i->roi_no+j].height + i->area[no][2] + i->area[no][3]};
i->roi [no][j] = r;
i->size[no][j] = r.width * r.height;
}
if (i->flags[no] == GEGL_CL_BUFFER_READ)
{
for (j=0; j < i->n; j++)
{
gpointer data;
/* un-tile */
switch (i->conv[no])
{
case GEGL_CL_COLOR_NOT_SUPPORTED:
{
gegl_buffer_cl_cache_flush (i->buffer[no], &i->roi[no][j]);
g_assert (i->tex_op[no][j] == NULL);
i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY,
i->size[no][j] * i->op_cl_format_size [no],
NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* pre-pinned memory */
data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_op[no][j], CL_TRUE,
CL_MAP_WRITE,
0, i->size[no][j] * i->op_cl_format_size [no],
0, NULL, NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* color conversion using BABL */
gegl_buffer_get (i->buffer[no], &i->roi[no][j], 1.0, i->format[no], data, GEGL_AUTO_ROWSTRIDE, GEGL_ABYSS_NONE);
cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_op[no][j], data,
0, NULL, NULL);
if (cl_err != CL_SUCCESS) CL_ERROR;
i->tex[no][j] = i->tex_op[no][j];
break;
}
case GEGL_CL_COLOR_EQUAL:
{
i->tex_buf[no][j] = gegl_buffer_cl_cache_get (i->buffer[no], &i->roi[no][j]);
if (i->tex_buf[no][j])
i->tex_buf_from_cache [no][j] = TRUE; /* don't free texture from cache */
else
{
gegl_buffer_cl_cache_flush (i->buffer[no], &i->roi[no][j]);
g_assert (i->tex_buf[no][j] == NULL);
i->tex_buf[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY,
i->size[no][j] * i->buf_cl_format_size [no],
NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* pre-pinned memory */
data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_buf[no][j], CL_TRUE,
CL_MAP_WRITE,
0, i->size[no][j] * i->buf_cl_format_size [no],
0, NULL, NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* color conversion will be performed in the GPU later */
gegl_buffer_cl_worker_transf (i->buffer[no], data, i->buf_cl_format_size [no], i->roi[no][j], FALSE);
cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_buf[no][j], data,
0, NULL, NULL);
if (cl_err != CL_SUCCESS) CL_ERROR;
}
i->tex[no][j] = i->tex_buf[no][j];
break;
}
case GEGL_CL_COLOR_CONVERT:
{
i->tex_buf[no][j] = gegl_buffer_cl_cache_get (i->buffer[no], &i->roi[no][j]);
if (i->tex_buf[no][j])
i->tex_buf_from_cache [no][j] = TRUE; /* don't free texture from cache */
else
{
gegl_buffer_cl_cache_flush (i->buffer[no], &i->roi[no][j]);
g_assert (i->tex_buf[no][j] == NULL);
i->tex_buf[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY,
i->size[no][j] * i->buf_cl_format_size [no],
NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* pre-pinned memory */
data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_buf[no][j], CL_TRUE,
CL_MAP_WRITE,
0, i->size[no][j] * i->buf_cl_format_size [no],
0, NULL, NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* color conversion will be performed in the GPU later */
/* get buffer data using multiple worker threads
to increase bandwidth */
gegl_buffer_cl_worker_transf (i->buffer[no], data, i->buf_cl_format_size [no], i->roi[no][j], FALSE);
cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_buf[no][j], data,
0, NULL, NULL);
if (cl_err != CL_SUCCESS) CL_ERROR;
}
g_assert (i->tex_op[no][j] == NULL);
i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_READ_WRITE,
i->size[no][j] * i->op_cl_format_size [no],
NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* color conversion in the GPU (input) */
g_assert (i->tex_buf[no][j] && i->tex_op[no][j]);
cl_err = gegl_cl_color_conv (i->tex_buf[no][j], i->tex_op[no][j], i->size[no][j],
i->buffer[no]->soft_format, i->format[no]);
if (cl_err == FALSE) CL_ERROR;
i->tex[no][j] = i->tex_op[no][j];
break;
}
}
}
/* Wait Processing */
cl_err = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
if (cl_err != CL_SUCCESS) CL_ERROR;
}
else if (i->flags[no] == GEGL_CL_BUFFER_WRITE)
{
for (j=0; j < i->n; j++)
{
switch (i->conv[no])
{
case GEGL_CL_COLOR_NOT_SUPPORTED:
{
g_assert (i->tex_op[no][j] == NULL);
i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_WRITE_ONLY,
i->size[no][j] * i->op_cl_format_size [no],
NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
i->tex[no][j] = i->tex_op[no][j];
break;
}
case GEGL_CL_COLOR_EQUAL:
{
g_assert (i->tex_buf[no][j] == NULL);
i->tex_buf[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_WRITE, /* cache */
i->size[no][j] * i->buf_cl_format_size [no],
NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
i->tex[no][j] = i->tex_buf[no][j];
break;
}
case GEGL_CL_COLOR_CONVERT:
{
g_assert (i->tex_buf[no][j] == NULL);
i->tex_buf[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_WRITE, /* cache */
i->size[no][j] * i->buf_cl_format_size [no],
NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
g_assert (i->tex_op[no][j] == NULL);
i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_READ_WRITE,
i->size[no][j] * i->op_cl_format_size [no],
NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
i->tex[no][j] = i->tex_op[no][j];
break;
}
}
}
}
else if (i->flags[no] == GEGL_CL_BUFFER_AUX)
{
for (j=0; j < i->n; j++)
{
g_assert (i->tex_op[no][j] == NULL);
i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_READ_WRITE,
i->size[no][j] * i->op_cl_format_size [no],
NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
i->tex[no][j] = i->tex_op[no][j];
}
}
}
i->roi_no += i->n;
i->iteration_no++;
if (result == FALSE)
{
for (no=0; no<i->iterators;no++)
{
if (i->buffer[no])
{
gint j;
gboolean found = FALSE;
for (j=0; j<no; j++)
if (i->buffer[no]==i->buffer[j])
{
found = TRUE;
break;
}
if (!found)
gegl_buffer_unlock (i->buffer[no]);
g_object_unref (i->buffer[no]);
}
}
i->is_finished = TRUE;
g_free (i->roi_all);
g_slice_free (GeglBufferClIterators, i);
}
*err = FALSE;
return result;
error:
for (no=0; no<i->iterators;no++)
for (j=0; j < i->n; j++)
{
if (i->tex_buf[no][j]) gegl_clReleaseMemObject (i->tex_buf[no][j]);
if (i->tex_op [no][j]) gegl_clReleaseMemObject (i->tex_op [no][j]);
i->tex [no][j] = NULL;
i->tex_buf[no][j] = NULL;
i->tex_op [no][j] = NULL;
}
*err = TRUE;
return FALSE;
}
static gboolean
cl_bilateral (cl_mem in_tex,
cl_mem out_tex,
const GeglRectangle *roi,
const GeglRectangle *src_rect,
gint s_sigma,
gfloat r_sigma)
{
cl_int cl_err = 0;
gint c;
const gint width = src_rect->width;
const gint height = src_rect->height;
const gint sw = (width -1) / s_sigma + 1;
const gint sh = (height-1) / s_sigma + 1;
const gint depth = (int)(1.0f / r_sigma) + 1;
size_t global_ws[2];
size_t local_ws[2];
cl_mem grid = NULL;
cl_mem blur[4] = {NULL, NULL, NULL, NULL};
if (!cl_data)
{
const char *kernel_name[] = {"bilateral_downsample",
"bilateral_blur",
"bilateral_interpolate",
NULL};
cl_data = gegl_cl_compile_and_build (bilateral_filter_fast_cl_source, kernel_name);
}
if (!cl_data)
return 1;
grid = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_READ_WRITE,
sw * sh * depth * sizeof(cl_float8),
NULL,
&cl_err);
CL_CHECK;
for(c = 0; c < 4; c++)
{
blur[c] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_READ_WRITE,
sw * sh * depth * sizeof(cl_float2),
NULL, &cl_err);
CL_CHECK;
}
local_ws[0] = 8;
local_ws[1] = 8;
global_ws[0] = ((sw + local_ws[0] - 1)/local_ws[0])*local_ws[0];
global_ws[1] = ((sh + local_ws[1] - 1)/local_ws[1])*local_ws[1];
gegl_cl_set_kernel_args (cl_data->kernel[0],
sizeof(cl_mem), &in_tex,
sizeof(cl_mem), &grid,
sizeof(cl_int), &width,
sizeof(cl_int), &height,
sizeof(cl_int), &sw,
sizeof(cl_int), &sh,
sizeof(cl_int), &depth,
sizeof(cl_int), &s_sigma,
sizeof(cl_float), &r_sigma,
NULL);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
cl_data->kernel[0], 2,
NULL, global_ws, local_ws,
0, NULL, NULL);
CL_CHECK;
local_ws[0] = 16;
local_ws[1] = 16;
global_ws[0] = ((sw + local_ws[0] - 1)/local_ws[0])*local_ws[0];
global_ws[1] = ((sh + local_ws[1] - 1)/local_ws[1])*local_ws[1];
gegl_cl_set_kernel_args (cl_data->kernel[1],
sizeof(cl_mem), &grid,
sizeof(cl_mem), &blur[0],
sizeof(cl_mem), &blur[1],
sizeof(cl_mem), &blur[2],
sizeof(cl_mem), &blur[3],
sizeof(cl_int), &sw,
sizeof(cl_int), &sh,
sizeof(cl_int), &depth,
NULL);
cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
cl_data->kernel[1], 2,
NULL, global_ws, local_ws,
0, NULL, NULL);
CL_CHECK;
global_ws[0] = width;
global_ws[1] = height;
gegl_cl_set_kernel_args (cl_data->kernel[2],
sizeof(cl_mem), &in_tex,
sizeof(cl_mem), &blur[0],
sizeof(cl_mem), &blur[1],
sizeof(cl_mem), &blur[2],
sizeof(cl_mem), &blur[3],
sizeof(cl_mem), &out_tex,
sizeof(cl_int), &width,
sizeof(cl_int), &sw,
sizeof(cl_int), &sh,
sizeof(cl_int), &depth,
sizeof(cl_int), &s_sigma,
sizeof(cl_float), &r_sigma,
NULL);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[2], 2,
NULL, global_ws, NULL,
0, NULL, NULL);
CL_CHECK;
cl_err = gegl_clFinish (gegl_cl_get_command_queue ());
CL_CHECK;
cl_err = gegl_clReleaseMemObject (grid);
CL_CHECK_ONLY (cl_err);
for(c = 0; c < 4; c++)
{
cl_err = gegl_clReleaseMemObject (blur[c]);
CL_CHECK_ONLY (cl_err);
}
return FALSE;
error:
if (grid)
gegl_clReleaseMemObject (grid);
for (c = 0; c < 4; c++)
{
if (blur[c])
gegl_clReleaseMemObject (blur[c]);
}
return TRUE;
}
static gboolean
cl_process (GeglOperation *operation,
cl_mem out_tex,
const GeglRectangle *roi)
{
GeglProperties *o = GEGL_PROPERTIES (operation);
const size_t gbl_size[] = { roi->width, roi->height };
cl_int cl_err = 0;
cl_int offset_x;
cl_int offset_y;
cl_int width;
cl_int height;
cl_float3 sedges;
cl_float3 contours;
cl_float3 frequency;
cl_float brightness;
cl_float polarization;
cl_float scattering;
cl_int iterations;
cl_float weird_factor;
if (!cl_data)
{
const char *kernel_name[] = { "cl_diffraction_patterns", NULL };
cl_data = gegl_cl_compile_and_build (diffraction_patterns_cl_source,
kernel_name);
if (!cl_data)
return TRUE;
}
offset_x = roi->x;
offset_y = roi->y;
width = o->width;
height = o->height;
sedges.s[0] = o->red_sedges;
sedges.s[1] = o->green_sedges;
sedges.s[2] = o->blue_sedges;
contours.s[0] = o->red_contours;
contours.s[1] = o->green_contours;
contours.s[2] = o->blue_contours;
frequency.s[0] = o->red_frequency;
frequency.s[1] = o->green_frequency;
frequency.s[2] = o->blue_frequency;
brightness = o->brightness;
polarization = o->polarization;
scattering = o->scattering;
iterations = ITERATIONS;
weird_factor = WEIRD_FACTOR;
cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
sizeof(cl_mem), &out_tex,
sizeof(cl_int), &offset_x,
sizeof(cl_int), &offset_y,
sizeof(cl_int), &width,
sizeof(cl_int), &height,
sizeof(cl_float3), &sedges,
sizeof(cl_float3), &contours,
sizeof(cl_float3), &frequency,
sizeof(cl_float), &brightness,
sizeof(cl_float), &polarization,
sizeof(cl_float), &scattering,
sizeof(cl_int), &iterations,
sizeof(cl_float), &weird_factor,
NULL);
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;
cl_err = gegl_clFinish (gegl_cl_get_command_queue ());
CL_CHECK;
return FALSE;
error:
return TRUE;
}
static gboolean
cl_box_blur (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], global_ws[2];
size_t local_ws_hor[2], local_ws_ver[2], local_ws[2];
size_t step_size ;
if (!cl_data)
{
const char *kernel_name[] = { "kernel_blur_hor", "kernel_blur_ver","kernel_box_blur_fast", NULL};
cl_data = gegl_cl_compile_and_build (box_blur_cl_source, kernel_name);
}
if (!cl_data)
return TRUE;
step_size = 64;
local_ws[0]=256;
local_ws[1]=1;
if( radius <=110 )
{
global_ws[0] = (roi->width + local_ws[0] - 2 * radius - 1) / ( local_ws[0] - 2 * radius ) * local_ws[0];
global_ws[1] = (roi->height + step_size - 1) / step_size;
cl_err = gegl_cl_set_kernel_args(cl_data->kernel[2],
sizeof(cl_mem), (void *)&in_tex,
sizeof(cl_mem), (void *)&out_tex,
sizeof(cl_float4)*local_ws[0], (void *)NULL,
sizeof(cl_int), (void *)&roi->width,
sizeof(cl_int), (void *)&roi->height,
sizeof(cl_int), (void *)&radius,
sizeof(cl_int), (void *)&step_size, NULL);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(),
cl_data->kernel[2], 2,
NULL, global_ws, local_ws, 0, NULL, NULL );
CL_CHECK;
}
else
{
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_cl_set_kernel_args (cl_data->kernel[0],
sizeof(cl_mem), (void*)&in_tex,
sizeof(cl_mem), (void*)&aux_tex,
sizeof(cl_int), (void*)&roi->width,
sizeof(cl_int), (void*)&radius,
NULL);
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_cl_set_kernel_args (cl_data->kernel[1],
sizeof(cl_mem), (void*)&aux_tex,
sizeof(cl_mem), (void*)&out_tex,
sizeof(cl_int), (void*)&roi->width,
sizeof(cl_int), (void*)&radius,
NULL);
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;
}
