Blame operations/common/c2g.c

Packit Service 2781ba
/* STRESS, Spatio Temporal Retinex Envelope with Stochastic Sampling
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
 * Copyright 2007,2009 Øyvind Kolås     <pippin@gimp.org>
Packit Service 2781ba
 *                     Ivar Farup       <ivarf@hig.no>
Packit Service 2781ba
 *                     Allesandro Rizzi <rizzi@dti.unimi.it>
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_int_ui (radius, _("Radius"), 2, 3000, 300, 2, 3000, 1.6,
Packit Service 2781ba
                _("Neighborhood taken into account, this is the radius in pixels taken into account when deciding which colors map to which gray values"))
Packit Service 2781ba
gegl_chant_int_ui (samples, _("Samples"), 1, 1000, 4, 1, 20, 1.0,
Packit Service 2781ba
                _("Number of samples to do per iteration looking for the range of colors"))
Packit Service 2781ba
gegl_chant_int_ui (iterations, _("Iterations"), 1, 1000, 10, 1, 20, 1.0,
Packit Service 2781ba
                _("Number of iterations, a higher number of iterations provides less noisy results at a computational cost"))
Packit Service 2781ba
Packit Service 2781ba
/*
Packit Service 2781ba
gegl_chant_double (rgamma, _("Radial Gamma"), 0.0, 8.0, 2.0,
Packit Service 2781ba
                _("Gamma applied to radial distribution"))
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       "c2g.c"
Packit Service 2781ba
Packit Service 2781ba
#include "gegl-chant.h"
Packit Service 2781ba
#include <math.h>
Packit Service 2781ba
#include <stdlib.h>
Packit Service 2781ba
#include "envelopes.h"
Packit Service 2781ba
Packit Service 2781ba
#define RGAMMA 2.0
Packit Service 2781ba
Packit Service 2781ba
static void c2g (GeglBuffer          *src,
Packit Service 2781ba
                 const GeglRectangle *src_rect,
Packit Service 2781ba
                 GeglBuffer          *dst,
Packit Service 2781ba
                 const GeglRectangle *dst_rect,
Packit Service 2781ba
                 gint                 radius,
Packit Service 2781ba
                 gint                 samples,
Packit Service 2781ba
                 gint                 iterations,
Packit Service 2781ba
                 gdouble              rgamma)
Packit Service 2781ba
{
Packit Service 2781ba
  gint x,y;
Packit Service 2781ba
  gint    dst_offset=0;
Packit Service 2781ba
  gfloat *src_buf;
Packit Service 2781ba
  gfloat *dst_buf;
Packit Service 2781ba
  gint    inw = src_rect->width;
Packit Service 2781ba
  gint    outw = dst_rect->width;
Packit Service 2781ba
  gint    inh = src_rect->height;
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 * 2);
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
  for (y=radius; y<dst_rect->height+radius; y++)
Packit Service 2781ba
    {
Packit Service 2781ba
      gint src_offset = (inw*y+radius)*4;
Packit Service 2781ba
      for (x=radius; x
Packit Service 2781ba
        {
Packit Service 2781ba
          gfloat *pixel= src_buf + src_offset;
Packit Service 2781ba
          gfloat  min[4];
Packit Service 2781ba
          gfloat  max[4];
Packit Service 2781ba
Packit Service 2781ba
          compute_envelopes (src_buf,
Packit Service 2781ba
                             inw, inh,
Packit Service 2781ba
                             x, y,
Packit Service 2781ba
                             radius, samples,
Packit Service 2781ba
                             iterations,
Packit Service 2781ba
                             FALSE, /* same spray */
Packit Service 2781ba
                             rgamma,
Packit Service 2781ba
                             min, max);
