Пример #1
0
static gboolean
cl_stretch_contrast (cl_mem               in_tex,
                     cl_mem               out_tex,
                     size_t               global_worksize,
                     const GeglRectangle *roi,
                     cl_float4            min,
                     cl_float4            diff)
{
  cl_int cl_err  = 0;

  cl_err = gegl_clSetKernelArg(cl_data->kernel[2], 0, sizeof(cl_mem),
                               (void*)&in_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[2], 1, sizeof(cl_mem),
                               (void*)&out_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[2], 2, sizeof(cl_float4),
                               (void*)&min);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[2], 3, sizeof(cl_float4),
                               (void*)&diff);
  CL_CHECK;

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

  return FALSE;

 error:
  return TRUE;
}
Пример #2
0
/* OpenCL processing function */
static cl_int
cl_process (GeglOperation       *op,
            cl_mem               in_tex,
            cl_mem               out_tex,
            size_t               global_worksize,
            const GeglRectangle *roi,
            gint                 level)
{
  /* Retrieve a pointer to GeglChantO structure which contains all the
   * chanted properties
   */

  GeglChantO *o = GEGL_CHANT_PROPERTIES (op);

  gfloat      in_range;
  gfloat      out_range;
  gfloat      in_offset;
  gfloat      out_offset;
  gfloat      scale;

  cl_int cl_err = 0;

  in_offset = o->in_low * 1.0;
  out_offset = o->out_low * 1.0;
  in_range = o->in_high-o->in_low;
  out_range = o->out_high-o->out_low;

  if (in_range == 0.0)
    in_range = 0.00000001;

  scale = out_range/in_range;

  if (!cl_data)
    {
      const char *kernel_name[] = {"kernel_levels", NULL};
      cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
    }
  if (!cl_data) return 1;

  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem),   (void*)&in_tex);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem),   (void*)&out_tex);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&in_offset);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&out_offset);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float), (void*)&scale);
  if (cl_err != CL_SUCCESS) return cl_err;

  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
                                        cl_data->kernel[0], 1,
                                        NULL, &global_worksize, NULL,
                                        0, NULL, NULL);
  if (cl_err != CL_SUCCESS) return cl_err;

  return cl_err;
}
Пример #3
0
/* OpenCL processing function */
static gboolean
cl_process (GeglOperation       *op,
            cl_mem               in_tex,
            cl_mem               out_tex,
            size_t               global_worksize,
            const GeglRectangle *roi,
            int                  level)
{
    /* Retrieve a pointer to GeglProperties structure which contains all the
     * chanted properties
     */

    GeglProperties   *o         = GEGL_PROPERTIES (op);
    const gfloat *coeffs    = o->user_data;

    cl_int cl_err = 0;

    if (! coeffs)
    {
        coeffs = o->user_data = preprocess (o);
    }

    if (!cl_data)
    {
        const char *kernel_name[] = {"gegl_color_temperature", NULL};
        cl_data = gegl_cl_compile_and_build (color_temperature_cl_source, kernel_name);
    }

    if (!cl_data) return 1;

    cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem),    (void*)&in_tex);
    CL_CHECK;
    cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem),    (void*)&out_tex);
    CL_CHECK;
    cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float),  (void*)&coeffs[0]);
    CL_CHECK;
    cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float),  (void*)&coeffs[1]);
    CL_CHECK;
    cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float),  (void*)&coeffs[2]);
    CL_CHECK;

    cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
                                         cl_data->kernel[0], 1,
                                         NULL, &global_worksize, NULL,
                                         0, NULL, NULL);
    CL_CHECK;

    return FALSE;

