globaltonemap.c 25.9 KB
Newer Older
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21
/*
    This file is part of darktable,
    copyright (c) 2012 Henrik Andersson.

    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.

    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.

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

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

41 42
#define REDUCESIZE 64

43
// NaN-safe clip: NaN compares false and will result in 0.0
44
#define CLIP(x) (((x) >= 0.0) ? ((x) <= 1.0 ? (x) : 1.0) : 0.0)
45
DT_MODULE_INTROSPECTION(3, dt_iop_global_tonemap_params_t)
46 47 48 49 50 51 52 53 54 55 56

typedef enum _iop_operator_t
{
  OPERATOR_REINHARD,
  OPERATOR_FILMIC,
  OPERATOR_DRAGO
} _iop_operator_t;

typedef struct dt_iop_global_tonemap_params_t
{
  _iop_operator_t operator;
57 58
  struct
  {
59 60 61
    float bias;
    float max_light; // cd/m2
  } drago;
62
  float detail;
63
} dt_iop_global_tonemap_params_t;
64

65 66 67 68 69 70 71 72 73
typedef struct dt_iop_global_tonemap_data_t
{
  _iop_operator_t operator;
  struct
  {
    float bias;
    float max_light; // cd/m2
  } drago;
  float detail;
74
} dt_iop_global_tonemap_data_t;
75

76 77 78
typedef struct dt_iop_global_tonemap_gui_data_t
{
  GtkWidget *operator;
79 80
  struct
  {
81 82 83
    GtkWidget *bias;
    GtkWidget *max_light;
  } drago;
84
  GtkWidget *detail;
85
  float lwmax;
86
  uint64_t hash;
87
  dt_pthread_mutex_t lock;
88
} dt_iop_global_tonemap_gui_data_t;
89

90 91 92 93 94 95 96
typedef struct dt_iop_global_tonemap_global_data_t
{
  int kernel_pixelmax_first;
  int kernel_pixelmax_second;
  int kernel_global_tonemap_reinhard;
  int kernel_global_tonemap_drago;
  int kernel_global_tonemap_filmic;
97
} dt_iop_global_tonemap_global_data_t;
98

99 100 101 102 103 104 105
const char *name()
{
  return _("global tonemap");
}

int flags()
{
106
  return IOP_FLAGS_INCLUDE_IN_STYLES | IOP_FLAGS_SUPPORTS_BLENDING | IOP_FLAGS_ALLOW_TILING;
107 108
}

109
int groups()
110 111 112 113
{
  return IOP_GROUP_TONE;
}

114 115
int legacy_params(dt_iop_module_t *self, const void *const old_params, const int old_version,
                  void *new_params, const int new_version)
116
{
117 118 119 120
  if(old_version < 3 && new_version == 3)
  {
    dt_iop_global_tonemap_params_t *o = (dt_iop_global_tonemap_params_t *)old_params;
    dt_iop_global_tonemap_params_t *n = (dt_iop_global_tonemap_params_t *)new_params;
121

122 123 124 125 126 127
    // only appended detail, 0 is no-op
    memcpy(n, o, sizeof(dt_iop_global_tonemap_params_t) - sizeof(float));
    n->detail = 0.0f;
    return 0;
  }
  return 1;
128 129
}

130 131 132
static inline void process_reinhard(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,
133
                                    dt_iop_global_tonemap_data_t *data)
134
{
135
  float *in = (float *)ivoid;
136 137 138 139
  float *out = (float *)ovoid;
  const int ch = piece->colors;

#ifdef _OPENMP
140
#pragma omp parallel for default(none) shared(in, out, data) schedule(static)
141
#endif
142
  for(size_t k = 0; k < (size_t)roi_out->width * roi_out->height; k++)
143
  {
144 145 146 147
    float *inp = in + ch * k;
    float *outp = out + ch * k;
    float l = inp[0] / 100.0;
    outp[0] = 100.0 * (l / (1.0f + l));
148 149 150
    outp[1] = inp[1];
    outp[2] = inp[2];
  }
151 152
}

153 154 155
static inline void process_drago(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, dt_iop_global_tonemap_data_t *data)
156
{
157
  dt_iop_global_tonemap_gui_data_t *g = (dt_iop_global_tonemap_gui_data_t *)self->gui_data;
158
  float *in = (float *)ivoid;
159 160 161 162
  float *out = (float *)ovoid;
  const int ch = piece->colors;

  /* precalcs */
