Blame operations/common/edge-sobel.c

Packit Service 2781ba
/* This file is an image processing operation for GEGL
Packit Service 2781ba
 *
Packit Service 2781ba
 * GEGL is free software; you can redistribute it and/or
Packit Service 2781ba
 * modify it under the terms of the GNU Lesser General Public
Packit Service 2781ba
 * License as published by the Free Software Foundation; either
Packit Service 2781ba
 * version 3 of the License, or (at your option) any later version.
Packit Service 2781ba
 *
Packit Service 2781ba
 * GEGL is distributed in the hope that it will be useful,
Packit Service 2781ba
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
Packit Service 2781ba
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
Packit Service 2781ba
 * Lesser General Public License for more details.
Packit Service 2781ba
 *
Packit Service 2781ba
 * You should have received a copy of the GNU Lesser General Public
Packit Service 2781ba
 * License along with GEGL; if not, see <http://www.gnu.org/licenses/>.
Packit Service 2781ba
 *
Packit Service 2781ba
 */
Packit Service 2781ba
Packit Service 2781ba
/*
Packit Service 2781ba
 * Copyright 2011 Victor Oliveira <victormatheus@gmail.com>
Packit Service 2781ba
 */
Packit Service 2781ba
Packit Service 2781ba
#include "config.h"
Packit Service 2781ba
#include <glib/gi18n-lib.h>
Packit Service 2781ba
Packit Service 2781ba
Packit Service 2781ba
#ifdef GEGL_CHANT_PROPERTIES
Packit Service 2781ba
Packit Service 2781ba
gegl_chant_boolean (horizontal,  _("Horizontal"),  TRUE,
Packit Service 2781ba
                    _("Horizontal"))
Packit Service 2781ba
Packit Service 2781ba
gegl_chant_boolean (vertical,  _("Vertical"),  TRUE,
Packit Service 2781ba
                    _("Vertical"))
Packit Service 2781ba
Packit Service 2781ba
gegl_chant_boolean (keep_signal,  _("Keep Signal"),  TRUE,
Packit Service 2781ba
                    _("Keep Signal"))
Packit Service 2781ba
Packit Service 2781ba
#else
Packit Service 2781ba
Packit Service 2781ba
#define GEGL_CHANT_TYPE_AREA_FILTER
Packit Service 2781ba
#define GEGL_CHANT_C_FILE       "edge-sobel.c"
Packit Service 2781ba
Packit Service 2781ba
#include "gegl-chant.h"
Packit Service 2781ba
#include <math.h>
Packit Service 2781ba
#include <stdio.h>
Packit Service 2781ba
Packit Service 2781ba
#define SOBEL_RADIUS 1
Packit Service 2781ba
Packit Service 2781ba
static void
Packit Service 2781ba
edge_sobel (GeglBuffer          *src,
Packit Service 2781ba
            const GeglRectangle *src_rect,
Packit Service 2781ba
            GeglBuffer          *dst,
Packit Service 2781ba
            const GeglRectangle *dst_rect,
Packit Service 2781ba
            gboolean            horizontal,
Packit Service 2781ba
            gboolean            vertical,
Packit Service 2781ba
            gboolean            keep_signal);
