highlights.c 36.6 KB
Newer Older
1
/*
2 3
   This file is part of darktable,
   copyright (c) 2009--2010 johannes hanika.
4

5 6 7 8
   darktable is free software: you can redistribute it and/or modify
   it under the terms of the GNU General Public License as published by
   the Free Software Foundation, either version 3 of the License, or
   (at your option) any later version.
9

10 11 12 13
   darktable 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 General Public License for more details.
14

15 16 17
   You should have received a copy of the GNU General Public License
   along with darktable.  If not, see <http://www.gnu.org/licenses/>.
 */
18
#ifdef HAVE_CONFIG_H
19
#include "config.h"
20 21
#endif
#include <assert.h>
22 23
#include <math.h>
#include <stdlib.h>
24
#include <string.h>
25
#if defined(__SSE__)
26
#include <xmmintrin.h>
27
#endif
28 29 30
#include "bauhaus/bauhaus.h"
#include "common/opencl.h"
#include "control/control.h"
31 32
#include "develop/develop.h"
#include "develop/imageop.h"
33
#include "develop/imageop_math.h"
34
#include "develop/tiling.h"
35
#include "gui/accelerators.h"
36
#include "gui/gtk.h"
37
#include "iop/iop_api.h"
38

39 40 41
#include <gtk/gtk.h>
#include <inttypes.h>

42

43
DT_MODULE_INTROSPECTION(2, dt_iop_highlights_params_t)
44 45 46

typedef enum dt_iop_highlights_mode_t
{
47
  DT_IOP_HIGHLIGHTS_CLIP = 0,
48 49
  DT_IOP_HIGHLIGHTS_LCH = 1,
  DT_IOP_HIGHLIGHTS_INPAINT = 2,
50
} dt_iop_highlights_mode_t;
51 52 53

typedef struct dt_iop_highlights_params_t
{
54
  dt_iop_highlights_mode_t mode;
55
  float blendL, blendC, blendh; // unused
56
  float clip;
57
} dt_iop_highlights_params_t;
58 59 60

typedef struct dt_iop_highlights_gui_data_t
{
61
  GtkWidget *clip;
62
  GtkWidget *mode;
63
} dt_iop_highlights_gui_data_t;
64

65
typedef dt_iop_highlights_params_t dt_iop_highlights_data_t;
66

67 68
typedef struct dt_iop_highlights_global_data_t
{
69
  int kernel_highlights_1f_clip;
70 71
  int kernel_highlights_1f_lch_bayer;
  int kernel_highlights_1f_lch_xtrans;
72
  int kernel_highlights_4f_clip;
73
} dt_iop_highlights_global_data_t;
74

75 76
const char *name()
{
77
  return _("highlight reconstruction");
78 79
}

80
int groups()
81
{
82
  return IOP_GROUP_BASIC;
83 84
}

85
int flags()
86
{
87
  return IOP_FLAGS_SUPPORTS_BLENDING | IOP_FLAGS_ALLOW_TILING | IOP_FLAGS_ONE_INSTANCE;
88
}
89

90 91
int legacy_params(dt_iop_module_t *self, const void *const old_params, const int old_version,
                  void *new_params, const int new_version)
92 93 94
{
  if(old_version == 1 && new_version == 2)
  {
95
    memcpy(new_params, old_params, sizeof(dt_iop_highlights_params_t) - sizeof(float));
96 97 98 99 100 101 102
    dt_iop_highlights_params_t *n = (dt_iop_highlights_params_t *)new_params;
    n->clip = 1.0f;
    return 0;
  }
  return 1;
}

103
#ifdef HAVE_OPENCL
104
int process_cl(struct dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, cl_mem dev_in, cl_mem dev_out,
105
               const dt_iop_roi_t *const roi_in, const dt_iop_roi_t *const roi_out)
106 107 108 109
{
  dt_iop_highlights_data_t *d = (dt_iop_highlights_data_t *)piece->data;
  dt_iop_highlights_global_data_t *gd = (dt_iop_highlights_global_data_t *)self->data;

110
  cl_int err = -999;
111 112
  cl_mem dev_xtrans = NULL;

113
  const int devid = piece->pipe->devid;
114 115 116
  const int width = roi_in->width;
  const int height = roi_in->height;

117 118 119
  const float clip = d->clip
                     * fminf(piece->pipe->dsc.processed_maximum[0],
                             fminf(piece->pipe->dsc.processed_maximum[1], piece->pipe->dsc.processed_maximum[2]));
120

121
  const uint32_t filters = piece->pipe->dsc.filters;
122

123
  if(!filters)
124
  {
125 126
    // non-raw images use dedicated kernel which just clips
    size_t sizes[] = { ROUNDUPWD(width), ROUNDUPHT(height), 1 };
127 128 129 130 131 132 133
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_4f_clip, 0, sizeof(cl_mem), (void *)&dev_in);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_4f_clip, 1, sizeof(cl_mem), (void *)&dev_out);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_4f_clip, 2, sizeof(int), (void *)&width);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_4f_clip, 3, sizeof(int), (void *)&height);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_4f_clip, 4, sizeof(int), (void *)&d->mode);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_4f_clip, 5, sizeof(float), (void *)&clip);
    err = dt_opencl_enqueue_kernel_2d(devid, gd->kernel_highlights_4f_clip, sizes);
134 135
    if(err != CL_SUCCESS) goto error;
  }