163
  const float eps = 0.0001f;
164 165
  float lwmax;
  float tmp_lwmax = NAN;
166

167 168 169 170
  // Drago needs the absolute Lmax value of the image. In pixelpipe FULL we can not reliably get this value
  // as the pixelpipe might only see part of the image (region of interest). Therefore we try to get lwmax from
  // the PREVIEW pixelpipe which luckily stores it for us.
  if(self->dev->gui_attached && g && piece->pipe->type == DT_DEV_PIXELPIPE_FULL)
171
  {
172 173 174 175 176 177 178 179
    dt_pthread_mutex_lock(&g->lock);
    const uint64_t hash = g->hash;
    dt_pthread_mutex_unlock(&g->lock);

    // note that the case 'hash == 0' on first invocation in a session implies that g->lwmax
    // is NAN which initiates special handling below to avoid inconsistent results. in all
    // other cases we make sure that the preview pipe has left us with proper readings for
    // lwmax. if data are not yet there we need to wait (with timeout).
180 181
    if(hash != 0 && !dt_dev_sync_pixelpipe_hash(self->dev, piece->pipe, 0, self->priority, &g->lock, &g->hash))
      dt_control_log(_("inconsistent output"));
182

183 184 185 186 187 188 189 190 191 192 193 194 195 196
    dt_pthread_mutex_lock(&g->lock);
    tmp_lwmax = g->lwmax;
    dt_pthread_mutex_unlock(&g->lock);
  }

  // in all other cases we calculate lwmax here
  if(isnan(tmp_lwmax))
  {
    lwmax = eps;
    for(size_t k = 0; k < (size_t)roi_out->width * roi_out->height; k++)
    {
      float *inp = in + ch * k;
      lwmax = fmaxf(lwmax, (inp[0] * 0.01f));
    }
197
  }
198 199 200 201 202 203 204 205
  else
  {
    lwmax = tmp_lwmax;
  }

  // PREVIEW pixelpipe stores lwmax
  if(self->dev->gui_attached && g && piece->pipe->type == DT_DEV_PIXELPIPE_PREVIEW)
  {
206
    uint64_t hash = dt_dev_hash_plus(self->dev, piece->pipe, 0, self->priority);
207 208
    dt_pthread_mutex_lock(&g->lock);
    g->lwmax = lwmax;
209
    g->hash = hash;
210 211 212
    dt_pthread_mutex_unlock(&g->lock);
  }

213
  const float ldc = data->drago.max_light * 0.01 / log10f(lwmax + 1);
214
  const float bl = logf(fmaxf(eps, data->drago.bias)) / logf(0.5);
215 216

#ifdef _OPENMP
217
#pragma omp parallel for default(none) shared(in, out, lwmax) schedule(static)
218
#endif
219
  for(size_t k = 0; k < (size_t)roi_out->width * roi_out->height; k++)
220
  {
221 222 223 224 225
    float *inp = in + ch * k;
    float *outp = out + ch * k;
    float lw = inp[0] * 0.01f;
    outp[0] = 100.0f
              * (ldc * logf(fmaxf(eps, lw + 1.0f)) / logf(fmaxf(eps, 2.0f + (powf(lw / lwmax, bl)) * 8.0f)));
226 227 228
    outp[1] = inp[1];
    outp[2] = inp[2];
  }
229 230
}

231 232 233
static inline void process_filmic(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,
234
                                  dt_iop_global_tonemap_data_t *data)
235
{
236
  float *in = (float *)ivoid;
237 238 239 240
  float *out = (float *)ovoid;
  const int ch = piece->colors;

#ifdef _OPENMP
241
#pragma omp parallel for default(none) shared(in, out, data) schedule(static)
242
#endif
243
  for(size_t k = 0; k < (size_t)roi_out->width * roi_out->height; k++)
244
  {
245 246 247 248 249
    float *inp = in + ch * k;
    float *outp = out + ch * k;
    float l = inp[0] / 100.0;
    float x = fmaxf(0.0f, l - 0.004f);
    outp[0] = 100.0 * ((x * (6.2 * x + .5)) / (x * (6.2 * x + 1.7) + 0.06));
250 251 252
    outp[1] = inp[1];
    outp[2] = inp[2];
  }
253 254
}

