/* 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 Dominik Ernst <dernst@gmx.de>
 *
 * Recursive Gauss IIR Filter as described by Young / van Vliet
 * in "Signal Processing 44 (1995) 139 - 151"
 *
 * NOTE: The IIR filter should not be used for radius < 0.5, since it
 *       becomes very inaccurate.
 */

#include "config.h"
#include <glib/gi18n-lib.h>

#ifdef GEGL_CHANT_PROPERTIES

gegl_chant_register_enum (gegl_gaussian_blur_filter)
  enum_value (GEGL_GAUSSIAN_BLUR_FILTER_AUTO, "Auto")
  enum_value (GEGL_GAUSSIAN_BLUR_FILTER_FIR,  "FIR")
  enum_value (GEGL_GAUSSIAN_BLUR_FILTER_IIR,  "IIR")
gegl_chant_register_enum_end (GeglGaussianBlurFilter)

gegl_chant_double_ui (std_dev_x, _("Size X"),
                      0.0, 10000.0, 4.0, 0.0, 1000.0, 1.5,
                      _("Standard deviation for the horizontal axis "
                        "(multiply by ~2 to get radius)"))

gegl_chant_double_ui (std_dev_y, _("Size Y"),
                      0.0, 10000.0, 4.0, 0.0, 1000.0, 1.5,
                      _("Standard deviation for the vertical axi. "
                        "(multiply by ~2 to get radius)"))

gegl_chant_enum      (filter, _("Filter"),
                      GeglGaussianBlurFilter, gegl_gaussian_blur_filter,
                      GEGL_GAUSSIAN_BLUR_FILTER_AUTO,
                      _("Optional parameter to override the automatic "
                        "selection of blur filter"))

#else

#define GEGL_CHANT_TYPE_AREA_FILTER
#define GEGL_CHANT_C_FILE "gaussian-blur.c"

#include "gegl-chant.h"
#include <math.h>
#include <stdio.h>

#define RADIUS_SCALE   4

static void
iir_young_find_constants (gfloat   radius,
                          gdouble *B,
                          gdouble *b);

static gint
fir_gen_convolve_matrix (gdouble   sigma,
                         gdouble **cmatrix_p);


static void
iir_young_find_constants (gfloat   sigma,
                          gdouble *B,
                          gdouble *b)
{
  gdouble q;

  if (sigma == 0.0)
    {
      /* to prevent unexpected ringing at tile boundaries,
         we define an (expensive) copy operation here */
      *B = 1.0;
      b[0] = 1.0;
      b[1] = b[2] = b[3] = 0.0;
      return;
    }

  if(sigma >= 2.5)
    q = 0.98711*sigma - 0.96330;
  else
    q = 3.97156 - 4.14554*sqrt(1-0.26891*sigma);

  b[0] = 1.57825 + (2.44413*q) + (1.4281*q*q) + (0.422205*q*q*q);
  b[1] = (2.44413*q) + (2.85619*q*q) + (1.26661*q*q*q);
  b[2] = -((1.4281*q*q) + (1.26661*q*q*q));
  b[3] = 0.422205*q*q*q;

  *B = 1 - ( (b[1]+b[2]+b[3])/b[0] );
}

static inline void
iir_young_blur_1D (gfloat  * buf,
                   gint      offset,
                   gint      delta_offset,
                   gdouble   B,
                   gdouble * b,
                   gfloat  * w,
                   gint      w_len)
{
  gint wcount, i;
  gdouble tmp;
  gdouble recip = 1.0 / b[0];

  /* forward filter */
  wcount = 0;

  while (wcount < w_len)
    {
      tmp = 0;

      for (i=1; i<4; i++)
        {
          if (wcount-i >= 0)
            tmp += b[i]*w[wcount-i];
        }

      tmp *= recip;
      tmp += B*buf[offset];
      w[wcount] = tmp;

      wcount++;
      offset += delta_offset;
    }

  /* backward filter */
  wcount = w_len - 1;
  offset -= delta_offset;

  while (wcount >= 0)
    {
      tmp = 0;

      for (i=1; i<4; i++)
        {
          if (wcount+i < w_len)
            tmp += b[i]*buf[offset+delta_offset*i];
        }

      tmp *= recip;
      tmp += B*w[wcount];
      buf[offset] = tmp;

      offset -= delta_offset;
      wcount--;
    }
}

