|
Packit |
bc1512 |
/* This file is an image processing operation for GEGL
|
|
Packit |
bc1512 |
*
|
|
Packit |
bc1512 |
* GEGL is free software; you can redistribute it and/or
|
|
Packit |
bc1512 |
* modify it under the terms of the GNU Lesser General Public
|
|
Packit |
bc1512 |
* License as published by the Free Software Foundation; either
|
|
Packit |
bc1512 |
* version 3 of the License, or (at your option) any later version.
|
|
Packit |
bc1512 |
*
|
|
Packit |
bc1512 |
* GEGL is distributed in the hope that it will be useful,
|
|
Packit |
bc1512 |
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
|
Packit |
bc1512 |
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
|
|
Packit |
bc1512 |
* Lesser General Public License for more details.
|
|
Packit |
bc1512 |
*
|
|
Packit |
bc1512 |
* You should have received a copy of the GNU Lesser General Public
|
|
Packit |
bc1512 |
* License along with GEGL; if not, see <http://www.gnu.org/licenses/>.
|
|
Packit |
bc1512 |
*
|
|
Packit |
bc1512 |
* Copyright 2005 Øyvind Kolås <pippin@gimp.org>,
|
|
Packit |
bc1512 |
* 2007 Øyvind Kolås <oeyvindk@hig.no>
|
|
Packit |
bc1512 |
*/
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
#include "config.h"
|
|
Packit |
bc1512 |
#include <glib/gi18n-lib.h>
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
#ifdef GEGL_CHANT_PROPERTIES
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
gegl_chant_int_ui (radius, _("Radius"), 0, 100, 8, 0, 40, 1.5,
|
|
Packit |
bc1512 |
_("Radius of square pixel region, (width and height will be radius*2+1)"))
|
|
Packit |
bc1512 |
gegl_chant_int (pairs, _("Pairs"), 1, 2, 2,
|
|
Packit |
bc1512 |
_("Number of pairs; higher number preserves more acute features"))
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
#else
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
#define GEGL_CHANT_TYPE_AREA_FILTER
|
|
Packit |
bc1512 |
#define GEGL_CHANT_C_FILE "snn-mean.c"
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
#include "gegl-chant.h"
|
|
Packit |
bc1512 |
#include <math.h>
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
static void
|
|
Packit |
bc1512 |
snn_mean (GeglBuffer *src,
|
|
Packit |
bc1512 |
GeglBuffer *dst,
|
|
Packit |
bc1512 |
const GeglRectangle *dst_rect,
|
|
Packit |
bc1512 |
gdouble radius,
|
|
Packit |
bc1512 |
gint pairs);
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
static void prepare (GeglOperation *operation)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
GeglOperationAreaFilter *area = GEGL_OPERATION_AREA_FILTER (operation);
|
|
Packit |
bc1512 |
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
area->left = area->right = area->top = area->bottom = ceil (o->radius);
|
|
Packit |
bc1512 |
gegl_operation_set_format (operation, "input", babl_format ("RGBA float"));
|
|
Packit |
bc1512 |
gegl_operation_set_format (operation, "output", babl_format ("RGBA float"));
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
static gboolean
|
|
Packit |
bc1512 |
cl_process (GeglOperation *operation,
|
|
Packit |
bc1512 |
GeglBuffer *input,
|
|
Packit |
bc1512 |
GeglBuffer *output,
|
|
Packit |
bc1512 |
const GeglRectangle *result);
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
static gboolean
|
|
Packit |
bc1512 |
process (GeglOperation *operation,
|
|
Packit |
bc1512 |
GeglBuffer *input,
|
|
Packit |
bc1512 |
GeglBuffer *output,
|
|
Packit |
bc1512 |
const GeglRectangle *result,
|
|
Packit |
bc1512 |
gint level)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
|
|
Packit |
bc1512 |
GeglBuffer *temp_in;
|
|
Packit |
bc1512 |
GeglRectangle compute;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
if (gegl_cl_is_accelerated ())
|
|
Packit |
bc1512 |
if (cl_process (operation, input, output, result))
|
|
Packit |
bc1512 |
return TRUE;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
compute = gegl_operation_get_required_for_output (
|
|
Packit |
bc1512 |
operation, "input", result);
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
if (o->radius < 1.0)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
output = g_object_ref (input);
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
else
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
temp_in = gegl_buffer_create_sub_buffer (input, &compute);
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
snn_mean (temp_in, output, result, o->radius, o->pairs);
|
|
Packit |
bc1512 |
g_object_unref (temp_in);
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
return TRUE;
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
#define RGB_LUMINANCE_RED (0.212671)
|
|
Packit |
bc1512 |
#define RGB_LUMINANCE_GREEN (0.715160)
|
|
Packit |
bc1512 |
#define RGB_LUMINANCE_BLUE (0.072169)
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
static inline gfloat rgb2luminance (gfloat *pix)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
return pix[0] * RGB_LUMINANCE_RED +
|
|
Packit |
bc1512 |
pix[1] * RGB_LUMINANCE_GREEN +
|
|
Packit |
bc1512 |
pix[2] * RGB_LUMINANCE_BLUE;
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
#define POW2(a)((a)*(a))
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
static inline gfloat colordiff (gfloat *pixA,
|
|
Packit |
bc1512 |
gfloat *pixB)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
return POW2(pixA[0]-pixB[0])+
|
|
Packit |
bc1512 |
POW2(pixA[1]-pixB[1])+
|
|
Packit |
bc1512 |
POW2(pixA[2]-pixB[2]);
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
static void
|
|
Packit |
bc1512 |
snn_mean (GeglBuffer *src,
|
|
Packit |
bc1512 |
GeglBuffer *dst,
|
|
Packit |
bc1512 |
const GeglRectangle *dst_rect,
|
|
Packit |
bc1512 |
gdouble dradius,
|
|
Packit |
bc1512 |
gint pairs)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
gint x,y;
|
|
Packit |
bc1512 |
gint offset;
|
|
Packit |
bc1512 |
gfloat *src_buf;
|
|
Packit |
bc1512 |
gfloat *dst_buf;
|
|
Packit |
bc1512 |
gint radius = dradius;
|
|
Packit |
bc1512 |
gint src_width = gegl_buffer_get_width (src);
|
|
Packit |
bc1512 |
gint src_height = gegl_buffer_get_height (src);
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
src_buf = g_new0 (gfloat, gegl_buffer_get_pixel_count (src) * 4);
|
|
Packit |
bc1512 |
dst_buf = g_new0 (gfloat, dst_rect->width * dst_rect->height * 4);
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
gegl_buffer_get (src, NULL, 1.0, babl_format ("RGBA float"), src_buf, GEGL_AUTO_ROWSTRIDE, GEGL_ABYSS_NONE);
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
offset = 0;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
for (y=0; y<dst_rect->height; y++)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
gfloat *center_pix;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
center_pix = src_buf + ((radius) + (y+radius)* src_width)*4;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
for (x=0; x<dst_rect->width; x++)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
gint u,v;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
gfloat accumulated[4]={0,};
|
|
Packit |
bc1512 |
gint count=0;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
/* iterate through the upper left quater of pixels */
|
|
Packit |
bc1512 |
for (v=-radius;v<=0;v++)
|
|
Packit |
bc1512 |
for (u=-radius;u<= (pairs==1?radius:0);u++)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
gfloat *selected_pix = center_pix;
|
|
Packit |
bc1512 |
gfloat best_diff = 1000.0;
|
|
Packit |
bc1512 |
gint i;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
/* skip computations for the center pixel */
|
|
Packit |
bc1512 |
if (u != 0 &&
|
|
Packit |
bc1512 |
v != 0)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
/* compute the coordinates of the symmetric pairs for
|
|
Packit |
bc1512 |
* this locaiton in the quadrant
|
|
Packit |
bc1512 |
*/
|
|
Packit |
bc1512 |
gint xs[4], ys[4];
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
xs[0] = x+u+radius;
|
|
Packit |
bc1512 |
xs[1] = x-u+radius;
|
|
Packit |
bc1512 |
xs[2] = x-u+radius;
|
|
Packit |
bc1512 |
xs[3] = x+u+radius;
|
|
Packit |
bc1512 |
ys[0] = y+v+radius;
|
|
Packit |
bc1512 |
ys[1] = y-v+radius;
|
|
Packit |
bc1512 |
ys[2] = y+v+radius;
|
|
Packit |
bc1512 |
ys[3] = y-v+radius;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
/* check which member of the symmetric quadruple to use */
|
|
Packit |
bc1512 |
for (i=0;i
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
if (xs[i] >= 0 && xs[i] < src_width &&
|
|
Packit |
bc1512 |
ys[i] >= 0 && ys[i] < src_height)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
gfloat *tpix = src_buf + (xs[i]+ys[i]* src_width)*4;
|
|
Packit |
bc1512 |
gfloat diff = colordiff (tpix, center_pix);
|
|
Packit |
bc1512 |
if (diff < best_diff)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
best_diff = diff;
|
|
Packit |
bc1512 |
selected_pix = tpix;
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
/* accumulate the components of the best sample from
|
|
Packit |
bc1512 |
* the symmetric quadruple
|
|
Packit |
bc1512 |
*/
|
|
Packit |
bc1512 |
for (i=0;i<4;i++)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
accumulated[i] += selected_pix[i];
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
count++;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
if (u==0 && v==0)
|
|
Packit |
bc1512 |
break; /* to avoid doubly processing when using only 1 pair */
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
for (u=0; u<4; u++)
|
|
Packit |
bc1512 |
dst_buf[offset*4+u] = accumulated[u]/count;
|
|
Packit |
bc1512 |
offset++;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
center_pix += 4;
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
gegl_buffer_set (dst, dst_rect, 0, babl_format ("RGBA float"), dst_buf,
|
|
Packit |
bc1512 |
GEGL_AUTO_ROWSTRIDE);
|
|
Packit |
bc1512 |
g_free (src_buf);
|
|
Packit |
bc1512 |
g_free (dst_buf);
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
#include "opencl/gegl-cl.h"
|
|
Packit |
bc1512 |
#include "buffer/gegl-buffer-cl-iterator.h"
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
static const char* kernel_source =
|
|
Packit |
bc1512 |
"float colordiff (float4 pixA, \n"
|
|
Packit |
bc1512 |
" float4 pixB) \n"
|
|
Packit |
bc1512 |
"{ \n"
|
|
Packit |
bc1512 |
" float4 pix = pixA-pixB; \n"
|
|
Packit |
bc1512 |
" pix *= pix; \n"
|
|
Packit |
bc1512 |
" return pix.x+pix.y+pix.z; \n"
|
|
Packit |
bc1512 |
"} \n"
|
|
Packit |
bc1512 |
" \n"
|
|
Packit |
bc1512 |
"__kernel void snn_mean_CL (__global const float4 *src_buf, \n"
|
|
Packit |
bc1512 |
" int src_width, \n"
|
|
Packit |
bc1512 |
" int src_height, \n"
|
|
Packit |
bc1512 |
" __global float4 *dst_buf, \n"
|
|
Packit |
bc1512 |
" int radius, \n"
|
|
Packit |
bc1512 |
" int pairs) \n"
|
|
Packit |
bc1512 |
"{ \n"
|
|
Packit |
bc1512 |
" int gidx =get_global_id(0); \n"
|
|
Packit |
bc1512 |
" int gidy =get_global_id(1); \n"
|
|
Packit |
bc1512 |
" int offset =gidy * get_global_size(0) + gidx; \n"
|
|
Packit |
bc1512 |
" \n"
|
|
Packit |
bc1512 |
" __global const float4 *center_pix= \n"
|
|
Packit |
bc1512 |
" src_buf + ((radius+gidx) + (gidy+radius)* src_width); \n"
|
|
Packit |
bc1512 |
" float4 accumulated=0; \n"
|
|
Packit |
bc1512 |
" \n"
|
|
Packit |
bc1512 |
" int count=0; \n"
|
|
Packit |
bc1512 |
" if(pairs==2) \n"
|
|
Packit |
bc1512 |
" { \n"
|
|
Packit |
bc1512 |
" for(int i=-radius;i<0;i++) \n"
|
|
Packit |
bc1512 |
" { \n"
|
|
Packit |
bc1512 |
" for(int j=-radius;j<0;j++) \n"
|
|
Packit |
bc1512 |
" { \n"
|
|
Packit |
bc1512 |
" __global const float4 *selected_pix = center_pix; \n"
|
|
Packit |
bc1512 |
" float best_diff = 1000.0f; \n"
|
|
Packit |
bc1512 |
" \n"
|
|
Packit |
bc1512 |
" int xs[4]={ \n"
|
|
Packit |
bc1512 |
" gidx+j+radius, gidx-j+radius, \n"
|
|
Packit |
bc1512 |
" gidx-j+radius, gidx+j+radius \n"
|
|
Packit |
bc1512 |
" }; \n"
|
|
Packit |
bc1512 |
" int ys[4]={ \n"
|
|
Packit |
bc1512 |
" gidy+i+radius, gidy-i+radius, \n"
|
|
Packit |
bc1512 |
" gidy+i+radius, gidy-i+radius}; \n"
|
|
Packit |
bc1512 |
" \n"
|
|
Packit |
bc1512 |
" for (int k=0;k<4;k++) \n"
|
|
Packit |
bc1512 |
" { \n"
|
|
Packit |
bc1512 |
" if (xs[k] >= 0 && xs[k] < src_width && \n"
|
|
Packit |
bc1512 |
" ys[k] >= 0 && ys[k] < src_height) \n"
|
|
Packit |
bc1512 |
" { \n"
|
|
Packit |
bc1512 |
" __global const float4 *tpix = \n"
|
|
Packit |
bc1512 |
" src_buf + (xs[k] + ys[k] * src_width);\n"
|
|
Packit |
bc1512 |
" float diff=colordiff(*tpix, *center_pix); \n"
|
|
Packit |
bc1512 |
" if (diff < best_diff) \n"
|
|
Packit |
bc1512 |
" { \n"
|
|
Packit |
bc1512 |
" best_diff = diff; \n"
|
|
Packit |
bc1512 |
" selected_pix = tpix; \n"
|
|
Packit |
bc1512 |
" } \n"
|
|
Packit |
bc1512 |
" } \n"
|
|
Packit |
bc1512 |
" } \n"
|
|
Packit |
bc1512 |
" \n"
|
|
Packit |
bc1512 |
" accumulated += *selected_pix; \n"
|
|
Packit |
bc1512 |
" \n"
|
|
Packit |
bc1512 |
" ++count; \n"
|
|
Packit |
bc1512 |
" if (i==0 && j==0) \n"
|
|
Packit |
bc1512 |
" break; \n"
|
|
Packit |
bc1512 |
" } \n"
|
|
Packit |
bc1512 |
" } \n"
|
|
Packit |
bc1512 |
" dst_buf[offset] = accumulated/count; \n"
|
|
Packit |
bc1512 |
" return; \n"
|
|
Packit |
bc1512 |
" } \n"
|
|
Packit |
bc1512 |
" else if(pairs==1) \n"
|
|
Packit |
bc1512 |
" { \n"
|
|
Packit |
bc1512 |
" for(int i=-radius;i<=0;i++) \n"
|
|
Packit |
bc1512 |
" { \n"
|
|
Packit |
bc1512 |
" for(int j=-radius;j<=radius;j++) \n"
|
|
Packit |
bc1512 |
" { \n"
|
|
Packit |
bc1512 |
" __global const float4 *selected_pix = center_pix; \n"
|
|
Packit |
bc1512 |
" float best_diff = 1000.0f; \n"
|
|
Packit |
bc1512 |
" \n"
|
|
Packit |
bc1512 |
" /* skip computations for the center pixel */ \n"
|
|
Packit |
bc1512 |
" if (i != 0 && j != 0) \n"
|
|
Packit |
bc1512 |
" { \n"
|
|
Packit |
bc1512 |
" int xs[4]={ \n"
|
|
Packit |
bc1512 |
" gidx+i+radius, gidx-i+radius, \n"
|
|
Packit |
bc1512 |
" gidx-i+radius, gidx+i+radius \n"
|
|
Packit |
bc1512 |
" }; \n"
|
|
Packit |
bc1512 |
" int ys[4]={ \n"
|
|
Packit |
bc1512 |
" gidy+j+radius, gidy-j+radius, \n"
|
|
Packit |
bc1512 |
" gidy+j+radius, gidy-j+radius \n"
|
|
Packit |
bc1512 |
" }; \n"
|
|
Packit |
bc1512 |
" \n"
|
|
Packit |
bc1512 |
" for (i=0;i<2;i++) \n"
|
|
Packit |
bc1512 |
" { \n"
|
|
Packit |
bc1512 |
" if (xs[i] >= 0 && xs[i] < src_width && \n"
|
|
Packit |
bc1512 |
" ys[i] >= 0 && ys[i] < src_height) \n"
|
|
Packit |
bc1512 |
" { \n"
|
|
Packit |
bc1512 |
" __global const float4 *tpix = \n"
|
|
Packit |
bc1512 |
" src_buf + (xs[i] + ys[i] * src_width);\n"
|
|
Packit |
bc1512 |
" float diff=colordiff (*tpix, *center_pix);\n"
|
|
Packit |
bc1512 |
" if (diff < best_diff) \n"
|
|
Packit |
bc1512 |
" { \n"
|
|
Packit |
bc1512 |
" best_diff = diff; \n"
|
|
Packit |
bc1512 |
" selected_pix = tpix; \n"
|
|
Packit |
bc1512 |
" } \n"
|
|
Packit |
bc1512 |
" } \n"
|
|
Packit |
bc1512 |
" } \n"
|
|
Packit |
bc1512 |
" } \n"
|
|
Packit |
bc1512 |
" accumulated += *selected_pix; \n"
|
|
Packit |
bc1512 |
" ++count; \n"
|
|
Packit |
bc1512 |
" if (i==0 && j==0) \n"
|
|
Packit |
bc1512 |
" break; \n"
|
|
Packit |
bc1512 |
" } \n"
|
|
Packit |
bc1512 |
" } \n"
|
|
Packit |
bc1512 |
" dst_buf[offset] = accumulated/count; \n"
|
|
Packit |
bc1512 |
" return; \n"
|
|
Packit |
bc1512 |
" } \n"
|
|
Packit |
bc1512 |
" return; \n"
|
|
Packit |
bc1512 |
"} \n";
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
static gegl_cl_run_data *cl_data = NULL;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
static cl_int
|
|
Packit |
bc1512 |
cl_snn_mean (cl_mem in_tex,
|
|
Packit |
bc1512 |
cl_mem out_tex,
|
|
Packit |
bc1512 |
const GeglRectangle *src_rect,
|
|
Packit |
bc1512 |
const GeglRectangle *roi,
|
|
Packit |
bc1512 |
gint radius,
|
|
Packit |
bc1512 |
gint pairs)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
cl_int cl_err = 0;
|
|
Packit |
bc1512 |
size_t global_ws[2];
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
if (!cl_data)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
const char *kernel_name[] = {"snn_mean_CL", NULL};
|
|
Packit |
bc1512 |
cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
if (!cl_data) return 1;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
global_ws[0] = roi->width;
|
|
Packit |
bc1512 |
global_ws[1] = roi->height;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex);
|
|
Packit |
bc1512 |
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_int), (void*)&src_rect->width);
|
|
Packit |
bc1512 |
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int), (void*)&src_rect->height);
|
|
Packit |
bc1512 |
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_mem), (void*)&out_tex);
|
|
Packit |
bc1512 |
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_int), (void*)&radius);
|
|
Packit |
bc1512 |
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_int), (void*)&pairs);
|
|
Packit |
bc1512 |
if (cl_err != CL_SUCCESS) return cl_err;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
|
|
Packit |
bc1512 |
cl_data->kernel[0], 2,
|
|
Packit |
bc1512 |
NULL, global_ws, NULL,
|
|
Packit |
bc1512 |
0, NULL, NULL);
|
|
Packit |
bc1512 |
if (cl_err != CL_SUCCESS) return cl_err;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
return cl_err;
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
static gboolean
|
|
Packit |
bc1512 |
cl_process (GeglOperation *operation,
|
|
Packit |
bc1512 |
GeglBuffer *input,
|
|
Packit |
bc1512 |
GeglBuffer *output,
|
|
Packit |
bc1512 |
const GeglRectangle *result)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
const Babl *in_format = gegl_operation_get_format (operation, "input");
|
|
Packit |
bc1512 |
const Babl *out_format = gegl_operation_get_format (operation, "output");
|
|
Packit |
bc1512 |
gint err;
|
|
Packit |
bc1512 |
gint j;
|
|
Packit |
bc1512 |
cl_int cl_err;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation);
|
|
Packit |
bc1512 |
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE, GEGL_ABYSS_NONE);
|
|
Packit |
bc1512 |
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);
|
|
Packit |
bc1512 |
while (gegl_buffer_cl_iterator_next (i, &err))
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
if (err) return FALSE;
|
|
Packit |
bc1512 |
for (j=0; j < i->n; j++)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
cl_err = cl_snn_mean(i->tex[read][j], i->tex[0][j], &i->roi[read][j], &i->roi[0][j], ceil(o->radius), o->pairs);
|
|
Packit |
bc1512 |
if (cl_err != CL_SUCCESS)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
g_warning("[OpenCL] Error in gegl:snn-mean: %s", gegl_cl_errstring(cl_err));
|
|
Packit |
bc1512 |
return FALSE;
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
return TRUE;
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
static void
|
|
Packit |
bc1512 |
gegl_chant_class_init (GeglChantClass *klass)
|
|
Packit |
bc1512 |
{
|
|
Packit |
bc1512 |
GeglOperationClass *operation_class;
|
|
Packit |
bc1512 |
GeglOperationFilterClass *filter_class;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
operation_class = GEGL_OPERATION_CLASS (klass);
|
|
Packit |
bc1512 |
filter_class = GEGL_OPERATION_FILTER_CLASS (klass);
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
filter_class->process = process;
|
|
Packit |
bc1512 |
operation_class->prepare = prepare;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
operation_class->opencl_support = TRUE;
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
gegl_operation_class_set_keys (operation_class,
|
|
Packit |
bc1512 |
"name" , "gegl:snn-mean",
|
|
Packit |
bc1512 |
"categories" , "misc",
|
|
Packit |
bc1512 |
"description",
|
|
Packit |
bc1512 |
_("Noise reducing edge enhancing blur filter based "
|
|
Packit |
bc1512 |
" on Symmetric Nearest Neighbours"),
|
|
Packit |
bc1512 |
NULL);
|
|
Packit |
bc1512 |
}
|
|
Packit |
bc1512 |
|
|
Packit |
bc1512 |
#endif
|