255 256
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)
257
{
258
  dt_iop_global_tonemap_data_t *data = (dt_iop_global_tonemap_data_t *)piece->data;
259
  const float scale = piece->iscale / roi_in->scale;
260
  const float sigma_r = 8.0f; // does not depend on scale
261 262 263
  const float iw = piece->buf_in.width / scale;
  const float ih = piece->buf_in.height / scale;
  const float sigma_s = fminf(iw, ih) * 0.03f;
264 265 266 267 268 269 270
  dt_bilateral_t *b = NULL;
  if(data->detail != 0.0f)
  {
    b = dt_bilateral_init(roi_in->width, roi_in->height, sigma_s, sigma_r);
    // get detail from unchanged input buffer
    dt_bilateral_splat(b, (float *)ivoid);
  }
271

272 273
  switch(data->operator)
  {
274 275 276 277 278 279 280 281 282
    case OPERATOR_REINHARD:
      process_reinhard(self, piece, ivoid, ovoid, roi_in, roi_out, data);
      break;
    case OPERATOR_DRAGO:
      process_drago(self, piece, ivoid, ovoid, roi_in, roi_out, data);
      break;
    case OPERATOR_FILMIC:
      process_filmic(self, piece, ivoid, ovoid, roi_in, roi_out, data);
      break;
283
  }
284

285 286 287 288
  if(data->detail != 0.0f)
  {
    dt_bilateral_blur(b);
    // and apply it to output buffer after logscale
289
    dt_bilateral_slice_to_output(b, (float *)ivoid, (float *)ovoid, data->detail);
290 291 292
    dt_bilateral_free(b);
  }

293
  if(piece->pipe->mask_display & DT_DEV_PIXELPIPE_DISPLAY_MASK) dt_iop_alpha_copy(ivoid, ovoid, roi_out->width, roi_out->height);
294 295
}

296
#ifdef HAVE_OPENCL
297
int process_cl(struct dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, cl_mem dev_in, cl_mem dev_out,
298
               const dt_iop_roi_t *const roi_in, const dt_iop_roi_t *const roi_out)
