/* 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 <http://www.gnu.org/licenses/>.
*
* Copyright 2006 Øyvind Kolås <pippin@gimp.org>
*/
#include "config.h"
#include <glib/gi18n-lib.h>
#ifdef GEGL_CHANT_PROPERTIES
/* no properties */
#else
#define GEGL_CHANT_TYPE_POINT_FILTER
#define GEGL_CHANT_C_FILE "invert.c"
#define GEGLV4
#include "gegl-chant.h"
static gboolean
process (GeglOperation *op,
void *in_buf,
void *out_buf,
glong samples,
const GeglRectangle *roi,
gint level)
{
glong i;
gfloat *in = in_buf;
gfloat *out = out_buf;
for (i=0; i<samples; i++)
{
int j;
for (j=0; j<3; j++)
{
gfloat c;
c = in[j];
c = 1.0 - c;
out[j] = c;
}
out[3]=in[3];
in += 4;
out+= 4;
}
return TRUE;
}
#include "opencl/gegl-cl.h"
static const char* kernel_source =
"__kernel void kernel_inv(__global const float4 *in, \n"
" __global float4 *out) \n"
"{ \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in[gid]; \n"
" float4 out_v; \n"
" out_v.xyz = (1.0f - in_v.xyz); \n"
" out_v.w = in_v.w; \n"
" out[gid] = out_v; \n"
"} \n";
static gegl_cl_run_data *cl_data = NULL;
/* OpenCL processing function */
static cl_int
cl_process (GeglOperation *op,
cl_mem in_tex,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *roi,
int level)
{
cl_int cl_err = 0;
if (!cl_data)
{
const char *kernel_name[] = {"kernel_inv", 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*)&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;
}
static void
gegl_chant_class_init (GeglChantClass *klass)
{
GeglOperationClass *operation_class;
GeglOperationPointFilterClass *point_filter_class;
operation_class = GEGL_OPERATION_CLASS (klass);
point_filter_class = GEGL_OPERATION_POINT_FILTER_CLASS (klass);
point_filter_class->process = process;
point_filter_class->cl_process = cl_process;
operation_class->opencl_support = TRUE;
gegl_operation_class_set_keys (operation_class,
"name" , "gegl:invert",
"categories" , "color",
"description",
_("Inverts the components (except alpha), the result is the "
"corresponding \"negative\" image."),
NULL);
}
#endif