1 /*
2 This file is part of darktable,
3 Copyright (C) 2010-2020 darktable developers.
4
5 darktable is free software: you can redistribute it and/or modify
6 it under the terms of the GNU General Public License as published by
7 the Free Software Foundation, either version 3 of the License, or
8 (at your option) any later version.
9
10 darktable is distributed in the hope that it will be useful,
11 but WITHOUT ANY WARRANTY; without even the implied warranty of
12 MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
13 GNU General Public License for more details.
14
15 You should have received a copy of the GNU General Public License
16 along with darktable. If not, see <http://www.gnu.org/licenses/>.
17 */
18 #ifdef HAVE_CONFIG_H
19 #include "config.h"
20 #endif
21 #include <assert.h>
22 #include <math.h>
23 #include <stdlib.h>
24 #include <string.h>
25 #if defined(__SSE__)
26 #include <xmmintrin.h>
27 #endif
28 #include "bauhaus/bauhaus.h"
29 #include "common/opencl.h"
30 #include "control/control.h"
31 #include "develop/develop.h"
32 #include "develop/imageop.h"
33 #include "develop/imageop_math.h"
34 #include "develop/imageop_gui.h"
35 #include "develop/tiling.h"
36 #include "gui/accelerators.h"
37 #include "gui/gtk.h"
38 #include "iop/iop_api.h"
39
40 #include <gtk/gtk.h>
41 #include <inttypes.h>
42
43
44 DT_MODULE_INTROSPECTION(2, dt_iop_highlights_params_t)
45
46 typedef enum dt_iop_highlights_mode_t
47 {
48 DT_IOP_HIGHLIGHTS_CLIP = 0, // $DESCRIPTION: "clip highlights"
49 DT_IOP_HIGHLIGHTS_LCH = 1, // $DESCRIPTION: "reconstruct in LCh"
50 DT_IOP_HIGHLIGHTS_INPAINT = 2, // $DESCRIPTION: "reconstruct color"
51 } dt_iop_highlights_mode_t;
52
53 typedef struct dt_iop_highlights_params_t
54 {
55 dt_iop_highlights_mode_t mode; // $DEFAULT: DT_IOP_HIGHLIGHTS_CLIP $DESCRIPTION: "method"
56 float blendL; // unused $DEFAULT: 1.0
57 float blendC; // unused $DEFAULT: 0.0
58 float blendh; // unused $DEFAULT: 0.0
59 float clip; // $MIN: 0.0 $MAX: 2.0 $DEFAULT: 1.0 $DESCRIPTION: "clipping threshold"
60 } dt_iop_highlights_params_t;
61
62 typedef struct dt_iop_highlights_gui_data_t
63 {
64 GtkWidget *clip;
65 GtkWidget *mode;
66 } dt_iop_highlights_gui_data_t;
67
68 typedef dt_iop_highlights_params_t dt_iop_highlights_data_t;
69
70 typedef struct dt_iop_highlights_global_data_t
71 {
72 int kernel_highlights_1f_clip;
73 int kernel_highlights_1f_lch_bayer;
74 int kernel_highlights_1f_lch_xtrans;
75 int kernel_highlights_4f_clip;
76 } dt_iop_highlights_global_data_t;
77
78
name()79 const char *name()
80 {
81 return _("highlight reconstruction");
82 }
83
description(struct dt_iop_module_t * self)84 const char *description(struct dt_iop_module_t *self)
85 {
86 return dt_iop_set_description(self, _("avoid magenta highlights and try to recover highlights colors"),
87 _("corrective"),
88 _("linear, raw, scene-referred"),
89 _("reconstruction, raw"),
90 _("linear, raw, scene-referred"));
91 }
92
default_group()93 int default_group()
94 {
95 return IOP_GROUP_BASIC | IOP_GROUP_TECHNICAL;
96 }
97
flags()98 int flags()
99 {
100 return IOP_FLAGS_SUPPORTS_BLENDING | IOP_FLAGS_ALLOW_TILING | IOP_FLAGS_ONE_INSTANCE;
101 }
102
default_colorspace(dt_iop_module_t * self,dt_dev_pixelpipe_t * pipe,dt_dev_pixelpipe_iop_t * piece)103 int default_colorspace(dt_iop_module_t *self, dt_dev_pixelpipe_t *pipe, dt_dev_pixelpipe_iop_t *piece)
104 {
105 return iop_cs_RAW;
106 }
107
legacy_params(dt_iop_module_t * self,const void * const old_params,const int old_version,void * new_params,const int new_version)108 int legacy_params(dt_iop_module_t *self, const void *const old_params, const int old_version,
109 void *new_params, const int new_version)
110 {
111 if(old_version == 1 && new_version == 2)
112 {
113 memcpy(new_params, old_params, sizeof(dt_iop_highlights_params_t) - sizeof(float));
114 dt_iop_highlights_params_t *n = (dt_iop_highlights_params_t *)new_params;
115 n->clip = 1.0f;
116 return 0;
117 }
118 return 1;
119 }
120
121 #ifdef HAVE_OPENCL
process_cl(struct dt_iop_module_t * self,dt_dev_pixelpipe_iop_t * piece,cl_mem dev_in,cl_mem dev_out,const dt_iop_roi_t * const roi_in,const dt_iop_roi_t * const roi_out)122 int process_cl(struct dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, cl_mem dev_in, cl_mem dev_out,
123 const dt_iop_roi_t *const roi_in, const dt_iop_roi_t *const roi_out)
124 {
125 dt_iop_highlights_data_t *d = (dt_iop_highlights_data_t *)piece->data;
126 dt_iop_highlights_global_data_t *gd = (dt_iop_highlights_global_data_t *)self->global_data;
127
128 cl_int err = -999;
129 cl_mem dev_xtrans = NULL;
130
131 const int devid = piece->pipe->devid;
132 const int width = roi_in->width;
133 const int height = roi_in->height;
134
135 const float clip = d->clip
136 * fminf(piece->pipe->dsc.processed_maximum[0],
137 fminf(piece->pipe->dsc.processed_maximum[1], piece->pipe->dsc.processed_maximum[2]));
138
139 const uint32_t filters = piece->pipe->dsc.filters;
140
141 if(!filters)
142 {
143 // non-raw images use dedicated kernel which just clips
144 size_t sizes[] = { ROUNDUPWD(width), ROUNDUPHT(height), 1 };
145 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_4f_clip, 0, sizeof(cl_mem), (void *)&dev_in);
146 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_4f_clip, 1, sizeof(cl_mem), (void *)&dev_out);
147 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_4f_clip, 2, sizeof(int), (void *)&width);
148 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_4f_clip, 3, sizeof(int), (void *)&height);
149 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_4f_clip, 4, sizeof(int), (void *)&d->mode);
150 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_4f_clip, 5, sizeof(float), (void *)&clip);
151 err = dt_opencl_enqueue_kernel_2d(devid, gd->kernel_highlights_4f_clip, sizes);
152 if(err != CL_SUCCESS) goto error;
153 }
154 else if(d->mode == DT_IOP_HIGHLIGHTS_CLIP)
155 {
156 // raw images with clip mode (both bayer and xtrans)
157 size_t sizes[] = { ROUNDUPWD(width), ROUNDUPHT(height), 1 };
158 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_clip, 0, sizeof(cl_mem), (void *)&dev_in);
159 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_clip, 1, sizeof(cl_mem), (void *)&dev_out);
160 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_clip, 2, sizeof(int), (void *)&width);
161 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_clip, 3, sizeof(int), (void *)&height);
162 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_clip, 4, sizeof(float), (void *)&clip);
163 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_clip, 5, sizeof(int), (void *)&roi_out->x);
164 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_clip, 6, sizeof(int), (void *)&roi_out->y);
165 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_clip, 7, sizeof(int), (void *)&filters);
166 err = dt_opencl_enqueue_kernel_2d(devid, gd->kernel_highlights_1f_clip, sizes);
167 if(err != CL_SUCCESS) goto error;
168 }
169 else if(d->mode == DT_IOP_HIGHLIGHTS_LCH && filters != 9u)
170 {
171 // bayer sensor raws with LCH mode
172 size_t sizes[] = { ROUNDUPWD(width), ROUNDUPHT(height), 1 };
173 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_bayer, 0, sizeof(cl_mem), (void *)&dev_in);
174 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_bayer, 1, sizeof(cl_mem), (void *)&dev_out);
175 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_bayer, 2, sizeof(int), (void *)&width);
176 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_bayer, 3, sizeof(int), (void *)&height);
177 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_bayer, 4, sizeof(float), (void *)&clip);
178 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_bayer, 5, sizeof(int), (void *)&roi_out->x);
179 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_bayer, 6, sizeof(int), (void *)&roi_out->y);
180 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_bayer, 7, sizeof(int), (void *)&filters);
181 err = dt_opencl_enqueue_kernel_2d(devid, gd->kernel_highlights_1f_lch_bayer, sizes);
182 if(err != CL_SUCCESS) goto error;
183 }
184 else if(d->mode == DT_IOP_HIGHLIGHTS_LCH && filters == 9u)
185 {
186 // xtrans sensor raws with LCH mode
187 int blocksizex, blocksizey;
188
189 dt_opencl_local_buffer_t locopt
190 = (dt_opencl_local_buffer_t){ .xoffset = 2 * 2, .xfactor = 1, .yoffset = 2 * 2, .yfactor = 1,
191 .cellsize = sizeof(float), .overhead = 0,
192 .sizex = 1 << 8, .sizey = 1 << 8 };
193
194 if(dt_opencl_local_buffer_opt(devid, gd->kernel_highlights_1f_lch_xtrans, &locopt))
195 {
196 blocksizex = locopt.sizex;
197 blocksizey = locopt.sizey;
198 }
199 else
200 blocksizex = blocksizey = 1;
201
202 dev_xtrans
203 = dt_opencl_copy_host_to_device_constant(devid, sizeof(piece->pipe->dsc.xtrans), piece->pipe->dsc.xtrans);
204 if(dev_xtrans == NULL) goto error;
205
206 size_t sizes[] = { ROUNDUP(width, blocksizex), ROUNDUP(height, blocksizey), 1 };
207 size_t local[] = { blocksizex, blocksizey, 1 };
208 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_xtrans, 0, sizeof(cl_mem), (void *)&dev_in);
209 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_xtrans, 1, sizeof(cl_mem), (void *)&dev_out);
210 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_xtrans, 2, sizeof(int), (void *)&width);
211 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_xtrans, 3, sizeof(int), (void *)&height);
212 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_xtrans, 4, sizeof(float), (void *)&clip);
213 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_xtrans, 5, sizeof(int), (void *)&roi_out->x);
214 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_xtrans, 6, sizeof(int), (void *)&roi_out->y);
215 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_xtrans, 7, sizeof(cl_mem), (void *)&dev_xtrans);
216 dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_xtrans, 8,
217 sizeof(float) * (blocksizex + 4) * (blocksizey + 4), NULL);
218
219 err = dt_opencl_enqueue_kernel_2d_with_local(devid, gd->kernel_highlights_1f_lch_xtrans, sizes, local);
220 if(err != CL_SUCCESS) goto error;
221 }
222
223 // update processed maximum
224 const float m = fmaxf(fmaxf(piece->pipe->dsc.processed_maximum[0], piece->pipe->dsc.processed_maximum[1]),
225 piece->pipe->dsc.processed_maximum[2]);
226 for(int k = 0; k < 3; k++) piece->pipe->dsc.processed_maximum[k] = m;
227
228 dt_opencl_release_mem_object(dev_xtrans);
229 return TRUE;
230
231 error:
232 dt_opencl_release_mem_object(dev_xtrans);
233 dt_print(DT_DEBUG_OPENCL, "[opencl_highlights] couldn't enqueue kernel! %d\n", err);
234 return FALSE;
235 }
236 #endif
237
tiling_callback(struct dt_iop_module_t * self,struct dt_dev_pixelpipe_iop_t * piece,const dt_iop_roi_t * roi_in,const dt_iop_roi_t * roi_out,struct dt_develop_tiling_t * tiling)238 void tiling_callback(struct dt_iop_module_t *self, struct dt_dev_pixelpipe_iop_t *piece,
239 const dt_iop_roi_t *roi_in, const dt_iop_roi_t *roi_out,
240 struct dt_develop_tiling_t *tiling)
241 {
242 dt_iop_highlights_data_t *d = (dt_iop_highlights_data_t *)piece->data;
243 const uint32_t filters = piece->pipe->dsc.filters;
244
245 tiling->factor = 2.0f; // in + out
246 tiling->maxbuf = 1.0f;
247 tiling->overhead = 0;
248
249 if(filters == 9u)
250 {
251 // xtrans
252 tiling->xalign = 6;
253 tiling->yalign = 6;
254 tiling->overlap = (d->mode == DT_IOP_HIGHLIGHTS_LCH) ? 2 : 0;
255 }
256 else if(filters)
257 {
258 // bayer
259 tiling->xalign = 2;
260 tiling->yalign = 2;
261 tiling->overlap = (d->mode == DT_IOP_HIGHLIGHTS_LCH) ? 1 : 0;
262 }
263 else
264 {
265 // non-raw
266 tiling->xalign = 1;
267 tiling->yalign = 1;
268 tiling->overlap = 0;
269 }
270 }
271
272 /* interpolate value for a pixel, ideal via ratio to nearby pixel */
interp_pix_xtrans(const int ratio_next,const ssize_t offset_next,const float clip0,const float clip_next,const float * const in,const float * const ratios)273 static inline float interp_pix_xtrans(const int ratio_next,
274 const ssize_t offset_next,
275 const float clip0, const float clip_next,
276 const float *const in,
277 const float *const ratios)
278 {
279 assert(ratio_next != 0);
280 // it's OK to exceed clipping of current pixel's color based on a
281 // neighbor -- that is the purpose of interpolating highlight
282 // colors
283 const float clip_val = fmaxf(clip0, clip_next);
284 if(in[offset_next] >= clip_next - 1e-5f)
285 {
286 // next pixel is also clipped
287 return clip_val;
288 }
289 else
290 {
291 // set this pixel in ratio to the next
292 assert(ratio_next != 0);
293 if (ratio_next > 0)
294 return fminf(in[offset_next] / ratios[ratio_next], clip_val);
295 else
296 return fminf(in[offset_next] * ratios[-ratio_next], clip_val);
297 }
298 }
299
interpolate_color_xtrans(const void * const ivoid,void * const ovoid,const dt_iop_roi_t * const roi_in,const dt_iop_roi_t * const roi_out,int dim,int dir,int other,const float * const clip,const uint8_t (* const xtrans)[6],const int pass)300 static inline void interpolate_color_xtrans(const void *const ivoid, void *const ovoid,
301 const dt_iop_roi_t *const roi_in,
302 const dt_iop_roi_t *const roi_out,
303 int dim, int dir, int other,
304 const float *const clip,
305 const uint8_t (*const xtrans)[6],
306 const int pass)
307 {
308 // In Bayer each row/col has only green/red or green/blue
309 // transitions, hence can reconstruct color by single ratio per
310 // row. In x-trans there can be transitions between arbitrary colors
311 // in a row/col (and 2x2 green blocks which provide no color
312 // transition information). Hence calculate multiple color ratios
313 // for each row/col.
314
315 // Lookup for color ratios, e.g. red -> blue is roff[0][2] and blue
316 // -> red is roff[2][0]. Returned value is an index into ratios. If
317 // negative, then need to invert the ratio. Identity color
318 // transitions aren't used.
319 const int roff[3][3] = {{ 0, -1, -2},
320 { 1, 0, -3},
321 { 2, 3, 0}};
322 // record ratios of color transitions 0:unused, 1:RG, 2:RB, and 3:GB
323 float ratios[4] = {1.0f, 1.0f, 1.0f, 1.0f};
324
325 // passes are 0:+x, 1:-x, 2:+y, 3:-y
326 // dims are 0:traverse a row, 1:traverse a column
327 // dir is 1:left to right, -1: right to left
328 int i = (dim == 0) ? 0 : other;
329 int j = (dim == 0) ? other : 0;
330 const ssize_t offs = (ssize_t)(dim ? roi_out->width : 1) * ((dir < 0) ? -1 : 1);
331 const ssize_t offl = offs - (dim ? 1 : roi_out->width);
332 const ssize_t offr = offs + (dim ? 1 : roi_out->width);
333 int beg, end;
334 if(dir == 1)
335 {
336 beg = 0;
337 end = (dim == 0) ? roi_out->width : roi_out->height;
338 }
339 else
340 {
341 beg = ((dim == 0) ? roi_out->width : roi_out->height) - 1;
342 end = -1;
343 }
344
345 float *in, *out;
346 if(dim == 1)
347 {
348 out = (float *)ovoid + (size_t)i + (size_t)beg * roi_out->width;
349 in = (float *)ivoid + (size_t)i + (size_t)beg * roi_in->width;
350 }
351 else
352 {
353 out = (float *)ovoid + (size_t)beg + (size_t)j * roi_out->width;
354 in = (float *)ivoid + (size_t)beg + (size_t)j * roi_in->width;
355 }
356
357 for(int k = beg; k != end; k += dir)
358 {
359 if(dim == 1)
360 j = k;
361 else
362 i = k;
363
364 const uint8_t f0 = FCxtrans(j, i, roi_in, xtrans);
365 const uint8_t f1 = FCxtrans(dim ? (j + dir) : j, dim ? i : (i + dir), roi_in, xtrans);
366 const uint8_t fl = FCxtrans(dim ? (j + dir) : (j - 1), dim ? (i - 1) : (i + dir), roi_in, xtrans);
367 const uint8_t fr = FCxtrans(dim ? (j + dir) : (j + 1), dim ? (i + 1) : (i + dir), roi_in, xtrans);
368 const float clip0 = clip[f0];
369 const float clip1 = clip[f1];
370 const float clipl = clip[fl];
371 const float clipr = clip[fr];
372 const float clip_max = fmaxf(fmaxf(clip[0], clip[1]), clip[2]);
373
374 if(i == 0 || i == roi_out->width - 1 || j == 0 || j == roi_out->height - 1)
375 {
376 if(pass == 3) out[0] = fminf(clip_max, in[0]);
377 }
378 else
379 {
380 // ratio to next pixel if this & next are unclamped and not in
381 // 2x2 green block
382 if ((f0 != f1) &&
383 (in[0] < clip0 && in[0] > 1e-5f) &&
384 (in[offs] < clip1 && in[offs] > 1e-5f))
385 {
386 const int r = roff[f0][f1];
387 assert(r != 0);
388 if (r > 0)
389 ratios[r] = (3.f * ratios[r] + (in[offs] / in[0])) / 4.f;
390 else
391 ratios[-r] = (3.f * ratios[-r] + (in[0] / in[offs])) / 4.f;
392 }
393
394 if(in[0] >= clip0 - 1e-5f)
395 {
396 // interplate color for clipped pixel
397 float add;
398 if(f0 != f1)
399 // next pixel is different color
400 add =
401 interp_pix_xtrans(roff[f0][f1], offs, clip0, clip1, in, ratios);
402 else
403 // at start of 2x2 green block, look diagonally
404 add = (fl != f0) ?
405 interp_pix_xtrans(roff[f0][fl], offl, clip0, clipl, in, ratios) :
406 interp_pix_xtrans(roff[f0][fr], offr, clip0, clipr, in, ratios);
407
408 if(pass == 0)
409 out[0] = add;
410 else if(pass == 3)
411 out[0] = fminf(clip_max, (out[0] + add) / 4.0f);
412 else
413 out[0] += add;
414 }
415 else
416 {
417 // pixel is not clipped
418 if(pass == 3) out[0] = in[0];
419 }
420 }
421 out += offs;
422 in += offs;
423 }
424 }
425
interpolate_color(const void * const ivoid,void * const ovoid,const dt_iop_roi_t * const roi_out,int dim,int dir,int other,const float * clip,const uint32_t filters,const int pass)426 static inline void interpolate_color(const void *const ivoid, void *const ovoid,
427 const dt_iop_roi_t *const roi_out, int dim, int dir, int other,
428 const float *clip, const uint32_t filters, const int pass)
429 {
430 float ratio = 1.0f;
431 float *in, *out;
432
433 int i = 0, j = 0;
434 if(dim == 0)
435 j = other;
436 else
437 i = other;
438 ssize_t offs = dim ? roi_out->width : 1;
439 if(dir < 0) offs = -offs;
440 int beg, end;
441 if(dim == 0 && dir == 1)
442 {
443 beg = 0;
444 end = roi_out->width;
445 }
446 else if(dim == 0 && dir == -1)
447 {
448 beg = roi_out->width - 1;
449 end = -1;
450 }
451 else if(dim == 1 && dir == 1)
452 {
453 beg = 0;
454 end = roi_out->height;
455 }
456 else if(dim == 1 && dir == -1)
457 {
458 beg = roi_out->height - 1;
459 end = -1;
460 }
461 else
462 return;
463
464 if(dim == 1)
465 {
466 out = (float *)ovoid + i + (size_t)beg * roi_out->width;
467 in = (float *)ivoid + i + (size_t)beg * roi_out->width;
468 }
469 else
470 {
471 out = (float *)ovoid + beg + (size_t)j * roi_out->width;
472 in = (float *)ivoid + beg + (size_t)j * roi_out->width;
473 }
474 for(int k = beg; k != end; k += dir)
475 {
476 if(dim == 1)
477 j = k;
478 else
479 i = k;
480 const float clip0 = clip[FC(j, i, filters)];
481 const float clip1 = clip[FC(dim ? (j + 1) : j, dim ? i : (i + 1), filters)];
482 if(i == 0 || i == roi_out->width - 1 || j == 0 || j == roi_out->height - 1)
483 {
484 if(pass == 3) out[0] = in[0];
485 }
486 else
487 {
488 if(in[0] < clip0 && in[0] > 1e-5f)
489 { // both are not clipped
490 if(in[offs] < clip1 && in[offs] > 1e-5f)
491 { // update ratio, exponential decay. ratio = in[odd]/in[even]
492 if(k & 1)
493 ratio = (3.0f * ratio + in[0] / in[offs]) / 4.0f;
494 else
495 ratio = (3.0f * ratio + in[offs] / in[0]) / 4.0f;
496 }
497 }
498
499 if(in[0] >= clip0 - 1e-5f)
500 { // in[0] is clipped, restore it as in[1] adjusted according to ratio
501 float add = 0.0f;
502 if(in[offs] >= clip1 - 1e-5f)
503 add = fmaxf(clip0, clip1);
504 else if(k & 1)
505 add = in[offs] * ratio;
506 else
507 add = in[offs] / ratio;
508
509 if(pass == 0)
510 out[0] = add;
511 else if(pass == 3)
512 out[0] = (out[0] + add) / 4.0f;
513 else
514 out[0] += add;
515 }
516 else
517 {
518 if(pass == 3) out[0] = in[0];
519 }
520 }
521 out += offs;
522 in += offs;
523 }
524 }
525
526 /*
527 * these 2 constants were computed using following Sage code:
528 *
529 * sqrt3 = sqrt(3)
530 * sqrt12 = sqrt(12) # 2*sqrt(3)
531 *
532 * print 'sqrt3 = ', sqrt3, ' ~= ', RealField(128)(sqrt3)
533 * print 'sqrt12 = ', sqrt12, ' ~= ', RealField(128)(sqrt12)
534 */
535 #define SQRT3 1.7320508075688772935274463415058723669L
536 #define SQRT12 3.4641016151377545870548926830117447339L // 2*SQRT3
537
process_lch_bayer(dt_iop_module_t * self,dt_dev_pixelpipe_iop_t * piece,const void * const ivoid,void * const ovoid,const dt_iop_roi_t * const roi_in,const dt_iop_roi_t * const roi_out,const float clip)538 static void process_lch_bayer(dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, const void *const ivoid,
539 void *const ovoid, const dt_iop_roi_t *const roi_in,
540 const dt_iop_roi_t *const roi_out, const float clip)
541 {
542 const uint32_t filters = piece->pipe->dsc.filters;
543
544 #ifdef _OPENMP
545 #pragma omp parallel for default(none) \
546 dt_omp_firstprivate(clip, filters, ivoid, ovoid, roi_out) \
547 schedule(static) collapse(2)
548 #endif
549 for(int j = 0; j < roi_out->height; j++)
550 {
551 for(int i = 0; i < roi_out->width; i++)
552 {
553 float *const out = (float *)ovoid + (size_t)roi_out->width * j + i;
554 const float *const in = (float *)ivoid + (size_t)roi_out->width * j + i;
555
556 if(i == roi_out->width - 1 || j == roi_out->height - 1)
557 {
558 // fast path for border
559 out[0] = MIN(clip, in[0]);
560 }
561 else
562 {
563 int clipped = 0;
564
565 // sample 1 bayer block. thus we will have 2 green values.
566 float R = 0.0f, Gmin = FLT_MAX, Gmax = -FLT_MAX, B = 0.0f;
567 for(int jj = 0; jj <= 1; jj++)
568 {
569 for(int ii = 0; ii <= 1; ii++)
570 {
571 const float val = in[(size_t)jj * roi_out->width + ii];
572
573 clipped = (clipped || (val > clip));
574
575 const int c = FC(j + jj + roi_out->y, i + ii + roi_out->x, filters);
576 switch(c)
577 {
578 case 0:
579 R = val;
580 break;
581 case 1:
582 Gmin = MIN(Gmin, val);
583 Gmax = MAX(Gmax, val);
584 break;
585 case 2:
586 B = val;
587 break;
588 }
589 }
590 }
591
592 if(clipped)
593 {
594 const float Ro = MIN(R, clip);
595 const float Go = MIN(Gmin, clip);
596 const float Bo = MIN(B, clip);
597
598 const float L = (R + Gmax + B) / 3.0f;
599
600 float C = SQRT3 * (R - Gmax);
601 float H = 2.0f * B - Gmax - R;
602
603 const float Co = SQRT3 * (Ro - Go);
604 const float Ho = 2.0f * Bo - Go - Ro;
605
606 if(R != Gmax && Gmax != B)
607 {
608 const float ratio = sqrtf((Co * Co + Ho * Ho) / (C * C + H * H));
609 C *= ratio;
610 H *= ratio;
611 }
612
613 float RGB[3] = { 0.0f, 0.0f, 0.0f };
614
615 /*
616 * backtransform proof, sage:
617 *
618 * R,G,B,L,C,H = var('R,G,B,L,C,H')
619 * solve([L==(R+G+B)/3, C==sqrt(3)*(R-G), H==2*B-G-R], R, G, B)
620 *
621 * result:
622 * [[R == 1/6*sqrt(3)*C - 1/6*H + L, G == -1/6*sqrt(3)*C - 1/6*H + L, B == 1/3*H + L]]
623 */
624 RGB[0] = L - H / 6.0f + C / SQRT12;
625 RGB[1] = L - H / 6.0f - C / SQRT12;
626 RGB[2] = L + H / 3.0f;
627
628 out[0] = RGB[FC(j + roi_out->y, i + roi_out->x, filters)];
629 }
630 else
631 {
632 out[0] = in[0];
633 }
634 }
635 }
636 }
637 }
638
process_lch_xtrans(dt_iop_module_t * self,dt_dev_pixelpipe_iop_t * piece,const void * const ivoid,void * const ovoid,const dt_iop_roi_t * const roi_in,const dt_iop_roi_t * const roi_out,const float clip)639 static void process_lch_xtrans(dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, const void *const ivoid,
640 void *const ovoid, const dt_iop_roi_t *const roi_in,
641 const dt_iop_roi_t *const roi_out, const float clip)
642 {
643 const uint8_t(*const xtrans)[6] = (const uint8_t(*const)[6])piece->pipe->dsc.xtrans;
644
645 #ifdef _OPENMP
646 #pragma omp parallel for default(none) \
647 dt_omp_firstprivate(clip, ivoid, ovoid, roi_in, roi_out, xtrans) \
648 schedule(static)
649 #endif
650 for(int j = 0; j < roi_out->height; j++)
651 {
652 float *out = (float *)ovoid + (size_t)roi_out->width * j;
653 float *in = (float *)ivoid + (size_t)roi_in->width * j;
654
655 // bit vector used as ring buffer to remember clipping of current
656 // and last two columns, checking current pixel and its vertical
657 // neighbors
658 int cl = 0;
659
660 for(int i = 0; i < roi_out->width; i++)
661 {
662 // update clipping ring buffer
663 cl = (cl << 1) & 6;
664 if(j >= 2 && j <= roi_out->height - 3)
665 {
666 cl |= (in[-roi_in->width] > clip) | (in[0] > clip) | (in[roi_in->width] > clip);
667 }
668
669 if(i < 2 || i > roi_out->width - 3 || j < 2 || j > roi_out->height - 3)
670 {
671 // fast path for border
672 out[0] = MIN(clip, in[0]);
673 }
674 else
675 {
676 // if current pixel is clipped, always reconstruct
677 int clipped = (in[0] > clip);
678 if(!clipped)
679 {
680 clipped = cl;
681 if(clipped)
682 {
683 // If the ring buffer can't show we are in an obviously
684 // unclipped region, this is the slow case: check if there
685 // is any 3x3 block touching the current pixel which has
686 // no clipping, as then don't need to reconstruct the
687 // current pixel. This avoids zippering in edge
688 // transitions from clipped to unclipped areas. The
689 // X-Trans sensor seems prone to this, unlike Bayer, due
690 // to its irregular pattern.
691 for(int offset_j = -2; offset_j <= 0; offset_j++)
692 {
693 for(int offset_i = -2; offset_i <= 0; offset_i++)
694 {
695 if(clipped)
696 {
697 clipped = 0;
698 for(int jj = offset_j; jj <= offset_j + 2; jj++)
699 {
700 for(int ii = offset_i; ii <= offset_i + 2; ii++)
701 {
702 const float val = in[(ssize_t)jj * roi_in->width + ii];
703 clipped = (clipped || (val > clip));
704 }
705 }
706 }
707 }
708 }
709 }
710 }
711
712 if(clipped)
713 {
714 float mean[3] = { 0.0f, 0.0f, 0.0f };
715 int cnt[3] = { 0, 0, 0 };
716 float RGBmax[3] = { -FLT_MAX, -FLT_MAX, -FLT_MAX };
717
718 for(int jj = -1; jj <= 1; jj++)
719 {
720 for(int ii = -1; ii <= 1; ii++)
721 {
722 const float val = in[(ssize_t)jj * roi_in->width + ii];
723 const int c = FCxtrans(j+jj, i+ii, roi_in, xtrans);
724 mean[c] += val;
725 cnt[c]++;
726 RGBmax[c] = MAX(RGBmax[c], val);
727 }
728 }
729
730 const float Ro = MIN(mean[0]/cnt[0], clip);
731 const float Go = MIN(mean[1]/cnt[1], clip);
732 const float Bo = MIN(mean[2]/cnt[2], clip);
733
734 const float R = RGBmax[0];
735 const float G = RGBmax[1];
736 const float B = RGBmax[2];
737
738 const float L = (R + G + B) / 3.0f;
739
740 float C = SQRT3 * (R - G);
741 float H = 2.0f * B - G - R;
742
743 const float Co = SQRT3 * (Ro - Go);
744 const float Ho = 2.0f * Bo - Go - Ro;
745
746 if(R != G && G != B)
747 {
748 const float ratio = sqrtf((Co * Co + Ho * Ho) / (C * C + H * H));
749 C *= ratio;
750 H *= ratio;
751 }
752
753 float RGB[3] = { 0.0f, 0.0f, 0.0f };
754
755 RGB[0] = L - H / 6.0f + C / SQRT12;
756 RGB[1] = L - H / 6.0f - C / SQRT12;
757 RGB[2] = L + H / 3.0f;
758
759 out[0] = RGB[FCxtrans(j, i, roi_out, xtrans)];
760 }
761 else
762 out[0] = in[0];
763 }
764 out++;
765 in++;
766 }
767 }
768 }
769
770 #undef SQRT3
771 #undef SQRT12
772
process_clip_plain(dt_dev_pixelpipe_iop_t * piece,const void * const ivoid,void * const ovoid,const dt_iop_roi_t * const roi_in,const dt_iop_roi_t * const roi_out,const float clip)773 static void process_clip_plain(dt_dev_pixelpipe_iop_t *piece, const void *const ivoid, void *const ovoid,
774 const dt_iop_roi_t *const roi_in, const dt_iop_roi_t *const roi_out,
775 const float clip)
776 {
777 const float *const in = (const float *const)ivoid;
778 float *const out = (float *const)ovoid;
779
780 if(piece->pipe->dsc.filters)
781 { // raw mosaic
782 #ifdef _OPENMP
783 #pragma omp parallel for SIMD() default(none) \
784 dt_omp_firstprivate(clip, in, out, roi_out) \
785 schedule(static)
786 #endif
787 for(size_t k = 0; k < (size_t)roi_out->width * roi_out->height; k++)
788 {
789 out[k] = MIN(clip, in[k]);
790 }
791 }
792 else
793 {
794 const int ch = piece->colors;
795
796 #ifdef _OPENMP
797 #pragma omp parallel for SIMD() default(none) \
798 dt_omp_firstprivate(ch, clip, in, out, roi_out) \
799 schedule(static)
800 #endif
801 for(size_t k = 0; k < (size_t)ch * roi_out->width * roi_out->height; k++)
802 {
803 out[k] = MIN(clip, in[k]);
804 }
805 }
806 }
807
808 #if defined(__SSE__)
process_clip_sse2(dt_dev_pixelpipe_iop_t * piece,const void * const ivoid,void * const ovoid,const dt_iop_roi_t * const roi_in,const dt_iop_roi_t * const roi_out,const float clip)809 static void process_clip_sse2(dt_dev_pixelpipe_iop_t *piece, const void *const ivoid, void *const ovoid,
810 const dt_iop_roi_t *const roi_in, const dt_iop_roi_t *const roi_out,
811 const float clip)
812 {
813 if(piece->pipe->dsc.filters)
814 { // raw mosaic
815 const __m128 clipm = _mm_set1_ps(clip);
816 const size_t n = (size_t)roi_out->height * roi_out->width;
817 float *const out = (float *)ovoid;
818 float *const in = (float *)ivoid;
819 #ifdef _OPENMP
820 #pragma omp parallel for default(none) \
821 dt_omp_firstprivate(clipm, in, n, out) \
822 schedule(static)
823 #endif
824 for(size_t j = 0; j < (n & ~3u); j += 4) _mm_stream_ps(out + j, _mm_min_ps(clipm, _mm_load_ps(in + j)));
825 _mm_sfence();
826 // lets see if there's a non-multiple of four rest to process:
827 if(n & 3)
828 for(size_t j = n & ~3u; j < n; j++) out[j] = MIN(clip, in[j]);
829 }
830 else
831 {
832 const __m128 clipm = _mm_set1_ps(clip);
833 const int ch = piece->colors;
834
835 #ifdef _OPENMP
836 #pragma omp parallel for default(none) \
837 dt_omp_firstprivate(ch, clipm, ivoid, ovoid, roi_in, roi_out) \
838 schedule(static)
839 #endif
840 for(int j = 0; j < roi_out->height; j++)
841 {
842 float *out = (float *)ovoid + (size_t)ch * roi_out->width * j;
843 float *in = (float *)ivoid + (size_t)ch * roi_in->width * j;
844 for(int i = 0; i < roi_out->width; i++, in += ch, out += ch)
845 {
846 _mm_stream_ps(out, _mm_min_ps(clipm, _mm_set_ps(in[3], in[2], in[1], in[0])));
847 }
848 }
849 _mm_sfence();
850 }
851 }
852 #endif
853
process_clip(dt_dev_pixelpipe_iop_t * piece,const void * const ivoid,void * const ovoid,const dt_iop_roi_t * const roi_in,const dt_iop_roi_t * const roi_out,const float clip)854 static void process_clip(dt_dev_pixelpipe_iop_t *piece, const void *const ivoid, void *const ovoid,
855 const dt_iop_roi_t *const roi_in, const dt_iop_roi_t *const roi_out,
856 const float clip)
857 {
858 if(darktable.codepath.OPENMP_SIMD) process_clip_plain(piece, ivoid, ovoid, roi_in, roi_out, clip);
859 #if defined(__SSE__)
860 else if(darktable.codepath.SSE2)
861 process_clip_sse2(piece, ivoid, ovoid, roi_in, roi_out, clip);
862 #endif
863 else
864 dt_unreachable_codepath();
865 }
866
process(struct dt_iop_module_t * self,dt_dev_pixelpipe_iop_t * piece,const void * const ivoid,void * const ovoid,const dt_iop_roi_t * const roi_in,const dt_iop_roi_t * const roi_out)867 void process(struct dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, const void *const ivoid,
868 void *const ovoid, const dt_iop_roi_t *const roi_in, const dt_iop_roi_t *const roi_out)
869 {
870 const uint32_t filters = piece->pipe->dsc.filters;
871 dt_iop_highlights_data_t *data = (dt_iop_highlights_data_t *)piece->data;
872
873 const float clip
874 = data->clip * fminf(piece->pipe->dsc.processed_maximum[0],
875 fminf(piece->pipe->dsc.processed_maximum[1], piece->pipe->dsc.processed_maximum[2]));
876 // const int ch = piece->colors;
877 if(!filters)
878 {
879 process_clip(piece, ivoid, ovoid, roi_in, roi_out, clip);
880 for(int k=0;k<3;k++)
881 piece->pipe->dsc.processed_maximum[k]
882 = fminf(piece->pipe->dsc.processed_maximum[0],
883 fminf(piece->pipe->dsc.processed_maximum[1], piece->pipe->dsc.processed_maximum[2]));
884 return;
885 }
886
887 switch(data->mode)
888 {
889 case DT_IOP_HIGHLIGHTS_INPAINT: // a1ex's (magiclantern) idea of color inpainting:
890 {
891 const float clips[4] = { 0.987 * data->clip * piece->pipe->dsc.processed_maximum[0],
892 0.987 * data->clip * piece->pipe->dsc.processed_maximum[1],
893 0.987 * data->clip * piece->pipe->dsc.processed_maximum[2], clip };
894
895 if(filters == 9u)
896 {
897 const uint8_t(*const xtrans)[6] = (const uint8_t(*const)[6])piece->pipe->dsc.xtrans;
898 #ifdef _OPENMP
899 #pragma omp parallel for default(none) \
900 dt_omp_firstprivate(clips, filters, ivoid, ovoid, roi_in, roi_out, \
901 xtrans) \
902 schedule(static)
903 #endif
904 for(int j = 0; j < roi_out->height; j++)
905 {
906 interpolate_color_xtrans(ivoid, ovoid, roi_in, roi_out, 0, 1, j, clips, xtrans, 0);
907 interpolate_color_xtrans(ivoid, ovoid, roi_in, roi_out, 0, -1, j, clips, xtrans, 1);
908 }
909 #ifdef _OPENMP
910 #pragma omp parallel for default(none) \
911 dt_omp_firstprivate(clips, filters, ivoid, ovoid, roi_in, roi_out, \
912 xtrans) \
913 schedule(static)
914 #endif
915 for(int i = 0; i < roi_out->width; i++)
916 {
917 interpolate_color_xtrans(ivoid, ovoid, roi_in, roi_out, 1, 1, i, clips, xtrans, 2);
918 interpolate_color_xtrans(ivoid, ovoid, roi_in, roi_out, 1, -1, i, clips, xtrans, 3);
919 }
920 }
921 else
922 {
923 #ifdef _OPENMP
924 #pragma omp parallel for default(none) \
925 dt_omp_firstprivate(clips, filters, ivoid, ovoid, roi_out) \
926 shared(data, piece) \
927 schedule(static)
928 #endif
929 for(int j = 0; j < roi_out->height; j++)
930 {
931 interpolate_color(ivoid, ovoid, roi_out, 0, 1, j, clips, filters, 0);
932 interpolate_color(ivoid, ovoid, roi_out, 0, -1, j, clips, filters, 1);
933 }
934
935 // up/down directions
936 #ifdef _OPENMP
937 #pragma omp parallel for default(none) \
938 dt_omp_firstprivate(clips, filters, ivoid, ovoid, roi_out) \
939 shared(data, piece) \
940 schedule(static)
941 #endif
942 for(int i = 0; i < roi_out->width; i++)
943 {
944 interpolate_color(ivoid, ovoid, roi_out, 1, 1, i, clips, filters, 2);
945 interpolate_color(ivoid, ovoid, roi_out, 1, -1, i, clips, filters, 3);
946 }
947 }
948 break;
949 }
950 case DT_IOP_HIGHLIGHTS_LCH:
951 if(filters == 9u)
952 process_lch_xtrans(self, piece, ivoid, ovoid, roi_in, roi_out, clip);
953 else
954 process_lch_bayer(self, piece, ivoid, ovoid, roi_in, roi_out, clip);
955 break;
956 default:
957 case DT_IOP_HIGHLIGHTS_CLIP:
958 process_clip(piece, ivoid, ovoid, roi_in, roi_out, clip);
959 break;
960 }
961
962 // update processed maximum
963 const float m = fmaxf(fmaxf(piece->pipe->dsc.processed_maximum[0], piece->pipe->dsc.processed_maximum[1]),
964 piece->pipe->dsc.processed_maximum[2]);
965 for(int k = 0; k < 3; k++) piece->pipe->dsc.processed_maximum[k] = m;
966
967 if(piece->pipe->mask_display & DT_DEV_PIXELPIPE_DISPLAY_MASK) dt_iop_alpha_copy(ivoid, ovoid, roi_out->width, roi_out->height);
968 }
969
commit_params(struct dt_iop_module_t * self,dt_iop_params_t * p1,dt_dev_pixelpipe_t * pipe,dt_dev_pixelpipe_iop_t * piece)970 void commit_params(struct dt_iop_module_t *self, dt_iop_params_t *p1, dt_dev_pixelpipe_t *pipe,
971 dt_dev_pixelpipe_iop_t *piece)
972 {
973 dt_iop_highlights_params_t *p = (dt_iop_highlights_params_t *)p1;
974 dt_iop_highlights_data_t *d = (dt_iop_highlights_data_t *)piece->data;
975 memcpy(d, p, sizeof(*p));
976
977 piece->process_cl_ready = 1;
978
979 // no OpenCL for DT_IOP_HIGHLIGHTS_INPAINT yet.
980 if(d->mode == DT_IOP_HIGHLIGHTS_INPAINT) piece->process_cl_ready = 0;
981 }
982
init_global(dt_iop_module_so_t * module)983 void init_global(dt_iop_module_so_t *module)
984 {
985 const int program = 2; // basic.cl, from programs.conf
986 dt_iop_highlights_global_data_t *gd
987 = (dt_iop_highlights_global_data_t *)malloc(sizeof(dt_iop_highlights_global_data_t));
988 module->data = gd;
989 gd->kernel_highlights_1f_clip = dt_opencl_create_kernel(program, "highlights_1f_clip");
990 gd->kernel_highlights_1f_lch_bayer = dt_opencl_create_kernel(program, "highlights_1f_lch_bayer");
991 gd->kernel_highlights_1f_lch_xtrans = dt_opencl_create_kernel(program, "highlights_1f_lch_xtrans");
992 gd->kernel_highlights_4f_clip = dt_opencl_create_kernel(program, "highlights_4f_clip");
993 }
994
cleanup_global(dt_iop_module_so_t * module)995 void cleanup_global(dt_iop_module_so_t *module)
996 {
997 dt_iop_highlights_global_data_t *gd = (dt_iop_highlights_global_data_t *)module->data;
998 dt_opencl_free_kernel(gd->kernel_highlights_4f_clip);
999 dt_opencl_free_kernel(gd->kernel_highlights_1f_lch_bayer);
1000 dt_opencl_free_kernel(gd->kernel_highlights_1f_lch_xtrans);
1001 dt_opencl_free_kernel(gd->kernel_highlights_1f_clip);
1002 free(module->data);
1003 module->data = NULL;
1004 }
1005
init_pipe(struct dt_iop_module_t * self,dt_dev_pixelpipe_t * pipe,dt_dev_pixelpipe_iop_t * piece)1006 void init_pipe(struct dt_iop_module_t *self, dt_dev_pixelpipe_t *pipe, dt_dev_pixelpipe_iop_t *piece)
1007 {
1008 piece->data = malloc(sizeof(dt_iop_highlights_data_t));
1009 }
1010
cleanup_pipe(struct dt_iop_module_t * self,dt_dev_pixelpipe_t * pipe,dt_dev_pixelpipe_iop_t * piece)1011 void cleanup_pipe(struct dt_iop_module_t *self, dt_dev_pixelpipe_t *pipe, dt_dev_pixelpipe_iop_t *piece)
1012 {
1013 free(piece->data);
1014 piece->data = NULL;
1015 }
1016
gui_update(struct dt_iop_module_t * self)1017 void gui_update(struct dt_iop_module_t *self)
1018 {
1019 dt_iop_highlights_gui_data_t *g = (dt_iop_highlights_gui_data_t *)self->gui_data;
1020 dt_iop_highlights_params_t *p = (dt_iop_highlights_params_t *)self->params;
1021 dt_bauhaus_slider_set(g->clip, p->clip);
1022 dt_bauhaus_combobox_set(g->mode, p->mode);
1023 }
1024
reload_defaults(dt_iop_module_t * module)1025 void reload_defaults(dt_iop_module_t *module)
1026 {
1027 // enable this per default if raw or sraw,
1028 module->default_enabled = dt_image_is_rawprepare_supported(&(module->dev->image_storage));
1029 }
1030
gui_init(struct dt_iop_module_t * self)1031 void gui_init(struct dt_iop_module_t *self)
1032 {
1033 dt_iop_highlights_gui_data_t *g = IOP_GUI_ALLOC(highlights);
1034
1035 g->mode = dt_bauhaus_combobox_from_params(self, "mode");
1036 gtk_widget_set_tooltip_text(g->mode, _("highlight reconstruction method"));
1037
1038 g->clip = dt_bauhaus_slider_from_params(self, "clip");
1039 dt_bauhaus_slider_set_digits(g->clip, 3);
1040 gtk_widget_set_tooltip_text(g->clip, _("manually adjust the clipping threshold against "
1041 "magenta highlights (you shouldn't ever need to touch this)"));
1042 }
1043
1044 // modelines: These editor modelines have been set for all relevant files by tools/update_modelines.sh
1045 // vim: shiftwidth=2 expandtab tabstop=2 cindent
1046 // kate: tab-indents: off; indent-width 2; replace-tabs on; indent-mode cstyle; remove-trailing-spaces modified;
1047