299
{
300
  dt_iop_global_tonemap_data_t *d = (dt_iop_global_tonemap_data_t *)piece->data;
301
  dt_iop_global_tonemap_global_data_t *gd = (dt_iop_global_tonemap_global_data_t *)self->data;
302
  dt_iop_global_tonemap_gui_data_t *g = (dt_iop_global_tonemap_gui_data_t *)self->gui_data;
303
  dt_bilateral_cl_t *b = NULL;
304 305 306 307

  cl_int err = -999;
  cl_mem dev_m = NULL;
  cl_mem dev_r = NULL;
308
  float *maximum = NULL;
309 310 311 312 313 314 315
  const int devid = piece->pipe->devid;
  int gtkernel = -1;

  const int width = roi_out->width;
  const int height = roi_out->height;
  float parameters[4] = { 0.0f };

316 317 318 319 320 321 322 323 324 325 326
  switch(d->operator)
  {
    case OPERATOR_REINHARD:
      gtkernel = gd->kernel_global_tonemap_reinhard;
      break;
    case OPERATOR_DRAGO:
      gtkernel = gd->kernel_global_tonemap_drago;
      break;
    case OPERATOR_FILMIC:
      gtkernel = gd->kernel_global_tonemap_filmic;
      break;
327 328
  }

329
  if(d->operator== OPERATOR_DRAGO)
330
  {
331 332
    const float eps = 0.0001f;
    float tmp_lwmax = NAN;
333

334 335 336
    // see comments in process() about lwmax value
    if(self->dev->gui_attached && g && piece->pipe->type == DT_DEV_PIXELPIPE_FULL)
    {
337 338 339 340
      dt_pthread_mutex_lock(&g->lock);
      const uint64_t hash = g->hash;
      dt_pthread_mutex_unlock(&g->lock);

341 342
      if(hash != 0 && !dt_dev_sync_pixelpipe_hash(self->dev, piece->pipe, 0, self->priority, &g->lock, &g->hash))
        dt_control_log(_("inconsistent output"));
343

344 345 346 347
      dt_pthread_mutex_lock(&g->lock);
      tmp_lwmax = g->lwmax;
      dt_pthread_mutex_unlock(&g->lock);
    }
348

349 350
    if(isnan(tmp_lwmax))
    {
351 352 353 354 355 356 357 358 359 360
      dt_opencl_local_buffer_t flocopt
        = (dt_opencl_local_buffer_t){ .xoffset = 0, .xfactor = 1, .yoffset = 0, .yfactor = 1,
                                      .cellsize = sizeof(float), .overhead = 0,
                                      .sizex = 1 << 4, .sizey = 1 << 4 };

      if(!dt_opencl_local_buffer_opt(devid, gd->kernel_pixelmax_first, &flocopt))
        goto error;

      const size_t bwidth = ROUNDUP(width, flocopt.sizex);
      const size_t bheight = ROUNDUP(height, flocopt.sizey);
361

362
      const int bufsize = (bwidth / flocopt.sizex) * (bheight / flocopt.sizey);
363

364 365 366 367 368 369 370 371 372 373 374 375
      dt_opencl_local_buffer_t slocopt
        = (dt_opencl_local_buffer_t){ .xoffset = 0, .xfactor = 1, .yoffset = 0, .yfactor = 1,
                                      .cellsize = sizeof(float), .overhead = 0,
                                      .sizex = 1 << 16, .sizey = 1 };

      if(!dt_opencl_local_buffer_opt(devid, gd->kernel_pixelmax_second, &slocopt))
        goto error;

      const int reducesize = MIN(REDUCESIZE, ROUNDUP(bufsize, slocopt.sizex) / slocopt.sizex);

      size_t sizes[3];
      size_t local[3];
376 377 378 379 380 381 382 383 384 385

      dev_m = dt_opencl_alloc_device_buffer(devid, (size_t)bufsize * sizeof(float));
      if(dev_m == NULL) goto error;

      dev_r = dt_opencl_alloc_device_buffer(devid, (size_t)reducesize * sizeof(float));
      if(dev_r == NULL) goto error;

      sizes[0] = bwidth;
      sizes[1] = bheight;
      sizes[2] = 1;
386 387
      local[0] = flocopt.sizex;
      local[1] = flocopt.sizey;
388 389 390 391 392
      local[2] = 1;
      dt_opencl_set_kernel_arg(devid, gd->kernel_pixelmax_first, 0, sizeof(cl_mem), &dev_in);
      dt_opencl_set_kernel_arg(devid, gd->kernel_pixelmax_first, 1, sizeof(int), &width);
      dt_opencl_set_kernel_arg(devid, gd->kernel_pixelmax_first, 2, sizeof(int), &height);
      dt_opencl_set_kernel_arg(devid, gd->kernel_pixelmax_first, 3, sizeof(cl_mem), &dev_m);
393
      dt_opencl_set_kernel_arg(devid, gd->kernel_pixelmax_first, 4, flocopt.sizex * flocopt.sizey * sizeof(float), NULL);
394 395 396
      err = dt_opencl_enqueue_kernel_2d_with_local(devid, gd->kernel_pixelmax_first, sizes, local);
      if(err != CL_SUCCESS) goto error;

397
      sizes[0] = reducesize * slocopt.sizex;
398 399
      sizes[1] = 1;
      sizes[2] = 1;
400
      local[0] = slocopt.sizex;
401 402 403 404 405
      local[1] = 1;
      local[2] = 1;
      dt_opencl_set_kernel_arg(devid, gd->kernel_pixelmax_second, 0, sizeof(cl_mem), &dev_m);
      dt_opencl_set_kernel_arg(devid, gd->kernel_pixelmax_second, 1, sizeof(cl_mem), &dev_r);
      dt_opencl_set_kernel_arg(devid, gd->kernel_pixelmax_second, 2, sizeof(int), &bufsize);
406
      dt_opencl_set_kernel_arg(devid, gd->kernel_pixelmax_second, 3, slocopt.sizex * sizeof(float), NULL);
407 408 409
      err = dt_opencl_enqueue_kernel_2d_with_local(devid, gd->kernel_pixelmax_second, sizes, local);
      if(err != CL_SUCCESS) goto error;

410
      maximum = dt_alloc_align(16, reducesize * sizeof(float));
411
      err = dt_opencl_read_buffer_from_device(devid, (void *)maximum, dev_r, 0,
412
                                            (size_t)reducesize * sizeof(float), CL_TRUE);
413
      if(err != CL_SUCCESS) goto error;
414

415 416 417
      dt_opencl_release_mem_object(dev_r);
      dt_opencl_release_mem_object(dev_m);
      dev_r = dev_m = NULL;
418

419 420 421 422 423 424 425 426
      for(int k = 1; k < reducesize; k++)
      {
        float mine = maximum[0];
        float other = maximum[k];
        maximum[0] = (other > mine) ? other : mine;
      }

      tmp_lwmax = MAX(eps, (maximum[0] * 0.01f));
427

428 429
      dt_free_align(maximum);
      maximum = NULL;
430 431
    }

432
    const float lwmax = tmp_lwmax;
433
    const float ldc = d->drago.max_light * 0.01f / log10f(lwmax + 1.0f);
434 435 436 437 438 439
    const float bl = logf(MAX(eps, d->drago.bias)) / logf(0.5f);

    parameters[0] = eps;
    parameters[1] = ldc;
    parameters[2] = bl;
    parameters[3] = lwmax;
440 441 442

    if(self->dev->gui_attached && g && piece->pipe->type == DT_DEV_PIXELPIPE_PREVIEW)
    {
443
      uint64_t hash = dt_dev_hash_plus(self->dev, piece->pipe, 0, self->priority);
444 445
      dt_pthread_mutex_lock(&g->lock);
      g->lwmax = lwmax;
446
      g->hash = hash;
447 448
      dt_pthread_mutex_unlock(&g->lock);
    }
449 450
  }

451
  const float scale = piece->iscale / roi_in->scale;
452
  const float sigma_r = 8.0f; // does not depend on scale
453 454 455
  const float iw = piece->buf_in.width / scale;
  const float ih = piece->buf_in.height / scale;
  const float sigma_s = fminf(iw, ih) * 0.03f;
456 457 458 459 460 461 462 463 464 465

  if(d->detail != 0.0f)
  {
    b = dt_bilateral_init_cl(devid, roi_in->width, roi_in->height, sigma_s, sigma_r);
    if(!b) goto error;
    // get detail from unchanged input buffer
    err = dt_bilateral_splat_cl(b, dev_in);
    if(err != CL_SUCCESS) goto error;
  }

466 467 468 469 470
  size_t sizes[2] = { ROUNDUPWD(width), ROUNDUPHT(height) };
  dt_opencl_set_kernel_arg(devid, gtkernel, 0, sizeof(cl_mem), &dev_in);
  dt_opencl_set_kernel_arg(devid, gtkernel, 1, sizeof(cl_mem), &dev_out);
  dt_opencl_set_kernel_arg(devid, gtkernel, 2, sizeof(int), &width);
  dt_opencl_set_kernel_arg(devid, gtkernel, 3, sizeof(int), &height);
471
  dt_opencl_set_kernel_arg(devid, gtkernel, 4, 4 * sizeof(float), &parameters);
472 473
  err = dt_opencl_enqueue_kernel_2d(devid, gtkernel, sizes);
  if(err != CL_SUCCESS) goto error;
474 475 476 477

  if(d->detail != 0.0f)
  {
    err = dt_bilateral_blur_cl(b);
478
    if(err != CL_SUCCESS) goto error;
479 480
    // and apply it to output buffer after logscale
    err = dt_bilateral_slice_to_output_cl(b, dev_in, dev_out, d->detail);
481
    if(err != CL_SUCCESS) goto error;
482 483 484
    dt_bilateral_free_cl(b);
  }

485 486 487
  return TRUE;

error:
488
  if(b) dt_bilateral_free_cl(b);
489 490
  dt_opencl_release_mem_object(dev_m);
  dt_opencl_release_mem_object(dev_r);
491
  dt_free_align(maximum);
492 493 494 495 496 497
  dt_print(DT_DEBUG_OPENCL, "[opencl_global_tonemap] couldn't enqueue kernel! %d\n", err);
  return FALSE;
}
#endif


