static gboolean
cl_process (GeglOperation *operation,
GeglBuffer *input,
GeglBuffer *output,
const GeglRectangle *result)
{
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *out_format = gegl_operation_get_format (operation, "output");
gint err;
gint j;
cl_int cl_err;
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE, GEGL_ABYSS_NONE);
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))
{
if (err) return FALSE;
for (j=0; j < i->n; j++)
{
cl_err = cl_mono_mixer(i->tex[read][j], i->tex[0][j], i->size[0][j], &i->roi[0][j], o->red ,o->green , o->blue);
if (cl_err != CL_SUCCESS)
{
g_warning("[OpenCL] Error in gegl:mono-mixer: %s", gegl_cl_errstring(cl_err));
return FALSE;
}
}
}
return TRUE;
}
GeglBufferClIterator *
gegl_buffer_cl_iterator_new (GeglBuffer *buffer,
const GeglRectangle *roi,
const Babl *format,
guint flags,
GeglAbyssPolicy abyss_policy)
{
GeglBufferClIterator *i = (gpointer)g_slice_new0 (GeglBufferClIterators);
/* Because the iterator is nulled above, we can forgo explicitly setting
* i->is_finished to FALSE. */
gegl_buffer_cl_iterator_add (i, buffer, roi, format, flags, abyss_policy);
return i;
}
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
process (GeglOperation *operation,
GeglBuffer *input,
const GeglRectangle *result,
gint level)
{
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
if (o->buffer)
{
GeglBuffer *output = GEGL_BUFFER (o->buffer);
const Babl *in_format = gegl_buffer_get_format (input);
const Babl *out_format = gegl_buffer_get_format (output);
if (gegl_operation_use_opencl (operation)
&& gegl_cl_color_supported (in_format, out_format) == GEGL_CL_COLOR_CONVERT)
{
size_t size;
gboolean err;
cl_int cl_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,
out_format,
GEGL_CL_BUFFER_READ,
GEGL_ABYSS_NONE);
gegl_cl_color_babl (out_format, &size);
GEGL_NOTE (GEGL_DEBUG_OPENCL,
"write-buffer: "
"%p %p %s %s {%d %d %d %d}",
input,
output,
babl_get_name (in_format),
babl_get_name (out_format),
result->x,
result->y,
result->width,
result->height);
while (gegl_buffer_cl_iterator_next (i, &err))
{
if (err) break;
cl_err = gegl_clEnqueueCopyBuffer (gegl_cl_get_command_queue (),
i->tex[read],
i->tex[0],
0,
0,
i->size[0] * size,
0,
NULL,
NULL);
if (cl_err != CL_SUCCESS)
{
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", gegl_cl_errstring (cl_err));
break;
}
}
if (cl_err || err)
gegl_buffer_copy (input, result, output, result);
}
else
gegl_buffer_copy (input, result, output, result);
gegl_buffer_flush (output);
}
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;
}
