gegl-cl-color.c 10 KB
Newer Older
1 2 3 4 5 6 7 8 9 10 11 12 13
/* This file is part of GEGL
 *
 * GEGL is free software; you can redistribute it and/or
 * modify it under the terms of the GNU Lesser General Public
 * License as published by the Free Software Foundation; either
 * version 3 of the License, or (at your option) any later version.
 *
 * GEGL is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 * Lesser General Public License for more details.
 *
 * You should have received a copy of the GNU Lesser General Public
14
 * License along with GEGL; if not, see <https://www.gnu.org/licenses/>.
15 16 17 18
 *
 * Copyright 2012 Victor Oliveira (victormatheus@gmail.com)
 */

Victor Oliveira's avatar
Victor Oliveira committed
19 20 21 22 23 24

/* Here we have implemented in OpenCL a subset of color space conversions provided by BABL
   that are commonly used. For now there is no support for conversions that need a intermediate
   representation (ex: A->B, B->C, C->D), just among two color spaces.
*/

25
#include "config.h"
26

27
#include "gegl.h"
28
#include "gegl/gegl-debug.h"
29
#include "gegl-cl.h"
30 31
#include "gegl-cl-color.h"

32
#include "opencl/colors.cl.h"
33
#include "opencl/colors-8bit-lut.cl.h"
34

35 36 37
static GHashTable  *color_kernels_hash = NULL;

typedef struct
38
{
39 40 41 42 43 44 45 46 47 48 49 50
  const Babl *from_fmt;
  const Babl *to_fmt;
  const char *kernel_name;
  cl_kernel   kernel;
} ColorConversionInfo;

static guint
color_kernels_hash_hashfunc (gconstpointer key)
{
  const ColorConversionInfo *ekey = key;
  return GPOINTER_TO_INT (ekey->from_fmt) ^ GPOINTER_TO_INT (ekey->to_fmt);
}
51

52 53 54 55 56 57
static gboolean
color_kernels_hash_equalfunc (gconstpointer a,
                              gconstpointer b)
{
  const ColorConversionInfo *ea = a;
  const ColorConversionInfo *eb = b;
58

59 60 61 62 63 64 65
  if (ea->from_fmt == eb->from_fmt &&
      ea->to_fmt   == eb->to_fmt)
    {
      return TRUE;
    }
  return FALSE;
}
66

67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99 100 101 102
static gboolean
gegl_cl_color_load_conversion_set (ColorConversionInfo  *conversions,
                                   gint                  num_conversions,
                                   GeglClRunData       **kernels,
                                   const gchar          *source)
{
  const char *kernel_names[num_conversions + 1];

  for (int i = 0; i < num_conversions; ++i)
    {
      kernel_names[i] = conversions[i].kernel_name;
    }

  kernel_names[num_conversions] = NULL;

  *kernels = gegl_cl_compile_and_build (source, kernel_names);

  if (!(*kernels))
    {
      return FALSE;
    }

  for (int i = 0; i < num_conversions; ++i)
    {
      ColorConversionInfo *info = g_new (ColorConversionInfo, 1);

      conversions[i].kernel = (*kernels)->kernel[i];
      *info = conversions[i];

      g_hash_table_insert (color_kernels_hash, info, info);
    }

  return TRUE;
}

