GeglNode *
gegl_operation_detect (GeglOperation *operation,
gint x,
gint y)
{
GeglNode *node = NULL;
GeglOperationClass *klass;
if (!operation)
return NULL;
g_return_val_if_fail (GEGL_IS_OPERATION (operation), NULL);
node = operation->node;
klass = GEGL_OPERATION_GET_CLASS (operation);
if (klass->detect)
{
return klass->detect (operation, x, y);
}
if (x >= node->have_rect.x &&
x < node->have_rect.x + node->have_rect.width &&
y >= node->have_rect.y &&
y < node->have_rect.y + node->have_rect.height)
{
return node;
}
return NULL;
}
void
gegl_operation_context_get_property (GeglOperationContext *context,
const gchar *property_name,
GValue *value)
{
GParamSpec *pspec;
GValue *storage;
pspec = g_object_class_find_property (G_OBJECT_GET_CLASS (context->operation),
property_name);
if (!pspec)
{
g_warning ("%s: node %s has no pad|property named '%s'",
G_STRFUNC,
GEGL_OPERATION_GET_CLASS (context->operation)->name,
property_name);
}
storage = gegl_operation_context_get_value (context, property_name);
if (storage != NULL)
{
g_value_copy (storage, value);
}
}
gboolean
gegl_operation_process (GeglOperation *operation,
GeglOperationContext *context,
const gchar *output_pad,
const GeglRectangle *result,
gint level)
{
GeglOperationClass *klass;
g_return_val_if_fail (GEGL_IS_OPERATION (operation), FALSE);
klass = GEGL_OPERATION_GET_CLASS (operation);
if (!strcmp (output_pad, "output") &&
(result->width == 0 || result->height == 0))
{
GeglBuffer *output = gegl_buffer_new (NULL, NULL);
g_warning ("%s Eeek: processing 0px rectangle", G_STRLOC);
/* when this case is hit.. we've done something bad.. */
gegl_operation_context_take_object (context, "output", G_OBJECT (output));
return TRUE;
}
return klass->process (operation, context, output_pad, result, level);
}
/* Calls the prepare function on the operation that extends this base class */
void
gegl_operation_prepare (GeglOperation *self)
{
GeglOperationClass *klass;
g_return_if_fail (GEGL_IS_OPERATION (self));
klass = GEGL_OPERATION_GET_CLASS (self);
/* build OpenCL kernel */
if (!klass->cl_data)
{
const gchar *cl_source = gegl_operation_class_get_key (klass, "cl-source");
if (cl_source)
{
char *name = strdup (klass->name);
const char *kernel_name[] = {name, NULL};
char *k;
for (k=name; *k; k++)
switch (*k)
{
case ' ': case ':': case '-':
*k = '_';
break;
}
klass->cl_data = gegl_cl_compile_and_build (cl_source, kernel_name);
free (name);
}
}
if (klass->prepare)
klass->prepare (self);
}
static gboolean
gegl_operation_point_filter_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");
GeglOperationPointFilterClass *point_filter_class;
GeglOperationClass *operation_class = GEGL_OPERATION_GET_CLASS (operation);
point_filter_class = GEGL_OPERATION_POINT_FILTER_GET_CLASS (operation);
if ((result->width > 0) && (result->height > 0))
{
if (gegl_operation_use_opencl (operation) &&
(operation_class->cl_data || point_filter_class->cl_process))
{
if (gegl_operation_point_filter_cl_process (operation, input, output, result, level))
return TRUE;
}
{
GeglBufferIterator *i = gegl_buffer_iterator_new (output, result, level, out_format, GEGL_BUFFER_WRITE, GEGL_ABYSS_NONE);
gint read = /*output == input ? 0 :*/ gegl_buffer_iterator_add (i, input, result, level, in_format, GEGL_BUFFER_READ, GEGL_ABYSS_NONE);
/* using separate read and write iterators for in-place ideally a single
* readwrite indice would be sufficient
*/
while (gegl_buffer_iterator_next (i))
point_filter_class->process (operation, i->data[read], i->data[0], i->length, &i->roi[0], level);
}
}
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;
}
void
gegl_operation_context_set_property (GeglOperationContext *context,
const gchar *property_name,
const GValue *value)
{
GParamSpec *pspec;
GValue *storage;
g_return_if_fail (context != NULL);
pspec = g_object_class_find_property (G_OBJECT_GET_CLASS (G_OBJECT (context->operation)), property_name);
if (!pspec)
{
g_warning ("%s: node %s has no pad|property named '%s'",
G_STRFUNC,
GEGL_OPERATION_GET_CLASS (context->operation)->name,
property_name);
}
/* if the value already exists in the context it will be reused */
storage = gegl_operation_context_add_value (context, property_name, G_PARAM_SPEC_VALUE_TYPE(pspec));
/* storage needs to have the correct type */
g_value_copy (value, storage);
}
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
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;
}
/* Calls an extending class' get_bound_box method if defined otherwise
* just returns a zero-initiliased bouding box
*/
GeglRectangle
gegl_operation_get_bounding_box (GeglOperation *self)
{
GeglOperationClass *klass = GEGL_OPERATION_GET_CLASS (self);
GeglRectangle rect = { 0, 0, 0, 0 };
if (klass->get_bounding_box)
return klass->get_bounding_box (self);
return rect;
}
const gchar *
gegl_operation_get_name (GeglOperation *operation)
{
GeglOperationClass *klass;
g_return_val_if_fail (GEGL_IS_OPERATION (operation), NULL);
klass = GEGL_OPERATION_GET_CLASS (operation);
return klass->name;
}
GeglRectangle
gegl_operation_get_cached_region (GeglOperation *operation,
const GeglRectangle *roi)
{
GeglOperationClass *klass = GEGL_OPERATION_GET_CLASS (operation);
if (!klass->get_cached_region)
{
return *roi;
}
return klass->get_cached_region (operation, roi);
}
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;
}
GeglRectangle
gegl_operation_get_required_for_output (GeglOperation *operation,
const gchar *input_pad,
const GeglRectangle *roi)
{
GeglOperationClass *klass = GEGL_OPERATION_GET_CLASS (operation);
if (roi->width == 0 ||
roi->height == 0)
return *roi;
g_assert (klass->get_required_for_output);
return klass->get_required_for_output (operation, input_pad, roi);
}
void
gegl_operation_attach (GeglOperation *self,
GeglNode *node)
{
GeglOperationClass *klass;
g_return_if_fail (GEGL_IS_OPERATION (self));
g_return_if_fail (GEGL_IS_NODE (node));
klass = GEGL_OPERATION_GET_CLASS (self);
g_assert (klass->attach);
self->node = node;
klass->attach (self);
}
void
gegl_operation_context_remove_property (GeglOperationContext *self,
const gchar *property_name)
{
Property *property = NULL;
GSList *found;
found = g_slist_find_custom (self->property, property_name, lookup_property);
if (found)
property = found->data;
if (!property)
{
g_warning ("didn't find property %s for %s", property_name,
GEGL_OPERATION_GET_CLASS (self->operation)->name);
return;
}
self->property = g_slist_remove (self->property, property);
property_destroy (property);
}
GeglBuffer *
gegl_operation_context_get_output_maybe_in_place (GeglOperation *operation,
GeglOperationContext *context,
GeglBuffer *input,
const GeglRectangle *roi)
{
GeglOperationClass *klass = GEGL_OPERATION_GET_CLASS (operation);
GeglBuffer *output;
if (klass->want_in_place &&
gegl_can_do_inplace_processing (operation, input, roi))
{
output = g_object_ref (input);
gegl_operation_context_take_object (context, "output", G_OBJECT (output));
}
else
{
output = gegl_operation_context_get_target (context, "output");
}
return output;
}
GeglRectangle
gegl_operation_get_invalidated_by_change (GeglOperation *self,
const gchar *input_pad,
const GeglRectangle *input_region)
{
GeglOperationClass *klass;
GeglRectangle retval = { 0, };
g_return_val_if_fail (GEGL_IS_OPERATION (self), retval);
g_return_val_if_fail (input_pad != NULL, retval);
g_return_val_if_fail (input_region != NULL, retval);
klass = GEGL_OPERATION_GET_CLASS (self);
if (input_region->width == 0 ||
input_region->height == 0)
return *input_region;
if (klass->get_invalidated_by_change)
return klass->get_invalidated_by_change (self, input_pad, input_region);
return *input_region;
}
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;
}
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;
}