error:
    return TRUE;
}
Пример #4
0
static gboolean
cl_process (GeglOperation       *operation,
            cl_mem              in,
            cl_mem              out,
            size_t              global_worksize,
            const GeglRectangle *roi,
            gint                level)
{
  GeglProperties *o = GEGL_PROPERTIES (operation);
  gfloat      color[4];
  gegl_color_get_pixel (o->color, babl_format ("R'G'B'A float"), color);

  if (!cl_data)
    {
      const char *kernel_name[] = {"cl_color_to_alpha",NULL};
      cl_data = gegl_cl_compile_and_build (color_to_alpha_cl_source, kernel_name);
    }
  if (!cl_data)
    return TRUE;
  else
  {
    cl_int cl_err = 0;
    cl_float4 f_color;
    f_color.s[0] = color[0];
    f_color.s[1] = color[1];
    f_color.s[2] = color[2];
    f_color.s[3] = color[3];

    cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0,  sizeof(cl_mem),   (void*)&in);
    CL_CHECK;
    cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1,  sizeof(cl_mem),   (void*)&out);
    CL_CHECK;
    cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2,  sizeof(cl_float4),(void*)&f_color);
    CL_CHECK;

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

  return  FALSE;

 error:
  return TRUE;
}
Пример #5
0
/* OpenCL processing function */
static cl_int
cl_process (GeglOperation       *op,
            cl_mem               in_tex,
            cl_mem               out_tex,
            size_t               global_worksize,
            const GeglRectangle *roi,
            gint                 level)
{
  /* Retrieve a pointer to GeglProperties structure which contains all the
   * chanted properties
   */

  GeglProperties *o = GEGL_PROPERTIES (op);

  gfloat      black_level = (gfloat) o->black_level;
  gfloat      diff;
  gfloat      exposure_negated = (gfloat) -o->exposure;
  gfloat      gain;
  gfloat      white;
  
  cl_int cl_err = 0;

  if (!cl_data)
    {
      const char *kernel_name[] = {"kernel_exposure", NULL};
      cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
    }
  if (!cl_data) return 1;

  white = exp2f (exposure_negated);
  diff = MAX (white - black_level, 0.01);
  gain = 1.0f / diff;

  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem),   (void*)&in_tex);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem),   (void*)&out_tex);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&black_level);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&gain);
  if (cl_err != CL_SUCCESS) return cl_err;

  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
                                        cl_data->kernel[0], 1,
                                        NULL, &global_worksize, NULL,
                                        0, NULL, NULL);
  if (cl_err != CL_SUCCESS) return cl_err;

  return cl_err;
}
Пример #6
0
static cl_int
cl_edge_laplace (cl_mem                in_tex,
                 cl_mem                aux_tex,
                 cl_mem                out_tex,
                 const GeglRectangle  *src_rect,
                 const GeglRectangle  *roi,
                 gint                  radius)
{
  cl_int cl_err = 0;
  size_t global_ws[2];
  if (!cl_data)
  {
    const char *kernel_name[] = {"pre_edgelaplace", "knl_edgelaplace", NULL};
    cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
  }

  if (!cl_data) return 1;

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

  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem),   (void*)&in_tex);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem),   (void*)&aux_tex);
  if (cl_err != CL_SUCCESS) return cl_err;

  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
                                       cl_data->kernel[0], 2,
                                       NULL, global_ws, NULL,
                                       0, NULL, NULL);
  if (cl_err != CL_SUCCESS) return cl_err;

  cl_err = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
  if (CL_SUCCESS != cl_err) return cl_err;

  cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem),   (void*)&aux_tex);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_mem),   (void*)&out_tex);
  if (cl_err != CL_SUCCESS) return cl_err;

  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
                                       cl_data->kernel[1], 2,
                                       NULL, global_ws, NULL,
                                       0, NULL, NULL);
  if (cl_err != CL_SUCCESS) return cl_err;
  return cl_err;
}
Пример #7
0
/* OpenCL processing function */
static cl_int
cl_process (GeglOperation       *op,
            cl_mem               in_tex,
            cl_mem               out_tex,
            size_t               global_worksize,
            const GeglRectangle *roi,
            int                  level)
{
  /* Retrieve a pointer to GeglChantO structure which contains all the
   * chanted properties
   */

  GeglChantO   *o         = GEGL_CHANT_PROPERTIES (op);
  const gfloat *coeffs    = o->chant_data;

  cl_int cl_err = 0;

  if (! coeffs)
    {
      coeffs = o->chant_data = preprocess (o);
    }

  if (!cl_data)
    {
      const char *kernel_name[] = {"kernel_temp", NULL};
      cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
    }

  if (!cl_data) return 1;

  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem),    (void*)&in_tex);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem),    (void*)&out_tex);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float),  (void*)&coeffs[0]);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float),  (void*)&coeffs[1]);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float),  (void*)&coeffs[2]);
  if (cl_err != CL_SUCCESS) return cl_err;

  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
                                        cl_data->kernel[0], 1,
                                        NULL, &global_worksize, NULL,
                                        0, NULL, NULL);
  if (cl_err != CL_SUCCESS) return cl_err;

  return cl_err;
}
Пример #8
0
static gboolean
cl_process (GeglOperation       *op,
            cl_mem               in_tex,
            cl_mem               aux_tex,
            cl_mem               out_tex,
            size_t               global_worksize,
            const GeglRectangle *roi,
            gint                 level)
{
  cl_int cl_err = 0;
  int kernel;
  gfloat value;

  if (!cl_data)
    {
      const char *kernel_name[] = {"gegl_opacity_RaGaBaA_float", "gegl_opacity_RGBA_float", NULL};
      cl_data = gegl_cl_compile_and_build (opacity_cl_source, kernel_name);
    }
  if (!cl_data) return TRUE;

  value = GEGL_CHANT_PROPERTIES (op)->value;

  kernel = (GEGL_CHANT_PROPERTIES (op)->chant_data == NULL)? 0 : 1;

  cl_err = gegl_clSetKernelArg(cl_data->kernel[kernel], 0, sizeof(cl_mem),   (void*)&in_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[kernel], 1, sizeof(cl_mem),   (aux_tex)? (void*)&aux_tex : NULL);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[kernel], 2, sizeof(cl_mem),   (void*)&out_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[kernel], 3, sizeof(cl_float), (void*)&value);
  CL_CHECK;

  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
                                        cl_data->kernel[kernel], 1,
                                        NULL, &global_worksize, NULL,
                                        0, NULL, NULL);
  CL_CHECK;

  return FALSE;