136 137 138 139 140 141 142 143 144 145 146 147 148 149 150 151 152 153 154 155 156 157 158 159 160 161 162 163 164 165 166
  else if(d->mode == DT_IOP_HIGHLIGHTS_CLIP)
  {
    // raw images with clip mode (both bayer and xtrans)
    size_t sizes[] = { ROUNDUPWD(width), ROUNDUPHT(height), 1 };
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_clip, 0, sizeof(cl_mem), (void *)&dev_in);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_clip, 1, sizeof(cl_mem), (void *)&dev_out);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_clip, 2, sizeof(int), (void *)&width);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_clip, 3, sizeof(int), (void *)&height);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_clip, 4, sizeof(float), (void *)&clip);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_clip, 5, sizeof(int), (void *)&roi_out->x);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_clip, 6, sizeof(int), (void *)&roi_out->y);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_clip, 7, sizeof(int), (void *)&filters);
    err = dt_opencl_enqueue_kernel_2d(devid, gd->kernel_highlights_1f_clip, sizes);
    if(err != CL_SUCCESS) goto error;
  }
  else if(d->mode == DT_IOP_HIGHLIGHTS_LCH && filters != 9u)
  {
    // bayer sensor raws with LCH mode
    size_t sizes[] = { ROUNDUPWD(width), ROUNDUPHT(height), 1 };
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_bayer, 0, sizeof(cl_mem), (void *)&dev_in);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_bayer, 1, sizeof(cl_mem), (void *)&dev_out);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_bayer, 2, sizeof(int), (void *)&width);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_bayer, 3, sizeof(int), (void *)&height);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_bayer, 4, sizeof(float), (void *)&clip);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_bayer, 5, sizeof(int), (void *)&roi_out->x);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_bayer, 6, sizeof(int), (void *)&roi_out->y);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_bayer, 7, sizeof(int), (void *)&filters);
    err = dt_opencl_enqueue_kernel_2d(devid, gd->kernel_highlights_1f_lch_bayer, sizes);
    if(err != CL_SUCCESS) goto error;
  }
  else if(d->mode == DT_IOP_HIGHLIGHTS_LCH && filters == 9u)
167
  {
168
    // xtrans sensor raws with LCH mode
169
    int blocksizex, blocksizey;
170

171 172 173 174
    dt_opencl_local_buffer_t locopt
      = (dt_opencl_local_buffer_t){ .xoffset = 2 * 2, .xfactor = 1, .yoffset = 2 * 2, .yfactor = 1,
                                    .cellsize = sizeof(float), .overhead = 0,
                                    .sizex = 1 << 8, .sizey = 1 << 8 };
175

176
    if(dt_opencl_local_buffer_opt(devid, gd->kernel_highlights_1f_lch_xtrans, &locopt))
177
    {
178 179
      blocksizex = locopt.sizex;
      blocksizey = locopt.sizey;
180 181
    }
    else
182
      blocksizex = blocksizey = 1;
183 184 185 186 187 188 189 190 191 192 193 194 195 196 197 198 199 200 201

    dev_xtrans
        = dt_opencl_copy_host_to_device_constant(devid, sizeof(piece->pipe->dsc.xtrans), piece->pipe->dsc.xtrans);
    if(dev_xtrans == NULL) goto error;

    size_t sizes[] = { ROUNDUP(width, blocksizex), ROUNDUP(height, blocksizey), 1 };
    size_t local[] = { blocksizex, blocksizey, 1 };
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_xtrans, 0, sizeof(cl_mem), (void *)&dev_in);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_xtrans, 1, sizeof(cl_mem), (void *)&dev_out);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_xtrans, 2, sizeof(int), (void *)&width);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_xtrans, 3, sizeof(int), (void *)&height);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_xtrans, 4, sizeof(float), (void *)&clip);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_xtrans, 5, sizeof(int), (void *)&roi_out->x);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_xtrans, 6, sizeof(int), (void *)&roi_out->y);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_xtrans, 7, sizeof(cl_mem), (void *)&dev_xtrans);
    dt_opencl_set_kernel_arg(devid, gd->kernel_highlights_1f_lch_xtrans, 8,
                               (blocksizex + 4) * (blocksizey + 4) * sizeof(float), NULL);

    err = dt_opencl_enqueue_kernel_2d_with_local(devid, gd->kernel_highlights_1f_lch_xtrans, sizes, local);
202 203
    if(err != CL_SUCCESS) goto error;
  }
204 205

  // update processed maximum
206 207 208
  const float m = fmaxf(fmaxf(piece->pipe->dsc.processed_maximum[0], piece->pipe->dsc.processed_maximum[1]),
                        piece->pipe->dsc.processed_maximum[2]);
  for(int k = 0; k < 3; k++) piece->pipe->dsc.processed_maximum[k] = m;
209

210
  dt_opencl_release_mem_object(dev_xtrans);
211 212 213
  return TRUE;

error:
214
  dt_opencl_release_mem_object(dev_xtrans);
215 216
  dt_print(DT_DEBUG_OPENCL, "[opencl_highlights] couldn't enqueue kernel! %d\n", err);
  return FALSE;
217 218
}
#endif
219

220 221 222 223 224 225 226 227 228 229 230 231 232 233 234 235 236 237 238 239 240 241 242 243 244 245 246 247 248 249 250 251 252 253
void 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)
{
  dt_iop_highlights_data_t *d = (dt_iop_highlights_data_t *)piece->data;
  const uint32_t filters = piece->pipe->dsc.filters;

  tiling->factor = 2.0f;  // in + out
  tiling->maxbuf = 1.0f;
  tiling->overhead = 0;

  if(filters == 9u)
  {
    // xtrans
    tiling->xalign = 6;
    tiling->yalign = 6;
    tiling->overlap = (d->mode == DT_IOP_HIGHLIGHTS_LCH) ? 2 : 0;
  }
  else if(filters)
  {
    // bayer
    tiling->xalign = 2;
    tiling->yalign = 2;
    tiling->overlap = (d->mode == DT_IOP_HIGHLIGHTS_LCH) ? 1 : 0;
  }
  else
  {
    // non-raw
    tiling->xalign = 1;
    tiling->yalign = 1;
    tiling->overlap = 0;
  }
}