498 499 500
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)
501
{
502 503
  dt_iop_global_tonemap_data_t *d = (dt_iop_global_tonemap_data_t *)piece->data;

504 505 506 507
  const float scale = piece->iscale / roi_in->scale;
  const float iw = piece->buf_in.width / scale;
  const float ih = piece->buf_in.height / scale;
  const float sigma_s = fminf(iw, ih) * 0.03f;
508
  const float sigma_r = 8.0f;
509
  const int detail = (d->detail != 0.0f);
510 511 512 513 514

  const int width = roi_in->width;
  const int height = roi_in->height;
  const int channels = piece->colors;

515
  const size_t basebuffer = width * height * channels * sizeof(float);
516

517
  tiling->factor = 2.0f + (detail ? (float)dt_bilateral_memory_use2(width, height, sigma_s, sigma_r) / basebuffer : 0.0f);
518
  tiling->maxbuf
519
      = (detail ? MAX(1.0f, (float)dt_bilateral_singlebuffer_size2(width, height, sigma_s, sigma_r) / basebuffer) : 1.0f);
520
  tiling->overhead = 0;
521
  tiling->overlap = (detail ? ceilf(4 * sigma_s) : 0);
522 523 524 525 526
  tiling->xalign = 1;
  tiling->yalign = 1;
  return;
}