void print_buf_checksums(const char *descr, gfloat *buf, int channels, const GeglRectangle *rect)
{
  float sum[4] = {0,0,0,0};
  int h,v,c;
  for (v=0; v<rect->height; v++)
    for (h=0; h<rect->width; h++)
      for (c=0; c<channels; c++)
        sum[c] += buf[v*rect->width*channels +
                      h*channels+c];
  printf("%s: sums = %f %f %f %f\n", descr, sum[0],sum[1],sum[2],sum[3]);
}


/* expects src and dst buf to have the same height and no y-offset */
static void
iir_young_hor_blur (GeglBuffer          *src,
                    const GeglRectangle *src_rect,
                    GeglBuffer          *dst,
                    const GeglRectangle *dst_rect,
                    gdouble              B,
                    gdouble             *b,
                    const Babl          *format)
{
  gint v;
  gint c;
  gint w_len;
  gfloat *buf;
  gfloat *w;
  gint channels = babl_format_get_n_components(format);
  printf("iir_hor: channels format= %d %s\n", channels,babl_get_name(format));

  buf = g_new0 (gfloat, src_rect->height * src_rect->width * channels);
  w   = g_new0 (gfloat, src_rect->width);

  gegl_buffer_get (src, src_rect, 1.0, format,
                   buf, GEGL_AUTO_ROWSTRIDE, GEGL_ABYSS_NONE);

  printf("Before: first pixels = %f %f %f\n", buf[0], buf[channels], buf[channels*2]);

  w_len = src_rect->width;
  for (v=0; v<src_rect->height; v++)
    {
      for (c = 0; c < channels; c++)
        {
          iir_young_blur_1D (buf,
                             v*src_rect->width*channels+c,
                             channels,
                             B,
                             b,
                             w,
                             w_len);
        }
    }

  print_buf_checksums("hor", buf, channels, src_rect);
  gegl_buffer_set (dst, src_rect, 0.0, format,
                   buf, GEGL_AUTO_ROWSTRIDE);
  g_free (buf);
  g_free (w);
}

/* expects src and dst buf to have the same width and no x-offset */
static void
iir_young_ver_blur (GeglBuffer          *src,
                    const GeglRectangle *src_rect,
                    GeglBuffer          *dst,
                    const GeglRectangle *dst_rect,
                    gdouble              B,
                    gdouble             *b,
                    const Babl          *format)
{
  gint u;
  gint c;
  gint w_len;
  gfloat *buf;
  gfloat *w;
  gint channels = babl_format_get_n_components(format);
  printf("iir_ver: channels format= %d %s\n", channels, babl_get_name(format));

  buf = g_new0 (gfloat, src_rect->height * src_rect->width * channels);
  w   = g_new0 (gfloat, src_rect->height);

  gegl_buffer_get (src, src_rect, 1.0, format,
                   buf, GEGL_AUTO_ROWSTRIDE, GEGL_ABYSS_NONE);

  w_len = src_rect->height;
  for (u=0; u<dst_rect->width; u++)
    {
      for (c = 0; c < channels; c++)
        {
          iir_young_blur_1D (buf,
                             u*channels + c,
                             src_rect->width*channels,
                             B,
                             b,
                             w,
                             w_len);
        }
    }

  print_buf_checksums("ver", buf, channels, src_rect);
  gegl_buffer_set (dst, src_rect, 0,
                   format, buf, GEGL_AUTO_ROWSTRIDE);
  g_free (buf);
  g_free (w);
}


static gint
fir_calc_convolve_matrix_length (gdouble sigma)
{
  return sigma ? ceil (sigma)*6.0+1.0 : 1;
}

static gint
fir_gen_convolve_matrix (gdouble   sigma,
                         gdouble **cmatrix_p)
{
  gint     matrix_length;
  gdouble *cmatrix;

  matrix_length = fir_calc_convolve_matrix_length (sigma);
  cmatrix = g_new (gdouble, matrix_length);
  if (!cmatrix)
    return 0;

  if (matrix_length == 1)
    {
      cmatrix[0] = 1;
    }
  else
    {
      gint i,x;
      gdouble sum = 0;

      for (i=0; i<matrix_length/2+1; i++)
        {
          gdouble y;

          x = i - (matrix_length/2);
          y = (1.0/(sigma*sqrt(2.0*G_PI))) * exp(-(x*x) / (2.0*sigma*sigma));

          cmatrix[i] = y;
          sum += cmatrix[i];
        }

      for (i=matrix_length/2 + 1; i<matrix_length; i++)
        {
          cmatrix[i] = cmatrix[matrix_length-i-1];
          sum += cmatrix[i];
        }

      for (i=0; i<matrix_length; i++)
      {
        cmatrix[i] /= sum;
      }
    }

  *cmatrix_p = cmatrix;
  return matrix_length;
}