Packit Service 2781ba
          {
Packit Service 2781ba
            /* this should be replaced with a better/faster projection of
Packit Service 2781ba
             * pixel onto the vector spanned by min -> max, currently
Packit Service 2781ba
             * computed by comparing the distance to min with the sum
Packit Service 2781ba
             * of the distance to min/max.
Packit Service 2781ba
             */
Packit Service 2781ba
Packit Service 2781ba
            gfloat nominator = 0;
Packit Service 2781ba
            gfloat denominator = 0;
Packit Service 2781ba
            gint c;
Packit Service 2781ba
            for (c=0; c<3; c++)
Packit Service 2781ba
              {
Packit Service 2781ba
                nominator   += (pixel[c] - min[c]) * (pixel[c] - min[c]);
Packit Service 2781ba
                denominator += (pixel[c] - max[c]) * (pixel[c] - max[c]);
Packit Service 2781ba
              }
Packit Service 2781ba
Packit Service 2781ba
            nominator = sqrt (nominator);
Packit Service 2781ba
            denominator = sqrt (denominator);
Packit Service 2781ba
            denominator = nominator + denominator;
Packit Service 2781ba
Packit Service 2781ba
            if (denominator>0.000)
Packit Service 2781ba
              {
Packit Service 2781ba
                dst_buf[dst_offset+0] = nominator/denominator;
Packit Service 2781ba
              }
Packit Service 2781ba
            else
Packit Service 2781ba
              {
Packit Service 2781ba
                /* shouldn't happen */
Packit Service 2781ba
                dst_buf[dst_offset+0] = 0.5;
Packit Service 2781ba
              }
Packit Service 2781ba
            dst_buf[dst_offset+1] = src_buf[src_offset+3];
Packit Service 2781ba
Packit Service 2781ba
            src_offset+=4;
Packit Service 2781ba
            dst_offset+=2;
Packit Service 2781ba
          }
Packit Service 2781ba
        }
Packit Service 2781ba
    }