254 255 256 257 258 259
/* interpolate value for a pixel, ideal via ratio to nearby pixel */
static inline float 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)
260
{
261 262 263 264 265 266
  assert(ratio_next != 0);
  // it's OK to exceed clipping of current pixel's color based on a
  // neighbor -- that is the purpose of interpolating highlight
  // colors
  const float clip_val = fmaxf(clip0, clip_next);
  if(in[offset_next] >= clip_next - 1e-5f)
267
  {
268 269
    // next pixel is also clipped
    return clip_val;
270
  }
271
  else
272
  {
273 274 275 276 277 278
    // set this pixel in ratio to the next
    assert(ratio_next != 0);
    if (ratio_next > 0)
      return fminf(in[offset_next] / ratios[ratio_next], clip_val);
    else
      return fminf(in[offset_next] * ratios[-ratio_next], clip_val);
279
  }
280 281 282 283 284 285 286 287 288 289 290 291 292 293 294 295 296 297 298 299 300 301 302 303 304 305 306 307 308 309 310 311 312 313 314 315 316
}

static inline void 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)
{
  // In Bayer each row/col has only green/red or green/blue
  // transitions, hence can reconstruct color by single ratio per
  // row. In x-trans there can be transitions between arbitary colors
  // in a row/col (and 2x2 green blocks which provide no color
  // transition information). Hence calculate multiple color ratios
  // for each row/col.

  // Lookup for color ratios, e.g. red -> blue is roff[0][2] and blue
  // -> red is roff[2][0]. Returned value is an index into ratios. If
  // negative, then need to invert the ratio. Identity color
  // transitions aren't used.
  const int roff[3][3] = {{ 0, -1, -2},
                          { 1,  0, -3},
                          { 2,  3,  0}};
  // record ratios of color transitions 0:unused, 1:RG, 2:RB, and 3:GB
  float ratios[4] = {1.0f, 1.0f, 1.0f, 1.0f};

  // passes are 0:+x, 1:-x, 2:+y, 3:-y
  // dims are 0:traverse a row, 1:traverse a column
  // dir is 1:left to right, -1: right to left
  int i = (dim == 0) ? 0 : other;
  int j = (dim == 0) ? other : 0;
  const ssize_t offs = (dim ? roi_out->width : 1) * ((dir < 0) ? -1 : 1);
  const ssize_t offl = offs - (dim ? 1 : roi_out->width);
  const ssize_t offr = offs + (dim ? 1 : roi_out->width);
  int beg, end;
  if(dir == 1)
317 318
  {
    beg = 0;
319
    end = (dim == 0) ? roi_out->width : roi_out->height;
320
  }
321
  else
322
  {
323
    beg = ((dim == 0) ? roi_out->width : roi_out->height) - 1;
324 325
    end = -1;
  }
326

327
  float *in, *out;
328 329
  if(dim == 1)
  {
330 331
    out = (float *)ovoid + (size_t)i + (size_t)beg * roi_out->width;
    in = (float *)ivoid + (size_t)i + (size_t)beg * roi_in->width;
332 333 334
  }
  else
  {
335 336
    out = (float *)ovoid + (size_t)beg + (size_t)j * roi_out->width;
    in = (float *)ivoid + (size_t)beg + (size_t)j * roi_in->width;
337 338
  }

339
  for(int k = beg; k != end; k += dir)
340
  {
341 342 343 344
    if(dim == 1)
      j = k;
    else
      i = k;
345 346 347 348 349 350 351 352 353 354 355 356

    const uint8_t f0 = FCxtrans(j, i, roi_in, xtrans);
    const uint8_t f1 = FCxtrans(dim ? (j + dir) : j, dim ? i : (i + dir), roi_in, xtrans);
    const uint8_t fl = FCxtrans(dim ? (j + dir) : (j - 1), dim ? (i - 1) : (i + dir), roi_in, xtrans);
    const uint8_t fr = FCxtrans(dim ? (j + dir) : (j + 1), dim ? (i + 1) : (i + dir), roi_in, xtrans);
    const float clip0 = clip[f0];
    const float clip1 = clip[f1];
    const float clipl = clip[fl];
    const float clipr = clip[fr];
    const float clip_max = fmaxf(fmaxf(clip[0], clip[1]), clip[2]);

    if(i == 0 || i == roi_out->width - 1 || j == 0 || j == roi_out->height - 1)
357
    {
358
      if(pass == 3) out[0] = fminf(clip_max, in[0]);
359 360 361
    }
    else
    {
362 363 364 365 366
      // ratio to next pixel if this & next are unclamped and not in
      // 2x2 green block
      if ((f0 != f1) &&
          (in[0] < clip0 && in[0] > 1e-5f) &&
          (in[offs] < clip1 && in[offs] > 1e-5f))
367
      {
368 369 370 371 372 373
        const int r = roff[f0][f1];
        assert(r != 0);
        if (r > 0)
          ratios[r] = (3.f * ratios[r] + (in[offs] / in[0])) / 4.f;
        else
          ratios[-r] = (3.f * ratios[-r] + (in[0] / in[offs])) / 4.f;
374 375
      }

376
      if(in[0] >= clip0 - 1e-5f)
377
      {
378 379 380 381 382 383
        // interplate color for clipped pixel
        float add;
        if(f0 != f1)
          // next pixel is different color
          add =
            interp_pix_xtrans(roff[f0][f1], offs, clip0, clip1, in, ratios);
384
        else
385 386 387 388
          // at start of 2x2 green block, look diagonally
          add = (fl != f0) ?
            interp_pix_xtrans(roff[f0][fl], offl, clip0, clipl, in, ratios) :
            interp_pix_xtrans(roff[f0][fr], offr, clip0, clipr, in, ratios);
389

390 391 392
        if(pass == 0)
          out[0] = add;
        else if(pass == 3)
393
          out[0] = fminf(clip_max, (out[0] + add) / 4.0f);
394 395
        else
          out[0] += add;
396 397 398
      }
      else
      {
399
        // pixel is not clipped
400
        if(pass == 3) out[0] = in[0];
401 402 403 404 405 406 407
      }
    }
    out += offs;
    in += offs;
  }
}

