/* This file is an image processing operation for GEGL * * GEGL is free software; you can redistribute it and/or * modify it under the terms of the GNU Lesser General Public * License as published by the Free Software Foundation; either * version 3 of the License, or (at your option) any later version. * * GEGL is distributed in the hope that it will be useful, * but WITHOUT ANY WARRANTY; without even the implied warranty of * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU * Lesser General Public License for more details. * * You should have received a copy of the GNU Lesser General Public * License along with GEGL; if not, see . * * Copyright 2006-2009 Øyvind Kolås */ #include "config.h" #include #ifdef GEGL_CHANT_PROPERTIES /* no properties */ #else #define GEGL_CHANT_TYPE_POINT_COMPOSER #define GEGL_CHANT_C_FILE "over.c" #include "gegl-chant.h" static void prepare (GeglOperation *operation) { const Babl *format = babl_format ("RaGaBaA float"); gegl_operation_set_format (operation, "input", format); gegl_operation_set_format (operation, "aux", format); gegl_operation_set_format (operation, "output", format); } static gboolean process (GeglOperation *op, void *in_buf, void *aux_buf, void *out_buf, glong n_pixels, const GeglRectangle *roi, gint level) { gfloat * GEGL_ALIGNED in = in_buf; gfloat * GEGL_ALIGNED aux = aux_buf; gfloat * GEGL_ALIGNED out = out_buf; if (aux==NULL) return TRUE; while (n_pixels--) { out[0] = aux[0] + in[0] * (1.0f - aux[3]); out[1] = aux[1] + in[1] * (1.0f - aux[3]); out[2] = aux[2] + in[2] * (1.0f - aux[3]); out[3] = aux[3] + in[3] - aux[3] * in[3]; in += 4; aux += 4; out += 4; } return TRUE; } #include "opencl/gegl-cl.h" static const char* kernel_source = "__kernel void kernel_over(__global const float4 *in, \n" " __global const float4 *aux, \n" " __global float4 *out) \n" "{ \n" " int gid = get_global_id(0); \n" " float4 in_v = in [gid]; \n" " float4 aux_v = aux[gid]; \n" " float4 out_v; \n" " out_v.xyz = aux_v.xyz + in_v.xyz * (1.0f - aux_v.w); \n" " out_v.w = aux_v.w + in_v.w - aux_v.w * in_v.w; \n" " out[gid] = out_v; \n" "} \n"; static gegl_cl_run_data *cl_data = NULL; static cl_int 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) { /* Retrieve a pointer to GeglChantO structure which contains all the * chanted properties */ cl_int cl_err = 0; if (aux_tex == NULL) return 0; if (!cl_data) { const char *kernel_name[] = {"kernel_over", 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*)&aux_tex); 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); if (cl_err != CL_SUCCESS) return cl_err; return cl_err; } /* Fast paths */ static gboolean operation_process (GeglOperation *operation, GeglOperationContext *context, const gchar *output_prop, const GeglRectangle *result, gint level) { GeglOperationClass *operation_class; gpointer input, aux; operation_class = GEGL_OPERATION_CLASS (gegl_chant_parent_class); /* get the raw values this does not increase the reference count */ input = gegl_operation_context_get_object (context, "input"); aux = gegl_operation_context_get_object (context, "aux"); /* pass the input/aux buffers directly through if they are alone*/ { const GeglRectangle *in_extent = NULL; const GeglRectangle *aux_extent = NULL; if (input) in_extent = gegl_buffer_get_abyss (input); if ((!input || (aux && !gegl_rectangle_intersect (NULL, in_extent, result)))) { gegl_operation_context_take_object (context, "output", g_object_ref (aux)); return TRUE; } if (aux) aux_extent = gegl_buffer_get_abyss (aux); if (!aux || (input && !gegl_rectangle_intersect (NULL, aux_extent, result))) { gegl_operation_context_take_object (context, "output", g_object_ref (input)); return TRUE; } } /* chain up, which will create the needed buffers for our actual * process function */ return operation_class->process (operation, context, output_prop, result, level); } static void gegl_chant_class_init (GeglChantClass *klass) { GeglOperationClass *operation_class; GeglOperationPointComposerClass *point_composer_class; operation_class = GEGL_OPERATION_CLASS (klass); point_composer_class = GEGL_OPERATION_POINT_COMPOSER_CLASS (klass); operation_class->prepare = prepare; operation_class->process = operation_process; point_composer_class->process = process; point_composer_class->cl_process = cl_process; operation_class->opencl_support = TRUE; operation_class->compat_name = "gegl:over"; gegl_operation_class_set_keys (operation_class, "name" , "svg:src-over", "categories" , "compositors:porter-duff", "description", _("Porter Duff operation over (d = cA + cB * (1 - aA))"), NULL); } #endif