static gboolean
cl_process (GeglOperation *operation,
GeglBuffer *input,
GeglBuffer *output,
const GeglRectangle *result)
{
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *out_format = gegl_operation_get_format (operation, "output");
gint err = 0;
GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation);
GeglProperties *o = GEGL_PROPERTIES (operation);
GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output,
result,
out_format,
GEGL_CL_BUFFER_WRITE);
gint read = gegl_buffer_cl_iterator_add_2 (i,
input,
result,
in_format,
GEGL_CL_BUFFER_READ,
op_area->left,
op_area->right,
op_area->top,
op_area->bottom,
GEGL_ABYSS_CLAMP);
gint aux = gegl_buffer_cl_iterator_add_aux (i,
result,
in_format,
0,
0,
op_area->top,
op_area->bottom);
while (gegl_buffer_cl_iterator_next (i, &err) && !err)
{
err = cl_box_blur (i->tex[read],
i->tex[aux],
i->tex[0],
i->size[0],
&i->roi[0],
o->radius);
if (err)
{
gegl_buffer_cl_iterator_stop (i);
break;
}
}
return !err;
}
static gboolean
operation_source_process (GeglOperation *operation,
GeglBuffer *output,
const GeglRectangle *result,
gint level)
{
const Babl *out_format = gegl_operation_get_format (operation, "output");
if ((result->width > 0) && (result->height > 0))
{
GeglBufferIterator *iter;
if (gegl_operation_use_opencl (operation) &&
babl_format_get_n_components (out_format) == 4 &&
babl_format_get_type (out_format, 0) == babl_type ("float"))
{
GeglBufferClIterator *cl_iter;
gboolean err;
GEGL_NOTE (GEGL_DEBUG_OPENCL, "GEGL_OPERATION_POINT_RENDER: %s", GEGL_OPERATION_GET_CLASS (operation)->name);
cl_iter = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE);
while (gegl_buffer_cl_iterator_next (cl_iter, &err) && !err)
{
err = checkerboard_cl_process (operation, cl_iter->tex[0], cl_iter->size[0], &cl_iter->roi[0], level);
if (err)
{
gegl_buffer_cl_iterator_stop (cl_iter);
break;
}
}
if (err)
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", GEGL_OPERATION_GET_CLASS (operation)->name);
else
return TRUE;
}
iter = gegl_buffer_iterator_new (output, result, level, out_format,
GEGL_ACCESS_WRITE, GEGL_ABYSS_NONE);
while (gegl_buffer_iterator_next (iter))
checkerboard_process (operation, iter->data[0], iter->length, &iter->roi[0], level);
}
return TRUE;
}
static gboolean
process (GeglOperation *operation,
GeglBuffer *out_buf,
const GeglRectangle *roi,
gint level)
{
GeglBufferIterator *iter;
const Babl *out_format = gegl_operation_get_format (operation,
"output");
if (gegl_operation_use_opencl (operation))
{
GeglBufferClIterator *cl_iter;
gboolean err;
GEGL_NOTE (GEGL_DEBUG_OPENCL, "GEGL_OPERATION_POINT_RENDER: %s",
GEGL_OPERATION_GET_CLASS (operation)->name);
cl_iter = gegl_buffer_cl_iterator_new (out_buf, roi, out_format,
GEGL_CL_BUFFER_WRITE);
while (gegl_buffer_cl_iterator_next (cl_iter, &err) && !err)
{
err = cl_process (operation, cl_iter->tex[0], cl_iter->roi);
if (err)
{
gegl_buffer_cl_iterator_stop (cl_iter);
break;
}
}
if (err)
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s",
GEGL_OPERATION_GET_CLASS (operation)->name);
else
return TRUE;
}
iter = gegl_buffer_iterator_new (out_buf, roi, level, out_format,
GEGL_ACCESS_WRITE, GEGL_ABYSS_NONE);
while (gegl_buffer_iterator_next (iter))
c_process (operation, iter->data[0], &iter->roi[0]);
return TRUE;
}
static gboolean
bilateral_cl_process (GeglOperation *operation,
GeglBuffer *input,
GeglBuffer *output,
const GeglRectangle *result,
gint s_sigma,
gfloat r_sigma)
{
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *out_format = gegl_operation_get_format (operation, "output");
gint err = 0;
GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output,
result,
out_format,
GEGL_CL_BUFFER_WRITE);
gint read = gegl_buffer_cl_iterator_add (i,
input,
result,
in_format,
GEGL_CL_BUFFER_READ,
GEGL_ABYSS_NONE);
while (gegl_buffer_cl_iterator_next (i, &err) && !err)
{
err = cl_bilateral (i->tex[read],
i->tex[0],
&i->roi[0],
&i->roi[read],
s_sigma,
r_sigma);
if (err)
{
gegl_buffer_cl_iterator_stop (i);
break;
}
}
return !err;
}
static gboolean
cl_process (GeglOperation *operation,
GeglBuffer *input,
GeglBuffer *output,
const GeglRectangle *roi)
{
const Babl *in_format = babl_format ("RaGaBaA float");
const Babl *out_format = babl_format ("RaGaBaA float");
gint err;
gfloat bg_color[4];
gint norm;
GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation);
GeglProperties *o = GEGL_PROPERTIES (operation);
GeglRectangle *image_extent;
GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output,
roi,
out_format,
GEGL_CL_BUFFER_WRITE);
gint read = gegl_buffer_cl_iterator_add_2 (i,
input,
roi,
in_format,
GEGL_CL_BUFFER_READ,
op_area->left,
op_area->right,
op_area->top,
op_area->bottom,
GEGL_ABYSS_CLAMP);
gint aux = gegl_buffer_cl_iterator_add_aux (i,
roi,
in_format,
op_area->left,
op_area->right,
op_area->top,
op_area->bottom);
gegl_color_get_pixel (o->background, babl_format ("RaGaBaA float"), bg_color);
norm = 0;
norm = o->norm == GEGL_PIXELIZE_NORM_EUCLIDEAN ? 1 : norm;
norm = o->norm == GEGL_PIXELIZE_NORM_INFINITY ? 2 : norm;
image_extent = gegl_operation_source_get_bounding_box (operation, "input");
while (gegl_buffer_cl_iterator_next (i, &err) && !err)
{
err = cl_pixelize(i->tex[read],
i->tex[aux],
i->tex[0],
&i->roi[read],
&i->roi[0],
o->size_x,
o->size_y,
o->ratio_x,
o->ratio_y,
bg_color,
norm,
image_extent);
if (err)
{
gegl_buffer_cl_iterator_stop (i);
break;
}
}
return !err;
}
static gboolean
cl_process (GeglOperation *operation,
GeglBuffer *input,
GeglBuffer *output,
const GeglRectangle *result)
{
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *out_format = gegl_operation_get_format (operation, "output");
gfloat min[] = {1.0f, 1.0f, 1.0f, 1.0f};
gfloat max[] = {0.0f, 0.0f, 0.0f, 0.0f};
gfloat i_min[4], i_max[4], diff[4];
cl_int err = 0;
gint read, c;
GeglBufferClIterator *i;
GeglProperties *o;
cl_float4 cl_min, cl_diff;
o = GEGL_PROPERTIES (operation);
if (cl_build_kernels ())
return FALSE;
i = gegl_buffer_cl_iterator_new (input,
result,
in_format,
GEGL_CL_BUFFER_READ);
while (gegl_buffer_cl_iterator_next (i, &err) && !err)
{
err = cl_buffer_get_min_max (i->tex[0],
i->size[0],
&i->roi[0],
i_min,
i_max);
if (err)
{
gegl_buffer_cl_iterator_stop (i);
break;
}
for (c = 0; c < 3; c++)
{
if (i_min[c] < min[c])
min[c] = i_min[c];
if (i_max[c] > max[c])
max[c] = i_max[c];
}
}
if (err)
return FALSE;
if (o->keep_colors)
reduce_min_max_global (min, max);
for (c = 0; c < 3; c ++)
{
diff[c] = max[c] - min[c];
/* Avoid a divide by zero error if the image is a solid color */
if (diff[c] < 1e-3)
{
min[c] = 0.0;
diff[c] = 1.0;
}
}
cl_diff.x = diff[0];
cl_diff.y = diff[1];
cl_diff.z = diff[2];
cl_diff.w = 1.0;
cl_min.x = min[0];
cl_min.y = min[1];
cl_min.z = min[2];
cl_min.w = 0.0;
i = gegl_buffer_cl_iterator_new (output,
result,
out_format,
GEGL_CL_BUFFER_WRITE);
read = gegl_buffer_cl_iterator_add_2 (i,
input,
result,
in_format,
GEGL_CL_BUFFER_READ,
0,
0,
0,
0,
GEGL_ABYSS_NONE);
while (gegl_buffer_cl_iterator_next (i, &err) && !err)
{
err = cl_stretch_contrast (i->tex[read],
i->tex[0],
i->size[0],
&i->roi[0],
cl_min,
cl_diff);
if (err)
{
gegl_buffer_cl_iterator_stop (i);
break;
}
}
return !err;
}
static gboolean
fir_cl_process (GeglBuffer *input,
GeglBuffer *output,
const GeglRectangle *result,
const Babl *format,
gfloat *cmatrix,
gint clen,
GeglOrientation orientation,
GeglAbyssPolicy abyss)
{
gboolean err = FALSE;
cl_int cl_err;
cl_mem cl_cmatrix = NULL;
GeglBufferClIterator *i;
gint read;
gint left, right, top, bottom;
if (orientation == GEGL_ORIENTATION_HORIZONTAL)
{
right = left = clen / 2;
top = bottom = 0;
}
else
{
right = left = 0;
top = bottom = clen / 2;
}
i = gegl_buffer_cl_iterator_new (output,
result,
format,
GEGL_CL_BUFFER_WRITE);
read = gegl_buffer_cl_iterator_add_2 (i,
input,
result,
format,
GEGL_CL_BUFFER_READ,
left, right,
top, bottom,
abyss);
cl_cmatrix = gegl_clCreateBuffer (gegl_cl_get_context(),
CL_MEM_COPY_HOST_PTR | CL_MEM_READ_ONLY,
clen * sizeof(cl_float), cmatrix, &cl_err);
CL_CHECK;
while (gegl_buffer_cl_iterator_next (i, &err) && !err)
{
err = cl_gaussian_blur(i->tex[read],
i->tex[0],
&i->roi[0],
cl_cmatrix,
clen,
orientation);
if (err)
{
gegl_buffer_cl_iterator_stop (i);
break;
}
}
cl_err = gegl_clReleaseMemObject (cl_cmatrix);
CL_CHECK;
cl_cmatrix = NULL;
return !err;
error:
if (cl_cmatrix)
gegl_clReleaseMemObject (cl_cmatrix);
return FALSE;
}
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;
}