408 409 410
static inline void 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)
411 412 413 414 415
{
  float ratio = 1.0f;
  float *in, *out;

  int i = 0, j = 0;
416 417 418 419
  if(dim == 0)
    j = other;
  else
    i = other;
420
  ssize_t offs = dim ? roi_out->width : 1;
421
  if(dir < 0) offs = -offs;
422
  int beg, end;
423 424 425 426 427 428 429 430 431 432 433 434 435 436 437 438 439 440 441 442 443 444
  if(dim == 0 && dir == 1)
  {
    beg = 0;
    end = roi_out->width;
  }
  else if(dim == 0 && dir == -1)
  {
    beg = roi_out->width - 1;
    end = -1;
  }
  else if(dim == 1 && dir == 1)
  {
    beg = 0;
    end = roi_out->height;
  }
  else if(dim == 1 && dir == -1)
  {
    beg = roi_out->height - 1;
    end = -1;
  }
  else
    return;
445 446 447

  if(dim == 1)
  {
448 449
    out = (float *)ovoid + i + (size_t)beg * roi_out->width;
    in = (float *)ivoid + i + (size_t)beg * roi_out->width;
450 451 452
  }
  else
  {
453 454
    out = (float *)ovoid + beg + (size_t)j * roi_out->width;
    in = (float *)ivoid + beg + (size_t)j * roi_out->width;
455
  }
456
  for(int k = beg; k != end; k += dir)
457
  {
458 459 460 461
    if(dim == 1)
      j = k;
    else
      i = k;
462
    const float clip0 = clip[FC(j, i, filters)];
463 464
    const float clip1 = clip[FC(dim ? (j + 1) : j, dim ? i : (i + 1), filters)];
    if(i == 0 || i == roi_out->width - 1 || j == 0 || j == roi_out->height - 1)
465
    {
466
      if(pass == 3) out[0] = in[0];
467 468 469
    }
    else
    {
470
      if(in[0] < clip0 && in[0] > 1e-5f)
471
      { // both are not clipped
472
        if(in[offs] < clip1 && in[offs] > 1e-5f)
473
        { // update ratio, exponential decay. ratio = in[odd]/in[even]
474 475 476 477
          if(k & 1)
            ratio = (3.0f * ratio + in[0] / in[offs]) / 4.0f;
          else
            ratio = (3.0f * ratio + in[offs] / in[0]) / 4.0f;
478 479 480
        }
      }

481
      if(in[0] >= clip0 - 1e-5f)
482 483
      { // in[0] is clipped, restore it as in[1] adjusted according to ratio
        float add = 0.0f;
484 485 486 487 488 489
        if(in[offs] >= clip1 - 1e-5f)
          add = fmaxf(clip0, clip1);
        else if(k & 1)
          add = in[offs] * ratio;
        else
          add = in[offs] / ratio;
490

491 492 493 494 495 496
        if(pass == 0)
          out[0] = add;
        else if(pass == 3)
          out[0] = (out[0] + add) / 4.0f;
        else
          out[0] += add;
497 498 499
      }
      else
      {
500
        if(pass == 3) out[0] = in[0];
501 502 503 504 505 506
      }
    }
    out += offs;
    in += offs;
  }
}
507

508 509 510 511 512 513 514 515 516 517 518 519 520 521 522
/*
 * these 2 constants were computed using following Sage code:
 *
 * sqrt3 = sqrt(3)
 * sqrt12 = sqrt(12) # 2*sqrt(3)
 *
 * print 'sqrt3 = ', sqrt3, ' ~= ', RealField(128)(sqrt3)
 * print 'sqrt12 = ', sqrt12, ' ~= ', RealField(128)(sqrt12)
 */
#define SQRT3 1.7320508075688772935274463415058723669L
#define SQRT12 3.4641016151377545870548926830117447339L // 2*SQRT3