527 528
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)
529 530 531 532
{
  dt_iop_global_tonemap_params_t *p = (dt_iop_global_tonemap_params_t *)p1;
  dt_iop_global_tonemap_data_t *d = (dt_iop_global_tonemap_data_t *)piece->data;

533
  d->operator= p->operator;
534 535 536 537
  d->drago.bias = p->drago.bias;
  d->drago.max_light = p->drago.max_light;
  d->detail = p->detail;

538 539 540
  // drago needs the maximum L-value of the whole image so it must not use tiling
  if(d->operator == OPERATOR_DRAGO) piece->process_tiling_ready = 0;

541 542 543 544 545 546
#ifdef HAVE_OPENCL
  if(d->detail != 0.0f)
    piece->process_cl_ready = (piece->process_cl_ready && !(darktable.opencl->avoid_atomics));
#endif
}

547
void init_pipe(struct dt_iop_module_t *self, dt_dev_pixelpipe_t *pipe, dt_dev_pixelpipe_iop_t *piece)
548
{
549
  piece->data = calloc(1, sizeof(dt_iop_global_tonemap_data_t));
550 551
  self->commit_params(self, self->default_params, pipe, piece);
}
552

553
void cleanup_pipe(struct dt_iop_module_t *self, dt_dev_pixelpipe_t *pipe, dt_dev_pixelpipe_iop_t *piece)
554 555
{
  free(piece->data);
556
  piece->data = NULL;
557
}
558

559 560 561
void init_global(dt_iop_module_so_t *module)
{
  const int program = 8; // extended.cl from programs.conf
562 563
  dt_iop_global_tonemap_global_data_t *gd
      = (dt_iop_global_tonemap_global_data_t *)malloc(sizeof(dt_iop_global_tonemap_global_data_t));
564 565 566 567 568 569 570 571 572 573 574 575 576 577 578 579 580 581 582 583 584 585 586
  module->data = gd;
  gd->kernel_pixelmax_first = dt_opencl_create_kernel(program, "pixelmax_first");
  gd->kernel_pixelmax_second = dt_opencl_create_kernel(program, "pixelmax_second");
  gd->kernel_global_tonemap_reinhard = dt_opencl_create_kernel(program, "global_tonemap_reinhard");
  gd->kernel_global_tonemap_drago = dt_opencl_create_kernel(program, "global_tonemap_drago");
  gd->kernel_global_tonemap_filmic = dt_opencl_create_kernel(program, "global_tonemap_filmic");
}


void cleanup_global(dt_iop_module_so_t *module)
{
  dt_iop_global_tonemap_global_data_t *gd = (dt_iop_global_tonemap_global_data_t *)module->data;
  dt_opencl_free_kernel(gd->kernel_pixelmax_first);
  dt_opencl_free_kernel(gd->kernel_pixelmax_second);
  dt_opencl_free_kernel(gd->kernel_global_tonemap_reinhard);
  dt_opencl_free_kernel(gd->kernel_global_tonemap_drago);
  dt_opencl_free_kernel(gd->kernel_global_tonemap_filmic);
  free(module->data);
  module->data = NULL;
}



