gegl-cl-color.c 6.32 KB
Newer Older
1 2 3 4 5 6 7 8
#include "gegl.h"
#include "gegl-cl-color.h"
#include "gegl-cl-init.h"

#include "gegl-cl-color-kernel.h"

static gegl_cl_run_data *kernels_color = NULL;

9 10 11
#define CL_FORMAT_N 8

static const Babl *format[CL_FORMAT_N];
12 13 14 15 16 17 18 19 20 21

void
gegl_cl_color_compile_kernels(void)
{
  const char *kernel_name[] = {"non_premultiplied_to_premultiplied", /* 0 */
                               "premultiplied_to_non_premultiplied", /* 1 */
                               "rgba2rgba_gamma_2_2",                /* 2 */
                               "rgba_gamma_2_22rgba",                /* 3 */
                               "rgba2rgba_gamma_2_2_premultiplied",  /* 4 */
                               "rgba_gamma_2_2_premultiplied2rgba",  /* 5 */
22 23
                               "rgbaf_to_rgbau8",                    /* 6 */
                               "rgbau8_to_rgbaf",                    /* 7 */
24 25 26 27 28 29 30 31
                               NULL};

  format[0] = babl_format ("RaGaBaA float"),
  format[1] = babl_format ("RGBA float"),
  format[2] = babl_format ("R'G'B'A float"),
  format[3] = babl_format ("RGBA float"),
  format[4] = babl_format ("R'aG'aB'aA float"),
  format[5] = babl_format ("RGBA float"),
32 33
  format[6] = babl_format ("RGBA u8"),
  format[7] = babl_format ("RGBA float"),
34 35 36 37

  kernels_color = gegl_cl_compile_and_build (kernel_color_source, kernel_name);
}

38
gegl_cl_color_op
39 40 41 42 43 44
gegl_cl_color_supported (const Babl *in_format, const Babl *out_format)
{
  int i;
  gboolean supported_format_in  = FALSE;
  gboolean supported_format_out = FALSE;

45 46 47 48
  if (in_format == out_format)
    return CL_COLOR_EQUAL;

  for (i = 0; i < CL_FORMAT_N; i++)
49 50 51 52 53
    {
      if (format[i] == in_format)  supported_format_in  = TRUE;
      if (format[i] == out_format) supported_format_out = TRUE;
    }

54 55 56 57
  if (supported_format_in && supported_format_out)
    return CL_COLOR_CONVERT;
  else
    return CL_COLOR_NOT_SUPPORTED;
58 59 60 61 62 63 64 65
}

#define CONV_1(x)   {conv[0] = x; conv[1] = -1;}
#define CONV_2(x,y) {conv[0] = x; conv[1] =  y;}

//#define CL_ERROR {g_assert(0);}
#define CL_ERROR {g_printf("[OpenCL] Error in %s:%d@%s - %s\n", __FILE__, __LINE__, __func__, gegl_cl_errstring(errcode)); return FALSE;}

66 67
/* in_tex and aux_tex may be destroyed to keep intermediate results,
   converted result will be stored in in_tex */