static void 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)
Dan Torop's avatar
Dan Torop committed
523
{
524
  const uint32_t filters = piece->pipe->dsc.filters;
525

Dan Torop's avatar
Dan Torop committed
526 527 528
#ifdef _OPENMP
#pragma omp parallel for schedule(dynamic) default(none)
#endif
529
  for(int j = 0; j < roi_out->height; j++)
Dan Torop's avatar
Dan Torop committed
530
  {
531 532 533 534
    float *out = (float *)ovoid + (size_t)roi_out->width * j;
    float *in = (float *)ivoid + (size_t)roi_out->width * j;

    for(int i = 0; i < roi_out->width; i++, in++, out++)
Dan Torop's avatar
Dan Torop committed
535
    {
536
      if(i == roi_out->width - 1 || j == roi_out->height - 1)
Dan Torop's avatar
Dan Torop committed
537 538
      {
        // fast path for border
539
        out[0] = MIN(clip, in[0]);
Dan Torop's avatar
Dan Torop committed
540 541 542
      }
      else
      {
543 544 545 546
        int clipped = 0;

        // sample 1 bayer block. thus we will have 2 green values.
        float R = 0.0f, Gmin = FLT_MAX, Gmax = -FLT_MAX, B = 0.0f;
Dan Torop's avatar
Dan Torop committed
547 548 549 550
        for(int jj = 0; jj <= 1; jj++)
        {
          for(int ii = 0; ii <= 1; ii++)
          {
551 552 553 554 555 556 557 558 559 560 561 562 563 564 565 566 567 568
            const float val = in[(size_t)jj * roi_out->width + ii];

            clipped = (clipped || (val > clip));

            const int c = FC(j + jj + roi_out->y, i + ii + roi_out->x, filters);
            switch(c)
            {
              case 0:
                R = val;
                break;
              case 1:
                Gmin = MIN(Gmin, val);
                Gmax = MAX(Gmax, val);
                break;
              case 2:
                B = val;
                break;
            }
Dan Torop's avatar
Dan Torop committed
569 570
          }
        }
571 572

        if(clipped)
Dan Torop's avatar
Dan Torop committed
573
        {
574 575 576 577 578 579 580 581 582 583 584 585 586 587 588 589 590 591 592 593 594 595 596 597 598 599 600 601 602 603 604 605 606 607 608
          const float Ro = MIN(R, clip);
          const float Go = MIN(Gmin, clip);
          const float Bo = MIN(B, clip);

          const float L = (R + Gmax + B) / 3.0f;

          float C = SQRT3 * (R - Gmax);
          float H = 2.0f * B - Gmax - R;

          const float Co = SQRT3 * (Ro - Go);
          const float Ho = 2.0f * Bo - Go - Ro;

          if(R != Gmax && Gmax != B)
          {
            const float ratio = sqrtf((Co * Co + Ho * Ho) / (C * C + H * H));
            C *= ratio;
            H *= ratio;
          }

          float RGB[3] = { 0.0f, 0.0f, 0.0f };

          /*
           * backtransform proof, sage:
           *
           * R,G,B,L,C,H = var('R,G,B,L,C,H')
           * solve([L==(R+G+B)/3, C==sqrt(3)*(R-G), H==2*B-G-R], R, G, B)
           *
           * result:
           * [[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]]
           */
          RGB[0] = L - H / 6.0f + C / SQRT12;
          RGB[1] = L - H / 6.0f - C / SQRT12;
          RGB[2] = L + H / 3.0f;

          out[0] = RGB[FC(j + roi_out->y, i + roi_out->x, filters)];
Dan Torop's avatar
Dan Torop committed
609 610
        }
        else
611
        {
Dan Torop's avatar
Dan Torop committed
612
          out[0] = in[0];
613
        }
Dan Torop's avatar
Dan Torop committed
614 615 616 617 618
      }
    }
  }
}

619
static void process_lch_xtrans(dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, const void *const ivoid,
620 621
                               void *const ovoid, const dt_iop_roi_t *const roi_in,
                               const dt_iop_roi_t *const roi_out, const float clip)
622
{
623
  const uint8_t(*const xtrans)[6] = (const uint8_t(*const)[6])piece->pipe->dsc.xtrans;
624

625
#ifdef _OPENMP
626
#pragma omp parallel for schedule(dynamic) default(none)
627
#endif
628
  for(int j = 0; j < roi_out->height; j++)
629
  {
630 631 632
    float *out = (float *)ovoid + (size_t)roi_out->width * j;
    float *in = (float *)ivoid + (size_t)roi_in->width * j;

633 634 635 636 637
    // bit vector used as ring buffer to remember clipping of current
    // and last two columns, checking current pixel and its vertical
    // neighbors
    int cl = 0;

638
    for(int i = 0; i < roi_out->width; i++)
639
    {
640 641 642 643 644 645 646
      // update clipping ring buffer
      cl = (cl << 1) & 6;
      if(j >= 2 && j <= roi_out->height - 3)
      {
        cl |= (in[-roi_in->width] > clip) | (in[0] > clip) | (in[roi_in->width] > clip);
      }

647
      if(i < 2 || i > roi_out->width - 3 || j < 2 || j > roi_out->height - 3)
648 649
      {
        // fast path for border
650
        out[0] = MIN(clip, in[0]);
651 652 653
      }
      else
      {
654 655 656
        // if current pixel is clipped, always reconstruct
        int clipped = (in[0] > clip);
        if(!clipped)
657
        {
658 659
          clipped = cl;
          if(clipped)
660
          {
661 662 663 664 665 666 667 668 669
            // If the ring buffer can't show we are in an obviously
            // unclipped region, this is the slow case: check if there
            // is any 3x3 block touching the current pixel which has
            // no clipping, as then don't need to reconstruct the
            // current pixel. This avoids zippering in edge
            // transitions from clipped to unclipped areas. The
            // X-Trans sensor seems prone to this, unlike Bayer, due
            // to its irregular pattern.
            for(int offset_j = -2; offset_j <= 0; offset_j++)
670
            {
671
              for(int offset_i = -2; offset_i <= 0; offset_i++)
672
              {
673
                if(clipped)
674
                {
675 676
                  clipped = 0;
                  for(int jj = offset_j; jj <= offset_j + 2; jj++)
677
                  {
678 679 680 681 682
                    for(int ii = offset_i; ii <= offset_i + 2; ii++)
                    {
                      const float val = in[(ssize_t)jj * roi_in->width + ii];
                      clipped = (clipped || (val > clip));
                    }
683 684 685
                  }
                }
              }
686
            }
687 688
          }
        }
689 690 691

        if(clipped)
        {
692 693 694
          float mean[3] = { 0.0f, 0.0f, 0.0f };
          int cnt[3] = { 0, 0, 0 };
          float RGBmax[3] = { -FLT_MAX, -FLT_MAX, -FLT_MAX };
695

696 697 698 699
          for(int jj = -1; jj <= 1; jj++)
          {
            for(int ii = -1; ii <= 1; ii++)
            {
700
              const float val = in[(ssize_t)jj * roi_in->width + ii];
701 702 703 704 705 706
              const int c = FCxtrans(j+jj, i+ii, roi_in, xtrans);
              mean[c] += val;
              cnt[c]++;
              RGBmax[c] = MAX(RGBmax[c], val);
            }
          }
707

708 709 710 711 712 713 714 715 716 717 718 719
          const float Ro = MIN(mean[0]/cnt[0], clip);
          const float Go = MIN(mean[1]/cnt[1], clip);
          const float Bo = MIN(mean[2]/cnt[2], clip);

          const float R = RGBmax[0];
          const float G = RGBmax[1];
          const float B = RGBmax[2];

          const float L = (R + G + B) / 3.0f;

          float C = SQRT3 * (R - G);
          float H = 2.0f * B - G - R;
720

721 722
          const float Co = SQRT3 * (Ro - Go);
          const float Ho = 2.0f * Bo - Go - Ro;
723

724
          if(R != G && G != B)
725 726 727 728 729 730
          {
            const float ratio = sqrtf((Co * Co + Ho * Ho) / (C * C + H * H));
            C *= ratio;
            H *= ratio;
          }

731 732 733 734 735 736 737
          float RGB[3] = { 0.0f, 0.0f, 0.0f };

          RGB[0] = L - H / 6.0f + C / SQRT12;
          RGB[1] = L - H / 6.0f - C / SQRT12;
          RGB[2] = L + H / 3.0f;

          out[0] = RGB[FCxtrans(j, i, roi_out, xtrans)];
738
        }
739 740
        else
          out[0] = in[0];
741
      }
742 743
      out++;
      in++;
744 745 746 747
    }
  }
}