587
static void operator_callback(GtkWidget *combobox, gpointer user_data)
588 589 590 591 592
{
  dt_iop_module_t *self = (dt_iop_module_t *)user_data;
  dt_iop_global_tonemap_gui_data_t *g = (dt_iop_global_tonemap_gui_data_t *)self->gui_data;
  if(self->dt->gui->reset) return;
  dt_iop_global_tonemap_params_t *p = (dt_iop_global_tonemap_params_t *)self->params;
593
  p->operator= dt_bauhaus_combobox_get(combobox);
594

595 596
  gtk_widget_set_visible(g->drago.bias, FALSE);
  gtk_widget_set_visible(g->drago.max_light, FALSE);
597
  /* show ui for selected operator */
598
  if(p->operator== OPERATOR_DRAGO)
599 600 601 602
  {
    gtk_widget_set_visible(g->drago.bias, TRUE);
    gtk_widget_set_visible(g->drago.max_light, TRUE);
  }
603 604 605 606

  dt_dev_add_history_item(darktable.develop, self, TRUE);
}

607
static void _drago_bias_callback(GtkWidget *w, gpointer user_data)
608 609 610 611 612 613 614 615
{
  dt_iop_module_t *self = (dt_iop_module_t *)user_data;
  if(self->dt->gui->reset) return;
  dt_iop_global_tonemap_params_t *p = (dt_iop_global_tonemap_params_t *)self->params;
  p->drago.bias = dt_bauhaus_slider_get(w);
  dt_dev_add_history_item(darktable.develop, self, TRUE);
}

616
static void _drago_max_light_callback(GtkWidget *w, gpointer user_data)
617 618 619 620 621 622 623 624
{
  dt_iop_module_t *self = (dt_iop_module_t *)user_data;
  if(self->dt->gui->reset) return;
  dt_iop_global_tonemap_params_t *p = (dt_iop_global_tonemap_params_t *)self->params;
  p->drago.max_light = dt_bauhaus_slider_get(w);
  dt_dev_add_history_item(darktable.develop, self, TRUE);
}

625
static void detail_callback(GtkWidget *w, gpointer user_data)
626
{
627 628 629 630
  dt_iop_module_t *self = (dt_iop_module_t *)user_data;
  dt_iop_global_tonemap_params_t *p = (dt_iop_global_tonemap_params_t *)self->params;
  p->detail = dt_bauhaus_slider_get(w);
  dt_dev_add_history_item(darktable.develop, self, TRUE);
631 632 633 634 635 636 637 638
}

void gui_update(struct dt_iop_module_t *self)
{
  dt_iop_module_t *module = (dt_iop_module_t *)self;
  dt_iop_global_tonemap_gui_data_t *g = (dt_iop_global_tonemap_gui_data_t *)self->gui_data;
  dt_iop_global_tonemap_params_t *p = (dt_iop_global_tonemap_params_t *)module->params;
  dt_bauhaus_combobox_set(g->operator, p->operator);
639 640 641
  gtk_widget_set_visible(g->drago.bias, FALSE);
  gtk_widget_set_visible(g->drago.max_light, FALSE);
  /* show ui for selected operator */
642
  if(p->operator== OPERATOR_DRAGO)
643 644 645 646 647
  {
    gtk_widget_set_visible(g->drago.bias, TRUE);
    gtk_widget_set_visible(g->drago.max_light, TRUE);
  }

648 649 650
  /* drago */
  dt_bauhaus_slider_set(g->drago.bias, p->drago.bias);
  dt_bauhaus_slider_set(g->drago.max_light, p->drago.max_light);
651
  dt_bauhaus_slider_set(g->detail, p->detail);
652 653 654 655 656 657

  dt_pthread_mutex_lock(&g->lock);
  g->lwmax = NAN;
  g->hash = 0;
  dt_pthread_mutex_unlock(&g->lock);

658 659 660 661
}