Packit Service 2781ba
Packit Service 2781ba
Packit Service 2781ba
Packit Service 2781ba
static void prepare (GeglOperation *operation)
Packit Service 2781ba
{
Packit Service 2781ba
  GeglOperationAreaFilter *area = GEGL_OPERATION_AREA_FILTER (operation);
Packit Service 2781ba
  //GeglChantO              *o = GEGL_CHANT_PROPERTIES (operation);
Packit Service 2781ba
Packit Service 2781ba
  area->left = area->right = area->top = area->bottom = SOBEL_RADIUS;
Packit Service 2781ba
  gegl_operation_set_format (operation, "input", babl_format ("RGBA float"));
Packit Service 2781ba
  gegl_operation_set_format (operation, "output", babl_format ("RGBA float"));
Packit Service 2781ba
}
Packit Service 2781ba
Packit Service 2781ba
#include "opencl/gegl-cl.h"
Packit Service 2781ba
#include "buffer/gegl-buffer-cl-iterator.h"
Packit Service 2781ba
Packit Service 2781ba
static const char* kernel_source =
Packit Service 2781ba
"#define SOBEL_RADIUS 1                                                \n"
Packit Service 2781ba
"kernel void kernel_edgesobel(global float4 *in,                       \n"
Packit Service 2781ba
"                             global float4 *out,                      \n"
Packit Service 2781ba
"                             const int horizontal,                    \n"
Packit Service 2781ba
"                             const int vertical,                      \n"
Packit Service 2781ba
"                             const int keep_signal)                   \n"
Packit Service 2781ba
"{                                                                     \n"
Packit Service 2781ba
"    int gidx = get_global_id(0);                                      \n"
Packit Service 2781ba
"    int gidy = get_global_id(1);                                      \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    float4 hor_grad = 0.0f;                                           \n"
Packit Service 2781ba
"    float4 ver_grad = 0.0f;                                           \n"
Packit Service 2781ba
"    float4 gradient = 0.0f;                                           \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    int dst_width = get_global_size(0);                               \n"
Packit Service 2781ba
"    int src_width = dst_width + SOBEL_RADIUS * 2;                     \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    int i = gidx + SOBEL_RADIUS, j = gidy + SOBEL_RADIUS;             \n"
Packit Service 2781ba
"    int gid1d = i + j * src_width;                                    \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    float4 pix_fl = in[gid1d - 1 - src_width];                        \n"
Packit Service 2781ba
"    float4 pix_fm = in[gid1d     - src_width];                        \n"
Packit Service 2781ba
"    float4 pix_fr = in[gid1d + 1 - src_width];                        \n"
Packit Service 2781ba
"    float4 pix_ml = in[gid1d - 1            ];                        \n"
Packit Service 2781ba
"    float4 pix_mm = in[gid1d                ];                        \n"
Packit Service 2781ba
"    float4 pix_mr = in[gid1d + 1            ];                        \n"
Packit Service 2781ba
"    float4 pix_bl = in[gid1d - 1 + src_width];                        \n"
Packit Service 2781ba
"    float4 pix_bm = in[gid1d     + src_width];                        \n"
Packit Service 2781ba
"    float4 pix_br = in[gid1d + 1 + src_width];                        \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    if (horizontal)                                                   \n"
Packit Service 2781ba
"    {                                                                 \n"
Packit Service 2781ba
"        hor_grad +=                                                   \n"
Packit Service 2781ba
"            - 1.0f * pix_fl + 1.0f * pix_fr                           \n"
Packit Service 2781ba
"            - 2.0f * pix_ml + 2.0f * pix_mr                           \n"
Packit Service 2781ba
"            - 1.0f * pix_bl + 1.0f * pix_br;                          \n"
Packit Service 2781ba
"    }                                                                 \n"
Packit Service 2781ba
"    if (vertical)                                                     \n"
Packit Service 2781ba
"    {                                                                 \n"
Packit Service 2781ba
"        ver_grad +=                                                   \n"
Packit Service 2781ba
"            - 1.0f * pix_fl - 2.0f * pix_fm                           \n"
Packit Service 2781ba
"            - 1.0f * pix_fr + 1.0f * pix_bl                           \n"
Packit Service 2781ba
"            + 2.0f * pix_bm + 1.0f * pix_br;                          \n"
Packit Service 2781ba
"    }                                                                 \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    if (horizontal && vertical)                                       \n"
Packit Service 2781ba
"    {                                                                 \n"
Packit Service 2781ba
"        gradient = sqrt(                                              \n"
Packit Service 2781ba
"            hor_grad * hor_grad +                                     \n"
Packit Service 2781ba
"            ver_grad * ver_grad) / 1.41f;                             \n"
Packit Service 2781ba
"    }                                                                 \n"
Packit Service 2781ba
"    else                                                              \n"
Packit Service 2781ba
"    {                                                                 \n"
Packit Service 2781ba
"        if (keep_signal)                                              \n"
Packit Service 2781ba
"            gradient = hor_grad + ver_grad;                           \n"
Packit Service 2781ba
"        else                                                          \n"
Packit Service 2781ba
"            gradient = fabs(hor_grad + ver_grad);                     \n"
Packit Service 2781ba
"    }                                                                 \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    gradient.w = pix_mm.w;                                            \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    out[gidx + gidy * dst_width] = gradient;                          \n"
Packit Service 2781ba
"}                                                                     \n";
Packit Service 2781ba
Packit Service 2781ba
static gegl_cl_run_data *cl_data = NULL;
Packit Service 2781ba
Packit Service 2781ba
static cl_int
Packit Service 2781ba
cl_edge_sobel (cl_mem              in_tex,
Packit Service 2781ba
               cl_mem              out_tex,
Packit Service 2781ba
               size_t              global_worksize,
Packit Service 2781ba
               const GeglRectangle *roi,
Packit Service 2781ba
               gboolean            horizontal,
Packit Service 2781ba
               gboolean            vertical,
Packit Service 2781ba
               gboolean            keep_signal)