748 749 750
#undef SQRT3
#undef SQRT12

751 752 753
static void 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)
754
{
755 756
  const float *const in = (const float *const)ivoid;
  float *const out = (float *const)ovoid;
757

758
  if(piece->pipe->dsc.filters)
759 760
  { // raw mosaic
#ifdef _OPENMP
761
#pragma omp parallel for SIMD() default(none) schedule(static)
762 763 764 765 766 767 768
#endif
    for(size_t k = 0; k < (size_t)roi_out->width * roi_out->height; k++)
    {
      out[k] = MIN(clip, in[k]);
    }
  }
  else
769
  {
770 771 772
    const int ch = piece->colors;

#ifdef _OPENMP
773
#pragma omp parallel for SIMD() default(none) schedule(static)
774 775 776 777 778 779 780 781 782 783 784 785 786
#endif
    for(size_t k = 0; k < (size_t)ch * roi_out->width * roi_out->height; k++)
    {
      out[k] = MIN(clip, in[k]);
    }
  }
}

#if defined(__SSE__)
static void 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)
{
787
  if(piece->pipe->dsc.filters)
788
  { // raw mosaic
789
    const __m128 clipm = _mm_set1_ps(clip);
790 791 792
    const size_t n = (size_t)roi_out->height * roi_out->width;
    float *const out = (float *)ovoid;
    float *const in = (float *)ivoid;
793
#ifdef _OPENMP
794 795 796 797 798 799 800 801 802 803 804 805 806 807 808
#pragma omp parallel for schedule(static) default(none)
#endif
    for(size_t j = 0; j < (n & ~3u); j += 4) _mm_stream_ps(out + j, _mm_min_ps(clipm, _mm_load_ps(in + j)));
    _mm_sfence();
    // lets see if there's a non-multiple of four rest to process:
    if(n & 3)
      for(size_t j = n & ~3u; j < n; j++) out[j] = MIN(clip, in[j]);
  }
  else
  {
    const __m128 clipm = _mm_set1_ps(clip);
    const int ch = piece->colors;

#ifdef _OPENMP
#pragma omp parallel for default(none) schedule(static)
809
#endif
810
    for(int j = 0; j < roi_out->height; j++)
811
    {
812 813 814
      float *out = (float *)ovoid + (size_t)ch * roi_out->width * j;
      float *in = (float *)ivoid + (size_t)ch * roi_in->width * j;
      for(int i = 0; i < roi_out->width; i++, in += ch, out += ch)
815
      {
816
        _mm_stream_ps(out, _mm_min_ps(clipm, _mm_set_ps(in[3], in[2], in[1], in[0])));
817 818 819
      }
    }
    _mm_sfence();
820 821 822 823 824 825 826 827 828 829 830 831 832 833 834 835 836
  }
}
#endif

static void 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)
{
  if(darktable.codepath.OPENMP_SIMD) process_clip_plain(piece, ivoid, ovoid, roi_in, roi_out, clip);
#if defined(__SSE__)
  else if(darktable.codepath.SSE2)
    process_clip_sse2(piece, ivoid, ovoid, roi_in, roi_out, clip);
#endif
  else
    dt_unreachable_codepath();
}

