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