68
gboolean
69
gegl_cl_color_conv (cl_mem *in_tex, cl_mem *aux_tex, const size_t size[2],
70 71 72 73 74 75
                    const Babl *in_format, const Babl *out_format)
{
  int i;
  int errcode;
  int conv[2] = {-1, -1};

76 77 78 79
  cl_mem ping_tex = *in_tex, pong_tex = *aux_tex;

  if (gegl_cl_color_supported (in_format, out_format) == CL_COLOR_NOT_SUPPORTED)
    return FALSE;
80 81 82 83 84 85 86 87

  if (in_format == out_format)
    {
      const size_t origin[3] = {0, 0, 0};
      const size_t region[3] = {size[0], size[1], 1};

      /* just copy in_tex to out_tex */
      errcode = gegl_clEnqueueCopyImage (gegl_cl_get_command_queue(),
88
                                         *in_tex, *aux_tex, origin, origin, region,
89 90 91 92 93 94 95 96 97 98 99 100 101
                                         0, NULL, NULL);
      if (errcode != CL_SUCCESS) CL_ERROR

      errcode = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
      if (errcode != CL_SUCCESS) CL_ERROR
    }
  else
    {
      if      (in_format == babl_format ("RGBA float"))
        {
          if      (out_format == babl_format ("RaGaBaA float"))    CONV_1(0)
          else if (out_format == babl_format ("R'G'B'A float"))    CONV_1(2)
          else if (out_format == babl_format ("R'aG'aB'aA float")) CONV_1(4)
102
          else if (out_format == babl_format ("RGBA u8"))          CONV_1(6)
103 104 105 106 107 108
        }
      else if (in_format == babl_format ("RaGaBaA float"))
        {
          if      (out_format == babl_format ("RGBA float"))       CONV_1(1)
          else if (out_format == babl_format ("R'G'B'A float"))    CONV_2(1, 2)
          else if (out_format == babl_format ("R'aG'aB'aA float")) CONV_2(1, 4)
109
          else if (out_format == babl_format ("RGBA u8"))          CONV_2(1, 6)
110 111 112 113 114 115
        }
      else if (in_format == babl_format ("R'G'B'A float"))
        {
          if      (out_format == babl_format ("RGBA float"))       CONV_1(3)
          else if (out_format == babl_format ("RaGaBaA float"))    CONV_2(3, 0)
          else if (out_format == babl_format ("R'aG'aB'aA float")) CONV_2(3, 4)
116
          else if (out_format == babl_format ("RGBA u8"))          CONV_2(3, 6)
117 118 119 120 121 122
        }
      else if (in_format == babl_format ("R'aG'aB'aA float"))
        {
          if      (out_format == babl_format ("RGBA float"))       CONV_1(5)
          else if (out_format == babl_format ("RaGaBaA float"))    CONV_2(5, 0)
          else if (out_format == babl_format ("R'G'B'A float"))    CONV_2(5, 2)
123 124 125 126 127 128 129 130
          else if (out_format == babl_format ("RGBA u8"))          CONV_2(5, 6)
        }
      else if (in_format == babl_format ("RGBA u8"))
        {
          if      (out_format == babl_format ("RGBA float"))       CONV_1(7)
          else if (out_format == babl_format ("RaGaBaA float"))    CONV_2(7, 0)
          else if (out_format == babl_format ("R'G'B'A float"))    CONV_2(7, 2)
          else if (out_format == babl_format ("RGBA u8"))          CONV_2(7, 6)
131 132
        }

133 134
      /* XXX: maybe there are precision problems if a 8-bit texture is used as intermediate */
      for (i=0; conv[i] >= 0 && i<2; i++)
135
        {
136 137 138 139 140 141 142 143 144 145 146 147 148 149 150 151 152 153 154 155
            cl_mem tmp_tex;

            errcode = gegl_clSetKernelArg(kernels_color->kernel[conv[i]], 0, sizeof(cl_mem), (void*)&ping_tex);
            if (errcode != CL_SUCCESS) CL_ERROR

            errcode = gegl_clSetKernelArg(kernels_color->kernel[conv[i]], 1, sizeof(cl_mem), (void*)&pong_tex);
            if (errcode != CL_SUCCESS) CL_ERROR

            errcode = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
                                                  kernels_color->kernel[conv[i]], 2,
                                                  NULL, size, NULL,
                                                  0, NULL, NULL);
            if (errcode != CL_SUCCESS) CL_ERROR

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

            tmp_tex = ping_tex;
            ping_tex = pong_tex;
            pong_tex = tmp_tex;
156
        }
157 158 159 160 161 162 163 164 165 166 167 168

      if (i % 2 == 0)
        {
          *in_tex  = ping_tex;
          *aux_tex = pong_tex;
        }
      else
        {
          *in_tex  = pong_tex;
          *aux_tex = ping_tex;
        }

169 170 171 172 173 174
    }

  return TRUE;
}

#undef CL_ERROR