837 838
void 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)
839
{
840
  const uint32_t filters = piece->pipe->dsc.filters;
841 842 843
  dt_iop_highlights_data_t *data = (dt_iop_highlights_data_t *)piece->data;

  const float clip
844 845
      = data->clip * fminf(piece->pipe->dsc.processed_maximum[0],
                           fminf(piece->pipe->dsc.processed_maximum[1], piece->pipe->dsc.processed_maximum[2]));
846
  // const int ch = piece->colors;
847
  if(!filters)
848 849
  {
    process_clip(piece, ivoid, ovoid, roi_in, roi_out, clip);
850
    for(int k=0;k<3;k++)
851 852 853
      piece->pipe->dsc.processed_maximum[k]
          = fminf(piece->pipe->dsc.processed_maximum[0],
                  fminf(piece->pipe->dsc.processed_maximum[1], piece->pipe->dsc.processed_maximum[2]));
854 855
    return;
  }
856

857
  switch(data->mode)
858
  {
859 860
    case DT_IOP_HIGHLIGHTS_INPAINT: // a1ex's (magiclantern) idea of color inpainting:
    {
861 862 863
      const float clips[4] = { 0.987 * data->clip * piece->pipe->dsc.processed_maximum[0],
                               0.987 * data->clip * piece->pipe->dsc.processed_maximum[1],
                               0.987 * data->clip * piece->pipe->dsc.processed_maximum[2], clip };
864

865
      if(filters == 9u)
866
      {
867
        const uint8_t(*const xtrans)[6] = (const uint8_t(*const)[6])piece->pipe->dsc.xtrans;
868
#ifdef _OPENMP
869
#pragma omp parallel for schedule(dynamic) default(none)
870
#endif
871
        for(int j = 0; j < roi_out->height; j++)
872
        {
873 874
          interpolate_color_xtrans(ivoid, ovoid, roi_in, roi_out, 0, 1, j, clips, xtrans, 0);
          interpolate_color_xtrans(ivoid, ovoid, roi_in, roi_out, 0, -1, j, clips, xtrans, 1);
875 876
        }
#ifdef _OPENMP
877
#pragma omp parallel for schedule(dynamic) default(none)
878
#endif
879
        for(int i = 0; i < roi_out->width; i++)
880
        {
881 882
          interpolate_color_xtrans(ivoid, ovoid, roi_in, roi_out, 1, 1, i, clips, xtrans, 2);
          interpolate_color_xtrans(ivoid, ovoid, roi_in, roi_out, 1, -1, i, clips, xtrans, 3);
883 884
        }
      }
Dan Torop's avatar
Dan Torop committed
885 886
      else
      {
887
#ifdef _OPENMP
888
#pragma omp parallel for schedule(dynamic) default(none) shared(data, piece)
889
#endif
Dan Torop's avatar
Dan Torop committed
890 891
        for(int j = 0; j < roi_out->height; j++)
        {
892 893
          interpolate_color(ivoid, ovoid, roi_out, 0, 1, j, clips, filters, 0);
          interpolate_color(ivoid, ovoid, roi_out, 0, -1, j, clips, filters, 1);
Dan Torop's avatar
Dan Torop committed
894
        }
895

896
// up/down directions
897
#ifdef _OPENMP
898
#pragma omp parallel for schedule(dynamic) default(none) shared(data, piece)
899
#endif
Dan Torop's avatar
Dan Torop committed
900 901
        for(int i = 0; i < roi_out->width; i++)
        {
902 903
          interpolate_color(ivoid, ovoid, roi_out, 1, 1, i, clips, filters, 2);
          interpolate_color(ivoid, ovoid, roi_out, 1, -1, i, clips, filters, 3);
Dan Torop's avatar
Dan Torop committed
904
        }
905 906 907
      }
      break;
    }
908
    case DT_IOP_HIGHLIGHTS_LCH:
909
      if(filters == 9u)
910
        process_lch_xtrans(self, piece, ivoid, ovoid, roi_in, roi_out, clip);
Dan Torop's avatar
Dan Torop committed
911
      else
912
        process_lch_bayer(self, piece, ivoid, ovoid, roi_in, roi_out, clip);
913
      break;
914 915
    default:
    case DT_IOP_HIGHLIGHTS_CLIP:
916
      process_clip(piece, ivoid, ovoid, roi_in, roi_out, clip);
917 918
      break;
  }
919

920
  // update processed maximum
921 922 923
  const float m = fmaxf(fmaxf(piece->pipe->dsc.processed_maximum[0], piece->pipe->dsc.processed_maximum[1]),
                        piece->pipe->dsc.processed_maximum[2]);
  for(int k = 0; k < 3; k++) piece->pipe->dsc.processed_maximum[k] = m;
924

925
  if(piece->pipe->mask_display & DT_DEV_PIXELPIPE_DISPLAY_MASK) dt_iop_alpha_copy(ivoid, ovoid, roi_out->width, roi_out->height);
926 927
}

928
static void clip_callback(GtkWidget *slider, dt_iop_module_t *self)
929 930 931 932 933 934 935
{
  if(self->dt->gui->reset) return;
  dt_iop_highlights_params_t *p = (dt_iop_highlights_params_t *)self->params;
  p->clip = dt_bauhaus_slider_get(slider);
  dt_dev_add_history_item(darktable.develop, self, TRUE);
}

936
static void mode_changed(GtkWidget *combo, dt_iop_module_t *self)
937
{
938
  dt_iop_highlights_params_t *p = (dt_iop_highlights_params_t *)self->params;
939
  p->mode = dt_bauhaus_combobox_get(combo);
940
  if(p->mode > DT_IOP_HIGHLIGHTS_INPAINT) p->mode = DT_IOP_HIGHLIGHTS_INPAINT;
941
  dt_dev_add_history_item(darktable.develop, self, TRUE);
942 943
}

944 945
void commit_params(struct dt_iop_module_t *self, dt_iop_params_t *p1, dt_dev_pixelpipe_t *pipe,
                   dt_dev_pixelpipe_iop_t *piece)
946
{
947 948
  dt_iop_highlights_params_t *p = (dt_iop_highlights_params_t *)p1;
  dt_iop_highlights_data_t *d = (dt_iop_highlights_data_t *)piece->data;
949
  memcpy(d, p, sizeof(*p));
950 951 952 953

  piece->process_cl_ready = 1;

  // no OpenCL for DT_IOP_HIGHLIGHTS_INPAINT yet.
954
  if(d->mode == DT_IOP_HIGHLIGHTS_INPAINT) piece->process_cl_ready = 0;
955 956 957 958 959
}

void init_global(dt_iop_module_so_t *module)
{
  const int program = 2; // basic.cl, from programs.conf
960 961
  dt_iop_highlights_global_data_t *gd
      = (dt_iop_highlights_global_data_t *)malloc(sizeof(dt_iop_highlights_global_data_t));
962
  module->data = gd;
963
  gd->kernel_highlights_1f_clip = dt_opencl_create_kernel(program, "highlights_1f_clip");
964 965
  gd->kernel_highlights_1f_lch_bayer = dt_opencl_create_kernel(program, "highlights_1f_lch_bayer");
  gd->kernel_highlights_1f_lch_xtrans = dt_opencl_create_kernel(program, "highlights_1f_lch_xtrans");
966
  gd->kernel_highlights_4f_clip = dt_opencl_create_kernel(program, "highlights_4f_clip");
967 968 969 970 971
}