Packit Service 2781ba
{
Packit Service 2781ba
  if (!cl_data)
Packit Service 2781ba
    {
Packit Service 2781ba
      const char *kernel_name[] = {"kernel_edgesobel", NULL};
Packit Service 2781ba
      cl_data = gegl_cl_compile_and_build(kernel_source, kernel_name);
Packit Service 2781ba
    }
Packit Service 2781ba
  if (!cl_data)  return 0;
Packit Service 2781ba
Packit Service 2781ba
  {
Packit Service 2781ba
  const size_t gbl_size[2] = {roi->width,roi->height};
Packit Service 2781ba
  cl_int n_horizontal  = horizontal;
Packit Service 2781ba
  cl_int n_vertical    = vertical;
Packit Service 2781ba
  cl_int n_keep_signal = keep_signal;
Packit Service 2781ba
  cl_int cl_err = 0;
Packit Service 2781ba
Packit Service 2781ba
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex);
Packit Service 2781ba
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex);
Packit Service 2781ba
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int), (void*)&n_horizontal);
Packit Service 2781ba
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_int), (void*)&n_vertical);
Packit Service 2781ba
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_int), (void*)&n_keep_signal);
Packit Service 2781ba
  if (cl_err != CL_SUCCESS) return cl_err;
Packit Service 2781ba
Packit Service 2781ba
  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(),
Packit Service 2781ba
                                       cl_data->kernel[0], 2,
Packit Service 2781ba
                                       NULL, gbl_size, NULL,
Packit Service 2781ba
                                       0, NULL, NULL);
Packit Service 2781ba
  if (cl_err != CL_SUCCESS) return cl_err;
Packit Service 2781ba
  }
Packit Service 2781ba
Packit Service 2781ba
  return CL_SUCCESS;
Packit Service 2781ba
}
Packit Service 2781ba
Packit Service 2781ba
static gboolean
Packit Service 2781ba
cl_process (GeglOperation       *operation,
Packit Service 2781ba
      GeglBuffer          *input,
Packit Service 2781ba
      GeglBuffer          *output,
Packit Service 2781ba
      const GeglRectangle *result)
Packit Service 2781ba
{
Packit Service 2781ba
  const Babl *in_format  = gegl_operation_get_format (operation, "input");
Packit Service 2781ba
  const Babl *out_format = gegl_operation_get_format (operation, "output");
Packit Service 2781ba
  gint err;
Packit Service 2781ba
  gint j;
Packit Service 2781ba
  cl_int cl_err;
Packit Service 2781ba
Packit Service 2781ba
  GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation);
Packit Service 2781ba
  GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
Packit Service 2781ba
Packit Service 2781ba
  GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output,result, out_format, GEGL_CL_BUFFER_WRITE, GEGL_ABYSS_NONE);
Packit Service 2781ba
                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 Service 2781ba
  while (gegl_buffer_cl_iterator_next (i, &err))
Packit Service 2781ba
  {
Packit Service 2781ba
    if (err) return FALSE;
Packit Service 2781ba
    for (j=0; j < i->n; j++)
Packit Service 2781ba
    {
Packit Service 2781ba
      cl_err = cl_edge_sobel(i->tex[read][j], i->tex[0][j], i->size[0][j],&i->roi[0][j], o->horizontal, o->vertical, o->keep_signal);
Packit Service 2781ba
      if (cl_err != CL_SUCCESS)
Packit Service 2781ba
      {
Packit Service 2781ba
        g_warning("[OpenCL] Error in gegl:edge-sobel: %s", gegl_cl_errstring(cl_err));
Packit Service 2781ba
        return FALSE;
Packit Service 2781ba
      }
Packit Service 2781ba
    }
Packit Service 2781ba
  }
Packit Service 2781ba
  return TRUE;