gboolean
103 104 105
gegl_cl_color_compile_kernels (void)
{
  static GeglClRunData *float_kernels = NULL;
106 107
  static GeglClRunData *lut8_kernels = NULL;
  gboolean result = TRUE;
Victor Oliveira's avatar
Victor Oliveira committed
108

109 110 111
  ColorConversionInfo float_conversions[] = {
    { babl_format ("RGBA u8"), babl_format ("RGBA float"), "rgbau8_to_rgbaf", NULL },
    { babl_format ("RGBA float"), babl_format ("RGBA u8"), "rgbaf_to_rgbau8", NULL },
112

113 114 115 116
    { babl_format ("RGBA float"), babl_format("RaGaBaA float"), "rgbaf_to_ragabaf", NULL },
    { babl_format ("RaGaBaA float"), babl_format("RGBA float"), "ragabaf_to_rgbaf", NULL },
    { babl_format ("RGBA u8"), babl_format("RaGaBaA float"), "rgbau8_to_ragabaf", NULL },
    { babl_format ("RaGaBaA float"), babl_format("RGBA u8"), "ragabaf_to_rgbau8", NULL },
117

118 119 120 121
    { babl_format ("RGBA float"), babl_format("R'G'B'A float"), "rgbaf_to_rgba_gamma_f", NULL },
    { babl_format ("R'G'B'A float"), babl_format("RGBA float"), "rgba_gamma_f_to_rgbaf", NULL },
    { babl_format ("RGBA u8"), babl_format("R'G'B'A float"), "rgbau8_to_rgba_gamma_f", NULL },
    { babl_format ("R'G'B'A float"), babl_format("RGBA u8"), "rgba_gamma_f_to_rgbau8", NULL },
122

123 124 125 126
    { babl_format ("RGBA float"), babl_format("Y'CbCrA float"), "rgbaf_to_ycbcraf", NULL },
    { babl_format ("Y'CbCrA float"), babl_format("RGBA float"), "ycbcraf_to_rgbaf", NULL },
    { babl_format ("RGBA u8"), babl_format("Y'CbCrA float"), "rgbau8_to_ycbcraf", NULL },
    { babl_format ("Y'CbCrA float"), babl_format("RGBA u8"), "ycbcraf_to_rgbau8", NULL },
127

128 129
    { babl_format ("RGB u8"), babl_format("RGBA float"), "rgbu8_to_rgbaf", NULL },
    { babl_format ("RGBA float"), babl_format("RGB u8"), "rgbaf_to_rgbu8", NULL },
130

131
    { babl_format ("Y u8"), babl_format("Y float"), "yu8_to_yf", NULL },
132

133 134 135 136
    { babl_format ("RGBA float"), babl_format("YA float"), "rgbaf_to_yaf", NULL },
    { babl_format ("YA float"), babl_format("RGBA float"), "yaf_to_rgbaf", NULL },
    { babl_format ("RGBA u8"), babl_format("YA float"), "rgbau8_to_yaf", NULL },
    { babl_format ("YA float"), babl_format("RGBA u8"), "yaf_to_rgbau8", NULL },
137
    { babl_format ("RaGaBaA float"), babl_format("YA float"), "ragabaf_to_yaf", NULL },
138

139 140
    { babl_format ("RGBA float"), babl_format("R'G'B'A u8"), "rgbaf_to_rgba_gamma_u8", NULL },
    { babl_format ("RGBA float"), babl_format("R'G'B' u8"), "rgbaf_to_rgb_gamma_u8", NULL },
141

142 143
    { babl_format ("RaGaBaA float"), babl_format("R'G'B'A u8"), "ragabaf_to_rgba_gamma_u8", NULL },
    { babl_format ("RaGaBaA float"), babl_format("R'G'B' u8"), "ragabaf_to_rgb_gamma_u8", NULL },
144

145
    { babl_format ("YA float"), babl_format("R'G'B'A u8"), "yaf_to_rgba_gamma_u8", NULL },
146 147 148 149 150 151 152 153
    { babl_format ("YA float"), babl_format("R'G'B' u8"), "yaf_to_rgb_gamma_u8", NULL },

    { babl_format ("YA float"), babl_format ("RaGaBaA float"), "yaf_to_ragabaf", NULL },

    { babl_format ("Y float"), babl_format ("RaGaBaA float"), "yf_to_ragabaf", NULL },

    { babl_format ("RGBA float"), babl_format ("RGB float"), "rgbaf_to_rgbf", NULL },

154 155
    { babl_format ("R'G'B' float"), babl_format ("RGBA float"), "rgb_gamma_f_to_rgbaf", NULL },

156 157 158
    /* Reuse some conversions */
    { babl_format ("R'G'B' u8"), babl_format ("R'G'B'A float"), "rgbu8_to_rgbaf", NULL },
    { babl_format ("R'G'B'A u8"), babl_format ("R'G'B'A float"), "rgbau8_to_rgbaf", NULL },
159 160

    { babl_format ("R'G'B' float"), babl_format ("RaGaBaA float"), "rgb_gamma_f_to_rgbaf", NULL },
161
  };
162

163 164 165 166 167 168 169 170
  ColorConversionInfo lut8_conversions[] = {
    { babl_format ("R'G'B'A u8"), babl_format("RGBA float"), "rgba_gamma_u8_to_rgbaf", NULL },
    { babl_format ("R'G'B'A u8"), babl_format("RaGaBaA float"), "rgba_gamma_u8_to_ragabaf", NULL },
    { babl_format ("R'G'B'A u8"), babl_format("YA float"), "rgba_gamma_u8_to_yaf", NULL },
    { babl_format ("R'G'B' u8"), babl_format("RGBA float"), "rgb_gamma_u8_to_rgbaf", NULL },
    { babl_format ("R'G'B' u8"), babl_format("RaGaBaA float"), "rgb_gamma_u8_to_ragabaf", NULL },
    { babl_format ("R'G'B' u8"), babl_format("YA float"), "rgb_gamma_u8_to_yaf", NULL },
  };
171

172 173 174
  /* Fail if the kernels are already compiled, meaning this must have been called twice */
  g_return_val_if_fail (!float_kernels, FALSE);
  g_return_val_if_fail (!color_kernels_hash, FALSE);
175

176 177 178 179
  color_kernels_hash = g_hash_table_new_full (color_kernels_hash_hashfunc,
                                              color_kernels_hash_equalfunc,
                                              NULL,
                                              NULL);
180

181 182 183 184 185 186
  gegl_cl_color_load_conversion_set (float_conversions,
                                     G_N_ELEMENTS (float_conversions),
                                     &float_kernels,
                                     colors_cl_source);

  if (!float_kernels)
187
    {
Daniel Sabo's avatar
Daniel Sabo committed
188
      g_warning ("Failed to compile color conversions (float_kernels)");
189 190
      result = FALSE;
    }
191

192 193 194 195
  gegl_cl_color_load_conversion_set (lut8_conversions,
                                     G_N_ELEMENTS (lut8_conversions),
                                     &lut8_kernels,
                                     colors_8bit_lut_cl_source);
196

197 198
  if (!lut8_kernels)
    {
Daniel Sabo's avatar
Daniel Sabo committed
199
      g_warning ("Failed to compile color conversions (lut8_kernels)");
200
      result = FALSE;
201
    }
202

203
  return result;
204 205
}