void cleanup_global(dt_iop_module_so_t *module)
{
  dt_iop_highlights_global_data_t *gd = (dt_iop_highlights_global_data_t *)module->data;
972
  dt_opencl_free_kernel(gd->kernel_highlights_4f_clip);
973 974
  dt_opencl_free_kernel(gd->kernel_highlights_1f_lch_bayer);
  dt_opencl_free_kernel(gd->kernel_highlights_1f_lch_xtrans);
975
  dt_opencl_free_kernel(gd->kernel_highlights_1f_clip);
976 977
  free(module->data);
  module->data = NULL;
978 979
}

980
void init_pipe(struct dt_iop_module_t *self, dt_dev_pixelpipe_t *pipe, dt_dev_pixelpipe_iop_t *piece)
981
{
982 983
  piece->data = malloc(sizeof(dt_iop_highlights_data_t));
  self->commit_params(self, self->default_params, pipe, piece);
984 985
}

986
void cleanup_pipe(struct dt_iop_module_t *self, dt_dev_pixelpipe_t *pipe, dt_dev_pixelpipe_iop_t *piece)
987
{
988
  free(piece->data);
989
  piece->data = NULL;
990 991 992 993
}

void gui_update(struct dt_iop_module_t *self)
{
994 995 996
  dt_iop_module_t *module = (dt_iop_module_t *)self;
  dt_iop_highlights_gui_data_t *g = (dt_iop_highlights_gui_data_t *)self->gui_data;
  dt_iop_highlights_params_t *p = (dt_iop_highlights_params_t *)module->params;
997
  dt_bauhaus_slider_set(g->clip, p->clip);
998
  dt_bauhaus_combobox_set(g->mode, p->mode);
999 1000
}

1001
void reload_defaults(dt_iop_module_t *module)
1002
{
1003 1004
  dt_iop_highlights_params_t tmp = (dt_iop_highlights_params_t){
    .mode = DT_IOP_HIGHLIGHTS_CLIP, .blendL = 1.0, .blendC = 0.0, .blendh = 0.0, .clip = 1.0
1005 1006 1007 1008 1009
  };

  // we might be called from presets update infrastructure => there is no image
  if(!module->dev) goto end;

1010
  // only on for raw images:
1011
  if(dt_image_is_raw(&module->dev->image_storage))
1012
    module->default_enabled = 1;
1013
  else
1014
    module->default_enabled = 0;
1015

1016
end:
1017 1018
  memcpy(module->params, &tmp, sizeof(dt_iop_highlights_params_t));
  memcpy(module->default_params, &tmp, sizeof(dt_iop_highlights_params_t));
1019 1020
}

1021 1022
void init(dt_iop_module_t *module)
{
1023
  // module->data = malloc(sizeof(dt_iop_highlights_data_t));
1024 1025
  module->params = calloc(1, sizeof(dt_iop_highlights_params_t));
  module->default_params = calloc(1, sizeof(dt_iop_highlights_params_t));
Heiko Bauke's avatar
Heiko Bauke committed
1026
  module->priority = 58; // module order created by iop_dependencies.py, do not edit!
1027
  module->default_enabled = 1;
1028 1029
  module->params_size = sizeof(dt_iop_highlights_params_t);
  module->gui_data = NULL;
1030 1031
}

1032 1033
void cleanup(dt_iop_module_t *module)
{
1034 1035
  free(module->params);
  module->params = NULL;
1036 1037 1038 1039
}

void gui_init(struct dt_iop_module_t *self)
{
1040 1041 1042 1043
  self->gui_data = malloc(sizeof(dt_iop_highlights_gui_data_t));
  dt_iop_highlights_gui_data_t *g = (dt_iop_highlights_gui_data_t *)self->gui_data;
  dt_iop_highlights_params_t *p = (dt_iop_highlights_params_t *)self->params;

1044
  self->widget = gtk_box_new(GTK_ORIENTATION_VERTICAL, DT_BAUHAUS_SPACE);
1045

1046 1047
  g->mode = dt_bauhaus_combobox_new(self);
  gtk_box_pack_start(GTK_BOX(self->widget), g->mode, TRUE, TRUE, 0);
1048
  dt_bauhaus_widget_set_label(g->mode, NULL, _("method"));
1049 1050
  dt_bauhaus_combobox_add(g->mode, _("clip highlights"));
  dt_bauhaus_combobox_add(g->mode, _("reconstruct in LCh"));
1051
  dt_bauhaus_combobox_add(g->mode, _("reconstruct color"));
1052
  gtk_widget_set_tooltip_text(g->mode, _("highlight reconstruction method"));
1053

1054
  g->clip = dt_bauhaus_slider_new_with_range(self, 0.0, 2.0, 0.01, p->clip, 3);
1055 1056
  gtk_widget_set_tooltip_text(g->clip, _("manually adjust the clipping threshold against "
                                         "magenta highlights (you shouldn't ever need to touch this)"));
1057
  dt_bauhaus_widget_set_label(g->clip, NULL, _("clipping threshold"));
1058 1059
  gtk_box_pack_start(GTK_BOX(self->widget), g->clip, TRUE, TRUE, 0);

1060

1061 1062
  g_signal_connect(G_OBJECT(g->clip), "value-changed", G_CALLBACK(clip_callback), self);
  g_signal_connect(G_OBJECT(g->mode), "value-changed", G_CALLBACK(mode_changed), self);
1063 1064 1065 1066
}

void gui_cleanup(struct dt_iop_module_t *self)
{
1067 1068
  free(self->gui_data);
  self->gui_data = NULL;
1069 1070
}

Richard Wonka's avatar
Richard Wonka committed
1071
// modelines: These editor modelines have been set for all relevant files by tools/update_modelines.sh
1072
// vim: shiftwidth=2 expandtab tabstop=2 cindent
Tobias Ellinghaus's avatar
Tobias Ellinghaus committed
1073
// kate: tab-indents: off; indent-width 2; replace-tabs on; indent-mode cstyle; remove-trailing-spaces modified;