Packit Service 2781ba
}
Packit Service 2781ba
Packit Service 2781ba
static gboolean
Packit Service 2781ba
process (GeglOperation       *operation,
Packit Service 2781ba
         GeglBuffer          *input,
Packit Service 2781ba
         GeglBuffer          *output,
Packit Service 2781ba
         const GeglRectangle *result,
Packit Service 2781ba
         gint                 level)
Packit Service 2781ba
{
Packit Service 2781ba
  GeglChantO   *o = GEGL_CHANT_PROPERTIES (operation);
Packit Service 2781ba
  GeglRectangle compute;
Packit Service 2781ba
Packit Service 2781ba
  compute = gegl_operation_get_required_for_output (operation, "input",result);
Packit Service 2781ba
Packit Service 2781ba
  if (gegl_cl_is_accelerated ())
Packit Service 2781ba
    if(cl_process(operation, input, output, result))
Packit Service 2781ba
      return TRUE;
Packit Service 2781ba
Packit Service 2781ba
  edge_sobel (input, &compute, output, result, o->horizontal, o->vertical, o->keep_signal);
Packit Service 2781ba
  return  TRUE;
Packit Service 2781ba
}
Packit Service 2781ba
Packit Service 2781ba
inline static gfloat
Packit Service 2781ba
RMS(gfloat a, gfloat b)
Packit Service 2781ba
{
Packit Service 2781ba
  return sqrtf(a*a+b*b);
Packit Service 2781ba
}
Packit Service 2781ba
Packit Service 2781ba
static void
Packit Service 2781ba
edge_sobel (GeglBuffer          *src,
Packit Service 2781ba
            const GeglRectangle *src_rect,
Packit Service 2781ba
            GeglBuffer          *dst,
Packit Service 2781ba
            const GeglRectangle *dst_rect,
Packit Service 2781ba
            gboolean            horizontal,
Packit Service 2781ba
            gboolean            vertical,
Packit Service 2781ba
            gboolean            keep_signal)
