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;
GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (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_2 (i, input, result, in_format, GEGL_CL_BUFFER_READ, op_area->left, op_area->right, op_area->top, op_area->bottom, GEGL_ABYSS_NONE);
gint aux = gegl_buffer_cl_iterator_add_2 (i, NULL, result, in_format, GEGL_CL_BUFFER_AUX, op_area->left, op_area->right, op_area->top, op_area->bottom, 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_edge_laplace(i->tex[read][j], i->tex[aux][j], i->tex[0][j], &i->roi[read][j], &i->roi[0][j], LAPLACE_RADIUS);
if (cl_err != CL_SUCCESS)
{
g_warning("[OpenCL] Error in gegl:edge-laplace: %s", gegl_cl_errstring(cl_err));
return FALSE;
}
}
}
return TRUE;
}
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;
GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation);
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_2 (i, input, result, in_format, GEGL_CL_BUFFER_READ, op_area->left, op_area->right, op_area->top, op_area->bottom, 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_bilateral_filter(i->tex[read][j], i->tex[0][j], i->size[0][j], &i->roi[0][j], ceil(o->blur_radius), o->edge_preservation);
if (cl_err != CL_SUCCESS)
{
g_warning("[OpenCL] Error in gegl:bilateral-filter: %s", gegl_cl_errstring(cl_err));
return FALSE;
}
}
}
return TRUE;
}
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;
}
/* XXX: same program_source with different kernel_name[], context or device
* will retrieve the same key
*/
gegl_cl_run_data *
gegl_cl_compile_and_build (const char *program_source, const char *kernel_name[])
{
gint errcode;
gegl_cl_run_data *cl_data = NULL;
if ((cl_data = (gegl_cl_run_data *)g_hash_table_lookup(cl_program_hash, program_source)) == NULL)
{
size_t length = strlen(program_source);
gint i;
guint kernel_n = 0;
while (kernel_name[++kernel_n] != NULL);
cl_data = (gegl_cl_run_data *) g_malloc(sizeof(gegl_cl_run_data)+sizeof(cl_kernel)*kernel_n);
CL_SAFE_CALL( cl_data->program = gegl_clCreateProgramWithSource(gegl_cl_get_context(), 1, &program_source,
&length, &errcode) );
errcode = gegl_clBuildProgram(cl_data->program, 0, NULL, NULL, NULL, NULL);
if (errcode != CL_SUCCESS)
{
char buffer[2000];
CL_SAFE_CALL( errcode = gegl_clGetProgramBuildInfo(cl_data->program,
gegl_cl_get_device_id(),
CL_PROGRAM_BUILD_LOG,
sizeof(buffer), buffer, NULL) );
g_warning("OpenCL Build Error:%s\n%s", gegl_cl_errstring(errcode), buffer);
return NULL;
}
else
{
g_printf("[OpenCL] Compiling successful\n");
}
for (i=0; i<kernel_n; i++)
CL_SAFE_CALL( cl_data->kernel[i] =
gegl_clCreateKernel(cl_data->program, kernel_name[i], &errcode) );
g_hash_table_insert(cl_program_hash, g_strdup (program_source), (void*)cl_data);
}
return cl_data;
}
static gboolean
cl_process (GeglOperation *operation,
GeglBuffer *input,
GeglBuffer *output,
const GeglRectangle *result,
const GeglRectangle *src_rect)
{
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;
GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation);
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
gdouble theta = o->angle * G_PI / 180.0;
gfloat offset_x = (gfloat)(o->length * cos(theta));
gfloat offset_y = (gfloat)(o->length * sin(theta));
gint num_steps = (gint)ceil(o->length) + 1;
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_2 (i, input, result, in_format, GEGL_CL_BUFFER_READ,
op_area->left, op_area->right, op_area->top, op_area->bottom, 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_motion_blur(i->tex[read][j], i->tex[0][j], i->size[0][j], &i->roi[0][j], &i->roi[read][j], num_steps, offset_x, offset_y);
if (cl_err != CL_SUCCESS)
{
g_warning("[OpenCL] Error in gegl:motion-blur: %s", gegl_cl_errstring(cl_err));
return FALSE;
}
}
}
return TRUE;
}
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;
}