Packit Service 2781ba
  gegl_buffer_set (dst, dst_rect, 0, babl_format ("YA float"), dst_buf, 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
static void prepare (GeglOperation *operation)
Packit Service 2781ba
{
Packit Service 2781ba
  GeglOperationAreaFilter *area = GEGL_OPERATION_AREA_FILTER (operation);
Packit Service 2781ba
  area->left = area->right = area->top = area->bottom =
Packit Service 2781ba
      ceil (GEGL_CHANT_PROPERTIES (operation)->radius);
Packit Service 2781ba
Packit Service 2781ba
  gegl_operation_set_format (operation, "input", babl_format ("RGBA float"));
Packit Service 2781ba
  gegl_operation_set_format (operation, "output", babl_format ("YA float"));
Packit Service 2781ba
}
Packit Service 2781ba
Packit Service 2781ba
static GeglRectangle
Packit Service 2781ba
get_bounding_box (GeglOperation *operation)
Packit Service 2781ba
{
Packit Service 2781ba
  GeglRectangle  result = {0,0,0,0};
Packit Service 2781ba
  GeglRectangle *in_rect = gegl_operation_source_get_bounding_box (operation,
Packit Service 2781ba
                                                                     "input");
Packit Service 2781ba
  if (!in_rect)
Packit Service 2781ba
    return result;
Packit Service 2781ba
  return *in_rect;
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 TRUE true                                                     \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"#define FALSE false                                                   \n"
Packit Service 2781ba
"#define ANGLE_PRIME 95273                                             \n"
Packit Service 2781ba
"#define RADIUS_PRIME 29537                                            \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"void sample_min_max(const __global   float4 *src_buf,                 \n"
Packit Service 2781ba
"                                     int     src_width,               \n"
Packit Service 2781ba
"                                     int     src_height,              \n"
Packit Service 2781ba
"                    const __global   float  *radiuses,                \n"
Packit Service 2781ba
"                    const __global   float  *lut_cos,                 \n"
Packit Service 2781ba
"                    const __global   float  *lut_sin,                 \n"
Packit Service 2781ba
"                                     int     x,                       \n"
Packit Service 2781ba
"                                     int     y,                       \n"
Packit Service 2781ba
"                                     int     radius,                  \n"
Packit Service 2781ba
"                                     int     samples,                 \n"
Packit Service 2781ba
"                                     float4 *min,                     \n"
Packit Service 2781ba
"                                     float4 *max,                     \n"
Packit Service 2781ba
"                                     int     j,                       \n"
Packit Service 2781ba
"                                     int     iterations)              \n"
Packit Service 2781ba
"{                                                                     \n"
Packit Service 2781ba
"    float4 best_min;                                                  \n"
Packit Service 2781ba
"    float4 best_max;                                                  \n"
Packit Service 2781ba
"    float4 center_pix = *(src_buf + src_width * y + x);               \n"
Packit Service 2781ba
"    int i;                                                            \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    best_min = center_pix;                                            \n"
Packit Service 2781ba
"    best_max = center_pix;                                            \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    int angle_no  = (src_width * y + x) * (iterations) *              \n"
Packit Service 2781ba
"                       samples + j * samples;                         \n"
Packit Service 2781ba
"    int radius_no = angle_no;                                         \n"
Packit Service 2781ba
"    angle_no  %= ANGLE_PRIME;                                         \n"
Packit Service 2781ba
"    radius_no %= RADIUS_PRIME;                                        \n"
Packit Service 2781ba
"    for(i=0; i
Packit Service 2781ba
"    {                                                                 \n"
Packit Service 2781ba
"        int angle;                                                    \n"
Packit Service 2781ba
"        float rmag;                                                   \n"
Packit Service 2781ba
"        /* if we've sampled outside the valid image                   \n"
Packit Service 2781ba
"           area, we grab another sample instead, this                 \n"
Packit Service 2781ba
"           should potentially work better than mirroring              \n"
Packit Service 2781ba
"           or extending the image */                                  \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"         angle = angle_no++;                                          \n"
Packit Service 2781ba
"         rmag  = radiuses[radius_no++] * radius;                      \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"         if( angle_no  >= ANGLE_PRIME)                                \n"
Packit Service 2781ba
"             angle_no   = 0;                                          \n"
Packit Service 2781ba
"         if( radius_no >= RADIUS_PRIME)                               \n"
Packit Service 2781ba
"             radius_no  = 0;                                          \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"         int u = x + rmag * lut_cos[angle];                           \n"
Packit Service 2781ba
"         int v = y + rmag * lut_sin[angle];                           \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"         if(u>=src_width || u <0 || v>=src_height || v<0)             \n"
Packit Service 2781ba
"         {                                                            \n"
Packit Service 2781ba
"             //--i;                                                   \n"
Packit Service 2781ba
"             continue;                                                \n"
Packit Service 2781ba
"         }                                                            \n"
Packit Service 2781ba
"         float4 pixel = *(src_buf + (src_width * v + u));             \n"
Packit Service 2781ba
"         if(pixel.w<=0.0f)                                            \n"
Packit Service 2781ba
"         {                                                            \n"
Packit Service 2781ba
"             //--i;                                                   \n"
Packit Service 2781ba
"             continue;                                                \n"
Packit Service 2781ba
"         }                                                            \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"         best_min = pixel < best_min ? pixel : best_min;              \n"
Packit Service 2781ba
"         best_max = pixel > best_max ? pixel : best_max;              \n"
Packit Service 2781ba
"    }                                                                 \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    (*min).xyz = best_min.xyz;                                        \n"
Packit Service 2781ba
"    (*max).xyz = best_max.xyz;                                        \n"
Packit Service 2781ba
"}                                                                     \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"void compute_envelopes_CL(const __global  float4 *src_buf,            \n"
Packit Service 2781ba
"                                          int     src_width,          \n"
Packit Service 2781ba
"                                          int     src_height,         \n"
Packit Service 2781ba
"                          const __global  float  *radiuses,           \n"
Packit Service 2781ba
"                          const __global  float  *lut_cos,            \n"
Packit Service 2781ba
"                          const __global  float  *lut_sin,            \n"
Packit Service 2781ba
"                                          int     x,                  \n"
Packit Service 2781ba
"                                          int     y,                  \n"
Packit Service 2781ba
"                                          int     radius,             \n"
Packit Service 2781ba
"                                          int     samples,            \n"
Packit Service 2781ba
"                                          int     iterations,         \n"
Packit Service 2781ba
"                                          float4 *min_envelope,       \n"
Packit Service 2781ba
"                                          float4 *max_envelope)       \n"
Packit Service 2781ba
"{                                                                     \n"
Packit Service 2781ba
"    float4 range_sum = 0;                                             \n"
Packit Service 2781ba
"    float4 relative_brightness_sum = 0;                               \n"
Packit Service 2781ba
"    float4 pixel = *(src_buf + src_width * y + x);                    \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    int i;                                                            \n"
Packit Service 2781ba
"    for(i =0; i
Packit Service 2781ba
"    {                                                                 \n"
Packit Service 2781ba
"        float4 min,max;                                               \n"
Packit Service 2781ba
"        float4 range, relative_brightness;                            \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"        sample_min_max(src_buf, src_width, src_height,                \n"
Packit Service 2781ba
"                        radiuses, lut_cos, lut_sin, x, y,             \n"
Packit Service 2781ba
"                        radius,samples,&min,&max,i,iterations);       \n"
Packit Service 2781ba
"        range = max - min;                                            \n"
Packit Service 2781ba
"        relative_brightness = range <= 0.0f ?                         \n"
Packit Service 2781ba
"                               0.5f : (pixel - min) / range;          \n"
Packit Service 2781ba
"        relative_brightness_sum += relative_brightness;               \n"
Packit Service 2781ba
"        range_sum += range;                                           \n"
Packit Service 2781ba
"    }                                                                 \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    float4 relative_brightness = relative_brightness_sum / iterations;\n"
Packit Service 2781ba
"    float4 range = range_sum / iterations;                            \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    if(max_envelope)                                                  \n"
Packit Service 2781ba
"        *max_envelope = pixel + (1.0f - relative_brightness) * range; \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    if(min_envelope)                                                  \n"
Packit Service 2781ba
"        *min_envelope = pixel - relative_brightness * range;          \n"
Packit Service 2781ba
"}                                                                     \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"__kernel void C2g_CL(const __global float4 *src_buf,                  \n"
Packit Service 2781ba
"                                    int     src_width,                \n"
Packit Service 2781ba
"                                    int     src_height,               \n"
Packit Service 2781ba
"                     const __global float  *radiuses,                 \n"
Packit Service 2781ba
"                     const __global float  *lut_cos,                  \n"
Packit Service 2781ba
"                     const __global float  *lut_sin,                  \n"
Packit Service 2781ba
"                           __global float2 *dst_buf,                  \n"
Packit Service 2781ba
"                                    int     radius,                   \n"
Packit Service 2781ba
"                                    int     samples,                  \n"
Packit Service 2781ba
"                                    int     iterations)               \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
"    int x = gidx + radius;                                            \n"
Packit Service 2781ba
"    int y = gidy + radius;                                            \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    int src_offset = (src_width * y + x);                             \n"
Packit Service 2781ba
"    int dst_offset = gidx + get_global_size(0) * gidy;                \n"
Packit Service 2781ba
"    float4 min,max;                                                   \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    compute_envelopes_CL(src_buf, src_width, src_height,              \n"
Packit Service 2781ba
"                         radiuses, lut_cos, lut_sin, x, y,            \n"
Packit Service 2781ba
"                         radius, samples, iterations, &min, &max;;    \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    float4 pixel = *(src_buf + src_offset);                           \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    float nominator=0, denominator=0;                                 \n"
Packit Service 2781ba
"    float4 t1 = (pixel - min) * (pixel - min);                        \n"
Packit Service 2781ba
"    float4 t2 = (pixel - max) * (pixel - max);                        \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    nominator   = t1.x + t1.y + t1.z;                                 \n"
Packit Service 2781ba
"    denominator = t2.x + t2.y + t2.z;                                 \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    nominator   = sqrt(nominator);                                    \n"
Packit Service 2781ba
"    denominator = sqrt(denominator);                                  \n"
Packit Service 2781ba
"    denominator+= nominator + denominator;                            \n"
Packit Service 2781ba
"                                                                      \n"
Packit Service 2781ba
"    dst_buf[dst_offset].x = (denominator > 0.000f)                    \n"
Packit Service 2781ba
"                             ? (nominator / denominator) : 0.5f;      \n"
Packit Service 2781ba
"    dst_buf[dst_offset].y =  src_buf[src_offset].w;                   \n"
Packit Service 2781ba
"}                                                                     \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_c2g (cl_mem                in_tex,
Packit Service 2781ba
    cl_mem                    out_tex,
Packit Service 2781ba
    size_t                    global_worksize,
Packit Service 2781ba
    const GeglRectangle      *src_roi,
Packit Service 2781ba
    const GeglRectangle      *roi,
Packit Service 2781ba
    gint                      radius,
Packit Service 2781ba
    gint                      samples,
Packit Service 2781ba
    gint                      iterations,
Packit Service 2781ba
    gdouble                   rgamma)
Packit Service 2781ba
{
Packit Service 2781ba
  cl_int cl_err = 0;
Packit Service 2781ba
  cl_mem cl_lut_cos, cl_lut_sin, cl_radiuses;
Packit Service 2781ba
  const size_t gbl_size[2] = {roi->width, roi->height};
Packit Service 2781ba
Packit Service 2781ba
  if (!cl_data)
Packit Service 2781ba
    {
Packit Service 2781ba
      const char *kernel_name[] ={"C2g_CL", 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
  compute_luts(rgamma);
Packit Service 2781ba
Packit Service 2781ba
  cl_lut_cos = gegl_clCreateBuffer(gegl_cl_get_context(),
Packit Service 2781ba
                                   CL_MEM_READ_ONLY,
Packit Service 2781ba
                                   ANGLE_PRIME * sizeof(cl_float), NULL, &cl_err);
Packit Service 2781ba
Packit Service 2781ba
  cl_err |= gegl_clEnqueueWriteBuffer(gegl_cl_get_command_queue(), cl_lut_cos,
Packit Service 2781ba
                                      CL_TRUE, 0, ANGLE_PRIME * sizeof(cl_float), lut_cos,
Packit Service 2781ba
                                      0, NULL, NULL);
Packit Service 2781ba
  if (CL_SUCCESS != cl_err)   return cl_err;
Packit Service 2781ba
Packit Service 2781ba
  cl_lut_sin = gegl_clCreateBuffer(gegl_cl_get_context(),
Packit Service 2781ba
                                   CL_MEM_READ_ONLY,
Packit Service 2781ba
                                   ANGLE_PRIME * sizeof(cl_float), NULL, &cl_err);
Packit Service 2781ba
Packit Service 2781ba
  cl_err |= gegl_clEnqueueWriteBuffer(gegl_cl_get_command_queue(), cl_lut_sin,
Packit Service 2781ba
                                      CL_TRUE, 0, ANGLE_PRIME * sizeof(cl_float), lut_sin,
Packit Service 2781ba
                                      0, NULL, NULL);
Packit Service 2781ba
  if (CL_SUCCESS != cl_err)    return cl_err;
Packit Service 2781ba
Packit Service 2781ba
  cl_radiuses = gegl_clCreateBuffer(gegl_cl_get_context(),
Packit Service 2781ba
                                    CL_MEM_READ_ONLY,
Packit Service 2781ba
                                    RADIUS_PRIME * sizeof(cl_float), NULL, &cl_err);
Packit Service 2781ba
Packit Service 2781ba
  cl_err |= gegl_clEnqueueWriteBuffer(gegl_cl_get_command_queue(), cl_radiuses,
Packit Service 2781ba
                                      CL_TRUE, 0, RADIUS_PRIME * sizeof(cl_float), radiuses,
Packit Service 2781ba
                                      0, NULL, NULL);
Packit Service 2781ba
  if (CL_SUCCESS != cl_err)    return cl_err;
Packit Service 2781ba
Packit Service 2781ba
  {
Packit Service 2781ba
  cl_int cl_src_width  = src_roi->width;
Packit Service 2781ba
  cl_int cl_src_height = src_roi->height;
Packit Service 2781ba
  cl_int cl_radius     = radius;
Packit Service 2781ba
  cl_int cl_samples    = samples;
Packit Service 2781ba
  cl_int cl_iterations = iterations;
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_int), (void*)&cl_src_width);
Packit Service 2781ba
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int), (void*)&cl_src_height);
Packit Service 2781ba
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_mem), (void*)&cl_radiuses);
Packit Service 2781ba
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_mem), (void*)&cl_lut_cos);
Packit Service 2781ba
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_mem), (void*)&cl_lut_sin);
Packit Service 2781ba
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 6, sizeof(cl_mem), (void*)&out_tex);
Packit Service 2781ba
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 7, sizeof(cl_int), (void*)&cl_radius);
Packit Service 2781ba
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 8, sizeof(cl_int), (void*)&cl_samples);
Packit Service 2781ba
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 9, sizeof(cl_int), (void*)&cl_iterations);
Packit Service 2781ba
  if (cl_err != CL_SUCCESS) return cl_err;