error:
  return TRUE;
}
Пример #9
0
/* OpenCL processing function */
static cl_int
cl_process (GeglOperation       *op,
            cl_mem               in_tex,
            cl_mem               out_tex,
            size_t               global_worksize,
            const GeglRectangle *roi,
            gint                 level)
{
    /* Retrieve a pointer to GeglChantO structure which contains all the
     * chanted properties
     */

    GeglChantO *o = GEGL_CHANT_PROPERTIES (op);

    gfloat      gain = powf(2.0, o->exposure);
    gfloat      offset = o->offset;
    gfloat      gamma = 1.0 / o->gamma;

    cl_int cl_err = 0;

    if (!cl_data)
    {
        const char *kernel_name[] = {"kernel_exposure", NULL};
        cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
    }
    if (!cl_data) return 1;

    cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem),   (void*)&in_tex);
    cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem),   (void*)&out_tex);
    cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&gain);
    cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&offset);
    cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float), (void*)&gamma);
    if (cl_err != CL_SUCCESS) return cl_err;

    cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
                                         cl_data->kernel[0], 1,
                                         NULL, &global_worksize, NULL,
                                         0, NULL, NULL);
    if (cl_err != CL_SUCCESS) return cl_err;

    return cl_err;
}
Пример #10
0
static gboolean
cl_bilateral_filter (cl_mem                in_tex,
                     cl_mem                out_tex,
                     size_t                global_worksize,
                     const GeglRectangle  *roi,
                     gfloat                radius,
                     gfloat                preserve)
{
  cl_int cl_err = 0;
  size_t global_ws[2];

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

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

  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem),   (void*)&in_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem),   (void*)&out_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&radius);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&preserve);
  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;

  return FALSE;

error:
  return TRUE;
}
Пример #11
0
static cl_int
cl_mono_mixer(cl_mem                in_tex,
              cl_mem                out_tex,
              size_t                global_worksize,
              const GeglRectangle  *roi,
              gfloat                red,
              gfloat                green,
              gfloat                blue)
{
  cl_int cl_err = 0;
  if (!cl_data)
  {
    const char *kernel_name[] = {"Mono_mixer_cl", NULL};
    cl_data = gegl_cl_compile_and_build(kernel_source, kernel_name);
  }
  if (!cl_data) return 0;

  {
  cl_float4 color;
  color.s[0] = red;
  color.s[1] = green;
  color.s[2] = blue;
  color.s[3] = 1.0f;

  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem),    (void*)&in_tex);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_float4), (void*)&color);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_mem),    (void*)&out_tex);
  if (cl_err != CL_SUCCESS) return cl_err;

  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(), cl_data->kernel[0],
                                       1, NULL,
                                       &global_worksize, NULL,
                                       0, NULL, NULL);
  }

  return cl_err;
}
Пример #12
0
static gboolean
cl_edge_sobel (cl_mem              in_tex,
               cl_mem              out_tex,
               size_t              global_worksize,
               const GeglRectangle *roi,
               gboolean            horizontal,
               gboolean            vertical,
               gboolean            keep_signal,
               gboolean            has_alpha)
{
  const size_t gbl_size[2] = {roi->width, roi->height};
  cl_int n_horizontal  = horizontal;
  cl_int n_vertical    = vertical;
  cl_int n_keep_signal = keep_signal;
  cl_int n_has_alpha   = has_alpha;
  cl_int cl_err = 0;

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

  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int), (void*)&n_horizontal);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_int), (void*)&n_vertical);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_int), (void*)&n_keep_signal);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_int), (void*)&n_has_alpha);
  CL_CHECK;

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

  return FALSE;

error:
  return TRUE;
}
Пример #13
0
static gboolean
cl_snn_mean (cl_mem                in_tex,
             cl_mem                out_tex,
             const GeglRectangle  *src_rect,
             const GeglRectangle  *roi,
             gint                  radius,
             gint                  pairs)
{
  cl_int cl_err = 0;
  size_t global_ws[2];

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


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

  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem),   (void*)&in_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_int),   (void*)&src_rect->width);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int),   (void*)&src_rect->height);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_mem),   (void*)&out_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_int),   (void*)&radius);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_int),   (void*)&pairs);
  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;

  return FALSE;