Packit Service 2781ba
{
Packit Service 2781ba
Packit Service 2781ba
  gint x,y;
Packit Service 2781ba
  gint offset;
Packit Service 2781ba
  gfloat *src_buf;
Packit Service 2781ba
  gfloat *dst_buf;
Packit Service 2781ba
Packit Service 2781ba
  gint src_width = src_rect->width;
Packit Service 2781ba
Packit Service 2781ba
  src_buf = g_new0 (gfloat, src_rect->width * src_rect->height * 4);
Packit Service 2781ba
  dst_buf = g_new0 (gfloat, dst_rect->width * dst_rect->height * 4);
Packit Service 2781ba
Packit Service 2781ba
  gegl_buffer_get (src, src_rect, 1.0, babl_format ("RGBA float"), src_buf, GEGL_AUTO_ROWSTRIDE,
Packit Service 2781ba
                   GEGL_ABYSS_NONE);
Packit Service 2781ba
Packit Service 2781ba
  offset = 0;
Packit Service 2781ba
Packit Service 2781ba
  for (y=0; y<dst_rect->height; y++)
Packit Service 2781ba
    for (x=0; x<dst_rect->width; x++)
Packit Service 2781ba
      {
Packit Service 2781ba
Packit Service 2781ba
        gfloat hor_grad[3] = {0.0f, 0.0f, 0.0f};
Packit Service 2781ba
        gfloat ver_grad[3] = {0.0f, 0.0f, 0.0f};
Packit Service 2781ba
        gfloat gradient[4] = {0.0f, 0.0f, 0.0f, 0.0f};
Packit Service 2781ba
Packit Service 2781ba
        gfloat *center_pix = src_buf + ((x+SOBEL_RADIUS)+((y+SOBEL_RADIUS) * src_width)) * 4;
Packit Service 2781ba
Packit Service 2781ba
        gint c;
Packit Service 2781ba
Packit Service 2781ba
        if (horizontal)
Packit Service 2781ba
          {
Packit Service 2781ba
            gint i=x+SOBEL_RADIUS, j=y+SOBEL_RADIUS;
Packit Service 2781ba
            gfloat *src_pix = src_buf + (i + j * src_width) * 4;
Packit Service 2781ba
Packit Service 2781ba
            for (c=0;c<3;c++)
Packit Service 2781ba
                hor_grad[c] +=
Packit Service 2781ba
                    -1.0f*src_pix[c-4-src_width*4]+ src_pix[c+4-src_width*4] +
Packit Service 2781ba
                    -2.0f*src_pix[c-4] + 2.0f*src_pix[c+4] +
Packit Service 2781ba
                    -1.0f*src_pix[c-4+src_width*4]+ src_pix[c+4+src_width*4];
Packit Service 2781ba
          }
Packit Service 2781ba
Packit Service 2781ba
        if (vertical)
Packit Service 2781ba
          {
Packit Service 2781ba
            gint i=x+SOBEL_RADIUS, j=y+SOBEL_RADIUS;
Packit Service 2781ba
            gfloat *src_pix = src_buf + (i + j * src_width) * 4;
Packit Service 2781ba
Packit Service 2781ba
            for (c=0;c<3;c++)
Packit Service 2781ba
                ver_grad[c] +=
Packit Service 2781ba
                  -1.0f*src_pix[c-4-src_width*4]-2.0f*src_pix[c-src_width*4]-1.0f*src_pix[c+4-src_width*4] +
Packit Service 2781ba
                  src_pix[c-4+src_width*4]+2.0f*src_pix[c+src_width*4]+     src_pix[c+4+src_width*4];
Packit Service 2781ba
        }
Packit Service 2781ba
Packit Service 2781ba
        if (horizontal && vertical)
Packit Service 2781ba
          {
Packit Service 2781ba
            for (c=0;c<3;c++)
Packit Service 2781ba
              // normalization to [0, 1]
Packit Service 2781ba
              gradient[c] = RMS(hor_grad[c],ver_grad[c])/1.41f;
Packit Service 2781ba
          }
Packit Service 2781ba
        else
Packit Service 2781ba
          {
Packit Service 2781ba
            if (keep_signal)
Packit Service 2781ba
              {
Packit Service 2781ba
                for (c=0;c<3;c++)
Packit Service 2781ba
                  gradient[c] = hor_grad[c]+ver_grad[c];
Packit Service 2781ba
              }
Packit Service 2781ba
            else
Packit Service 2781ba
              {
Packit Service 2781ba
                for (c=0;c<3;c++)
Packit Service 2781ba
                  gradient[c] = fabsf(hor_grad[c]+ver_grad[c]);
Packit Service 2781ba
              }
Packit Service 2781ba
          }
Packit Service 2781ba
Packit Service 2781ba
        //alpha
Packit Service 2781ba
        gradient[3] = center_pix[3];
Packit Service 2781ba
Packit Service 2781ba
        for (c=0; c<4;c++)
Packit Service 2781ba
          dst_buf[offset*4+c] = gradient[c];
Packit Service 2781ba
Packit Service 2781ba
        offset++;
Packit Service 2781ba
      }
Packit Service 2781ba
Packit Service 2781ba
  gegl_buffer_set (dst, dst_rect, 0, babl_format ("RGBA float"), dst_buf,
Packit Service 2781ba
                   GEGL_AUTO_ROWSTRIDE);
Packit Service 2781ba
  g_free (src_buf);
Packit Service 2781ba
  g_free (dst_buf);
Packit Service 2781ba
}
Packit Service 2781ba
Packit Service 2781ba
Packit Service 2781ba
static void
Packit Service 2781ba
gegl_chant_class_init (GeglChantClass *klass)
Packit Service 2781ba
{
Packit Service 2781ba
  GeglOperationClass       *operation_class;
Packit Service 2781ba
  GeglOperationFilterClass *filter_class;
Packit Service 2781ba
Packit Service 2781ba
  operation_class  = GEGL_OPERATION_CLASS (klass);
Packit Service 2781ba
  filter_class     = GEGL_OPERATION_FILTER_CLASS (klass);
Packit Service 2781ba
Packit Service 2781ba
  filter_class->process   = process;
Packit Service 2781ba
  operation_class->prepare = prepare;
Packit Service 2781ba
Packit Service 2781ba
  operation_class->opencl_support = TRUE;
Packit Service 2781ba
Packit Service 2781ba
  gegl_operation_class_set_keys (operation_class,
Packit Service 2781ba
    "name"       , "gegl:edge-sobel",
Packit Service 2781ba
    "categories" , "edge-detect",
Packit Service 2781ba
    "description",
Packit Service 2781ba
          _("Specialized direction-dependent edge detection"),
Packit Service 2781ba
          NULL);
Packit Service 2781ba
}
Packit Service 2781ba
Packit Service 2781ba
#endif