1 /* This file is part of GEGL
2  *
3  * GEGL is free software; you can redistribute it and/or
4  * modify it under the terms of the GNU Lesser General Public
5  * License as published by the Free Software Foundation; either
6  * version 3 of the License, or (at your option) any later version.
7  *
8  * GEGL is distributed in the hope that it will be useful,
9  * but WITHOUT ANY WARRANTY; without even the implied warranty of
10  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
11  * Lesser General Public License for more details.
12  *
13  * 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  * Copyright 2012 Victor Oliveira (victormatheus@gmail.com)
17  */
18 
19 
20 /* Here we have implemented in OpenCL a subset of color space conversions provided by BABL
21    that are commonly used. For now there is no support for conversions that need a intermediate
22    representation (ex: A->B, B->C, C->D), just among two color spaces.
23 */
24 
25 #include "config.h"
26 
27 #include "gegl.h"
28 #include "gegl/gegl-debug.h"
29 #include "gegl-cl.h"
30 #include "gegl-cl-color.h"
31 
32 #include "opencl/colors.cl.h"
33 #include "opencl/colors-8bit-lut.cl.h"
34 
35 static GHashTable  *color_kernels_hash = NULL;
36 
37 typedef struct
38 {
39   const Babl *from_fmt;
40   const Babl *to_fmt;
41   const char *kernel_name;
42   cl_kernel   kernel;
43 } ColorConversionInfo;
44 
45 static guint
color_kernels_hash_hashfunc(gconstpointer key)46 color_kernels_hash_hashfunc (gconstpointer key)
47 {
48   const ColorConversionInfo *ekey = key;
49   return GPOINTER_TO_INT (ekey->from_fmt) ^ GPOINTER_TO_INT (ekey->to_fmt);
50 }
51 
52 static gboolean
color_kernels_hash_equalfunc(gconstpointer a,gconstpointer b)53 color_kernels_hash_equalfunc (gconstpointer a,
54                               gconstpointer b)
55 {
56   const ColorConversionInfo *ea = a;
57   const ColorConversionInfo *eb = b;
58 
59   if (ea->from_fmt == eb->from_fmt &&
60       ea->to_fmt   == eb->to_fmt)
61     {
62       return TRUE;
63     }
64   return FALSE;
65 }
66 
67 static gboolean
gegl_cl_color_load_conversion_set(ColorConversionInfo * conversions,gint num_conversions,GeglClRunData ** kernels,const gchar * source)68 gegl_cl_color_load_conversion_set (ColorConversionInfo  *conversions,
69                                    gint                  num_conversions,
70                                    GeglClRunData       **kernels,
71                                    const gchar          *source)
72 {
73   const char *kernel_names[num_conversions + 1];
74 
75   for (int i = 0; i < num_conversions; ++i)
76     {
77       kernel_names[i] = conversions[i].kernel_name;
78     }
79 
80   kernel_names[num_conversions] = NULL;
81 
82   *kernels = gegl_cl_compile_and_build (source, kernel_names);
83 
84   if (!(*kernels))
85     {
86       return FALSE;
87     }
88 
89   for (int i = 0; i < num_conversions; ++i)
90     {
91       ColorConversionInfo *info = g_new (ColorConversionInfo, 1);
92 
93       conversions[i].kernel = (*kernels)->kernel[i];
94       *info = conversions[i];
95 
96       g_hash_table_insert (color_kernels_hash, info, info);
97     }
98 
99   return TRUE;
100 }
101 
102 gboolean
gegl_cl_color_compile_kernels(void)103 gegl_cl_color_compile_kernels (void)
104 {
105   static GeglClRunData *float_kernels = NULL;
106   static GeglClRunData *lut8_kernels = NULL;
107   gboolean result = TRUE;
108 
109   ColorConversionInfo float_conversions[] = {
110     { babl_format ("RGBA u8"), babl_format ("RGBA float"), "rgbau8_to_rgbaf", NULL },
111     { babl_format ("RGBA float"), babl_format ("RGBA u8"), "rgbaf_to_rgbau8", NULL },
112 
113     { babl_format ("RGBA float"), babl_format("RaGaBaA float"), "rgbaf_to_ragabaf", NULL },
114     { babl_format ("RaGaBaA float"), babl_format("RGBA float"), "ragabaf_to_rgbaf", NULL },
115     { babl_format ("RGBA u8"), babl_format("RaGaBaA float"), "rgbau8_to_ragabaf", NULL },
116     { babl_format ("RaGaBaA float"), babl_format("RGBA u8"), "ragabaf_to_rgbau8", NULL },
117 
118     { babl_format ("RGBA float"), babl_format("R'G'B'A float"), "rgbaf_to_rgba_gamma_f", NULL },
119     { babl_format ("R'G'B'A float"), babl_format("RGBA float"), "rgba_gamma_f_to_rgbaf", NULL },
120     { babl_format ("RGBA u8"), babl_format("R'G'B'A float"), "rgbau8_to_rgba_gamma_f", NULL },
121     { babl_format ("R'G'B'A float"), babl_format("RGBA u8"), "rgba_gamma_f_to_rgbau8", NULL },
122 
123     { babl_format ("RGBA float"), babl_format("Y'CbCrA float"), "rgbaf_to_ycbcraf", NULL },
124     { babl_format ("Y'CbCrA float"), babl_format("RGBA float"), "ycbcraf_to_rgbaf", NULL },
125     { babl_format ("RGBA u8"), babl_format("Y'CbCrA float"), "rgbau8_to_ycbcraf", NULL },
126     { babl_format ("Y'CbCrA float"), babl_format("RGBA u8"), "ycbcraf_to_rgbau8", NULL },
127 
128     { babl_format ("RGB u8"), babl_format("RGBA float"), "rgbu8_to_rgbaf", NULL },
129     { 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     { babl_format ("RGBA float"), babl_format("YA float"), "rgbaf_to_yaf", NULL },
134     { babl_format ("YA float"), babl_format("RGBA float"), "yaf_to_rgbaf", NULL },
135     { babl_format ("RGBA u8"), babl_format("YA float"), "rgbau8_to_yaf", NULL },
136     { 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     { babl_format ("RGBA float"), babl_format("R'G'B'A u8"), "rgbaf_to_rgba_gamma_u8", NULL },
140     { babl_format ("RGBA float"), babl_format("R'G'B' u8"), "rgbaf_to_rgb_gamma_u8", NULL },
141 
142     { babl_format ("RaGaBaA float"), babl_format("R'G'B'A u8"), "ragabaf_to_rgba_gamma_u8", NULL },
143     { 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     { babl_format ("YA float"), babl_format("R'G'B' u8"), "yaf_to_rgb_gamma_u8", NULL },
147 
148     { babl_format ("YA float"), babl_format ("RaGaBaA float"), "yaf_to_ragabaf", NULL },
149 
150     { babl_format ("Y float"), babl_format ("RaGaBaA float"), "yf_to_ragabaf", NULL },
151 
152     { babl_format ("RGBA float"), babl_format ("RGB float"), "rgbaf_to_rgbf", NULL },
153 
154     { babl_format ("R'G'B' float"), babl_format ("RGBA float"), "rgb_gamma_f_to_rgbaf", NULL },
155 
156     /* Reuse some conversions */
157     { babl_format ("R'G'B' u8"), babl_format ("R'G'B'A float"), "rgbu8_to_rgbaf", NULL },
158     { 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   ColorConversionInfo lut8_conversions[] = {
164     { babl_format ("R'G'B'A u8"), babl_format("RGBA float"), "rgba_gamma_u8_to_rgbaf", NULL },
165     { babl_format ("R'G'B'A u8"), babl_format("RaGaBaA float"), "rgba_gamma_u8_to_ragabaf", NULL },
166     { babl_format ("R'G'B'A u8"), babl_format("YA float"), "rgba_gamma_u8_to_yaf", NULL },
167     { babl_format ("R'G'B' u8"), babl_format("RGBA float"), "rgb_gamma_u8_to_rgbaf", NULL },
168     { babl_format ("R'G'B' u8"), babl_format("RaGaBaA float"), "rgb_gamma_u8_to_ragabaf", NULL },
169     { babl_format ("R'G'B' u8"), babl_format("YA float"), "rgb_gamma_u8_to_yaf", NULL },
170   };
171 
172   /* Fail if the kernels are already compiled, meaning this must have been called twice */
173   g_return_val_if_fail (!float_kernels, FALSE);
174   g_return_val_if_fail (!color_kernels_hash, FALSE);
175 
176   color_kernels_hash = g_hash_table_new_full (color_kernels_hash_hashfunc,
177                                               color_kernels_hash_equalfunc,
178                                               NULL,
179                                               NULL);
180 
181   gegl_cl_color_load_conversion_set (float_conversions,
182                                      G_N_ELEMENTS (float_conversions),
183                                      &float_kernels,
184                                      colors_cl_source);
185 
186   if (!float_kernels)
187     {
188       g_warning ("OpenCL Failed to compile color conversions (float_kernels)");
189       result = FALSE;
190     }
191 
192   gegl_cl_color_load_conversion_set (lut8_conversions,
193                                      G_N_ELEMENTS (lut8_conversions),
194                                      &lut8_kernels,
195                                      colors_8bit_lut_cl_source);
196 
197   if (!lut8_kernels)
198     {
199       g_warning ("OpenCL Failed to compile color conversions (lut8_kernels)");
200       result = FALSE;
201     }
202 
203   return result;
204 }
205 
206 static cl_kernel
find_color_kernel(const Babl * in_format,const Babl * out_format)207 find_color_kernel (const Babl *in_format,
208                    const Babl *out_format)
209 {
210   ColorConversionInfo *info;
211   ColorConversionInfo search = { in_format, out_format, NULL, NULL };
212   info = g_hash_table_lookup (color_kernels_hash, &search);
213   if (info)
214     return info->kernel;
215   return NULL;
216 }
217 
218 gboolean
gegl_cl_color_babl(const Babl * buffer_format,size_t * bytes)219 gegl_cl_color_babl (const Babl *buffer_format,
220                     size_t     *bytes)
221 {
222   if (bytes)
223     *bytes = babl_format_get_bytes_per_pixel (buffer_format);
224 
225   return TRUE;
226 }
227 
228 GeglClColorOp
gegl_cl_color_supported(const Babl * in_format,const Babl * out_format)229 gegl_cl_color_supported (const Babl *in_format,
230                          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   GEGL_NOTE (GEGL_DEBUG_OPENCL, "Missing OpenCL conversion for %s -> %s",
239              babl_get_name (in_format),
240              babl_get_name (out_format));
241 
242   return GEGL_CL_COLOR_NOT_SUPPORTED;
243 }
244 
245 gboolean
gegl_cl_color_conv(cl_mem in_tex,cl_mem out_tex,const size_t size,const Babl * in_format,const Babl * out_format)246 gegl_cl_color_conv (cl_mem         in_tex,
247                     cl_mem         out_tex,
248                     const size_t   size,
249                     const Babl    *in_format,
250                     const Babl    *out_format)
251 {
252   cl_kernel kernel = NULL;
253   cl_int cl_err = 0;
254 
255   if (in_format == out_format)
256     {
257       size_t s = babl_format_get_bytes_per_pixel (in_format);
258 
259       /* just copy in_tex to out_tex */
260       cl_err = gegl_clEnqueueCopyBuffer (gegl_cl_get_command_queue(),
261                                          in_tex, out_tex, 0, 0, size * s,
262                                          0, NULL, NULL);
263       CL_CHECK;
264       return FALSE;
265     }
266 
267   kernel = find_color_kernel (in_format, out_format);
268   if (!kernel)
269     return FALSE;
270 
271   cl_err = gegl_cl_set_kernel_args (kernel,
272                                     sizeof(cl_mem), &in_tex,
273                                     sizeof(cl_mem), &out_tex,
274                                     NULL);
275   CL_CHECK;
276 
277   cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
278                                         kernel, 1,
279                                         NULL, &size, NULL,
280                                         0, NULL, NULL);
281   CL_CHECK;
282   return FALSE;
283 
284 error:
285   return TRUE;
286 }
287