206 207 208
static cl_kernel
find_color_kernel (const Babl *in_format,
                   const Babl *out_format)
209
{
210 211 212 213 214 215
  ColorConversionInfo *info;
  ColorConversionInfo search = { in_format, out_format, NULL, NULL };
  info = g_hash_table_lookup (color_kernels_hash, &search);
  if (info)
    return info->kernel;
  return NULL;
216 217
}

218
gboolean
219 220
gegl_cl_color_babl (const Babl *buffer_format,
                    size_t     *bytes)
221
{
Victor Oliveira's avatar
Victor Oliveira committed
222
  if (bytes)
223
    *bytes = babl_format_get_bytes_per_pixel (buffer_format);
224 225 226 227

  return TRUE;
}

228
GeglClColorOp
229 230
gegl_cl_color_supported (const Babl *in_format,
                         const Babl *out_format)
231
{
232
  if (in_format == out_format)
233
    return GEGL_CL_COLOR_EQUAL;
234

235
  if (color_kernels_hash && find_color_kernel (in_format, out_format))
236
    return GEGL_CL_COLOR_CONVERT;
237

238 239 240 241
  GEGL_NOTE (GEGL_DEBUG_OPENCL, "Missing OpenCL conversion for %s -> %s",
             babl_get_name (in_format),
             babl_get_name (out_format));

242
  return GEGL_CL_COLOR_NOT_SUPPORTED;
243 244 245
}

gboolean
246 247 248 249 250
gegl_cl_color_conv (cl_mem         in_tex,
                    cl_mem         out_tex,
                    const size_t   size,
                    const Babl    *in_format,
                    const Babl    *out_format)
251
{
252 253
  cl_kernel kernel = NULL;
  cl_int cl_err = 0;
254 255 256

  if (in_format == out_format)
    {
257
      size_t s = babl_format_get_bytes_per_pixel (in_format);
258 259

      /* just copy in_tex to out_tex */
260 261 262 263
      cl_err = gegl_clEnqueueCopyBuffer (gegl_cl_get_command_queue(),
                                         in_tex, out_tex, 0, 0, size * s,
                                         0, NULL, NULL);
      CL_CHECK;
264
      return FALSE;
265
    }
266

267 268 269
  kernel = find_color_kernel (in_format, out_format);
  if (!kernel)
    return FALSE;
270

271 272 273 274 275 276 277 278 279 280 281
  cl_err = gegl_cl_set_kernel_args (kernel,
                                    sizeof(cl_mem), &in_tex,
                                    sizeof(cl_mem), &out_tex,
                                    NULL);
  CL_CHECK;

  cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
                                        kernel, 1,
                                        NULL, &size, NULL,
                                        0, NULL, NULL);
  CL_CHECK;
282 283 284
  return FALSE;

error:
285 286
  return TRUE;
}