error:
  return TRUE;
}
Пример #14
0
static gboolean
cl_process (GeglOperation       *self,
            cl_mem               in_tex,
            cl_mem               out_tex,
            size_t               global_worksize,
            const GeglRectangle *roi,
            gint                 level)
{
  GeglChantO *o = GEGL_CHANT_PROPERTIES (self);
  gint       num_sampling_points;
  gdouble    *xs, *ys;
  gfloat     *ysf = NULL;
  cl_mem     cl_curve = NULL;
  cl_ulong   cl_max_constant_size;
  cl_int     cl_err = 0;

  num_sampling_points = o->sampling_points;

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

  if (num_sampling_points > 0)
    {
      xs = g_new (gdouble, num_sampling_points);
      ys = g_new (gdouble, num_sampling_points);

      gegl_curve_calc_values (o->curve, 0.0, 1.0, num_sampling_points, xs, ys);
      g_free (xs);

      /*We need to downscale the array to pass it to the GPU*/
      ysf = g_new (gfloat, num_sampling_points);
      copy_double_array_to_float_array (ys, ysf, num_sampling_points);
      g_free (ys);

      cl_err = gegl_clGetDeviceInfo (gegl_cl_get_device (),
                                     CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE,
                                     sizeof (cl_ulong),
                                     &cl_max_constant_size,
                                     NULL);
      CL_CHECK;

      GEGL_NOTE (GEGL_DEBUG_OPENCL, "Max Constant Mem Size: %lu bytes",
                 (unsigned long) cl_max_constant_size);

      if (sizeof (cl_float) * num_sampling_points < cl_max_constant_size)
        {
          cl_curve = gegl_clCreateBuffer (gegl_cl_get_context (),
                                          CL_MEM_USE_HOST_PTR | CL_MEM_READ_ONLY,
                                          num_sampling_points * sizeof (cl_float),
                                          ysf, &cl_err);
          CL_CHECK;
          cl_err = gegl_clSetKernelArg (cl_data->kernel[0], 0, sizeof (cl_mem),
                                        (void*) &in_tex);
          CL_CHECK;
          cl_err = gegl_clSetKernelArg (cl_data->kernel[0], 1, sizeof (cl_mem),
                                        (void*) &out_tex);
          CL_CHECK;
          cl_err = gegl_clSetKernelArg (cl_data->kernel[0], 2, sizeof (cl_mem),
                                        (void*) &cl_curve);
          CL_CHECK;
          cl_err = gegl_clSetKernelArg (cl_data->kernel[0], 3, sizeof (gint),
                                        (void*) &num_sampling_points);
          CL_CHECK;
          cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
                                                cl_data->kernel[0], 1,
                                                NULL, &global_worksize, NULL,
                                                0, NULL, NULL);
          CL_CHECK;

          cl_err = gegl_clFinish (gegl_cl_get_command_queue ());
          CL_CHECK;

          cl_err = gegl_clReleaseMemObject (cl_curve);
          CL_CHECK_ONLY (cl_err);
        }
      else
        {
          /*If the curve size doesn't fit constant memory is better to use CPU*/
          GEGL_NOTE (GEGL_DEBUG_OPENCL,
                     "Not enough constant memory for the curve");
          g_free (ysf);
          return TRUE;
        }

      g_free (ysf);
      return FALSE;
error:
      if (ysf)
        g_free (ysf);
      if (cl_curve)
        gegl_clReleaseMemObject (cl_curve);

      return TRUE;
    }
  else  /*If the curve doesn't have a lookup table is better to use CPU*/
    {
      GEGL_NOTE (GEGL_DEBUG_OPENCL,
                 "Curve not suitable to be computed in the GPU");
      return TRUE;
    }
}
Пример #15
0
static void
motion_blur_cl  (GeglBuffer  *src,
                 const GeglRectangle  *src_rect,
                 GeglBuffer  *dst,
                 const GeglRectangle  *dst_rect,
                 const int num_steps,
                 const float offset_x,
                 const float offset_y)
{
    const Babl  * in_format = babl_format("RaGaBaA float");
    const Babl  *out_format = babl_format("RaGaBaA float");
    /* AreaFilter general processing flow.
       Loading data and making the necessary color space conversion. */
#include "gegl-cl-operation-area-filter-fw1.h"
    ///////////////////////////////////////////////////////////////////////////
    /* Algorithm specific processing flow.
       Build kernels, setting parameters, and running them. */

    if (!cl_data)
    {
        const char *kernel_name[] = {
            "motion_blur_CL", NULL
        };
        cl_data = gegl_cl_compile_and_build(kernel_source, kernel_name);
    }
    if (!cl_data) CL_ERROR;

    cl_int cl_src_width  = src_rect->width;
    cl_int cl_src_height = src_rect->height;
    cl_int cl_src_x      = src_rect->x;
    cl_int cl_src_y      = src_rect->y;
    cl_int cl_dst_x      = dst_rect->x;
    cl_int cl_dst_y      = dst_rect->y;
    cl_int cl_num_steps  = num_steps;
    cl_float cl_offset_x   = offset_x;
    cl_float cl_offset_y   = offset_y;

    CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
                               cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&src_mem));
    CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
                               cl_data->kernel[0], 1, sizeof(cl_int), (void*)&cl_src_width));
    CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
                               cl_data->kernel[0], 2, sizeof(cl_int), (void*)&cl_src_height));
    CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
                               cl_data->kernel[0], 3, sizeof(cl_int), (void*)&cl_src_x));
    CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
                               cl_data->kernel[0], 4, sizeof(cl_int), (void*)&cl_src_y));
    CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
                               cl_data->kernel[0], 5, sizeof(cl_mem), (void*)&dst_mem));
    CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
                               cl_data->kernel[0], 6, sizeof(cl_int), (void*)&cl_dst_x));
    CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
                               cl_data->kernel[0], 7, sizeof(cl_int), (void*)&cl_dst_y));
    CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
                               cl_data->kernel[0], 8, sizeof(cl_int), (void*)&cl_num_steps));
    CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
                               cl_data->kernel[0], 9, sizeof(cl_float), (void*)&cl_offset_x));
    CL_SAFE_CALL(errcode = gegl_clSetKernelArg(
                               cl_data->kernel[0], 10, sizeof(cl_float), (void*)&cl_offset_y));


    CL_SAFE_CALL(errcode = gegl_clEnqueueNDRangeKernel(
                               gegl_cl_get_command_queue(), cl_data->kernel[0],
                               2, NULL,
                               gbl_size, NULL,
                               0, NULL, NULL));

    errcode = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
    if (CL_SUCCESS != errcode) CL_ERROR;

    ///////////////////////////////////////////////////////////////////////////
    /* AreaFilter general processing flow.
       Making the necessary color space conversion and Saving data. */