void init(dt_iop_module_t *module)
{
662 663
  module->params = calloc(1, sizeof(dt_iop_global_tonemap_params_t));
  module->default_params = calloc(1, sizeof(dt_iop_global_tonemap_params_t));
664
  module->default_enabled = 0;
Heiko Bauke's avatar
Heiko Bauke committed
665
  module->priority = 544; // module order created by iop_dependencies.py, do not edit!
666 667
  module->params_size = sizeof(dt_iop_global_tonemap_params_t);
  module->gui_data = NULL;
668 669
  dt_iop_global_tonemap_params_t tmp
      = (dt_iop_global_tonemap_params_t){ OPERATOR_DRAGO, { 0.85f, 100.0f }, 0.0f };
670 671 672 673 674 675 676 677 678 679 680 681 682 683 684 685
  memcpy(module->params, &tmp, sizeof(dt_iop_global_tonemap_params_t));
  memcpy(module->default_params, &tmp, sizeof(dt_iop_global_tonemap_params_t));
}

void cleanup(dt_iop_module_t *module)
{
  free(module->params);
  module->params = NULL;
}

void gui_init(struct dt_iop_module_t *self)
{
  self->gui_data = malloc(sizeof(dt_iop_global_tonemap_gui_data_t));
  dt_iop_global_tonemap_gui_data_t *g = (dt_iop_global_tonemap_gui_data_t *)self->gui_data;
  dt_iop_global_tonemap_params_t *p = (dt_iop_global_tonemap_params_t *)self->params;

686 687
  dt_pthread_mutex_init(&g->lock, NULL);
  g->lwmax = NAN;
688
  g->hash = 0;
689

690
  self->widget = gtk_box_new(GTK_ORIENTATION_VERTICAL, DT_BAUHAUS_SPACE);
691 692

  /* operator */
693
  g->operator= dt_bauhaus_combobox_new(self);
694
  dt_bauhaus_widget_set_label(g->operator, NULL, _("operator"));
695 696 697 698 699

  dt_bauhaus_combobox_add(g->operator, "reinhard");
  dt_bauhaus_combobox_add(g->operator, "filmic");
  dt_bauhaus_combobox_add(g->operator, "drago");

700
  gtk_widget_set_tooltip_text(g->operator, _("the global tonemap operator"));
701
  g_signal_connect(G_OBJECT(g->operator), "value-changed", G_CALLBACK(operator_callback), self);
702 703 704
  gtk_box_pack_start(GTK_BOX(self->widget), GTK_WIDGET(g->operator), TRUE, TRUE, 0);

  /* drago bias */
705
  g->drago.bias = dt_bauhaus_slider_new_with_range(self, 0.5, 1.0, 0.05, p->drago.bias, 2);
706
  dt_bauhaus_widget_set_label(g->drago.bias, NULL, _("bias"));
707 708
  gtk_widget_set_tooltip_text(g->drago.bias, _("the bias for tonemapper controls the linearity, "
                                               "the higher the more details in blacks"));
709
  g_signal_connect(G_OBJECT(g->drago.bias), "value-changed", G_CALLBACK(_drago_bias_callback), self);
710
  gtk_box_pack_start(GTK_BOX(self->widget), g->drago.bias, TRUE, TRUE, 0);
711

712 713

  /* drago bias */
714
  g->drago.max_light = dt_bauhaus_slider_new_with_range(self, 1, 500, 10, p->drago.max_light, 2);
715
  dt_bauhaus_widget_set_label(g->drago.max_light, NULL, _("target"));
716
  gtk_widget_set_tooltip_text(g->drago.max_light, _("the target light for tonemapper specified as cd/m2"));
717
  g_signal_connect(G_OBJECT(g->drago.max_light), "value-changed", G_CALLBACK(_drago_max_light_callback), self);
718
  gtk_box_pack_start(GTK_BOX(self->widget), g->drago.max_light, TRUE, TRUE, 0);
719

720 721 722
  /* detail */
  g->detail = dt_bauhaus_slider_new_with_range(self, -1.0, 1.0, 0.01, 0.0, 3);
  gtk_box_pack_start(GTK_BOX(self->widget), g->detail, TRUE, TRUE, 0);
723
  dt_bauhaus_widget_set_label(g->detail, NULL, _("detail"));
724

725
  g_signal_connect(G_OBJECT(g->detail), "value-changed", G_CALLBACK(detail_callback), self);
726 727 728 729
}

void gui_cleanup(struct dt_iop_module_t *self)
{
730 731
  dt_iop_global_tonemap_gui_data_t *g = (dt_iop_global_tonemap_gui_data_t *)self->gui_data;
  dt_pthread_mutex_destroy(&g->lock);
732 733 734 735
  free(self->gui_data);
  self->gui_data = NULL;
}

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