static inline float
fir_get_mean_component_1D (gfloat  * buf,
                           gint      offset,
                           gint      delta_offset,
                           gdouble * cmatrix,
                           gint      matrix_length)
{
  gint    i;
  gdouble acc=0;

  for (i=0; i < matrix_length; i++)
    {
      acc += buf[offset] * cmatrix[i];
      offset += delta_offset;
    }

  return acc;
}

/* expects src and dst buf to have the same height and no y-offset */
static void
fir_hor_blur (GeglBuffer          *src,
              const GeglRectangle *src_rect,
              GeglBuffer          *dst,
              const GeglRectangle *dst_rect,
              gdouble             *cmatrix,
              gint                 matrix_length,
              gint                 xoff,     /* offset between src and dst */
              const Babl          *format) 
{
  gint        u,v;
  gint        offset;
  gfloat     *src_buf;
  gfloat     *dst_buf;
  const gint  radius = matrix_length/2;
  const gint  src_width = src_rect->width;
  gint        channels = babl_format_get_n_components(format);
  printf("fir_hor: channels = %d\n", channels);

  g_assert (xoff >= radius);

  src_buf = g_new0 (gfloat, src_rect->height * src_rect->width * channels);
  dst_buf = g_new0 (gfloat, dst_rect->height * dst_rect->width * channels);

  gegl_buffer_get (src, src_rect, 1.0, format,
                   src_buf, GEGL_AUTO_ROWSTRIDE, GEGL_ABYSS_NONE);

  offset = 0;
  for (v=0; v<dst_rect->height; v++)
    for (u=0; u<dst_rect->width; u++)
      {
        gint src_offset = (u-radius+xoff + v*src_width) * channels;
        gint c;
        for (c=0; c<channels; c++)
          dst_buf [offset++] = fir_get_mean_component_1D (src_buf,
                                                          src_offset + c,
                                                          channels,
                                                          cmatrix,
                                                          matrix_length);
      }

  gegl_buffer_set (dst, dst_rect, 0.0, format,
                   dst_buf, GEGL_AUTO_ROWSTRIDE);
  g_free (src_buf);
  g_free (dst_buf);
}

/* expects src and dst buf to have the same width and no x-offset */
static void
fir_ver_blur (GeglBuffer          *src,
              const GeglRectangle *src_rect,
              GeglBuffer          *dst,
              const GeglRectangle *dst_rect,
              gdouble             *cmatrix,
              gint                 matrix_length,
              gint                 yoff, /* offset between src and dst */
              const Babl          *format) 
{
  gint        u,v;
  gint        offset;
  gfloat     *src_buf;
  gfloat     *dst_buf;
  const gint  radius = matrix_length/2;
  const gint  src_width = src_rect->width;
  gint        channels = babl_format_get_n_components(format);
  printf("fir_ver: channels = %d\n", channels);

  g_assert (yoff >= radius);

  src_buf = g_new0 (gfloat, src_rect->width * src_rect->height * channels);
  dst_buf = g_new0 (gfloat, dst_rect->width * dst_rect->height * channels);

  gegl_buffer_get (src, src_rect, 1.0, format,
                   src_buf, GEGL_AUTO_ROWSTRIDE, GEGL_ABYSS_NONE);

  offset=0;
  for (v=0; v< dst_rect->height; v++)
    for (u=0; u< dst_rect->width; u++)
      {
        gint src_offset = (u + (v-radius+yoff)*src_width) * channels;
        gint c;
        for (c=0; c<channels; c++)
          dst_buf [offset++] = fir_get_mean_component_1D (src_buf,
                                                          src_offset + c,
                                                          src_width * channels,
                                                          cmatrix,
                                                          matrix_length);
      }

  gegl_buffer_set (dst, dst_rect, 0, format,
                   dst_buf, GEGL_AUTO_ROWSTRIDE);
  g_free (src_buf);
  g_free (dst_buf);
}