Packit Service 2781ba
  }
Packit Service 2781ba
Packit Service 2781ba
  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(), cl_data->kernel[0],
Packit Service 2781ba
                                       2, 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
  cl_err = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
Packit Service 2781ba
  if (CL_SUCCESS != cl_err) return cl_err;
Packit Service 2781ba
Packit Service 2781ba
  gegl_clFinish(gegl_cl_get_command_queue ());
Packit Service 2781ba
Packit Service 2781ba
  gegl_clReleaseMemObject(cl_radiuses);
Packit Service 2781ba
  gegl_clReleaseMemObject(cl_lut_cos);
Packit Service 2781ba
  gegl_clReleaseMemObject(cl_lut_sin);
Packit Service 2781ba
Packit Service 2781ba
  return cl_err;
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  = babl_format("RGBA float");
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,
Packit Service 2781ba
                                                           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_c2g(i->tex[read][j], i->tex[0][j],i->size[0][j], &i->roi[read][j],&i->roi[0][j],
Packit Service 2781ba
                      o->radius,o->samples,o->iterations,RGAMMA);
Packit Service 2781ba
      if (cl_err != CL_SUCCESS)
Packit Service 2781ba
      {
Packit Service 2781ba
        g_warning("[OpenCL] Error in gegl:c2g: %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
  compute = gegl_operation_get_required_for_output (operation, "input",result);
Packit Service 2781ba
Packit Service 2781ba
  if (o->radius < 500 && 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
  c2g (input, &compute, output, result,
Packit Service 2781ba
       o->radius,
Packit Service 2781ba
       o->samples,
Packit Service 2781ba
       o->iterations,
Packit Service 2781ba
       /*o->rgamma*/RGAMMA);
Packit Service 2781ba
Packit Service 2781ba
  return  TRUE;
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
  /* we override defined region to avoid growing the size of what is defined
Packit Service 2781ba
   * by the filter. This also allows the tricks used to treat alpha==0 pixels
Packit Service 2781ba
   * in the image as source data not to be skipped by the stochastic sampling
Packit Service 2781ba
   * yielding correct edge behavior.
Packit Service 2781ba
   */
Packit Service 2781ba
  operation_class->get_bounding_box = get_bounding_box;
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:c2g",
Packit Service 2781ba
      "categories",  "enhance",
Packit Service 2781ba
      "description", 
Packit Service 2781ba
     _("Color to grayscale conversion, uses envelopes formed from spatial "
Packit Service 2781ba
       " color differences to perform color-feature preserving grayscale "
Packit Service 2781ba
       " spatial contrast enhancement"),
Packit Service 2781ba
     NULL);
Packit Service 2781ba
}
Packit Service 2781ba
Packit Service 2781ba
#endif