From f5ee08af12efb96e92ea56a2b7e4baa984040162 Mon Sep 17 00:00:00 2001 From: Roman Lebedev Date: Tue, 2 Aug 2016 15:28:24 +0300 Subject: [PATCH 1/7] Add new raw overexposure indication iop Right now it is located right before highlight reconstruction iop, so it operates one the raw data, before demosaicing. As a result, for the areas with "clipped" colors, it produces flat areas with false color. 1. Not configurable. 2. Will be affected by all the iops that are after this module. --- data/kernels/basic.cl | 35 ++++++ po/POTFILES.in | 1 + src/iop/CMakeLists.txt | 1 + src/iop/rawoverexposed.c | 259 +++++++++++++++++++++++++++++++++++++++ 4 files changed, 296 insertions(+) create mode 100644 src/iop/rawoverexposed.c diff --git a/data/kernels/basic.cl b/data/kernels/basic.cl index 190dd85b7cd0..03a4b1b95086 100644 --- a/data/kernels/basic.cl +++ b/data/kernels/basic.cl @@ -163,6 +163,25 @@ highlights_4f_clip (read_only image2d_t in, write_only image2d_t out, const int write_imagef (out, (int2)(x, y), pixel); } +rawoverexposed_4f_mark (read_only image2d_t in, write_only image2d_t out, const int width, const int height, + const float threshold) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + + if(x >= width || y >= height) return; + + // 4f/pixel means that this has been debayered already. + // (this code path is just used for preview and non-raw images) + float4 pixel = read_imagef(in, sampleri, (int2)(x, y)); + + pixel.x = (pixel.x >= threshold) ? 0.0 : pixel.x; + pixel.y = (pixel.y >= threshold) ? 0.0 : pixel.y; + pixel.z = (pixel.z >= threshold) ? 0.0 : pixel.x; + + write_imagef (out, (int2)(x, y), pixel); +} + kernel void highlights_1f_clip (read_only image2d_t in, write_only image2d_t out, const int width, const int height, const float clip, const int rx, const int ry, const int filters) @@ -179,6 +198,22 @@ highlights_1f_clip (read_only image2d_t in, write_only image2d_t out, const int write_imagef (out, (int2)(x, y), pixel); } +kernel void +rawoverexposed_1f_mark (read_only image2d_t in, write_only image2d_t out, const int width, const int height, + const float threshold) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + + if(x >= width || y >= height) return; + + float4 pixel = read_imagef(in, sampleri, (int2)(x, y)); + + pixel.x = (pixel.x >= threshold) ? 0.0 : pixel.x; + + write_imagef (out, (int2)(x, y), pixel); +} + #define SQRT3 1.7320508075688772935274463415058723669f #define SQRT12 3.4641016151377545870548926830117447339f // 2*SQRT3 kernel void diff --git a/po/POTFILES.in b/po/POTFILES.in index 667e517fff9a..3140e7fe982d 100644 --- a/po/POTFILES.in +++ b/po/POTFILES.in @@ -115,6 +115,7 @@ src/iop/lowlight.c src/iop/lowpass.c src/iop/monochrome.c src/iop/nlmeans.c +src/iop/rawoverexposed.c src/iop/overexposed.c src/iop/profile_gamma.c src/iop/rawdenoise.c diff --git a/src/iop/CMakeLists.txt b/src/iop/CMakeLists.txt index ca00c9830882..b736a7f3ee46 100644 --- a/src/iop/CMakeLists.txt +++ b/src/iop/CMakeLists.txt @@ -77,6 +77,7 @@ add_iop(monochrome "monochrome.c" DEFAULT_VISIBLE) add_iop(basecurve "basecurve.c" DEFAULT_VISIBLE) add_iop(colorzones "colorzones.c") add_iop(highlights "highlights.c") +add_iop(rawoverexposed "rawoverexposed.c") add_iop(velvia "velvia.c") add_iop(vignette "vignette.c" DEFAULT_VISIBLE) add_iop(splittoning "splittoning.c") diff --git a/src/iop/rawoverexposed.c b/src/iop/rawoverexposed.c new file mode 100644 index 000000000000..5028611be01c --- /dev/null +++ b/src/iop/rawoverexposed.c @@ -0,0 +1,259 @@ +/* + This file is part of darktable, + copyright (c) 2016 Roman Lebedev. + + 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 . + */ + +#ifdef HAVE_CONFIG_H +#include "config.h" +#endif + +#include "bauhaus/bauhaus.h" // for dt_bauhaus_slider_get, dt_bauhaus_... +#include "common/darktable.h" // for darktable_t, dt_print, DT_MODULE_I... +#include "common/image.h" // for ::DT_IMAGE_RAW, dt_image_t +#include "common/opencl.h" // for dt_opencl_set_kernel_arg, dt_openc... +#include "develop/develop.h" // for dt_dev_add_history_item +#include "develop/imageop.h" // for dt_iop_module_t, dt_iop_roi_t, dt_... +#include "develop/imageop_math.h" // for dt_iop_alpha_copy +#include "develop/pixelpipe.h" // for dt_dev_pixelpipe_type_t::DT_DEV_PI... +#include "gui/gtk.h" // for dt_gui_gtk_t +#include "iop/iop_api.h" // for dt_iop_params_t +#include // for g_type_check_instance_cast, GCallback +#include // for TRUE, FALSE +#include // for _ +#include // for GtkWidget, gtk_box_get_type, gtk_b... +#include // for fminf +#include // for free, NULL, size_t, calloc, malloc +#include // for memcpy + +DT_MODULE_INTROSPECTION(1, dt_iop_rawoverexposed_params_t) + +typedef struct dt_iop_rawoverexposed_params_t +{ + float threshold; +} dt_iop_rawoverexposed_params_t; + +typedef struct dt_iop_rawoverexposed_gui_data_t +{ + GtkWidget *threshold; +} dt_iop_rawoverexposed_gui_data_t; + +typedef dt_iop_rawoverexposed_params_t dt_iop_rawoverexposed_data_t; + +typedef struct dt_iop_rawoverexposed_global_data_t +{ + int kernel_rawoverexposed_1f_mark; + int kernel_rawoverexposed_4f_mark; +} dt_iop_rawoverexposed_global_data_t; + +const char *name() +{ + return _("raw overexposed"); +} + +int groups() +{ + return IOP_GROUP_BASIC; +} + +int flags() +{ + return IOP_FLAGS_ALLOW_TILING | IOP_FLAGS_ONE_INSTANCE; +} + +#ifdef HAVE_OPENCL +int process_cl(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) +{ + dt_iop_rawoverexposed_data_t *d = piece->data; + dt_iop_rawoverexposed_global_data_t *gd = self->data; + + cl_int err = -999; + const int devid = piece->pipe->devid; + const int width = roi_in->width; + const int height = roi_in->height; + + size_t sizes[] = { ROUNDUPWD(width), ROUNDUPHT(height), 1 }; + const float threshold + = d->threshold * fminf(piece->pipe->dsc.processed_maximum[0], + fminf(piece->pipe->dsc.processed_maximum[1], piece->pipe->dsc.processed_maximum[2])); + + const int kernel = (!dt_dev_pixelpipe_uses_downsampled_input(piece->pipe) && piece->pipe->dsc.filters) + ? gd->kernel_rawoverexposed_1f_mark + : gd->kernel_rawoverexposed_4f_mark; + + dt_opencl_set_kernel_arg(devid, kernel, 0, sizeof(cl_mem), (void *)&dev_in); + dt_opencl_set_kernel_arg(devid, kernel, 1, sizeof(cl_mem), (void *)&dev_out); + dt_opencl_set_kernel_arg(devid, kernel, 2, sizeof(int), (void *)&width); + dt_opencl_set_kernel_arg(devid, kernel, 3, sizeof(int), (void *)&height); + dt_opencl_set_kernel_arg(devid, kernel, 4, sizeof(float), (void *)&threshold); + err = dt_opencl_enqueue_kernel_2d(devid, kernel, sizes); + if(err != CL_SUCCESS) goto error; + + return TRUE; + +error: + dt_print(DT_DEBUG_OPENCL, "[opencl_rawoverexposed] couldn't enqueue kernel! %d\n", err); + return FALSE; +} +#endif + +void process(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_rawoverexposed_data_t *data = piece->data; + + const float threshold + = data->threshold * fminf(piece->pipe->dsc.processed_maximum[0], + fminf(piece->pipe->dsc.processed_maximum[1], piece->pipe->dsc.processed_maximum[2])); + + const float *const in = (const float *const)ivoid; + float *const out = (float *const)ovoid; + + if(!dt_dev_pixelpipe_uses_downsampled_input(piece->pipe) && piece->pipe->dsc.filters) + { // raw mosaic +#ifdef _OPENMP +#pragma omp parallel for SIMD() default(none) schedule(static) +#endif + for(size_t k = 0; k < (size_t)roi_out->width * roi_out->height; k++) + { + out[k] = (in[k] >= threshold) ? 0.0 : in[k]; + } + } + else + { + const int ch = piece->colors; + +#ifdef _OPENMP +#pragma omp parallel for SIMD() default(none) schedule(static) +#endif + for(size_t k = 0; k < (size_t)ch * roi_out->width * roi_out->height; k++) + { + out[k] = (in[k] >= threshold) ? 0.0 : in[k]; + } + } + + if(piece->pipe->mask_display) dt_iop_alpha_copy(ivoid, ovoid, roi_out->width, roi_out->height); +} + +void commit_params(dt_iop_module_t *self, dt_iop_params_t *p1, dt_dev_pixelpipe_t *pipe, + dt_dev_pixelpipe_iop_t *piece) +{ + dt_iop_rawoverexposed_params_t *p = p1; + dt_iop_rawoverexposed_data_t *d = piece->data; + + *d = *p; + + if(pipe->type != DT_DEV_PIXELPIPE_FULL || /*!self->dev->overexposed.enabled ||*/ !self->dev->gui_attached) + piece->enabled = 0; +} + +void init_global(dt_iop_module_so_t *module) +{ + const int program = 2; // basic.cl, from programs.conf + module->data = malloc(sizeof(dt_iop_rawoverexposed_global_data_t)); + dt_iop_rawoverexposed_global_data_t *gd = module->data; + gd->kernel_rawoverexposed_1f_mark = dt_opencl_create_kernel(program, "rawoverexposed_1f_mark"); + gd->kernel_rawoverexposed_4f_mark = dt_opencl_create_kernel(program, "rawoverexposed_4f_mark"); +} + +void cleanup_global(dt_iop_module_so_t *module) +{ + dt_iop_rawoverexposed_global_data_t *gd = module->data; + dt_opencl_free_kernel(gd->kernel_rawoverexposed_4f_mark); + dt_opencl_free_kernel(gd->kernel_rawoverexposed_1f_mark); + free(module->data); + module->data = NULL; +} + +void init_pipe(dt_iop_module_t *self, dt_dev_pixelpipe_t *pipe, dt_dev_pixelpipe_iop_t *piece) +{ + piece->data = malloc(sizeof(dt_iop_rawoverexposed_data_t)); + self->commit_params(self, self->default_params, pipe, piece); +} + +void cleanup_pipe(dt_iop_module_t *self, dt_dev_pixelpipe_t *pipe, dt_dev_pixelpipe_iop_t *piece) +{ + free(piece->data); + piece->data = NULL; +} + +void reload_defaults(dt_iop_module_t *module) +{ + dt_iop_rawoverexposed_params_t tmp = (dt_iop_rawoverexposed_params_t){.threshold = 1.0 }; + + memcpy(module->params, &tmp, sizeof(dt_iop_rawoverexposed_params_t)); + memcpy(module->default_params, &tmp, sizeof(dt_iop_rawoverexposed_params_t)); +} + +void init(dt_iop_module_t *module) +{ + module->params = calloc(1, sizeof(dt_iop_rawoverexposed_params_t)); + module->default_params = calloc(1, sizeof(dt_iop_rawoverexposed_params_t)); + module->priority = 60; // module order created by iop_dependencies.py, do not edit! + module->default_enabled = 0; + module->params_size = sizeof(dt_iop_rawoverexposed_params_t); + module->gui_data = NULL; +} + +void cleanup(dt_iop_module_t *module) +{ + free(module->params); + module->params = NULL; +} + +void gui_update(dt_iop_module_t *self) +{ + dt_iop_rawoverexposed_gui_data_t *g = self->gui_data; + dt_iop_rawoverexposed_params_t *p = self->params; + dt_bauhaus_slider_set(g->threshold, p->threshold); +} + +static void threshold_callback(GtkWidget *slider, dt_iop_module_t *self) +{ + if(self->dt->gui->reset) return; + + dt_iop_rawoverexposed_params_t *p = self->params; + p->threshold = dt_bauhaus_slider_get(slider); + dt_dev_add_history_item(darktable.develop, self, TRUE); +} + +void gui_init(dt_iop_module_t *self) +{ + self->gui_data = malloc(sizeof(dt_iop_rawoverexposed_gui_data_t)); + + dt_iop_rawoverexposed_gui_data_t *g = self->gui_data; + dt_iop_rawoverexposed_params_t *p = self->params; + + self->widget = gtk_box_new(GTK_ORIENTATION_VERTICAL, DT_BAUHAUS_SPACE); + + g->threshold = dt_bauhaus_slider_new_with_range(self, 0.0, 2.0, 0.01, p->threshold, 3); + gtk_widget_set_tooltip_text(g->threshold, _("manually adjust the clipping threshold against " + "magenta highlights (you shouldn't ever need to touch this)")); + dt_bauhaus_widget_set_label(g->threshold, NULL, _("clipping threshold")); + g_signal_connect(g->threshold, "value-changed", G_CALLBACK(threshold_callback), self); + + gtk_box_pack_start(GTK_BOX(self->widget), g->threshold, TRUE, TRUE, 0); +} + +void gui_cleanup(dt_iop_module_t *self) +{ + free(self->gui_data); + self->gui_data = NULL; +} + +// modelines: These editor modelines have been set for all relevant files by tools/update_modelines.sh +// vim: shiftwidth=2 expandtab tabstop=2 cindent +// kate: tab-indents: off; indent-width 2; replace-tabs on; indent-mode cstyle; remove-trailing-spaces modified; From 99cf25e38dcd1726bc79c0a99d375ee8e1fd8330 Mon Sep 17 00:00:00 2001 From: Roman Lebedev Date: Sat, 6 Aug 2016 18:16:01 +0300 Subject: [PATCH 2/7] Temperature iop: store wb coeffs also in dsc --- src/develop/format.h | 6 ++++++ src/iop/temperature.c | 15 +++++++++++++++ 2 files changed, 21 insertions(+) diff --git a/src/develop/format.h b/src/develop/format.h index adff9776fab5..4f7bea7d6801 100644 --- a/src/develop/format.h +++ b/src/develop/format.h @@ -49,6 +49,12 @@ typedef struct dt_iop_buffer_dsc_t uint16_t raw_white_point; } rawprepare; + struct + { + int enabled; + float coeffs[4]; + } temperature; + /** sensor saturation, propagated through the operations */ float processed_maximum[4]; } dt_iop_buffer_dsc_t; diff --git a/src/iop/temperature.c b/src/iop/temperature.c index 2ccc205b27fb..549948895a76 100644 --- a/src/iop/temperature.c +++ b/src/iop/temperature.c @@ -446,8 +446,13 @@ void process(struct dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, const if(piece->pipe->mask_display) dt_iop_alpha_copy(ivoid, ovoid, roi_out->width, roi_out->height); } + + piece->pipe->dsc.temperature.enabled = 1; for(int k = 0; k < 4; k++) + { + piece->pipe->dsc.temperature.coeffs[k] = d->coeffs[k]; piece->pipe->dsc.processed_maximum[k] = d->coeffs[k] * piece->pipe->dsc.processed_maximum[k]; + } } #if defined(__SSE__) @@ -532,8 +537,13 @@ void process_sse2(struct dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, c if(piece->pipe->mask_display) dt_iop_alpha_copy(ivoid, ovoid, roi_out->width, roi_out->height); } + + piece->pipe->dsc.temperature.enabled = 1; for(int k = 0; k < 4; k++) + { + piece->pipe->dsc.temperature.coeffs[k] = d->coeffs[k]; piece->pipe->dsc.processed_maximum[k] = d->coeffs[k] * piece->pipe->dsc.processed_maximum[k]; + } } #endif @@ -592,8 +602,13 @@ int process_cl(struct dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, cl_m dt_opencl_release_mem_object(dev_coeffs); if(dev_xtrans != NULL) dt_opencl_release_mem_object(dev_xtrans); + + piece->pipe->dsc.temperature.enabled = 1; for(int k = 0; k < 3; k++) + { + piece->pipe->dsc.temperature.coeffs[k] = d->coeffs[k]; piece->pipe->dsc.processed_maximum[k] = d->coeffs[k] * piece->pipe->dsc.processed_maximum[k]; + } return TRUE; error: From 1fd61048415a5f51a77763e3ff0e98e3e4282cc4 Mon Sep 17 00:00:00 2001 From: Roman Lebedev Date: Sun, 4 Sep 2016 19:09:48 +0300 Subject: [PATCH 3/7] Rawoverexposed: second attempt, after overexposed, sample MIP_FULL. Much better, configurable since it works with 4 channels. --- data/darktableconfig.xml.in | 21 +++ data/kernels/basic.cl | 35 ----- src/develop/develop.c | 9 ++ src/develop/develop.h | 25 +++ src/dtgtk/paint.c | 40 +++++ src/dtgtk/paint.h | 2 + src/iop/rawoverexposed.c | 297 ++++++++++++++++++------------------ src/views/darkroom.c | 182 ++++++++++++++++++++++ tools/iop_dependencies.py | 14 +- 9 files changed, 441 insertions(+), 184 deletions(-) diff --git a/data/darktableconfig.xml.in b/data/darktableconfig.xml.in index 1a7df3ee5f59..4b9851df4ab8 100644 --- a/data/darktableconfig.xml.in +++ b/data/darktableconfig.xml.in @@ -1074,6 +1074,27 @@ do high quality resampling during export the image will first be processed in full resolution, and downscaled at the very end. this can result in better quality sometimes, but will always be slower. + + darkroom/ui/rawoverexposed/mode + int + 0 + + + + + darkroom/ui/rawoverexposed/colorscheme + int + 0 + + + + + darkroom/ui/rawoverexposed/threshold + float + 1.0 + + + darkroom/ui/overexposed/colorscheme int diff --git a/data/kernels/basic.cl b/data/kernels/basic.cl index 03a4b1b95086..190dd85b7cd0 100644 --- a/data/kernels/basic.cl +++ b/data/kernels/basic.cl @@ -163,25 +163,6 @@ highlights_4f_clip (read_only image2d_t in, write_only image2d_t out, const int write_imagef (out, (int2)(x, y), pixel); } -rawoverexposed_4f_mark (read_only image2d_t in, write_only image2d_t out, const int width, const int height, - const float threshold) -{ - const int x = get_global_id(0); - const int y = get_global_id(1); - - if(x >= width || y >= height) return; - - // 4f/pixel means that this has been debayered already. - // (this code path is just used for preview and non-raw images) - float4 pixel = read_imagef(in, sampleri, (int2)(x, y)); - - pixel.x = (pixel.x >= threshold) ? 0.0 : pixel.x; - pixel.y = (pixel.y >= threshold) ? 0.0 : pixel.y; - pixel.z = (pixel.z >= threshold) ? 0.0 : pixel.x; - - write_imagef (out, (int2)(x, y), pixel); -} - kernel void highlights_1f_clip (read_only image2d_t in, write_only image2d_t out, const int width, const int height, const float clip, const int rx, const int ry, const int filters) @@ -198,22 +179,6 @@ highlights_1f_clip (read_only image2d_t in, write_only image2d_t out, const int write_imagef (out, (int2)(x, y), pixel); } -kernel void -rawoverexposed_1f_mark (read_only image2d_t in, write_only image2d_t out, const int width, const int height, - const float threshold) -{ - const int x = get_global_id(0); - const int y = get_global_id(1); - - if(x >= width || y >= height) return; - - float4 pixel = read_imagef(in, sampleri, (int2)(x, y)); - - pixel.x = (pixel.x >= threshold) ? 0.0 : pixel.x; - - write_imagef (out, (int2)(x, y), pixel); -} - #define SQRT3 1.7320508075688772935274463415058723669f #define SQRT12 3.4641016151377545870548926830117447339f // 2*SQRT3 kernel void diff --git a/src/develop/develop.c b/src/develop/develop.c index 91699a812cbe..f35965eae14d 100644 --- a/src/develop/develop.c +++ b/src/develop/develop.c @@ -109,6 +109,11 @@ void dt_dev_init(dt_develop_t *dev, int32_t gui_attached) dev->proxy.exposure = NULL; + dev->rawoverexposed.enabled = FALSE; + dev->rawoverexposed.mode = dt_conf_get_int("darkroom/ui/rawoverexposed/mode"); + dev->rawoverexposed.colorscheme = dt_conf_get_int("darkroom/ui/rawoverexposed/colorscheme"); + dev->rawoverexposed.threshold = dt_conf_get_float("darkroom/ui/rawoverexposed/threshold"); + dev->overexposed.enabled = FALSE; dev->overexposed.colorscheme = dt_conf_get_int("darkroom/ui/overexposed/colorscheme"); dev->overexposed.lower = dt_conf_get_float("darkroom/ui/overexposed/lower"); @@ -152,6 +157,10 @@ void dt_dev_cleanup(dt_develop_t *dev) g_list_free_full(dev->proxy.exposure, g_free); + dt_conf_set_int("darkroom/ui/rawoverexposed/mode", dev->rawoverexposed.mode); + dt_conf_set_int("darkroom/ui/rawoverexposed/colorscheme", dev->rawoverexposed.colorscheme); + dt_conf_set_float("darkroom/ui/rawoverexposed/threshold", dev->rawoverexposed.threshold); + dt_conf_set_int("darkroom/ui/overexposed/colorscheme", dev->overexposed.colorscheme); dt_conf_set_float("darkroom/ui/overexposed/lower", dev->overexposed.lower); dt_conf_set_float("darkroom/ui/overexposed/upper", dev->overexposed.upper); diff --git a/src/develop/develop.h b/src/develop/develop.h index b581c253e967..340045308c42 100644 --- a/src/develop/develop.h +++ b/src/develop/develop.h @@ -49,6 +49,19 @@ typedef enum dt_dev_overexposed_colorscheme_t DT_DEV_OVEREXPOSED_PURPLEGREEN = 2 } dt_dev_overexposed_colorscheme_t; +typedef enum dt_dev_rawoverexposed_mode_t { + DT_DEV_RAWOVEREXPOSED_MODE_MARK_CFA = 0, + DT_DEV_RAWOVEREXPOSED_MODE_MARK_SOLID = 1, + DT_DEV_RAWOVEREXPOSED_MODE_FALSECOLOR = 2, +} dt_dev_rawoverexposed_mode_t; + +typedef enum dt_dev_rawoverexposed_colorscheme_t { + DT_DEV_RAWOVEREXPOSED_RED = 0, + DT_DEV_RAWOVEREXPOSED_GREEN = 1, + DT_DEV_RAWOVEREXPOSED_BLUE = 2, + DT_DEV_RAWOVEREXPOSED_BLACK = 3 +} dt_dev_rawoverexposed_colorscheme_t; + typedef enum dt_dev_histogram_type_t { DT_DEV_HISTOGRAM_LOGARITHMIC = 0, @@ -195,6 +208,18 @@ typedef struct dt_develop_t float upper; } overexposed; + // for the raw overexposure indicator + struct + { + guint timeout; + GtkWidget *floating_window, *button; // yes, having gtk stuff in here is ugly. live with it. + + gboolean enabled; + dt_dev_rawoverexposed_mode_t mode; + dt_dev_rawoverexposed_colorscheme_t colorscheme; + float threshold; + } rawoverexposed; + // the display profile related things (softproof, gamut check, profiles ...) struct { diff --git a/src/dtgtk/paint.c b/src/dtgtk/paint.c index 62c849a63725..6cc98484d15e 100644 --- a/src/dtgtk/paint.c +++ b/src/dtgtk/paint.c @@ -1260,6 +1260,46 @@ void dtgtk_cairo_paint_overexposed(cairo_t *cr, gint x, gint y, gint w, gint h, cairo_stroke(cr); } +void dtgtk_cairo_paint_rawoverexposed(cairo_t *cr, gint x, gint y, gint w, gint h, gint flags) +{ + gint s = w < h ? w : h; + cairo_translate(cr, x + (w / 2.0) - (s / 2.0), y + (h / 2.0) - (s / 2.0)); + cairo_scale(cr, s, s); + + float line_width = 0.15; + + cairo_set_line_cap(cr, CAIRO_LINE_CAP_ROUND); + cairo_set_line_width(cr, line_width); + + cairo_save(cr); + + const double step = ((line_width / 2.0) + (1.0 - line_width) / 2.0); + + // draw 4 CFA-like colored squares + + cairo_set_source_rgba(cr, 1.0, 0.0, 0.0, 1.0); // red + cairo_rectangle(cr, (line_width / 2.0), (line_width / 2.0), step, step); + cairo_fill(cr); + + cairo_set_source_rgba(cr, 0.0, 1.0, 0.0, 1.0); // green + cairo_rectangle(cr, step, (line_width / 2.0), step, step); + cairo_fill(cr); + + cairo_set_source_rgba(cr, 0.0, 1.0, 0.0, 1.0); // green + cairo_rectangle(cr, (line_width / 2.0), step, step, step); + cairo_fill(cr); + + cairo_set_source_rgba(cr, 0.0, 0.0, 1.0, 1.0); // blue + cairo_rectangle(cr, step, step, step, step); + cairo_fill(cr); + + cairo_restore(cr); + + /* outer rect */ + cairo_rectangle(cr, (line_width / 2.0), (line_width / 2.0), 1.0 - line_width, 1.0 - line_width); + cairo_stroke(cr); +} + void dtgtk_cairo_paint_gamut_check(cairo_t *cr, gint x, gint y, gint w, gint h, gint flags) { gint s = w < h ? w : h; diff --git a/src/dtgtk/paint.h b/src/dtgtk/paint.h index 7a082bf117be..13b8e8ed98c8 100644 --- a/src/dtgtk/paint.h +++ b/src/dtgtk/paint.h @@ -124,6 +124,8 @@ void dtgtk_cairo_paint_lock(cairo_t *cr, gint x, gint y, gint w, gint h, gint fl void dtgtk_cairo_paint_check_mark(cairo_t *cr, gint x, gint y, gint w, gint h, gint flags); /** paint an over/under exposure icon */ void dtgtk_cairo_paint_overexposed(cairo_t *cr, gint x, gint y, gint w, gint h, gint flags); +/** paint an raw over exposure icon */ +void dtgtk_cairo_paint_rawoverexposed(cairo_t *cr, gint x, gint y, gint w, gint h, gint flags); /** paint a gamut check icon */ void dtgtk_cairo_paint_gamut_check(cairo_t *cr, gint x, gint y, gint w, gint h, gint flags); /** paint a soft proofing icon */ diff --git a/src/iop/rawoverexposed.c b/src/iop/rawoverexposed.c index 5028611be01c..bc2872f712ee 100644 --- a/src/iop/rawoverexposed.c +++ b/src/iop/rawoverexposed.c @@ -20,37 +20,39 @@ #include "config.h" #endif -#include "bauhaus/bauhaus.h" // for dt_bauhaus_slider_get, dt_bauhaus_... -#include "common/darktable.h" // for darktable_t, dt_print, DT_MODULE_I... -#include "common/image.h" // for ::DT_IMAGE_RAW, dt_image_t -#include "common/opencl.h" // for dt_opencl_set_kernel_arg, dt_openc... -#include "develop/develop.h" // for dt_dev_add_history_item +#include "common/darktable.h" // for darktable, darktable_t, dt_alloc_a... +#include "common/image.h" // for dt_image_t, ::DT_IMAGE_4BAYER +#include "common/mipmap_cache.h" // for dt_mipmap_buffer_t, dt_mipmap_cach... +#include "control/control.h" // for dt_control_log +#include "develop/develop.h" // for dt_develop_t, dt_develop_t::(anony... #include "develop/imageop.h" // for dt_iop_module_t, dt_iop_roi_t, dt_... -#include "develop/imageop_math.h" // for dt_iop_alpha_copy +#include "develop/imageop_math.h" // for FC, FCxtrans #include "develop/pixelpipe.h" // for dt_dev_pixelpipe_type_t::DT_DEV_PI... -#include "gui/gtk.h" // for dt_gui_gtk_t #include "iop/iop_api.h" // for dt_iop_params_t -#include // for g_type_check_instance_cast, GCallback -#include // for TRUE, FALSE #include // for _ -#include // for GtkWidget, gtk_box_get_type, gtk_b... -#include // for fminf -#include // for free, NULL, size_t, calloc, malloc +#include // for GtkWidget +#include // for uint16_t, uint8_t, uint32_t +#include // for size_t, free, NULL, calloc, malloc #include // for memcpy -DT_MODULE_INTROSPECTION(1, dt_iop_rawoverexposed_params_t) +DT_MODULE(1) -typedef struct dt_iop_rawoverexposed_params_t +typedef struct dt_iop_rawoverexposed_t { - float threshold; -} dt_iop_rawoverexposed_params_t; + int dummy; +} dt_iop_rawoverexposed_t; -typedef struct dt_iop_rawoverexposed_gui_data_t -{ - GtkWidget *threshold; -} dt_iop_rawoverexposed_gui_data_t; +static const float dt_iop_rawoverexposed_colors[][4] __attribute__((aligned(16))) = { + { 1.0f, 0.0f, 0.0f, 1.0f }, // red + { 0.0f, 1.0f, 0.0f, 1.0f }, // green + { 0.0f, 0.0f, 1.0f, 1.0f }, // blue + { 0.0f, 0.0f, 0.0f, 1.0f } // black +}; -typedef dt_iop_rawoverexposed_params_t dt_iop_rawoverexposed_data_t; +typedef struct dt_iop_rawoverexposed_data_t +{ + unsigned int threshold[4]; +} dt_iop_rawoverexposed_data_t; typedef struct dt_iop_rawoverexposed_global_data_t { @@ -70,112 +72,158 @@ int groups() int flags() { - return IOP_FLAGS_ALLOW_TILING | IOP_FLAGS_ONE_INSTANCE; + return IOP_FLAGS_ALLOW_TILING | IOP_FLAGS_HIDDEN | IOP_FLAGS_ONE_INSTANCE | IOP_FLAGS_NO_HISTORY_STACK; } -#ifdef HAVE_OPENCL -int process_cl(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) +static void process_common_setup(dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece) { + dt_develop_t *dev = self->dev; dt_iop_rawoverexposed_data_t *d = piece->data; - dt_iop_rawoverexposed_global_data_t *gd = self->data; - - cl_int err = -999; - const int devid = piece->pipe->devid; - const int width = roi_in->width; - const int height = roi_in->height; - - size_t sizes[] = { ROUNDUPWD(width), ROUNDUPHT(height), 1 }; - const float threshold - = d->threshold * fminf(piece->pipe->dsc.processed_maximum[0], - fminf(piece->pipe->dsc.processed_maximum[1], piece->pipe->dsc.processed_maximum[2])); - - const int kernel = (!dt_dev_pixelpipe_uses_downsampled_input(piece->pipe) && piece->pipe->dsc.filters) - ? gd->kernel_rawoverexposed_1f_mark - : gd->kernel_rawoverexposed_4f_mark; - - dt_opencl_set_kernel_arg(devid, kernel, 0, sizeof(cl_mem), (void *)&dev_in); - dt_opencl_set_kernel_arg(devid, kernel, 1, sizeof(cl_mem), (void *)&dev_out); - dt_opencl_set_kernel_arg(devid, kernel, 2, sizeof(int), (void *)&width); - dt_opencl_set_kernel_arg(devid, kernel, 3, sizeof(int), (void *)&height); - dt_opencl_set_kernel_arg(devid, kernel, 4, sizeof(float), (void *)&threshold); - err = dt_opencl_enqueue_kernel_2d(devid, kernel, sizes); - if(err != CL_SUCCESS) goto error; - - return TRUE; - -error: - dt_print(DT_DEBUG_OPENCL, "[opencl_rawoverexposed] couldn't enqueue kernel! %d\n", err); - return FALSE; + + for(int k = 0; k < 4; k++) + { + // the clipping is detected as 1.0 in highlights iop + float threshold = dev->rawoverexposed.threshold; + + // but we check it on the raw input buffer, so we need backtransform thresholds + + // "undo" temperature iop + if(piece->pipe->dsc.temperature.enabled) threshold /= piece->pipe->dsc.temperature.coeffs[k]; + + // "undo" rawprepare iop + threshold *= piece->pipe->dsc.rawprepare.raw_white_point - piece->pipe->dsc.rawprepare.raw_black_level; + threshold += piece->pipe->dsc.rawprepare.raw_black_level; + + // and this is that 1.0 threshold, but in raw input buffer values + d->threshold[k] = (unsigned int)threshold; + } } -#endif void process(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_rawoverexposed_data_t *data = piece->data; + const dt_iop_rawoverexposed_data_t *const d = piece->data; + + process_common_setup(self, piece); + + dt_develop_t *dev = self->dev; + const dt_image_t *const image = &(dev->image_storage); - const float threshold - = data->threshold * fminf(piece->pipe->dsc.processed_maximum[0], - fminf(piece->pipe->dsc.processed_maximum[1], piece->pipe->dsc.processed_maximum[2])); + const int ch = piece->colors; + const int priority = self->priority; - const float *const in = (const float *const)ivoid; + const dt_dev_rawoverexposed_mode_t mode = dev->rawoverexposed.mode; + const int colorscheme = dev->rawoverexposed.colorscheme; + const float *const color = dt_iop_rawoverexposed_colors[colorscheme]; + + memcpy(ovoid, ivoid, (size_t)ch * roi_out->width * roi_out->height * sizeof(float)); + + dt_mipmap_buffer_t buf; + dt_mipmap_cache_get(darktable.mipmap_cache, &buf, image->id, DT_MIPMAP_FULL, DT_MIPMAP_BLOCKING, 'r'); + if(!buf.buf) + { + dt_control_log(_("failed to get raw buffer from image `%s'"), image->filename); + dt_mipmap_cache_release(darktable.mipmap_cache, &buf); + return; + } + +#if 0 + float pts[4] = {(float)(roi_out->x) / roi_in->scale, (float)(roi_out->y) / roi_in->scale, (float)(roi_out->x + roi_out->width) / roi_in->scale, (float)(roi_out->y + roi_out->height) / roi_in->scale}; + printf("in %f %f %f %f\n", pts[0], pts[1], pts[2], pts[3]); + dt_dev_distort_backtransform_plus(dev, dev->pipe, 0, priority, pts, 2); + printf("out %f %f %f %f\n\n", pts[0], pts[1], pts[2], pts[3]); +#endif + + const uint16_t *const raw = (const uint16_t *const)buf.buf; float *const out = (float *const)ovoid; - if(!dt_dev_pixelpipe_uses_downsampled_input(piece->pipe) && piece->pipe->dsc.filters) - { // raw mosaic + // NOT FROM THE PIPE !!! + const uint32_t filters = image->buf_dsc.filters; + const uint8_t(*const xtrans)[6] = (const uint8_t(*const)[6])image->buf_dsc.xtrans; + + // acquire temp memory for distorted pixel coords + const size_t coordbufsize = (size_t)roi_out->width * 2; + void *coordbuf = dt_alloc_align(16, coordbufsize * sizeof(float) * dt_get_num_threads()); + #ifdef _OPENMP -#pragma omp parallel for SIMD() default(none) schedule(static) +#pragma omp parallel for SIMD() default(none) shared(self, coordbuf, buf) schedule(static) #endif - for(size_t k = 0; k < (size_t)roi_out->width * roi_out->height; k++) + for(int j = 0; j < roi_out->height; j++) + { + float *bufptr = ((float *)coordbuf) + (size_t)coordbufsize * dt_get_thread_num(); + + // here are all the pixels of this row + for(int i = 0; i < roi_out->width; i++) { - out[k] = (in[k] >= threshold) ? 0.0 : in[k]; + bufptr[2 * i] = (float)(roi_out->x + i) / roi_in->scale; + bufptr[2 * i + 1] = (float)(roi_out->y + j) / roi_in->scale; } - } - else - { - const int ch = piece->colors; -#ifdef _OPENMP -#pragma omp parallel for SIMD() default(none) schedule(static) -#endif - for(size_t k = 0; k < (size_t)ch * roi_out->width * roi_out->height; k++) + // where did they come from? + dt_dev_distort_backtransform_plus(self->dev, self->dev->pipe, 0, priority, bufptr, roi_out->width); + + for(int i = 0; i < roi_out->width; i++) { - out[k] = (in[k] >= threshold) ? 0.0 : in[k]; + const size_t pout = (size_t)ch * (j * roi_out->width + i); + + // not sure which float -> int to use here + const int i_raw = (int)bufptr[2 * i]; + const int j_raw = (int)bufptr[2 * i + 1]; + + if(i_raw < 0 || j_raw < 0 || i_raw >= buf.width || j_raw >= buf.height) continue; + + int c; + if(filters == 9u) + { + c = FCxtrans(j_raw, i_raw, NULL, xtrans); + } + else // if(filters) + { + c = FC(j_raw, i_raw, filters); + } + + const size_t pin = (size_t)j_raw * buf.width + i_raw; + const float in = raw[pin]; + + // was the raw pixel clipped? + if(in < d->threshold[c]) continue; + + switch(mode) + { + case DT_DEV_RAWOVEREXPOSED_MODE_MARK_CFA: + memcpy(out + pout, dt_iop_rawoverexposed_colors[c], 4 * sizeof(float)); + break; + case DT_DEV_RAWOVEREXPOSED_MODE_MARK_SOLID: + memcpy(out + pout, color, 4 * sizeof(float)); + break; + case DT_DEV_RAWOVEREXPOSED_MODE_FALSECOLOR: + out[pout + c] = 0.0; + break; + } } } + dt_free_align(coordbuf); + + dt_mipmap_cache_release(darktable.mipmap_cache, &buf); + if(piece->pipe->mask_display) dt_iop_alpha_copy(ivoid, ovoid, roi_out->width, roi_out->height); } void commit_params(dt_iop_module_t *self, dt_iop_params_t *p1, dt_dev_pixelpipe_t *pipe, dt_dev_pixelpipe_iop_t *piece) { - dt_iop_rawoverexposed_params_t *p = p1; - dt_iop_rawoverexposed_data_t *d = piece->data; + dt_develop_t *dev = self->dev; - *d = *p; + if(pipe->type != DT_DEV_PIXELPIPE_FULL || !dev->rawoverexposed.enabled || !dev->gui_attached) piece->enabled = 0; - if(pipe->type != DT_DEV_PIXELPIPE_FULL || /*!self->dev->overexposed.enabled ||*/ !self->dev->gui_attached) - piece->enabled = 0; -} + const dt_image_t *const image = &(dev->image_storage); -void init_global(dt_iop_module_so_t *module) -{ - const int program = 2; // basic.cl, from programs.conf - module->data = malloc(sizeof(dt_iop_rawoverexposed_global_data_t)); - dt_iop_rawoverexposed_global_data_t *gd = module->data; - gd->kernel_rawoverexposed_1f_mark = dt_opencl_create_kernel(program, "rawoverexposed_1f_mark"); - gd->kernel_rawoverexposed_4f_mark = dt_opencl_create_kernel(program, "rawoverexposed_4f_mark"); -} + if(image->flags & DT_IMAGE_4BAYER) piece->enabled = 0; -void cleanup_global(dt_iop_module_so_t *module) -{ - dt_iop_rawoverexposed_global_data_t *gd = module->data; - dt_opencl_free_kernel(gd->kernel_rawoverexposed_4f_mark); - dt_opencl_free_kernel(gd->kernel_rawoverexposed_1f_mark); - free(module->data); - module->data = NULL; + if(image->buf_dsc.datatype != TYPE_UINT16 || !image->buf_dsc.filters) piece->enabled = 0; + + piece->process_cl_ready = 0; } void init_pipe(dt_iop_module_t *self, dt_dev_pixelpipe_t *pipe, dt_dev_pixelpipe_iop_t *piece) @@ -190,21 +238,14 @@ void cleanup_pipe(dt_iop_module_t *self, dt_dev_pixelpipe_t *pipe, dt_dev_pixelp piece->data = NULL; } -void reload_defaults(dt_iop_module_t *module) -{ - dt_iop_rawoverexposed_params_t tmp = (dt_iop_rawoverexposed_params_t){.threshold = 1.0 }; - - memcpy(module->params, &tmp, sizeof(dt_iop_rawoverexposed_params_t)); - memcpy(module->default_params, &tmp, sizeof(dt_iop_rawoverexposed_params_t)); -} - void init(dt_iop_module_t *module) { - module->params = calloc(1, sizeof(dt_iop_rawoverexposed_params_t)); - module->default_params = calloc(1, sizeof(dt_iop_rawoverexposed_params_t)); - module->priority = 60; // module order created by iop_dependencies.py, do not edit! - module->default_enabled = 0; - module->params_size = sizeof(dt_iop_rawoverexposed_params_t); + module->params = calloc(1, sizeof(dt_iop_rawoverexposed_t)); + module->default_params = calloc(1, sizeof(dt_iop_rawoverexposed_t)); + module->hide_enable_button = 1; + module->default_enabled = 1; + module->priority = 924; // module order created by iop_dependencies.py, do not edit! + module->params_size = sizeof(dt_iop_rawoverexposed_t); module->gui_data = NULL; } @@ -214,46 +255,6 @@ void cleanup(dt_iop_module_t *module) module->params = NULL; } -void gui_update(dt_iop_module_t *self) -{ - dt_iop_rawoverexposed_gui_data_t *g = self->gui_data; - dt_iop_rawoverexposed_params_t *p = self->params; - dt_bauhaus_slider_set(g->threshold, p->threshold); -} - -static void threshold_callback(GtkWidget *slider, dt_iop_module_t *self) -{ - if(self->dt->gui->reset) return; - - dt_iop_rawoverexposed_params_t *p = self->params; - p->threshold = dt_bauhaus_slider_get(slider); - dt_dev_add_history_item(darktable.develop, self, TRUE); -} - -void gui_init(dt_iop_module_t *self) -{ - self->gui_data = malloc(sizeof(dt_iop_rawoverexposed_gui_data_t)); - - dt_iop_rawoverexposed_gui_data_t *g = self->gui_data; - dt_iop_rawoverexposed_params_t *p = self->params; - - self->widget = gtk_box_new(GTK_ORIENTATION_VERTICAL, DT_BAUHAUS_SPACE); - - g->threshold = dt_bauhaus_slider_new_with_range(self, 0.0, 2.0, 0.01, p->threshold, 3); - gtk_widget_set_tooltip_text(g->threshold, _("manually adjust the clipping threshold against " - "magenta highlights (you shouldn't ever need to touch this)")); - dt_bauhaus_widget_set_label(g->threshold, NULL, _("clipping threshold")); - g_signal_connect(g->threshold, "value-changed", G_CALLBACK(threshold_callback), self); - - gtk_box_pack_start(GTK_BOX(self->widget), g->threshold, TRUE, TRUE, 0); -} - -void gui_cleanup(dt_iop_module_t *self) -{ - free(self->gui_data); - self->gui_data = NULL; -} - // modelines: These editor modelines have been set for all relevant files by tools/update_modelines.sh // vim: shiftwidth=2 expandtab tabstop=2 cindent // kate: tab-indents: off; indent-width 2; replace-tabs on; indent-mode cstyle; remove-trailing-spaces modified; diff --git a/src/views/darkroom.c b/src/views/darkroom.c index 48a94e165687..09a52268dfdb 100644 --- a/src/views/darkroom.c +++ b/src/views/darkroom.c @@ -1050,6 +1050,106 @@ static void upper_callback(GtkWidget *slider, gpointer user_data) dt_dev_reprocess_all(d); } +/* rawoverexposed */ +static void _rawoverexposed_quickbutton_clicked(GtkWidget *w, gpointer user_data) +{ + dt_develop_t *d = (dt_develop_t *)user_data; + d->rawoverexposed.enabled = !d->rawoverexposed.enabled; + // dt_dev_reprocess_center(d); + dt_dev_reprocess_all(d); +} + +static gboolean _rawoverexposed_close_popup(GtkWidget *widget, GdkEvent *event, gpointer user_data) +{ + dt_develop_t *d = (dt_develop_t *)user_data; + if(!gtk_widget_is_visible(darktable.bauhaus->popup_window)) gtk_widget_hide(d->rawoverexposed.floating_window); + return FALSE; +} + +static gboolean _rawoverexposed_show_popup(gpointer user_data) +{ + dt_develop_t *d = (dt_develop_t *)user_data; + /** finally move the window next to the button */ + gint x, y, wx, wy; + gint px, py, window_w, window_h; + GtkWidget *window = dt_ui_main_window(darktable.gui->ui); + gtk_widget_show_all(d->rawoverexposed.floating_window); + gdk_window_get_origin(gtk_widget_get_window(d->rawoverexposed.button), &px, &py); + + window_w = gdk_window_get_width(gtk_widget_get_window(d->rawoverexposed.floating_window)); + window_h = gdk_window_get_height(gtk_widget_get_window(d->rawoverexposed.floating_window)); + + gtk_widget_translate_coordinates(d->rawoverexposed.button, window, 0, 0, &wx, &wy); + x = px + wx - window_w + DT_PIXEL_APPLY_DPI(5); + y = py + wy - window_h - DT_PIXEL_APPLY_DPI(5); + gtk_window_move(GTK_WINDOW(d->rawoverexposed.floating_window), x, y); + + gtk_window_present(GTK_WINDOW(d->rawoverexposed.floating_window)); + + // when the mouse moves back over the main window we close the popup. + g_signal_connect(d->rawoverexposed.floating_window, "focus-out-event", G_CALLBACK(_rawoverexposed_close_popup), + user_data); + + return FALSE; +} + +static gboolean _rawoverexposed_quickbutton_pressed(GtkWidget *widget, GdkEvent *event, gpointer user_data) +{ + dt_develop_t *d = (dt_develop_t *)user_data; + const GdkEventButton *e = (GdkEventButton *)event; + if(e->button == 3) + { + _rawoverexposed_show_popup(user_data); + return TRUE; + } + else + { + d->rawoverexposed.timeout = g_timeout_add_seconds(1, _rawoverexposed_show_popup, user_data); + return FALSE; + } +} + +static gboolean _rawoverexposed_quickbutton_released(GtkWidget *widget, GdkEvent *event, gpointer user_data) +{ + dt_develop_t *d = (dt_develop_t *)user_data; + if(d->rawoverexposed.timeout > 0) g_source_remove(d->rawoverexposed.timeout); + d->rawoverexposed.timeout = 0; + return FALSE; +} + +static void rawoverexposed_mode_callback(GtkWidget *combo, gpointer user_data) +{ + dt_develop_t *d = (dt_develop_t *)user_data; + d->rawoverexposed.mode = dt_bauhaus_combobox_get(combo); + if(d->rawoverexposed.enabled == FALSE) + gtk_button_clicked(GTK_BUTTON(d->rawoverexposed.button)); + else + // dt_dev_reprocess_center(d); + dt_dev_reprocess_all(d); +} + +static void rawoverexposed_colorscheme_callback(GtkWidget *combo, gpointer user_data) +{ + dt_develop_t *d = (dt_develop_t *)user_data; + d->rawoverexposed.colorscheme = dt_bauhaus_combobox_get(combo); + if(d->rawoverexposed.enabled == FALSE) + gtk_button_clicked(GTK_BUTTON(d->rawoverexposed.button)); + else + // dt_dev_reprocess_center(d); + dt_dev_reprocess_all(d); +} + +static void rawoverexposed_threshold_callback(GtkWidget *slider, gpointer user_data) +{ + dt_develop_t *d = (dt_develop_t *)user_data; + d->rawoverexposed.threshold = dt_bauhaus_slider_get(slider); + if(d->rawoverexposed.enabled == FALSE) + gtk_button_clicked(GTK_BUTTON(d->rawoverexposed.button)); + else + // dt_dev_reprocess_center(d); + dt_dev_reprocess_all(d); +} + static gboolean _toolbox_toggle_callback(GtkAccelGroup *accel_group, GObject *acceleratable, guint keyval, GdkModifierType modifier, gpointer data) { @@ -1394,6 +1494,88 @@ void gui_init(dt_view_t *self) gtk_widget_set_tooltip_text(styles, _("quick access for applying any of your styles")); dt_view_manager_view_toolbox_add(darktable.view_manager, styles, DT_VIEW_DARKROOM); + /* create rawoverexposed popup tool */ + { + // the button + dev->rawoverexposed.button + = dtgtk_togglebutton_new(dtgtk_cairo_paint_rawoverexposed, CPF_STYLE_FLAT | CPF_DO_NOT_USE_BORDER); + gtk_widget_set_tooltip_text(dev->rawoverexposed.button, + _("toggle raw over exposed indication\nright click for options")); + g_signal_connect(G_OBJECT(dev->rawoverexposed.button), "clicked", + G_CALLBACK(_rawoverexposed_quickbutton_clicked), dev); + g_signal_connect(G_OBJECT(dev->rawoverexposed.button), "button-press-event", + G_CALLBACK(_rawoverexposed_quickbutton_pressed), dev); + g_signal_connect(G_OBJECT(dev->rawoverexposed.button), "button-release-event", + G_CALLBACK(_rawoverexposed_quickbutton_released), dev); + dt_view_manager_module_toolbox_add(darktable.view_manager, dev->rawoverexposed.button, DT_VIEW_DARKROOM); + + // and the popup window + const int panel_width = dt_conf_get_int("panel_width"); + + GtkWidget *window = dt_ui_main_window(darktable.gui->ui); + + dev->rawoverexposed.floating_window = gtk_window_new(GTK_WINDOW_TOPLEVEL); + gtk_window_set_default_size(GTK_WINDOW(dev->rawoverexposed.floating_window), panel_width, -1); + GtkWidget *frame = gtk_frame_new(NULL); + GtkWidget *event_box = gtk_event_box_new(); + GtkWidget *vbox = gtk_box_new(GTK_ORIENTATION_VERTICAL, 5); + gtk_widget_set_margin_start(vbox, DT_PIXEL_APPLY_DPI(8)); + gtk_widget_set_margin_end(vbox, DT_PIXEL_APPLY_DPI(8)); + gtk_widget_set_margin_top(vbox, DT_PIXEL_APPLY_DPI(8)); + gtk_widget_set_margin_bottom(vbox, DT_PIXEL_APPLY_DPI(8)); + + gtk_widget_set_can_focus(dev->rawoverexposed.floating_window, TRUE); + gtk_window_set_decorated(GTK_WINDOW(dev->rawoverexposed.floating_window), FALSE); + gtk_window_set_type_hint(GTK_WINDOW(dev->rawoverexposed.floating_window), GDK_WINDOW_TYPE_HINT_POPUP_MENU); + gtk_window_set_transient_for(GTK_WINDOW(dev->rawoverexposed.floating_window), GTK_WINDOW(window)); + gtk_widget_set_opacity(dev->rawoverexposed.floating_window, 0.9); + + gtk_widget_set_state_flags(frame, GTK_STATE_FLAG_SELECTED, TRUE); + gtk_frame_set_shadow_type(GTK_FRAME(frame), GTK_SHADOW_OUT); + + + gtk_container_add(GTK_CONTAINER(dev->rawoverexposed.floating_window), frame); + gtk_container_add(GTK_CONTAINER(frame), event_box); + gtk_container_add(GTK_CONTAINER(event_box), vbox); + + /** let's fill the encapsulating widgets */ + /* mode of operation */ + GtkWidget *mode = dt_bauhaus_combobox_new(NULL); + dt_bauhaus_widget_set_label(mode, NULL, _("mode")); + dt_bauhaus_combobox_add(mode, _("mark with CFA color")); + dt_bauhaus_combobox_add(mode, _("mark with solid color")); + dt_bauhaus_combobox_add(mode, _("false color")); + dt_bauhaus_combobox_set(mode, dev->rawoverexposed.mode); + gtk_widget_set_tooltip_text(mode, _("select how to mark the clipped pixels")); + g_signal_connect(G_OBJECT(mode), "value-changed", G_CALLBACK(rawoverexposed_mode_callback), dev); + gtk_box_pack_start(GTK_BOX(vbox), GTK_WIDGET(mode), TRUE, TRUE, 0); + gtk_widget_set_state_flags(mode, GTK_STATE_FLAG_SELECTED, TRUE); + + /* color scheme */ + GtkWidget *colorscheme = dt_bauhaus_combobox_new(NULL); + dt_bauhaus_widget_set_label(colorscheme, NULL, _("solid color scheme")); + dt_bauhaus_combobox_add(colorscheme, _("red")); + dt_bauhaus_combobox_add(colorscheme, _("green")); + dt_bauhaus_combobox_add(colorscheme, _("blue")); + dt_bauhaus_combobox_add(colorscheme, _("black")); + dt_bauhaus_combobox_set(colorscheme, dev->rawoverexposed.colorscheme); + gtk_widget_set_tooltip_text( + colorscheme, + _("select the solid color to indicate over exposure.\nwill only be used if mode = mark with solid color")); + g_signal_connect(G_OBJECT(colorscheme), "value-changed", G_CALLBACK(rawoverexposed_colorscheme_callback), dev); + gtk_box_pack_start(GTK_BOX(vbox), GTK_WIDGET(colorscheme), TRUE, TRUE, 0); + gtk_widget_set_state_flags(colorscheme, GTK_STATE_FLAG_SELECTED, TRUE); + + /* threshold */ + GtkWidget *threshold = dt_bauhaus_slider_new_with_range(NULL, 0.0, 2.0, 0.01, 1.0, 3); + dt_bauhaus_slider_set(threshold, dev->rawoverexposed.threshold); + dt_bauhaus_widget_set_label(threshold, NULL, _("clipping threshold")); + gtk_widget_set_tooltip_text( + threshold, _("threshold of what shall be considered overexposed\n1.0 - white level\n0.0 - black level")); + g_signal_connect(G_OBJECT(threshold), "value-changed", G_CALLBACK(rawoverexposed_threshold_callback), dev); + gtk_box_pack_start(GTK_BOX(vbox), GTK_WIDGET(threshold), TRUE, TRUE, 0); + } + /* create overexposed popup tool */ { // the button diff --git a/tools/iop_dependencies.py b/tools/iop_dependencies.py index 11c6a5dbe359..ef14ba5f49d8 100755 --- a/tools/iop_dependencies.py +++ b/tools/iop_dependencies.py @@ -282,6 +282,7 @@ def add_edges(gr): gr.add_edge(('gamma', 'splittoning')) gr.add_edge(('gamma', 'watermark')) gr.add_edge(('gamma', 'overexposed')) + gr.add_edge(('gamma', 'rawoverexposed')) gr.add_edge(('gamma', 'borders')) gr.add_edge(('gamma', 'dither')) gr.add_edge(('channelmixer', 'colorout')) @@ -292,6 +293,7 @@ def add_edges(gr): gr.add_edge(('splittoning', 'colorout')) gr.add_edge(('watermark', 'colorout')) gr.add_edge(('overexposed', 'colorout')) + gr.add_edge(('rawoverexposed', 'colorout')) gr.add_edge(('dither', 'colorout')) # borders should not change shape/color: @@ -304,6 +306,7 @@ def add_edges(gr): gr.add_edge(('borders', 'channelmixer')) # don't indicate borders as over/under exposed gr.add_edge(('borders', 'overexposed')) + gr.add_edge(('borders', 'rawoverexposed')) # can, but no need to # don't resample borders when scaling to the output dimensions gr.add_edge(('borders', 'finalscale')) @@ -315,7 +318,15 @@ def add_edges(gr): gr.add_edge(('finalscale', 'soften')) gr.add_edge(('finalscale', 'clahe')) gr.add_edge(('finalscale', 'channelmixer')) - gr.add_edge(('finalscale', 'overexposed')) + + # but can display overexposure after scaling + # NOTE: finalscale is only done in export pipe, + # while *overexposed is only done in full darkroom preview pipe + gr.add_edge(('overexposed', 'finalscale')) + gr.add_edge(('rawoverexposed', 'finalscale')) + + # let's display raw overexposure indication after usual overexposed + gr.add_edge(('rawoverexposed', 'overexposed')) # but watermark can be drawn on top of borders gr.add_edge(('watermark', 'borders')) @@ -504,6 +515,7 @@ def add_edges(gr): 'monochrome', 'nlmeans', 'overexposed', +'rawoverexposed', 'profile_gamma', 'rawdenoise', 'relight', From 994c844610615bd0886da63593ee25eeab5b4825 Mon Sep 17 00:00:00 2001 From: Roman Lebedev Date: Sun, 4 Sep 2016 19:13:28 +0300 Subject: [PATCH 4/7] Rawoverexposed iop: blind attempt at opencl kernel for false color mode --- data/kernels/basic.cl | 36 ++++++++++ src/iop/rawoverexposed.c | 139 +++++++++++++++++++++++++++++++++++++-- 2 files changed, 169 insertions(+), 6 deletions(-) diff --git a/data/kernels/basic.cl b/data/kernels/basic.cl index 190dd85b7cd0..02d307231afb 100644 --- a/data/kernels/basic.cl +++ b/data/kernels/basic.cl @@ -1805,6 +1805,42 @@ overexposed (read_only image2d_t in, write_only image2d_t out, const int width, } +/* kernel for the rawoverexposed plugin. */ +kernel void +rawoverexposed_falsecolor ( + read_only image2d_t in, write_only image2d_t out, global float *pi, + const int width, const int height, + read_only image2d_t raw, const int raw_width, const int raw_height, + const unsigned int filters, global unsigned int *threshold) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + + if(x >= width || y >= height) return; + + const int piwidth = 2*width; + global float *ppi = pi + mad24(y, piwidth, 2*x); + + const int raw_x = ppi[0]; + const int raw_y = ppi[1]; + + if(raw_x < 0 || raw_y < 0 || raw_x >= raw_width || raw_y >= raw_height) return; + + const uint raw_pixel = read_imageui(raw, sampleri, (int2)(raw_x, raw_y)).x; + + const int c = FC(raw_y, raw_x, filters); + + if(raw_pixel < threshold[c]) return; + + float4 pixel = read_imagef(in, sampleri, (int2)(x, y)); + + // falsecolor + pixel[c] = 0.0; + + write_imagef (out, (int2)(x, y), pixel); +} + + /* kernel for the lowlight plugin. */ kernel void lowlight (read_only image2d_t in, write_only image2d_t out, const int width, const int height, diff --git a/src/iop/rawoverexposed.c b/src/iop/rawoverexposed.c index bc2872f712ee..20ddf8b9b4ee 100644 --- a/src/iop/rawoverexposed.c +++ b/src/iop/rawoverexposed.c @@ -20,9 +20,10 @@ #include "config.h" #endif -#include "common/darktable.h" // for darktable, darktable_t, dt_alloc_a... -#include "common/image.h" // for dt_image_t, ::DT_IMAGE_4BAYER -#include "common/mipmap_cache.h" // for dt_mipmap_buffer_t, dt_mipmap_cach... +#include "common/darktable.h" // for darktable, darktable_t, dt_alloc_a... +#include "common/image.h" // for dt_image_t, ::DT_IMAGE_4BAYER +#include "common/mipmap_cache.h" // for dt_mipmap_buffer_t, dt_mipmap_cach... +#include "common/opencl.h" #include "control/control.h" // for dt_control_log #include "develop/develop.h" // for dt_develop_t, dt_develop_t::(anony... #include "develop/imageop.h" // for dt_iop_module_t, dt_iop_roi_t, dt_... @@ -56,8 +57,7 @@ typedef struct dt_iop_rawoverexposed_data_t typedef struct dt_iop_rawoverexposed_global_data_t { - int kernel_rawoverexposed_1f_mark; - int kernel_rawoverexposed_4f_mark; + int kernel_rawoverexposed_falsecolor; } dt_iop_rawoverexposed_global_data_t; const char *name() @@ -210,6 +210,114 @@ void process(dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, const void *c if(piece->pipe->mask_display) dt_iop_alpha_copy(ivoid, ovoid, roi_out->width, roi_out->height); } +#ifdef HAVE_OPENCL +int 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) +{ + const dt_iop_rawoverexposed_data_t *const d = piece->data; + dt_develop_t *dev = self->dev; + dt_iop_rawoverexposed_global_data_t *gd = (dt_iop_rawoverexposed_global_data_t *)self->data; + + dt_mipmap_buffer_t buf; + + cl_mem dev_raw = NULL; + float *coordbuf = NULL; + cl_mem dev_coord = NULL; + cl_mem dev_thresholds = NULL; + + cl_int err = -999; + + const int devid = piece->pipe->devid; + + const int width = roi_out->width; + const int height = roi_out->height; + + size_t origin[] = { 0, 0, 0 }; + size_t region[] = { width, height, 1 }; + + err = dt_opencl_enqueue_copy_image(devid, dev_in, dev_out, origin, origin, region); + if(err != CL_SUCCESS) goto error; + + const dt_image_t *const image = &(dev->image_storage); + + // NOT FROM THE PIPE !!! + const uint32_t filters = image->buf_dsc.filters; + + dt_mipmap_cache_get(darktable.mipmap_cache, &buf, image->id, DT_MIPMAP_FULL, DT_MIPMAP_BLOCKING, 'r'); + if(!buf.buf) + { + dt_control_log(_("failed to get raw buffer from image `%s'"), image->filename); + dt_mipmap_cache_release(darktable.mipmap_cache, &buf); + goto error; + } + + const int raw_width = buf.width; + const int raw_height = buf.height; + + dev_raw = dt_opencl_copy_host_to_device(devid, buf.buf, raw_width, raw_height, sizeof(uint16_t)); + if(dev_raw == NULL) goto error; + + const size_t coordbufsize = (size_t)height * width * 2 * sizeof(float); + + coordbuf = dt_alloc_align(16, coordbufsize); + if(coordbuf == NULL) goto error; + +#ifdef _OPENMP +#pragma omp parallel for SIMD() default(none) shared(self, coordbuf, buf) schedule(static) +#endif + for(int j = 0; j < height; j++) + { + float *bufptr = ((float *)coordbuf) + (size_t)2 * j * width; + + // here are all the pixels of this row + for(int i = 0; i < roi_out->width; i++) + { + bufptr[2 * i] = (float)(roi_out->x + i) / roi_in->scale; + bufptr[2 * i + 1] = (float)(roi_out->y + j) / roi_in->scale; + } + + // where did they come from? + dt_dev_distort_backtransform_plus(self->dev, self->dev->pipe, 0, self->priority, bufptr, roi_out->width); + } + + dev_coord = dt_opencl_alloc_device_buffer(devid, coordbufsize); + if(dev_coord == NULL) goto error; + + /* _blocking_ memory transfer: host coordbuf buffer -> opencl dev_coordbuf */ + err = dt_opencl_write_buffer_to_device(devid, coordbuf, dev_coord, 0, coordbufsize, CL_TRUE); + if(err != CL_SUCCESS) goto error; + + int kernel = gd->kernel_rawoverexposed_falsecolor; + + dev_thresholds = dt_opencl_copy_host_to_device_constant(devid, sizeof(unsigned int) * 4, (void *)d->threshold); + if(dev_thresholds == NULL) goto error; + + size_t sizes[2] = { ROUNDUPWD(width), ROUNDUPHT(height) }; + dt_opencl_set_kernel_arg(devid, kernel, 0, sizeof(cl_mem), &dev_in); + dt_opencl_set_kernel_arg(devid, kernel, 1, sizeof(cl_mem), &dev_out); + dt_opencl_set_kernel_arg(devid, kernel, 2, sizeof(cl_mem), &dev_coord); + dt_opencl_set_kernel_arg(devid, kernel, 3, sizeof(int), &width); + dt_opencl_set_kernel_arg(devid, kernel, 4, sizeof(int), &height); + dt_opencl_set_kernel_arg(devid, kernel, 5, sizeof(cl_mem), &dev_raw); + dt_opencl_set_kernel_arg(devid, kernel, 6, sizeof(int), &raw_width); + dt_opencl_set_kernel_arg(devid, kernel, 7, sizeof(int), &raw_height); + dt_opencl_set_kernel_arg(devid, kernel, 8, sizeof(uint32_t), &filters); + dt_opencl_set_kernel_arg(devid, kernel, 9, sizeof(cl_mem), &dev_thresholds); + err = dt_opencl_enqueue_kernel_2d(devid, kernel, sizes); + if(err != CL_SUCCESS) goto error; + return TRUE; + +error: + if(dev_thresholds != NULL) dt_opencl_release_mem_object(dev_thresholds); + if(dev_coord != NULL) dt_opencl_release_mem_object(dev_coord); + if(coordbuf != NULL) dt_free_align(coordbuf); + if(dev_raw != NULL) dt_opencl_release_mem_object(dev_raw); + dt_mipmap_cache_release(darktable.mipmap_cache, &buf); + dt_print(DT_DEBUG_OPENCL, "[opencl_rawoverexposed] couldn't enqueue kernel! %d\n", err); + return FALSE; +} +#endif + void commit_params(dt_iop_module_t *self, dt_iop_params_t *p1, dt_dev_pixelpipe_t *pipe, dt_dev_pixelpipe_iop_t *piece) { @@ -223,7 +331,26 @@ void commit_params(dt_iop_module_t *self, dt_iop_params_t *p1, dt_dev_pixelpipe_ if(image->buf_dsc.datatype != TYPE_UINT16 || !image->buf_dsc.filters) piece->enabled = 0; - piece->process_cl_ready = 0; + if(image->buf_dsc.filters == 9u) piece->process_cl_ready = 0; + + if(dev->rawoverexposed.mode != DT_DEV_RAWOVEREXPOSED_MODE_FALSECOLOR) piece->process_cl_ready = 0; +} + +void init_global(dt_iop_module_so_t *module) +{ + const int program = 2; // basic.cl from programs.conf + module->data = malloc(sizeof(dt_iop_rawoverexposed_global_data_t)); + dt_iop_rawoverexposed_global_data_t *gd = module->data; + gd->kernel_rawoverexposed_falsecolor = dt_opencl_create_kernel(program, "rawoverexposed_falsecolor"); +} + + +void cleanup_global(dt_iop_module_so_t *module) +{ + dt_iop_rawoverexposed_global_data_t *gd = (dt_iop_rawoverexposed_global_data_t *)module->data; + dt_opencl_free_kernel(gd->kernel_rawoverexposed_falsecolor); + free(module->data); + module->data = NULL; } void init_pipe(dt_iop_module_t *self, dt_dev_pixelpipe_t *pipe, dt_dev_pixelpipe_iop_t *piece) From 7717fa5693a36f533914c4d131b04610efd03f52 Mon Sep 17 00:00:00 2001 From: Roman Lebedev Date: Fri, 5 Aug 2016 22:47:04 +0300 Subject: [PATCH 5/7] Rawoverexposed iop: blind attempt at opencl kernel for mark solid mode --- data/kernels/basic.cl | 35 +++++++++++++++++++++++++++++++++++ src/iop/rawoverexposed.c | 23 +++++++++++++++++++++-- 2 files changed, 56 insertions(+), 2 deletions(-) diff --git a/data/kernels/basic.cl b/data/kernels/basic.cl index 02d307231afb..9d07ba47b6e0 100644 --- a/data/kernels/basic.cl +++ b/data/kernels/basic.cl @@ -1806,6 +1806,41 @@ overexposed (read_only image2d_t in, write_only image2d_t out, const int width, /* kernel for the rawoverexposed plugin. */ +kernel void +rawoverexposed_mark_solid ( + read_only image2d_t in, write_only image2d_t out, global float *pi, + const int width, const int height, + read_only image2d_t raw, const int raw_width, const int raw_height, + const unsigned int filters, global unsigned int *threshold, + const float4 solid_color) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + + if(x >= width || y >= height) return; + + const int piwidth = 2*width; + global float *ppi = pi + mad24(y, piwidth, 2*x); + + const int raw_x = ppi[0]; + const int raw_y = ppi[1]; + + if(raw_x < 0 || raw_y < 0 || raw_x >= raw_width || raw_y >= raw_height) return; + + const uint raw_pixel = read_imageui(raw, sampleri, (int2)(raw_x, raw_y)).x; + + const int c = FC(raw_y, raw_x, filters); + + if(raw_pixel < threshold[c]) return; + + float4 pixel = read_imagef(in, sampleri, (int2)(x, y)); + + // solid color + pixel.xyz = solid_color.xyz; + + write_imagef (out, (int2)(x, y), pixel); +} + kernel void rawoverexposed_falsecolor ( read_only image2d_t in, write_only image2d_t out, global float *pi, diff --git a/src/iop/rawoverexposed.c b/src/iop/rawoverexposed.c index 20ddf8b9b4ee..a9b4f5c93674 100644 --- a/src/iop/rawoverexposed.c +++ b/src/iop/rawoverexposed.c @@ -57,6 +57,7 @@ typedef struct dt_iop_rawoverexposed_data_t typedef struct dt_iop_rawoverexposed_global_data_t { + int kernel_rawoverexposed_mark_solid; int kernel_rawoverexposed_falsecolor; } dt_iop_rawoverexposed_global_data_t; @@ -240,6 +241,9 @@ int process_cl(struct dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, cl_m const dt_image_t *const image = &(dev->image_storage); + const int colorscheme = dev->rawoverexposed.colorscheme; + const float *const color = dt_iop_rawoverexposed_colors[colorscheme]; + // NOT FROM THE PIPE !!! const uint32_t filters = image->buf_dsc.filters; @@ -287,7 +291,16 @@ int process_cl(struct dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, cl_m err = dt_opencl_write_buffer_to_device(devid, coordbuf, dev_coord, 0, coordbufsize, CL_TRUE); if(err != CL_SUCCESS) goto error; - int kernel = gd->kernel_rawoverexposed_falsecolor; + int kernel; + switch(dev->rawoverexposed.mode) + { + case DT_DEV_RAWOVEREXPOSED_MODE_MARK_SOLID: + kernel = gd->kernel_rawoverexposed_mark_solid; + break; + default: + kernel = gd->kernel_rawoverexposed_falsecolor; + break; + } dev_thresholds = dt_opencl_copy_host_to_device_constant(devid, sizeof(unsigned int) * 4, (void *)d->threshold); if(dev_thresholds == NULL) goto error; @@ -303,6 +316,10 @@ int process_cl(struct dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, cl_m dt_opencl_set_kernel_arg(devid, kernel, 7, sizeof(int), &raw_height); dt_opencl_set_kernel_arg(devid, kernel, 8, sizeof(uint32_t), &filters); dt_opencl_set_kernel_arg(devid, kernel, 9, sizeof(cl_mem), &dev_thresholds); + + if(dev->rawoverexposed.mode == DT_DEV_RAWOVEREXPOSED_MODE_MARK_SOLID) + dt_opencl_set_kernel_arg(devid, kernel, 10, 4 * sizeof(float), color); + err = dt_opencl_enqueue_kernel_2d(devid, kernel, sizes); if(err != CL_SUCCESS) goto error; return TRUE; @@ -333,7 +350,7 @@ void commit_params(dt_iop_module_t *self, dt_iop_params_t *p1, dt_dev_pixelpipe_ if(image->buf_dsc.filters == 9u) piece->process_cl_ready = 0; - if(dev->rawoverexposed.mode != DT_DEV_RAWOVEREXPOSED_MODE_FALSECOLOR) piece->process_cl_ready = 0; + if(dev->rawoverexposed.mode == DT_DEV_RAWOVEREXPOSED_MODE_MARK_CFA) piece->process_cl_ready = 0; } void init_global(dt_iop_module_so_t *module) @@ -341,6 +358,7 @@ void init_global(dt_iop_module_so_t *module) const int program = 2; // basic.cl from programs.conf module->data = malloc(sizeof(dt_iop_rawoverexposed_global_data_t)); dt_iop_rawoverexposed_global_data_t *gd = module->data; + gd->kernel_rawoverexposed_mark_solid = dt_opencl_create_kernel(program, "rawoverexposed_mark_solid"); gd->kernel_rawoverexposed_falsecolor = dt_opencl_create_kernel(program, "rawoverexposed_falsecolor"); } @@ -349,6 +367,7 @@ void cleanup_global(dt_iop_module_so_t *module) { dt_iop_rawoverexposed_global_data_t *gd = (dt_iop_rawoverexposed_global_data_t *)module->data; dt_opencl_free_kernel(gd->kernel_rawoverexposed_falsecolor); + dt_opencl_free_kernel(gd->kernel_rawoverexposed_mark_solid); free(module->data); module->data = NULL; } From ca05f9cce026a2dc350fc23c56ed7143fb78a51b Mon Sep 17 00:00:00 2001 From: Roman Lebedev Date: Fri, 5 Aug 2016 22:47:18 +0300 Subject: [PATCH 6/7] Rawoverexposed iop: blind attempt at opencl kernel for mark cfa mode --- data/kernels/basic.cl | 39 +++++++++++++++++++++++++++++++++++++++ src/iop/rawoverexposed.c | 24 +++++++++++++++++++++--- 2 files changed, 60 insertions(+), 3 deletions(-) diff --git a/data/kernels/basic.cl b/data/kernels/basic.cl index 9d07ba47b6e0..5ac1ad779395 100644 --- a/data/kernels/basic.cl +++ b/data/kernels/basic.cl @@ -1806,6 +1806,45 @@ overexposed (read_only image2d_t in, write_only image2d_t out, const int width, /* kernel for the rawoverexposed plugin. */ +kernel void +rawoverexposed_mark_cfa ( + read_only image2d_t in, write_only image2d_t out, global float *pi, + const int width, const int height, + read_only image2d_t raw, const int raw_width, const int raw_height, + const unsigned int filters, global unsigned int *threshold, + global float *colors) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + + if(x >= width || y >= height) return; + + const int piwidth = 2*width; + global float *ppi = pi + mad24(y, piwidth, 2*x); + + const int raw_x = ppi[0]; + const int raw_y = ppi[1]; + + if(raw_x < 0 || raw_y < 0 || raw_x >= raw_width || raw_y >= raw_height) return; + + const uint raw_pixel = read_imageui(raw, sampleri, (int2)(raw_x, raw_y)).x; + + const int c = FC(raw_y, raw_x, filters); + + if(raw_pixel < threshold[c]) return; + + float4 pixel = read_imagef(in, sampleri, (int2)(x, y)); + + global float *color = colors + mad24(4, c, 0); + + // cfa color + pixel.x = color[0]; + pixel.y = color[1]; + pixel.z = color[2]; + + write_imagef (out, (int2)(x, y), pixel); +} + kernel void rawoverexposed_mark_solid ( read_only image2d_t in, write_only image2d_t out, global float *pi, diff --git a/src/iop/rawoverexposed.c b/src/iop/rawoverexposed.c index a9b4f5c93674..7e444d139dc8 100644 --- a/src/iop/rawoverexposed.c +++ b/src/iop/rawoverexposed.c @@ -57,6 +57,7 @@ typedef struct dt_iop_rawoverexposed_data_t typedef struct dt_iop_rawoverexposed_global_data_t { + int kernel_rawoverexposed_mark_cfa; int kernel_rawoverexposed_mark_solid; int kernel_rawoverexposed_falsecolor; } dt_iop_rawoverexposed_global_data_t; @@ -225,6 +226,7 @@ int process_cl(struct dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, cl_m float *coordbuf = NULL; cl_mem dev_coord = NULL; cl_mem dev_thresholds = NULL; + cl_mem dev_colors = NULL; cl_int err = -999; @@ -294,9 +296,22 @@ int process_cl(struct dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, cl_m int kernel; switch(dev->rawoverexposed.mode) { + case DT_DEV_RAWOVEREXPOSED_MODE_MARK_CFA: + kernel = gd->kernel_rawoverexposed_mark_cfa; + + dev_colors = dt_opencl_alloc_device_buffer(devid, sizeof(dt_iop_rawoverexposed_colors)); + if(dev_colors == NULL) goto error; + + /* _blocking_ memory transfer: host coordbuf buffer -> opencl dev_colors */ + err = dt_opencl_write_buffer_to_device(devid, (void *)dt_iop_rawoverexposed_colors, dev_colors, 0, + sizeof(dt_iop_rawoverexposed_colors), CL_TRUE); + if(err != CL_SUCCESS) goto error; + + break; case DT_DEV_RAWOVEREXPOSED_MODE_MARK_SOLID: kernel = gd->kernel_rawoverexposed_mark_solid; break; + case DT_DEV_RAWOVEREXPOSED_MODE_FALSECOLOR: default: kernel = gd->kernel_rawoverexposed_falsecolor; break; @@ -317,7 +332,9 @@ int process_cl(struct dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, cl_m dt_opencl_set_kernel_arg(devid, kernel, 8, sizeof(uint32_t), &filters); dt_opencl_set_kernel_arg(devid, kernel, 9, sizeof(cl_mem), &dev_thresholds); - if(dev->rawoverexposed.mode == DT_DEV_RAWOVEREXPOSED_MODE_MARK_SOLID) + if(dev->rawoverexposed.mode == DT_DEV_RAWOVEREXPOSED_MODE_MARK_CFA) + dt_opencl_set_kernel_arg(devid, kernel, 10, sizeof(cl_mem), &dev_colors); + else if(dev->rawoverexposed.mode == DT_DEV_RAWOVEREXPOSED_MODE_MARK_SOLID) dt_opencl_set_kernel_arg(devid, kernel, 10, 4 * sizeof(float), color); err = dt_opencl_enqueue_kernel_2d(devid, kernel, sizes); @@ -325,6 +342,7 @@ int process_cl(struct dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, cl_m return TRUE; error: + if(dev_colors != NULL) dt_opencl_release_mem_object(dev_colors); if(dev_thresholds != NULL) dt_opencl_release_mem_object(dev_thresholds); if(dev_coord != NULL) dt_opencl_release_mem_object(dev_coord); if(coordbuf != NULL) dt_free_align(coordbuf); @@ -349,8 +367,6 @@ void commit_params(dt_iop_module_t *self, dt_iop_params_t *p1, dt_dev_pixelpipe_ if(image->buf_dsc.datatype != TYPE_UINT16 || !image->buf_dsc.filters) piece->enabled = 0; if(image->buf_dsc.filters == 9u) piece->process_cl_ready = 0; - - if(dev->rawoverexposed.mode == DT_DEV_RAWOVEREXPOSED_MODE_MARK_CFA) piece->process_cl_ready = 0; } void init_global(dt_iop_module_so_t *module) @@ -358,6 +374,7 @@ void init_global(dt_iop_module_so_t *module) const int program = 2; // basic.cl from programs.conf module->data = malloc(sizeof(dt_iop_rawoverexposed_global_data_t)); dt_iop_rawoverexposed_global_data_t *gd = module->data; + gd->kernel_rawoverexposed_mark_cfa = dt_opencl_create_kernel(program, "rawoverexposed_mark_cfa"); gd->kernel_rawoverexposed_mark_solid = dt_opencl_create_kernel(program, "rawoverexposed_mark_solid"); gd->kernel_rawoverexposed_falsecolor = dt_opencl_create_kernel(program, "rawoverexposed_falsecolor"); } @@ -368,6 +385,7 @@ void cleanup_global(dt_iop_module_so_t *module) dt_iop_rawoverexposed_global_data_t *gd = (dt_iop_rawoverexposed_global_data_t *)module->data; dt_opencl_free_kernel(gd->kernel_rawoverexposed_falsecolor); dt_opencl_free_kernel(gd->kernel_rawoverexposed_mark_solid); + dt_opencl_free_kernel(gd->kernel_rawoverexposed_mark_cfa); free(module->data); module->data = NULL; } From 73c9faf3b0643244d8876ccd7c75d6941f5a6d1b Mon Sep 17 00:00:00 2001 From: Roman Lebedev Date: Sun, 4 Sep 2016 19:11:13 +0300 Subject: [PATCH 7/7] Rawoverexposed iop: explain why we account for wb. --- src/iop/rawoverexposed.c | 15 +++++++++++++++ 1 file changed, 15 insertions(+) diff --git a/src/iop/rawoverexposed.c b/src/iop/rawoverexposed.c index 7e444d139dc8..8fb09cafae78 100644 --- a/src/iop/rawoverexposed.c +++ b/src/iop/rawoverexposed.c @@ -92,6 +92,21 @@ static void process_common_setup(dt_iop_module_t *self, dt_dev_pixelpipe_iop_t * // "undo" temperature iop if(piece->pipe->dsc.temperature.enabled) threshold /= piece->pipe->dsc.temperature.coeffs[k]; + /* + * yes, technically, sensor clipping needs to be detected not accounting + * for white balance. + * + * but we are not after technical sensor clipping. + * + * pick some image that is overexposed, disable highlight clipping, apply + * negative exposure compensation. you'll see magenta highlight. + * if comment-out that ^ wb division, the module would not mark that + * area with magenta highlights as clipped, because technically + * the channels are not clipped, even though the colour is wrong. + * + * but we do want to see those magenta highlights marked... + */ + // "undo" rawprepare iop threshold *= piece->pipe->dsc.rawprepare.raw_white_point - piece->pipe->dsc.rawprepare.raw_black_level; threshold += piece->pipe->dsc.rawprepare.raw_black_level;