static void
prepare (GeglOperation *operation)
{

  GeglOperationAreaFilter *area = GEGL_OPERATION_AREA_FILTER (operation);
  GeglChantO              *o    = GEGL_CHANT_PROPERTIES (operation);
  const Babl *source_format = gegl_operation_get_source_format (operation, "input");

  gfloat fir_radius_x = fir_calc_convolve_matrix_length (o->std_dev_x) / 2;
  gfloat fir_radius_y = fir_calc_convolve_matrix_length (o->std_dev_y) / 2;
  gfloat iir_radius_x = o->std_dev_x * RADIUS_SCALE;
  gfloat iir_radius_y = o->std_dev_y * RADIUS_SCALE;
  gint n_components;
  const Babl *format;

  /* XXX: these should be calculated exactly considering o->filter, but we just
   * make sure there is enough space */
  area->left = area->right = ceil (MAX (fir_radius_x, iir_radius_x));
  area->top = area->bottom = ceil (MAX (fir_radius_y, iir_radius_y));

  if (source_format != NULL)
    n_components = babl_format_get_n_components(source_format);
  else
    n_components = 4; // Default

  switch (n_components)
    {
    case 1:
      format = babl_format("Y float");
      break;
    case 2:
      format = babl_format("YaA float");
      break;
    case 3:
      format = babl_format("RGB float");
      break;
    default:
      format = babl_format("RaGaBaA float");
      break;
    }

  gegl_operation_set_format (operation, "input", format);
  gegl_operation_set_format (operation, "output", format);
}

#include "opencl/gegl-cl.h"
#include "buffer/gegl-buffer-cl-iterator.h"

#include "opencl/gaussian-blur.cl.h"

static GeglClRunData *cl_data = NULL;

static gboolean
cl_gaussian_blur (cl_mem                in_tex,
                  cl_mem                out_tex,
                  cl_mem                aux_tex,
                  size_t                global_worksize,
                  const GeglRectangle  *roi,
                  const GeglRectangle  *src_rect,
                  const GeglRectangle  *aux_rect,
                  gfloat               *dmatrix_x,
                  gint                  matrix_length_x,
                  gint                  xoff,
                  gfloat               *dmatrix_y,
                  gint                  matrix_length_y,
                  gint                  yoff)
{
  cl_int cl_err = 0;

  size_t global_ws[2];

  cl_mem cl_matrix_x, cl_matrix_y;

  if (!cl_data)
    {
      const char *kernel_name[] = {"fir_ver_blur", "fir_hor_blur", NULL};
      cl_data = gegl_cl_compile_and_build (gaussian_blur_cl_source, kernel_name);
    }
  if (!cl_data) return TRUE;

  cl_matrix_x = gegl_clCreateBuffer(gegl_cl_get_context(),
                                    CL_MEM_USE_HOST_PTR | CL_MEM_READ_ONLY,
                                    matrix_length_x * sizeof(cl_float), dmatrix_x, &cl_err);
  CL_CHECK;

  cl_matrix_y = gegl_clCreateBuffer(gegl_cl_get_context(),
                                    CL_MEM_USE_HOST_PTR | CL_MEM_READ_ONLY,
                                    matrix_length_y * sizeof(cl_float), dmatrix_y, &cl_err);
  CL_CHECK;

  {
  global_ws[0] = aux_rect->width;
  global_ws[1] = aux_rect->height;

  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem), (void*)&in_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_int), (void*)&src_rect->width);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 2, sizeof(cl_mem), (void*)&aux_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 3, sizeof(cl_mem), (void*)&cl_matrix_x);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 4, sizeof(cl_int), (void*)&matrix_length_x);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 5, sizeof(cl_int), (void*)&xoff);
  CL_CHECK;

  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
                                       cl_data->kernel[1], 2,
                                       NULL, global_ws, NULL,
                                       0, NULL, NULL);
  CL_CHECK;
  }

  {
  global_ws[0] = roi->width;
  global_ws[1] = roi->height;

  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&aux_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_int), (void*)&aux_rect->width);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_mem), (void*)&out_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_mem), (void*)&cl_matrix_y);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_int), (void*)&matrix_length_y);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_int), (void*)&yoff);
  CL_CHECK;

  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
                                       cl_data->kernel[0], 2,
                                       NULL, global_ws, NULL,
                                       0, NULL, NULL);
  CL_CHECK;
  }

  cl_err = gegl_clFinish(gegl_cl_get_command_queue ());
  CL_CHECK;

  cl_err = gegl_clReleaseMemObject(cl_matrix_x);
  CL_CHECK;
  cl_err = gegl_clReleaseMemObject(cl_matrix_y);
  CL_CHECK;

  return FALSE;