#include "gegl-cl-operation-area-filter-fw2.h"
}
Пример #16
0
static cl_int
cl_motion_blur (cl_mem                in_tex,
                cl_mem                out_tex,
                size_t                global_worksize,
                const GeglRectangle  *roi,
                const GeglRectangle  *src_rect,
                gint                  num_steps,
                gfloat                offset_x,
                gfloat                offset_y)
{
  cl_int cl_err = 0;
  size_t global_ws[2];

  if (!cl_data)
  {
    const char *kernel_name[] = {"motion_blur_CL", NULL};
    cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
  }

  if (!cl_data) return 1;

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

  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0],  0, sizeof(cl_mem),   (void*)&in_tex);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0],  1, sizeof(cl_int),   (void*)&src_rect->width);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0],  2, sizeof(cl_int),   (void*)&src_rect->height);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0],  3, sizeof(cl_int),   (void*)&src_rect->x);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0],  4, sizeof(cl_int),   (void*)&src_rect->y);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0],  5, sizeof(cl_mem),   (void*)&out_tex);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0],  6, sizeof(cl_int),   (void*)&roi->x);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0],  7, sizeof(cl_int),   (void*)&roi->y);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0],  8, sizeof(cl_int),   (void*)&num_steps);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0],  9, sizeof(cl_float), (void*)&offset_x);
  cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 10, sizeof(cl_float), (void*)&offset_y);
  if (cl_err != CL_SUCCESS) return cl_err;

  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
                                       cl_data->kernel[0], 2,
                                       NULL, global_ws, NULL,
                                       0, NULL, NULL);
  if (cl_err != CL_SUCCESS) return cl_err;

  return cl_err;
}
Пример #17
0
static gboolean
cl_process (GeglOperation       *operation,
            cl_mem               in_tex,
            cl_mem               out_tex,
            size_t               global_worksize,
            const GeglRectangle *roi,
            gint                 level)
{
  GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
  gfloat      scale;
  gfloat      radius0, radius1;
  gint        roi_x, roi_y,x;
  gint        midx, midy;
  GeglRectangle *bounds = gegl_operation_source_get_bounding_box (operation, "input");

  gfloat length = hypot (bounds->width, bounds->height)/2;
  gfloat rdiff;
  gfloat cost, sint;
  gfloat      color[4];

  scale = bounds->width / (1.0 * bounds->height);
  scale = scale * (o->proportion) + 1.0 * (1.0-o->proportion);
  scale *= aspect_to_scale (o->squeeze);
  length = (bounds->width/2.0);

  if (scale > 1.0)
    length /= scale;

  gegl_color_get_pixel (o->color, babl_format ("RGBA float"), color);

  for (x=0; x<3; x++)   /* premultiply */
    color[x] *= color[3];

  radius0 = o->radius * (1.0-o->softness);
  radius1 = o->radius;
  rdiff = radius1-radius0;
  if (fabs (rdiff) < 0.0001)
    rdiff = 0.0001;

  midx = bounds->x + bounds->width * o->x;
  midy = bounds->y + bounds->height * o->y;

  roi_x = roi->x;
  roi_y = roi->y;

  /* constant for all pixels */
  cost = cos(-o->rotation * (G_PI*2/360.0));
  sint = sin(-o->rotation * (G_PI*2/360.0));

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

  {
  const size_t gbl_size[2] = {roi->width, roi->height};

  gint   shape = (gint) o->shape;
  gfloat gamma = o->gamma;

  cl_int cl_err = 0;
  cl_float4 f_color;
  f_color.s[0] = color[0];
  f_color.s[1] = color[1];
  f_color.s[2] = color[2];
  f_color.s[3] = color[3];

  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0,  sizeof(cl_mem),   (void*)&in_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1,  sizeof(cl_mem),   (void*)&out_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2,  sizeof(cl_float4),(void*)&f_color);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3,  sizeof(cl_float), (void*)&scale);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 4,  sizeof(cl_float), (void*)&cost);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 5,  sizeof(cl_float), (void*)&sint);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 6,  sizeof(cl_int),   (void*)&roi_x);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 7,  sizeof(cl_int),   (void*)&roi_y);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 8,  sizeof(cl_int),   (void*)&midx);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 9,  sizeof(cl_int),   (void*)&midy);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 10, sizeof(cl_int),   (void*)&shape);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 11, sizeof(cl_float), (void*)&gamma);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 12, sizeof(cl_float), (void*)&length);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 13, sizeof(cl_float), (void*)&radius0);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 14, sizeof(cl_float), (void*)&rdiff);
  CL_CHECK;

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

  return  FALSE;

