Example #1
0
static gboolean
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_in[2];
  size_t global_ws_aux[2];

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

  global_ws_in[0] = roi->width  + LAPLACE_RADIUS;
  global_ws_in[1] = roi->height + LAPLACE_RADIUS;

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

  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
                                    sizeof (cl_mem), &in_tex,
                                    sizeof (cl_mem), &aux_tex,
                                    NULL);
  CL_CHECK;

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


  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[1],
                                    sizeof (cl_mem), &aux_tex,
                                    sizeof (cl_mem), &out_tex,
                                    NULL);
  CL_CHECK;

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

  return FALSE;

 error:
  return TRUE;
}
Example #2
0
static gboolean
cl_process (GeglOperation       *operation,
            cl_mem               in_tex,
            cl_mem               aux_tex,
            cl_mem               out_tex,
            size_t               global_worksize,
            const GeglRectangle *roi,
            gint                 level)
{
  GeglOperationClass *operation_class = GEGL_OPERATION_GET_CLASS (operation);
  cl_int cl_err = 0;

  /* The kernel will have been compiled by our parent class */
  if (!operation_class->cl_data)
    return TRUE;

  cl_err = gegl_cl_set_kernel_args (operation_class->cl_data->kernel[0],
                                    sizeof(cl_mem), &in_tex,
                                    sizeof(cl_mem), &aux_tex,
                                    sizeof(cl_mem), &out_tex,
                                    NULL);
  CL_CHECK;

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

  return FALSE;

error:
  return TRUE;
}
Example #3
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;
}
Example #4
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;
}
Example #5
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);
  cl_float3   freq;
  cl_float3   phaseshift;
  cl_int3     keep;
  cl_int      cl_err = 0;

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

  if (!cl_data)
    return TRUE;

  freq.s[0] = o->cpn_1_frequency * G_PI;
  freq.s[1] = o->cpn_2_frequency * G_PI;
  freq.s[2] = o->cpn_3_frequency * G_PI;

  phaseshift.s[0] = G_PI * o->cpn_1_phaseshift / 180.0;
  phaseshift.s[1] = G_PI * o->cpn_2_phaseshift / 180.0;
  phaseshift.s[2] = G_PI * o->cpn_3_phaseshift / 180.0;

  keep.s[0] = (cl_int)o->cpn_1_keep;
  keep.s[1] = (cl_int)o->cpn_2_keep;
  keep.s[2] = (cl_int)o->cpn_3_keep;

  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
                                    sizeof(cl_mem),    &in,
                                    sizeof(cl_mem),    &out,
                                    sizeof(cl_float3), &freq,
                                    sizeof(cl_float3), &phaseshift,
                                    sizeof(cl_int3),   &keep,
                                    NULL);
  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;
}
/* 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;
}
Example #7
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);
  CeParamsType   *params = (CeParamsType*) o->user_data;
  cl_float3   color_diff;
  cl_float3   min;
  cl_float3   max;
  cl_int      cl_err = 0;
  gint        i;

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

  if (!cl_data)
    return TRUE;

  for (i = 0; i < 3; i++)
    {
      color_diff.s[i] = params->color_diff[i];
      min.s[i] = params->min[i];
      max.s[i] = params->max[i];
    }

  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
                                    sizeof(cl_mem),    &in,
                                    sizeof(cl_mem),    &out,
                                    sizeof(cl_float3), &color_diff,
                                    sizeof(cl_float3), &min,
                                    sizeof(cl_float3), &max,
                                    NULL);
  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;
}
Example #8
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;
}
Example #9
0
static gboolean
cl_gaussian_blur (cl_mem                 in_tex,
                  cl_mem                 out_tex,
                  const GeglRectangle   *roi,
                  cl_mem                 cl_cmatrix,
                  gint                   clen,
                  GeglOrientation        orientation)
{
  cl_int cl_err = 0;
  size_t global_ws[2];
  gint   kernel_num;

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

  if (!cl_data)
    return TRUE;

  if (orientation == GEGL_ORIENTATION_VERTICAL)
    kernel_num = 0;
  else
    kernel_num = 1;

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

  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[kernel_num],
                                    sizeof(cl_mem), (void*)&in_tex,
                                    sizeof(cl_mem), (void*)&out_tex,
                                    sizeof(cl_mem), (void*)&cl_cmatrix,
                                    sizeof(cl_int), (void*)&clen,
                                    NULL);
  CL_CHECK;

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

  cl_err = gegl_clFinish (gegl_cl_get_command_queue ());
  CL_CHECK;

  return FALSE;

error:
  return TRUE;
}
Example #10
0
static gboolean
cl_process (GeglOperation       *operation,
            cl_mem               in_buf,
            cl_mem               out_buf,
            const size_t         n_pixels,
            const GeglRectangle *roi,
            gint                 level)
{
  GeglOperationClass *operation_class = GEGL_OPERATION_GET_CLASS (operation);
  GeglClRunData *cl_data = operation_class->cl_data;
  GeglProperties *o = GEGL_PROPERTIES (operation);
  const size_t gbl_size[2] = {roi->width, roi->height};
  const size_t gbl_off[2]  = {roi->x, roi->y};
  cl_int cl_err = 0;
  cl_mem filter_pat = NULL;

  if (!cl_data)
    goto error;
  filter_pat = gegl_clCreateBuffer (gegl_cl_get_context (),
                                    CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                    pattern_width[o->pattern] *
                                    pattern_height[o->pattern] * sizeof(cl_int),
                                    (void*)pattern[o->pattern],
                                    &cl_err);
  CL_CHECK;
  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
                                    sizeof(cl_mem), &in_buf,
                                    sizeof(cl_mem), &out_buf,
                                    sizeof(cl_mem), &filter_pat,
                                    sizeof(cl_int), &pattern_width[o->pattern],
                                    sizeof(cl_int), &pattern_height[o->pattern],
                                    sizeof(cl_int), &o->additive,
                                    sizeof(cl_int), &o->rotated,
                                    NULL);
  CL_CHECK;
  cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
                                        cl_data->kernel[0], 2,
                                        gbl_off, gbl_size, NULL,
                                        0, NULL, NULL);
  CL_CHECK;
  cl_err = gegl_clFinish (gegl_cl_get_command_queue ());
  CL_CHECK;
  cl_err = gegl_clReleaseMemObject (filter_pat);
  CL_CHECK;
  return FALSE;
  error:
    if (filter_pat)
      gegl_clReleaseMemObject (filter_pat);
    return TRUE;
}
Example #11
0
static gboolean
checkerboard_cl_process (GeglOperation       *operation,
                         cl_mem               out_tex,
                         size_t               global_worksize,
                         const GeglRectangle *roi,
                         gint                 level)
{
  GeglProperties *o         = GEGL_PROPERTIES (operation);
  const Babl   *out_format  = gegl_operation_get_format (operation, "output");
  const size_t  gbl_size[2] = {roi->width, roi->height};
  const size_t  gbl_offs[2] = {roi->x, roi->y};
  cl_int        cl_err      = 0;
  float         color1[4];
  float         color2[4];

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

    if (!cl_data)
      return TRUE;
  }

  gegl_color_get_pixel (o->color1, out_format, color1);
  gegl_color_get_pixel (o->color2, out_format, color2);

  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
                                    sizeof(cl_mem), &out_tex,
                                    sizeof(color1), &color1,
                                    sizeof(color2), &color2,
                                    sizeof(cl_int), &o->x,
                                    sizeof(cl_int), &o->y,
                                    sizeof(cl_int), &o->x_offset,
                                    sizeof(cl_int), &o->y_offset,
                                    NULL);
  CL_CHECK;

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

  return FALSE;
error:
  return TRUE;
}
Example #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;
}
Example #13
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;
}
Example #14
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;
}
Example #15
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;
}
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;
}
Example #17
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;
}
Example #18
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;
}
Example #19
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);
  cl_float   threshold    = o->threshold;


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

  if (!cl_data)
    return TRUE;

  {
  cl_int cl_err = 0;

  gegl_cl_set_kernel_args (cl_data->kernel[0],
                           sizeof(cl_mem),   &in,
                           sizeof(cl_mem),   &out,
                           sizeof(cl_float), &threshold,
                           NULL);
  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;
}
Example #20
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;
}
Example #21
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;
}
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;
}
Example #23
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;
}
Example #24
0
static gboolean
cl_box_blur (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], global_ws[2];
  size_t local_ws_hor[2], local_ws_ver[2], local_ws[2];
  size_t step_size ;
  if (!cl_data)
    {
      const char *kernel_name[] = { "kernel_blur_hor", "kernel_blur_ver","kernel_box_blur_fast", NULL};
      cl_data = gegl_cl_compile_and_build (box_blur_cl_source, kernel_name);
    }

  if (!cl_data)
    return TRUE;
  step_size = 64;
  local_ws[0]=256;
  local_ws[1]=1;


  if( radius <=110 )
  {
    global_ws[0] = (roi->width + local_ws[0] - 2 * radius - 1) / ( local_ws[0] - 2 * radius ) * local_ws[0];
    global_ws[1] = (roi->height + step_size - 1) / step_size;
    cl_err = gegl_cl_set_kernel_args(cl_data->kernel[2],
                                     sizeof(cl_mem), (void *)&in_tex,
                                     sizeof(cl_mem), (void *)&out_tex,
                                     sizeof(cl_float4)*local_ws[0], (void *)NULL,
                                     sizeof(cl_int), (void *)&roi->width,
                                     sizeof(cl_int), (void *)&roi->height,
                                     sizeof(cl_int), (void *)&radius,
                                     sizeof(cl_int), (void *)&step_size, NULL);
    CL_CHECK;
    cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(),
                                         cl_data->kernel[2], 2,
                                         NULL, global_ws, local_ws, 0, NULL, NULL );
       CL_CHECK;

  }
  else
  {
    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_cl_set_kernel_args (cl_data->kernel[0],
                                      sizeof(cl_mem), (void*)&in_tex,
                                      sizeof(cl_mem), (void*)&aux_tex,
                                      sizeof(cl_int), (void*)&roi->width,
                                      sizeof(cl_int), (void*)&radius,
                                      NULL);
    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_cl_set_kernel_args (cl_data->kernel[1],
                                      sizeof(cl_mem), (void*)&aux_tex,
                                      sizeof(cl_mem), (void*)&out_tex,
                                      sizeof(cl_int), (void*)&roi->width,
                                      sizeof(cl_int), (void*)&radius,
                                      NULL);
    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;
}
Example #25
0
static gboolean
cl_pixelize (cl_mem               in_tex,
             cl_mem               aux_tex,
             cl_mem               out_tex,
             const GeglRectangle *src_rect,
             const GeglRectangle *roi,
             gint                 xsize,
             gint                 ysize,
             gfloat               xratio,
             gfloat               yratio,
             gfloat               bg_color[4],
             gint                 norm,
             GeglRectangle       *image_extent)
{
  cl_int cl_err = 0;
  const size_t gbl_size[2]= {roi->width, roi->height};

  gint cx0 = block_index (roi->x, xsize);
  gint cy0 = block_index (roi->y, ysize);
  gint block_count_x = block_index (roi->x + roi->width  + xsize - 1, xsize) - cx0;
  gint block_count_y = block_index (roi->y + roi->height + ysize - 1, ysize) - cy0;
  cl_int4 bbox = {{ image_extent->x, image_extent->y,
                    image_extent->x + image_extent->width,
                    image_extent->y + image_extent->height }};

  cl_int line_width = roi->width + 2 * xsize;

  size_t gbl_size_tmp[2] = {block_count_x, block_count_y};

  if (!cl_data)
  {
    const char *kernel_name[] = {"calc_block_color", "kernel_pixelize", NULL};
    cl_data = gegl_cl_compile_and_build (pixelize_cl_source, kernel_name);
  }

  if (!cl_data) return 1;

  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
                                    sizeof(cl_mem), (void*)&in_tex,
                                    sizeof(cl_mem), (void*)&aux_tex,
                                    sizeof(cl_int), (void*)&xsize,
                                    sizeof(cl_int), (void*)&ysize,
                                    sizeof(cl_int), (void*)&roi->x,
                                    sizeof(cl_int), (void*)&roi->y,
                                    sizeof(cl_int4), &bbox,
                                    sizeof(cl_int), (void*)&line_width,
                                    sizeof(cl_int), (void*)&block_count_x,
                                    NULL);
  CL_CHECK;

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

  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[1],
                                    sizeof(cl_mem),   (void*)&aux_tex,
                                    sizeof(cl_mem),   (void*)&out_tex,
                                    sizeof(cl_int),   (void*)&xsize,
                                    sizeof(cl_int),   (void*)&ysize,
                                    sizeof(cl_float), (void*)&xratio,
                                    sizeof(cl_float), (void*)&yratio,
                                    sizeof(cl_int),   (void*)&roi->x,
                                    sizeof(cl_int),   (void*)&roi->y,
                                    sizeof(cl_float4),(void*)bg_color,
                                    sizeof(cl_int),   (void*)&norm,
                                    sizeof(cl_int),   (void*)&block_count_x,
                                    NULL);
  CL_CHECK;

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

  return FALSE;

error:
  return TRUE;
}
Example #26
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;
}
Example #27
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"
}
Example #28
0
static gboolean
cl_bilateral (cl_mem                in_tex,
              cl_mem                out_tex,
              const GeglRectangle  *roi,
              const GeglRectangle  *src_rect,
              gint                  s_sigma,
              gfloat                r_sigma)
{
  cl_int cl_err = 0;

  gint c;

  const gint width    = src_rect->width;
  const gint height   = src_rect->height;

  const gint sw = (width -1) / s_sigma + 1;
  const gint sh = (height-1) / s_sigma + 1;
  const gint depth = (int)(1.0f / r_sigma) + 1;

  size_t global_ws[2];
  size_t local_ws[2];

  cl_mem grid = NULL;
  cl_mem blur[4] = {NULL, NULL, NULL, NULL};

  if (!cl_data)
    {
      const char *kernel_name[] = {"bilateral_downsample",
                                   "bilateral_blur",
                                   "bilateral_interpolate",
                                   NULL};
      cl_data = gegl_cl_compile_and_build (bilateral_filter_fast_cl_source, kernel_name);
    }

  if (!cl_data)
    return 1;


  grid = gegl_clCreateBuffer (gegl_cl_get_context (),
                              CL_MEM_READ_WRITE,
                              sw * sh * depth * sizeof(cl_float8),
                              NULL,
                              &cl_err);
  CL_CHECK;

  for(c = 0; c < 4; c++)
    {
      blur[c] = gegl_clCreateBuffer (gegl_cl_get_context (),
                                     CL_MEM_READ_WRITE,
                                     sw * sh * depth * sizeof(cl_float2),
                                     NULL, &cl_err);
      CL_CHECK;
    }

  local_ws[0] = 8;
  local_ws[1] = 8;

  global_ws[0] = ((sw + local_ws[0] - 1)/local_ws[0])*local_ws[0];
  global_ws[1] = ((sh + local_ws[1] - 1)/local_ws[1])*local_ws[1];

  gegl_cl_set_kernel_args (cl_data->kernel[0],
                           sizeof(cl_mem),   &in_tex,
                           sizeof(cl_mem),   &grid,
                           sizeof(cl_int),   &width,
                           sizeof(cl_int),   &height,
                           sizeof(cl_int),   &sw,
                           sizeof(cl_int),   &sh,
                           sizeof(cl_int),   &depth,
                           sizeof(cl_int),   &s_sigma,
                           sizeof(cl_float), &r_sigma,
                           NULL);
  CL_CHECK;

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

  local_ws[0] = 16;
  local_ws[1] = 16;

  global_ws[0] = ((sw + local_ws[0] - 1)/local_ws[0])*local_ws[0];
  global_ws[1] = ((sh + local_ws[1] - 1)/local_ws[1])*local_ws[1];

  gegl_cl_set_kernel_args (cl_data->kernel[1],
                           sizeof(cl_mem), &grid,
                           sizeof(cl_mem), &blur[0],
                           sizeof(cl_mem), &blur[1],
                           sizeof(cl_mem), &blur[2],
                           sizeof(cl_mem), &blur[3],
                           sizeof(cl_int), &sw,
                           sizeof(cl_int), &sh,
                           sizeof(cl_int), &depth,
                           NULL);

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

  global_ws[0] = width;
  global_ws[1] = height;

  gegl_cl_set_kernel_args (cl_data->kernel[2],
                           sizeof(cl_mem),   &in_tex,
                           sizeof(cl_mem),   &blur[0],
                           sizeof(cl_mem),   &blur[1],
                           sizeof(cl_mem),   &blur[2],
                           sizeof(cl_mem),   &blur[3],
                           sizeof(cl_mem),   &out_tex,
                           sizeof(cl_int),   &width,
                           sizeof(cl_int),   &sw,
                           sizeof(cl_int),   &sh,
                           sizeof(cl_int),   &depth,
                           sizeof(cl_int),   &s_sigma,
                           sizeof(cl_float), &r_sigma,
                           NULL);
  CL_CHECK;

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

  cl_err = gegl_clFinish (gegl_cl_get_command_queue ());
  CL_CHECK;

  cl_err = gegl_clReleaseMemObject (grid);
  CL_CHECK_ONLY (cl_err);

  for(c = 0; c < 4; c++)
    {
      cl_err = gegl_clReleaseMemObject (blur[c]);
      CL_CHECK_ONLY (cl_err);
    }

  return FALSE;

error:
  if (grid)
    gegl_clReleaseMemObject (grid);

  for (c = 0; c < 4; c++)
    {
      if (blur[c])
        gegl_clReleaseMemObject (blur[c]);
    }

  return TRUE;
}
Example #29
0
static gboolean
cl_process (GeglOperation       *operation,
            cl_mem               out_tex,
            const GeglRectangle *roi)
{
  GeglProperties   *o      = GEGL_PROPERTIES (operation);
  const size_t  gbl_size[] = { roi->width, roi->height };
  cl_int        cl_err     = 0;
  cl_int offset_x;
  cl_int offset_y;
  cl_int width;
  cl_int height;
  cl_float3 sedges;
  cl_float3 contours;
  cl_float3 frequency;
  cl_float brightness;
  cl_float polarization;
  cl_float scattering;
  cl_int iterations;
  cl_float weird_factor;

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

      if (!cl_data)
        return TRUE;
    }

  offset_x = roi->x;
  offset_y = roi->y;
  width = o->width;
  height = o->height;
  sedges.s[0] = o->red_sedges;
  sedges.s[1] = o->green_sedges;
  sedges.s[2] = o->blue_sedges;
  contours.s[0] = o->red_contours;
  contours.s[1] = o->green_contours;
  contours.s[2] = o->blue_contours;
  frequency.s[0] = o->red_frequency;
  frequency.s[1] = o->green_frequency;
  frequency.s[2] = o->blue_frequency;
  brightness = o->brightness;
  polarization = o->polarization;
  scattering = o->scattering;
  iterations = ITERATIONS;
  weird_factor = WEIRD_FACTOR;
  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
                                    sizeof(cl_mem),    &out_tex,
                                    sizeof(cl_int),    &offset_x,
                                    sizeof(cl_int),    &offset_y,
                                    sizeof(cl_int),    &width,
                                    sizeof(cl_int),    &height,
                                    sizeof(cl_float3), &sedges,
                                    sizeof(cl_float3), &contours,
                                    sizeof(cl_float3), &frequency,
                                    sizeof(cl_float),  &brightness,
                                    sizeof(cl_float),  &polarization,
                                    sizeof(cl_float),  &scattering,
                                    sizeof(cl_int),    &iterations,
                                    sizeof(cl_float),  &weird_factor,
                                    NULL);
  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;

  cl_err = gegl_clFinish (gegl_cl_get_command_queue ());
  CL_CHECK;

  return FALSE;

error:
  return TRUE;
}
Example #30
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;
    }
}