error:
  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);

  gdouble *cmatrix_x, *cmatrix_y;
  gint cmatrix_len_x, cmatrix_len_y;

  gfloat *fmatrix_x, *fmatrix_y;

  cmatrix_len_x = fir_gen_convolve_matrix (o->std_dev_x, &cmatrix_x);
  cmatrix_len_y = fir_gen_convolve_matrix (o->std_dev_y, &cmatrix_y);

  fmatrix_x = g_new (gfloat, cmatrix_len_x);
  fmatrix_y = g_new (gfloat, cmatrix_len_y);

  for(j=0; j<cmatrix_len_x; j++)
    fmatrix_x[j] = (gfloat) cmatrix_x[j];

  for(j=0; j<cmatrix_len_y; j++)
    fmatrix_y[j] = (gfloat) cmatrix_y[j];

  {
  GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE);
  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,
                                             0, 0, 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++)
        {
           err = cl_gaussian_blur(i->tex[read][j],
                                  i->tex[0][j],
                                  i->tex[aux][j],
                                  i->size[0][j],
                                  &i->roi[0][j],
                                  &i->roi[read][j],
                                  &i->roi[aux][j],
                                  fmatrix_x,
                                  cmatrix_len_x,
                                  op_area->left,
                                  fmatrix_y,
                                  cmatrix_len_y,
                                  op_area->top);
          if (err)
            {
              g_warning("[OpenCL] Error in gegl:gaussian-blur");
              return FALSE;
            }
        }
    }
  }

  g_free (fmatrix_x);
  g_free (fmatrix_y);

  g_free (cmatrix_x);
  g_free (cmatrix_y);
  return TRUE;
}

static gboolean
process (GeglOperation       *operation,
         GeglBuffer          *input,
         GeglBuffer          *output,
         const GeglRectangle *result,
         gint                 level)
{
  GeglRectangle rect;
  GeglBuffer *temp;
  GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation);
  GeglChantO              *o       = GEGL_CHANT_PROPERTIES (operation);

  GeglRectangle temp_extend;
  gdouble       B, b[4];
  gdouble      *cmatrix;
  gint          cmatrix_len;
  gboolean      force_iir;
  gboolean      force_fir;

  printf("process: result->x=%d\n", result->x);
  const Babl *format = gegl_operation_get_format(operation,"output");

  rect.x      = result->x - op_area->left;
  rect.width  = result->width + op_area->left + op_area->right;
  rect.y      = result->y - op_area->top;
  rect.height = result->height + op_area->top + op_area->bottom;

  force_iir = (o->filter == GEGL_GAUSSIAN_BLUR_FILTER_IIR);
  force_fir = (o->filter == GEGL_GAUSSIAN_BLUR_FILTER_FIR);

  if (gegl_cl_is_accelerated () && !force_iir)
    if (cl_process(operation, input, output, result))
      return TRUE;

  temp_extend = rect;
  temp_extend.x      = result->x;
  temp_extend.width  = result->width;
  temp = gegl_buffer_new (&temp_extend, format);
  printf("result->width = %d\n", result->width);

  if ((force_iir || o->std_dev_x > 1.0) && !force_fir)
    {
      iir_young_find_constants (o->std_dev_x, &B, b);
      iir_young_hor_blur (input, &rect, temp, &temp_extend,  B, b, format);
    }
  else
    {
      cmatrix_len = fir_gen_convolve_matrix (o->std_dev_x, &cmatrix);
      fir_hor_blur (input, &rect, temp, &temp_extend,
                    cmatrix, cmatrix_len, op_area->left, format);
      g_free (cmatrix);
    }


  if ((force_iir || o->std_dev_y > 1.0) && !force_fir)
    {
      iir_young_find_constants (o->std_dev_y, &B, b);
      iir_young_ver_blur (temp, &temp_extend, output, result, B, b, format);
    }
  else
    {
      cmatrix_len = fir_gen_convolve_matrix (o->std_dev_y, &cmatrix);
      fir_ver_blur (temp, &temp_extend, output, result, cmatrix, cmatrix_len,
                    op_area->top, format);
      g_free (cmatrix);
    }

  g_object_unref (temp);
  return  TRUE;
}


static void
gegl_chant_class_init (GeglChantClass *klass)
{
  GeglOperationClass       *operation_class;
  GeglOperationFilterClass *filter_class;

  operation_class = GEGL_OPERATION_CLASS (klass);
  filter_class    = GEGL_OPERATION_FILTER_CLASS (klass);

  operation_class->prepare        = prepare;
  operation_class->opencl_support = TRUE;

  filter_class->process           = process;

  gegl_operation_class_set_keys (operation_class,
    "name",        "gegl:gaussian-blur",
    "categories",  "blur",
    "description", _("Performs an averaging of neighboring pixels with the "
                     "normal distribution as weighting"),
    NULL);
}

#endif