error:
  return TRUE;
}
Пример #18
0
gboolean
gegl_operation_cl_set_kernel_args (GeglOperation *operation,
                                   cl_kernel      kernel,
                                   gint          *p,
                                   cl_int        *err)
{
  GParamSpec **self;
  GParamSpec **parent;
  guint n_self;
  guint n_parent;
  gint prop_no;

  self = g_object_class_list_properties (
            G_OBJECT_CLASS (g_type_class_ref (G_OBJECT_CLASS_TYPE (GEGL_OPERATION_GET_CLASS(operation)))),
            &n_self);

  parent = g_object_class_list_properties (
            G_OBJECT_CLASS (g_type_class_ref (GEGL_TYPE_OPERATION)),
            &n_parent);

  for (prop_no=0;prop_no<n_self;prop_no++)
    {
      gint parent_no;
      gboolean found=FALSE;
      for (parent_no=0;parent_no<n_parent;parent_no++)
        if (self[prop_no]==parent[parent_no])
          found=TRUE;
      /* only print properties if we are an addition compared to
       * GeglOperation
       */

      /* Removing pads */
      if (!strcmp(g_param_spec_get_name (self[prop_no]), "input") ||
          !strcmp(g_param_spec_get_name (self[prop_no]), "output") ||
          !strcmp(g_param_spec_get_name (self[prop_no]), "aux"))
        continue;

      if (!found)
        {
          if (g_type_is_a (G_PARAM_SPEC_VALUE_TYPE (self[prop_no]), G_TYPE_DOUBLE))
            {
              gdouble value;
              cl_float v;

              g_object_get (G_OBJECT (operation), g_param_spec_get_name (self[prop_no]), &value, NULL);

              v = value;
              *err = gegl_clSetKernelArg(kernel, (*p)++, sizeof(cl_float), (void*)&v);
            }
          else if (g_type_is_a (G_PARAM_SPEC_VALUE_TYPE (self[prop_no]), G_TYPE_FLOAT))
            {
              gfloat value;
              cl_float v;

              g_object_get (G_OBJECT (operation), g_param_spec_get_name (self[prop_no]), &value, NULL);

              v = value;
              *err = gegl_clSetKernelArg(kernel, (*p)++, sizeof(cl_float), (void*)&v);
            }
          else if (g_type_is_a (G_PARAM_SPEC_VALUE_TYPE (self[prop_no]), G_TYPE_INT))
            {
              gint value;
              cl_int v;

              g_object_get (G_OBJECT (operation), g_param_spec_get_name (self[prop_no]), &value, NULL);

              v = value;
              *err = gegl_clSetKernelArg(kernel, (*p)++, sizeof(cl_int), (void*)&v);
            }
          else if (g_type_is_a (G_PARAM_SPEC_VALUE_TYPE (self[prop_no]), G_TYPE_BOOLEAN))
            {
              gboolean value;
              cl_bool v;

              g_object_get (G_OBJECT (operation), g_param_spec_get_name (self[prop_no]), &value, NULL);

              v = value;
              *err = gegl_clSetKernelArg(kernel, (*p)++, sizeof(cl_bool), (void*)&v);
            }
          else
            {
              g_error ("Unsupported OpenCL kernel argument");
              return FALSE;
            }
        }
    }

  if (self)
    g_free (self);
  if (parent)
    g_free (parent);

  return TRUE;
}
Пример #19
0
static gboolean
cl_buffer_get_min_max (cl_mem               in_tex,
                       size_t               global_worksize,
                       const GeglRectangle *roi,
                       gfloat              min[4],
                       gfloat              max[4])
{
  cl_int   cl_err      = 0;
  size_t   local_ws, max_local_ws;
  size_t   work_groups;
  size_t   global_ws;
  cl_mem   cl_aux_min  = NULL;
  cl_mem   cl_aux_max  = NULL;
  cl_mem   cl_min_max  = NULL;
  cl_int   n_pixels    = (cl_int)global_worksize;
  cl_float4 min_max_buf[2];

  if (global_worksize < 1)
    {
      min[0] = min[1] = min[2] = min[3] =  G_MAXFLOAT;
      max[0] = max[1] = max[2] = max[3] = -G_MAXFLOAT;
      return FALSE;
    }

  cl_err = gegl_clGetDeviceInfo (gegl_cl_get_device (),
                                 CL_DEVICE_MAX_WORK_GROUP_SIZE,
                                 sizeof (size_t), &max_local_ws, NULL);
  CL_CHECK;

  max_local_ws = MIN (max_local_ws,
                      MIN (cl_data->work_group_size[0],
                           cl_data->work_group_size[1]));

  /* Needs to be a power of two */
  local_ws = 256;
  while (local_ws > max_local_ws)
    local_ws /= 2;

  work_groups = MIN ((global_worksize + local_ws - 1) / local_ws, local_ws);
  global_ws = work_groups * local_ws;


  cl_aux_min = gegl_clCreateBuffer (gegl_cl_get_context (),
                                    CL_MEM_READ_WRITE,
                                    local_ws * sizeof(cl_float4),
                                    NULL, &cl_err);
  CL_CHECK;
  cl_aux_max = gegl_clCreateBuffer (gegl_cl_get_context (),
                                    CL_MEM_READ_WRITE,
                                    local_ws * sizeof(cl_float4),
                                    NULL, &cl_err);
  CL_CHECK;
  cl_min_max = gegl_clCreateBuffer (gegl_cl_get_context (),
                                    CL_MEM_WRITE_ONLY,
                                    2 * sizeof(cl_float4),
                                    NULL, &cl_err);
  CL_CHECK;

  /* The full initialization is done in the two_stages_local_min_max_reduce
     kernel */
#if 0
  cl_err = gegl_clSetKernelArg(cl_data->kernel[3], 0, sizeof(cl_mem),
                               (void*)&cl_aux_min);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[3], 1, sizeof(cl_mem),
                               (void*)&cl_aux_max);
  CL_CHECK;

  cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
                                       cl_data->kernel[3], 1,
                                       NULL, &local_ws, &local_ws,
                                       0, NULL, NULL);
  CL_CHECK;
#endif

  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem),
                               (void*)&in_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem),
                               (void*)&cl_aux_min);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_mem),
                               (void*)&cl_aux_max);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3,
                               sizeof(cl_float4) * local_ws, NULL);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 4,
                               sizeof(cl_float4) * local_ws, NULL);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_int),
                               (void*)&n_pixels);
  CL_CHECK;

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

  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem),
                               (void*)&cl_aux_min);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_mem),
                               (void*)&cl_aux_max);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 2, sizeof(cl_mem),
                               (void*)&cl_min_max);
  CL_CHECK;

  /* Only one work group */
  cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
                                        cl_data->kernel[1], 1,
                                        NULL, &local_ws, &local_ws,
                                        0, NULL, NULL);
  CL_CHECK;

  /* Read the memory buffer, probably better to keep it in GPU memory */
  cl_err = gegl_clEnqueueReadBuffer (gegl_cl_get_command_queue (),
                                     cl_min_max, CL_TRUE, 0,
                                     2 * sizeof (cl_float4), &min_max_buf, 0,
                                     NULL, NULL);
  CL_CHECK;

  min[0] = min_max_buf[0].x;
  min[1] = min_max_buf[0].y;
  min[2] = min_max_buf[0].z;
  min[3] = min_max_buf[0].w;

  max[0] = min_max_buf[1].x;
  max[1] = min_max_buf[1].y;
  max[2] = min_max_buf[1].z;
  max[3] = min_max_buf[1].w;

  cl_err = gegl_clReleaseMemObject (cl_aux_min);
  CL_CHECK_ONLY (cl_err);

  cl_err = gegl_clReleaseMemObject (cl_aux_max);
  CL_CHECK_ONLY (cl_err);

  cl_err = gegl_clReleaseMemObject (cl_min_max);
  CL_CHECK_ONLY (cl_err);

  return FALSE;

error:
  if (cl_aux_min)
    gegl_clReleaseMemObject (cl_aux_min);
  if (cl_aux_max)
    gegl_clReleaseMemObject (cl_aux_max);
  if (cl_min_max)
    gegl_clReleaseMemObject (cl_min_max);
  return TRUE;
}
static gboolean
gegl_operation_point_filter_cl_process (GeglOperation       *operation,
                                        GeglBuffer          *input,
                                        GeglBuffer          *output,
                                        const GeglRectangle *result,
                                        gint                 level)
{
  const Babl *in_format  = gegl_operation_get_format (operation, "input");
  const Babl *out_format = gegl_operation_get_format (operation, "output");

  GeglOperationClass *operation_class = GEGL_OPERATION_GET_CLASS (operation);
  GeglOperationPointFilterClass *point_filter_class = GEGL_OPERATION_POINT_FILTER_GET_CLASS (operation);

  GeglBufferClIterator *iter = NULL;

  cl_int cl_err = 0;
  gboolean err;

  /* non-texturizable format! */
  if (!gegl_cl_color_babl (in_format,  NULL) ||
      !gegl_cl_color_babl (out_format, NULL))
    {
      GEGL_NOTE (GEGL_DEBUG_OPENCL, "Non-texturizable format!");
      return FALSE;
    }

  GEGL_NOTE (GEGL_DEBUG_OPENCL, "GEGL_OPERATION_POINT_FILTER: %s", operation_class->name);

  /* Process */
  iter = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE);

  gegl_buffer_cl_iterator_add (iter, input, result, in_format, GEGL_CL_BUFFER_READ, GEGL_ABYSS_NONE);

  while (gegl_buffer_cl_iterator_next (iter, &err))
    {
      if (err)
        return FALSE;

      if (point_filter_class->cl_process)
        {
          err = point_filter_class->cl_process (operation, iter->tex[1], iter->tex[0],
                                                iter->size[0], &iter->roi[0], level);
          if (err)
            {
              GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", operation_class->name);
              gegl_buffer_cl_iterator_stop (iter);
              return FALSE;
            }
        }
      else if (operation_class->cl_data)
        {
          gint p = 0;
          GeglClRunData *cl_data = operation_class->cl_data;

          cl_err = gegl_clSetKernelArg (cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&iter->tex[1]);
          CL_CHECK;
          cl_err = gegl_clSetKernelArg (cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&iter->tex[0]);
          CL_CHECK;

          gegl_operation_cl_set_kernel_args (operation, cl_data->kernel[0], &p, &cl_err);
          CL_CHECK;

          cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
                                                cl_data->kernel[0], 1,
                                                NULL, &iter->size[0], NULL,
                                                0, NULL, NULL);
          CL_CHECK;
        }
      else
        {
          g_warning ("OpenCL support enabled, but no way to execute");
          gegl_buffer_cl_iterator_stop (iter);
          return FALSE;
        }
    }

  return TRUE;

error:
  GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", gegl_cl_errstring (cl_err));
  if (iter)
    gegl_buffer_cl_iterator_stop (iter);
  return FALSE;
}
Пример #21
0
static gboolean
cl_box_max (cl_mem                in_tex,
            cl_mem                aux_tex,
            cl_mem                out_tex,
            size_t                global_worksize,
            const GeglRectangle  *roi,
            gint                  radius)
{
  cl_int cl_err = 0;
  size_t global_ws_hor[2], global_ws_ver[2];
  size_t local_ws_hor[2], local_ws_ver[2];

  if (!cl_data)
    {
      const char *kernel_name[] = {"kernel_max_hor", "kernel_max_ver", NULL};
      cl_data = gegl_cl_compile_and_build (box_max_cl_source, kernel_name);
    }
  if (!cl_data) return TRUE;

  local_ws_hor[0] = 1;
  local_ws_hor[1] = 256;
  global_ws_hor[0] = roi->height + 2 * radius;
  global_ws_hor[1] = ((roi->width + local_ws_hor[1] -1)/local_ws_hor[1]) * local_ws_hor[1];

  local_ws_ver[0] = 1;
  local_ws_ver[1] = 256;
  global_ws_ver[0] = roi->height;
  global_ws_ver[1] = ((roi->width + local_ws_ver[1] -1)/local_ws_ver[1]) * local_ws_ver[1];

  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem),   (void*)&in_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem),   (void*)&aux_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_int),   (void*)&roi->width);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_int),   (void*)&radius);
  CL_CHECK;

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

  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem),   (void*)&aux_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_mem),   (void*)&out_tex);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 2, sizeof(cl_int),   (void*)&roi->width);
  CL_CHECK;
  cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 3, sizeof(cl_int),   (void*)&radius);
  CL_CHECK;

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

  return FALSE;

error:
  return TRUE;
}