From 75bfb4e664242af7b115006c6c0e1a014313c5c0 Mon Sep 17 00:00:00 2001 From: Philipp Lutz Date: Sat, 29 Aug 2026 09:16:41 +0000 Subject: [PATCH 1/8] toneequal: add OpenCL implementation The tone equalizer was CPU only, so it pulled the pixelpipe back to the host on every zoom, pan and slider move. Add data/kernels/toneequal.cl and a process_cl() mirroring the CPU path: - luminance mask extraction for all seven estimators, - the fast guided filter (quantization, separable box average, linear and geometric-mean blending), - the exposure independent guided filter, reusing the existing dt_gaussian_*_cl() buffer blurs, - the exposure correction and the luminance mask display. The correction is evaluated as the sum of the eight octave gaussians rather than through d->correction_lut: the transcendentals are cheaper on a GPU than uploading a 320 kB table per invocation, and this is the reference formulation the lut approximates. The GUI reads the luminance mask from host memory to build its histogram and to report the luminance under the cursor, so process_cl() mirrors the caching logic of toneeq_process() and copies the mask back whenever the upstream pipe hash changes. Also add a tiling_callback() so the pixelpipe accounts for the intermediate buffers and falls back to the CPU rather than failing to allocate. The module stays untileable, the guided filter being global. Checked against the CPU path with the integration tests 0079, 0080 and 0129 to 0134, plus mask quantization and multi-iteration variants: the 8-bit outputs differ by at most 1 LSB on less than 0.1% of the pixels. --- data/kernels/programs.conf | 1 + data/kernels/toneequal.cl | 549 +++++++++++++++++++++++++++++ src/iop/toneequal.c | 689 ++++++++++++++++++++++++++++++++++++- 3 files changed, 1236 insertions(+), 3 deletions(-) create mode 100644 data/kernels/toneequal.cl diff --git a/data/kernels/programs.conf b/data/kernels/programs.conf index eec397a226c..2279fbf32ef 100644 --- a/data/kernels/programs.conf +++ b/data/kernels/programs.conf @@ -45,3 +45,4 @@ overlay.cl 41 bilinear.cl 42 spektrafilm.cl 43 satcurve.cl 44 +toneequal.cl 45 diff --git a/data/kernels/toneequal.cl b/data/kernels/toneequal.cl new file mode 100644 index 00000000000..833534c3de1 --- /dev/null +++ b/data/kernels/toneequal.cl @@ -0,0 +1,549 @@ +/* + This file is part of darktable, + Copyright (C) 2025 darktable developers. + + 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 . +*/ + +#include "common.h" + +// keep in sync with src/common/luminance_mask.h and src/iop/toneequal.c +#define TONEEQ_MIN_FLOAT 0x1p-16f // exp2f(-16.0f) +#define TONEEQ_MIN_EV (-8.0f) +#define TONEEQ_MAX_EV (0.0f) +#define TONEEQ_PIXEL_CHAN 8 + +// dt_iop_luminance_mask_method_t +#define TONEEQ_MEAN 0 +#define TONEEQ_LIGHTNESS 1 +#define TONEEQ_VALUE 2 +#define TONEEQ_NORM_1 3 +#define TONEEQ_NORM_2 4 +#define TONEEQ_NORM_POWER 5 +#define TONEEQ_GEOMEAN 6 + +// dt_iop_guided_filter_blending_t +#define TONEEQ_BLENDING_LINEAR 0 +#define TONEEQ_BLENDING_GEOMEAN 1 + + +static inline float _linear_contrast(const float pixel, + const float fulcrum, + const float contrast) +{ + // Increase the slope of the value around a fulcrum value + return fmax((pixel - fulcrum) * contrast + fulcrum, TONEEQ_MIN_FLOAT); +} + + +/* Flatten an RGB image into a luminance map, applying exposure and contrast + compensation on the fly. Mirrors luminance_mask() from common/luminance_mask.h */ +kernel void toneequal_luminance_mask(read_only image2d_t in, + global float *const luminance, + const int width, + const int height, + const int method, + const float exposure_boost, + const float fulcrum, + const float contrast_boost) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + if(x >= width || y >= height) return; + + const float4 pixel = read_imagef(in, samplerA, (int2)(x, y)); + float lum; + + switch(method) + { + case TONEEQ_LIGHTNESS: + { + // (max(RGB) + min(RGB)) / 2 is equivalent to HSL lightness + const float max_rgb = fmax(fmax(pixel.x, pixel.y), pixel.z); + const float min_rgb = fmin(fmin(pixel.x, pixel.y), pixel.z); + lum = exposure_boost * (max_rgb + min_rgb) / 2.0f; + break; + } + + case TONEEQ_VALUE: + // max(RGB) is equivalent to HSV value + lum = exposure_boost * fmax(fmax(pixel.x, pixel.y), pixel.z); + break; + + case TONEEQ_NORM_1: + // vector norm L1 + lum = exposure_boost * (fabs(pixel.x) + fabs(pixel.y) + fabs(pixel.z)); + break; + + case TONEEQ_NORM_2: + // vector norm L2 : euclidean norm + lum = exposure_boost + * dtcl_sqrt(pixel.x * pixel.x + pixel.y * pixel.y + pixel.z * pixel.z); + break; + + case TONEEQ_NORM_POWER: + { + // weird norm sort of perceptual. This is black magic really, but it looks good. + const float4 value = fabs(pixel); + const float4 RGB_square = value * value; + const float4 RGB_cubic = RGB_square * value; + const float numerator = RGB_cubic.x + RGB_cubic.y + RGB_cubic.z; + const float denominator = RGB_square.x + RGB_square.y + RGB_square.z; + lum = exposure_boost * numerator / denominator; + break; + } + + case TONEEQ_GEOMEAN: + { + // geometric_mean(RGB) + const float product = fabs(pixel.x) * fabs(pixel.y) * fabs(pixel.z); + lum = exposure_boost * dtcl_pow(product, 1.0f / 3.0f); + break; + } + + case TONEEQ_MEAN: + default: + // mean(RGB) is the intensity + lum = exposure_boost * (pixel.x + pixel.y + pixel.z) / 3.0f; + break; + } + + luminance[mad24(y, width, x)] = _linear_contrast(lum, fulcrum, contrast_boost); +} + + +/* Build the pixel correction as the sum of the contribution of each luminance + channel, weighted by the gaussian of the radial distance in EV. + This is the non-LUT reference implementation of apply_toneequalizer(): on GPU + the transcendentals are cheaper than a 320 kB lookup table upload. */ +static inline float _pixel_correction(const float exposure, + const float4 factors_low, + const float4 factors_high, + const float gauss_denom) +{ + // radial distances of the exposure octaves, see centers_ops[] in toneequal.c + const float4 centers_low = { -56.0f / 7.0f, -48.0f / 7.0f, + -40.0f / 7.0f, -32.0f / 7.0f }; + const float4 centers_high = { -24.0f / 7.0f, -16.0f / 7.0f, + -8.0f / 7.0f, 0.0f / 7.0f }; + + const float4 radius_low = exposure - centers_low; + const float4 radius_high = exposure - centers_high; + + const float4 gauss_low = + dtcl_exp(-radius_low * radius_low / gauss_denom) * factors_low; + const float4 gauss_high = + dtcl_exp(-radius_high * radius_high / gauss_denom) * factors_high; + + const float4 result = gauss_low + gauss_high; + + // the user-set correction is expected in [-2;+2] EV, so is the interpolated one + return clamp(result.x + result.y + result.z + result.w, 0.25f, 4.0f); +} + + +kernel void toneequal_apply(read_only image2d_t in, + global const float *const luminance, + write_only image2d_t out, + const int width, + const int height, + const int in_width, + const int offset_x, + const int offset_y, + const float4 factors_low, + const float4 factors_high, + const float gauss_denom) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + if(x >= width || y >= height) return; + + const int xin = x + offset_x; + const int yin = y + offset_y; + + // The radial-basis interpolation is valid in [-8; 0] EV and can quickly + // diverge outside + const float exposure = clamp(dtcl_log2(luminance[mad24(yin, in_width, xin)]), + TONEEQ_MIN_EV, TONEEQ_MAX_EV); + const float correction = + _pixel_correction(exposure, factors_low, factors_high, gauss_denom); + + const float4 pixel = read_imagef(in, samplerA, (int2)(xin, yin)); + write_imagef(out, (int2)(x, y), correction * pixel); +} + + +kernel void toneequal_display_mask(read_only image2d_t in, + global const float *const luminance, + write_only image2d_t out, + const int width, + const int height, + const int in_width, + const int offset_x, + const int offset_y) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + if(x >= width || y >= height) return; + + const int xin = x + offset_x; + const int yin = y + offset_y; + + // normalize the mask intensity between -8 EV and 0 EV for clarity, + // and add a "gamma" 2.0 for better legibility in shadows + const float lum = luminance[mad24(yin, in_width, xin)]; + const float intensity = + dtcl_sqrt(fmin(fmax(lum - 0.00390625f, 0.0f) / 0.99609375f, 1.0f)); + + // set gray level for the mask, copy alpha channel + const float4 pixel = read_imagef(in, samplerA, (int2)(xin, yin)); + write_imagef(out, (int2)(x, y), + (float4)(intensity, intensity, intensity, pixel.w)); +} + + +/* Quantize in exposure levels evenly spaced in log by sampling. + Mirrors quantize() from common/fast_guided_filter.h */ +kernel void toneequal_quantize(global const float *const in, + global float *const out, + const int width, + const int height, + const float sampling, + const float clip_min, + const float clip_max) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + if(x >= width || y >= height) return; + + const int k = mad24(y, width, x); + const float value = in[k]; + + if(sampling == 0.0f) + out[k] = value; // no-op + else if(sampling == 1.0f) // fast track + out[k] = clamp(dtcl_exp2(floor(dtcl_log2(value))), clip_min, clip_max); + else // slow track + out[k] = clamp(dtcl_exp2(floor(dtcl_log2(value) / sampling) * sampling), + clip_min, clip_max); +} + + +/* Separable box average over a window of size 2*radius + 1, mirroring + dt_box_mean() from common/box_filters.cc. The window is clipped at the + image borders and the average is taken over the samples actually seen, + so input and output buffers must be distinct. */ +#define TONEEQ_BOX_MEAN_LINE(TYPE, LENGTH, STRIDE) \ + TYPE sum = (TYPE)0.0f; \ + int hits = 0; \ + int i; \ + \ + /* add up the left half of the window */ \ + for(i = 0; i < min(radius, LENGTH); i++) \ + { \ + sum += in[(size_t)i * STRIDE]; \ + hits++; \ + } \ + \ + /* up to the point where we start removing values from the average */ \ + for(i = 0; i <= radius && i + radius < LENGTH; i++) \ + { \ + sum += in[(size_t)(i + radius) * STRIDE]; \ + hits++; \ + out[(size_t)i * STRIDE] = sum / (float)hits; \ + } \ + \ + /* if radius > LENGTH/2 we can neither add nor remove values */ \ + for(; i <= radius && i < LENGTH; i++) \ + out[(size_t)i * STRIDE] = sum / (float)hits; \ + \ + /* the bulk of the line */ \ + for(; i + radius < LENGTH; i++) \ + { \ + sum -= in[(size_t)(i - radius - 1) * STRIDE]; \ + sum += in[(size_t)(i + radius) * STRIDE]; \ + out[(size_t)i * STRIDE] = sum / (float)hits; \ + } \ + \ + /* the end of the line, no more values to add to the average */ \ + for(; i < LENGTH; i++) \ + { \ + sum -= in[(size_t)(i - radius - 1) * STRIDE]; \ + hits--; \ + out[(size_t)i * STRIDE] = sum / (float)hits; \ + } + +#define TONEEQ_BOX_MEAN_X(TYPE) \ + const int y = get_global_id(0); \ + if(y >= height) return; \ + \ + global const TYPE *const in = input + (size_t)y * width; \ + global TYPE *const out = output + (size_t)y * width; \ + TONEEQ_BOX_MEAN_LINE(TYPE, width, 1) + +#define TONEEQ_BOX_MEAN_Y(TYPE) \ + const int x = get_global_id(0); \ + if(x >= width) return; \ + \ + global const TYPE *const in = input + x; \ + global TYPE *const out = output + x; \ + TONEEQ_BOX_MEAN_LINE(TYPE, height, width) + +kernel void toneequal_box_mean_x_2c(global const float2 *const input, + global float2 *const output, + const int width, + const int height, + const int radius) +{ + TONEEQ_BOX_MEAN_X(float2) +} + +kernel void toneequal_box_mean_y_2c(global const float2 *const input, + global float2 *const output, + const int width, + const int height, + const int radius) +{ + TONEEQ_BOX_MEAN_Y(float2) +} + +kernel void toneequal_box_mean_x_4c(global const float4 *const input, + global float4 *const output, + const int width, + const int height, + const int radius) +{ + TONEEQ_BOX_MEAN_X(float4) +} + +kernel void toneequal_box_mean_y_4c(global const float4 *const input, + global float4 *const output, + const int width, + const int height, + const int radius) +{ + TONEEQ_BOX_MEAN_Y(float4) +} + + +/* Guided filter: pack guide and mask into an array of 4×1 vectors so that the + box average of all four terms can be done in a single pass. + Mirrors the first loop of variance_analyse() in common/fast_guided_filter.h */ +kernel void toneequal_gf_pack(global const float *const guide, + global const float *const mask, + global float4 *const packed, + const int width, + const int height) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + if(x >= width || y >= height) return; + + const int k = mad24(y, width, x); + const float g = guide[k]; + const float m = mask[k]; + + packed[k] = (float4)(g, m, g * g, g * m); +} + + +/* Guided filter: turn the averaged terms into the linear blending parameters + a and b such that mask = a * I + b */ +kernel void toneequal_gf_ab(global const float4 *const packed, + global float2 *const ab, + const int width, + const int height, + const float feathering) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + if(x >= width || y >= height) return; + + const int k = mad24(y, width, x); + const float4 avg = packed[k]; + + // avoid division by 0 + const float d = fmax((avg.z - avg.x * avg.x) + feathering, 1e-15f); + const float a = (avg.w - avg.x * avg.y) / d; + const float b = avg.y - a * avg.x; + + ab[k] = (float2)(a, b); +} + + +/* Guided filter: blend the guided image with the a and b parameters. + Mirrors apply_linear_blending[_w_geomean]() in common/fast_guided_filter.h */ +kernel void toneequal_gf_blend(global float *const image, + global const float2 *const ab, + const int width, + const int height, + const int filter) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + if(x >= width || y >= height) return; + + const int k = mad24(y, width, x); + const float pixel = image[k]; + const float2 blend = ab[k]; + + // Note : image[k] is positive at the outside of the luminance mask + const float blended = fmax(pixel * blend.x + blend.y, TONEEQ_MIN_FLOAT); + + image[k] = (filter == TONEEQ_BLENDING_GEOMEAN) + ? dtcl_sqrt(pixel * blended) + : blended; +} + + +/* Exposure independent guided filter: pack guide and mask so that a single + gaussian blur yields the averages of guide, guide², mask and mask × guide. + Mirrors eigf_variance_analysis() in common/eigf.h */ +kernel void toneequal_eigf_pack_4c(global const float *const guide, + global const float *const mask, + global float4 *const packed, + const int width, + const int height) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + if(x >= width || y >= height) return; + + const int k = mad24(y, width, x); + const float g = guide[k]; + const float m = mask[k]; + + packed[k] = (float4)(g, g * g, m, m * g); +} + + +kernel void toneequal_eigf_finish_4c(global float4 *const av, + const int width, + const int height) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + if(x >= width || y >= height) return; + + const int k = mad24(y, width, x); + const float4 avg = av[k]; + + // turn the averages of the squares into variance and covariance + av[k] = (float4)(avg.x, + avg.y - avg.x * avg.x, + avg.z, + avg.w - avg.x * avg.z); +} + + +/* same as above, but specialized for the case where guide == mask */ +kernel void toneequal_eigf_pack_2c(global const float *const guide, + global float2 *const packed, + const int width, + const int height) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + if(x >= width || y >= height) return; + + const int k = mad24(y, width, x); + const float g = guide[k]; + + packed[k] = (float2)(g, g * g); +} + + +kernel void toneequal_eigf_finish_2c(global float2 *const av, + const int width, + const int height) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + if(x >= width || y >= height) return; + + const int k = mad24(y, width, x); + const float2 avg = av[k]; + + av[k] = (float2)(avg.x, avg.y - avg.x * avg.x); +} + + +/* Exposure independent guided filter: blending step. + Mirrors eigf_blending() in common/eigf.h */ +kernel void toneequal_eigf_blend(global float *const image, + global const float *const mask, + global const float4 *const av, + const int width, + const int height, + const int filter, + const float feathering) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + if(x >= width || y >= height) return; + + const int k = mad24(y, width, x); + const float pixel = image[k]; + const float4 avg = av[k]; + + const float avg_g = avg.x; + const float var_g = avg.y; + const float avg_m = avg.z; + const float covar_mg = avg.w; + + const float norm_g = fmax(avg_g * pixel, 1E-6f); + const float norm_m = fmax(avg_m * mask[k], 1E-6f); + const float normalized_var_guide = var_g / norm_g; + const float normalized_covar = covar_mg / dtcl_sqrt(norm_g * norm_m); + const float a = normalized_covar / (normalized_var_guide + feathering); + const float b = avg_m - a * avg_g; + + const float blended = fmax(pixel * a + b, TONEEQ_MIN_FLOAT); + + image[k] = (filter == TONEEQ_BLENDING_GEOMEAN) + ? dtcl_sqrt(pixel * blended) + : blended; +} + + +/* same as above, but specialized for the case where guide == mask */ +kernel void toneequal_eigf_blend_no_mask(global float *const image, + global const float2 *const av, + const int width, + const int height, + const int filter, + const float feathering) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + if(x >= width || y >= height) return; + + const int k = mad24(y, width, x); + const float pixel = image[k]; + const float2 avg = av[k]; + + const float avg_g = avg.x; + const float var_g = avg.y; + + const float norm_g = fmax(avg_g * pixel, 1E-6f); + const float normalized_var_guide = var_g / norm_g; + const float a = normalized_var_guide / (normalized_var_guide + feathering); + const float b = avg_g - a * avg_g; + + const float blended = fmax(pixel * a + b, TONEEQ_MIN_FLOAT); + + image[k] = (filter == TONEEQ_BLENDING_GEOMEAN) + ? dtcl_sqrt(pixel * blended) + : blended; +} diff --git a/src/iop/toneequal.c b/src/iop/toneequal.c index 0f9081b7bb1..42ccfaeb8cb 100644 --- a/src/iop/toneequal.c +++ b/src/iop/toneequal.c @@ -95,6 +95,7 @@ #include "bauhaus/bauhaus.h" #include "common/darktable.h" +#include "common/bilinear.h" #include "common/fast_guided_filter.h" #include "common/eigf.h" #include "common/interpolation.h" @@ -109,6 +110,7 @@ #include "develop/imageop_math.h" #include "develop/imageop_gui.h" #include "develop/preview_data.h" +#include "develop/tiling.h" #include "dtgtk/drawingarea.h" #include "dtgtk/expander.h" #include "gui/accelerators.h" @@ -206,7 +208,23 @@ typedef struct dt_iop_toneequalizer_data_t typedef struct dt_iop_toneequalizer_global_data_t { - // TODO: put OpenCL kernels here at some point + int kernel_toneequal_luminance_mask; + int kernel_toneequal_apply; + int kernel_toneequal_display_mask; + int kernel_toneequal_quantize; + int kernel_toneequal_box_mean_x_2c; + int kernel_toneequal_box_mean_y_2c; + int kernel_toneequal_box_mean_x_4c; + int kernel_toneequal_box_mean_y_4c; + int kernel_toneequal_gf_pack; + int kernel_toneequal_gf_ab; + int kernel_toneequal_gf_blend; + int kernel_toneequal_eigf_pack_4c; + int kernel_toneequal_eigf_finish_4c; + int kernel_toneequal_eigf_pack_2c; + int kernel_toneequal_eigf_finish_2c; + int kernel_toneequal_eigf_blend; + int kernel_toneequal_eigf_blend_no_mask; } dt_iop_toneequalizer_global_data_t; @@ -813,8 +831,9 @@ static inline void apply_toneequalizer(const float *const restrict in, #else -// we keep this version for further reference (e.g. for implementing -// a gpu version) +// we keep this version for further reference; it is the one the OpenCL +// path implements, as evaluating the gaussians is cheaper on a GPU than +// uploading the correction lut __DT_CLONE_TARGETS__ static inline void apply_toneequalizer(const float *const restrict in, const float *const restrict luminance, @@ -1180,6 +1199,613 @@ void process(dt_iop_module_t *self, } +#ifdef HAVE_OPENCL + +/*** + * OpenCL implementation + * + * The GPU path mirrors the CPU code above: extract the luminance mask from the + * input, optionally refine it with the (exposure independent) guided filter, + * then correct each pixel exposure from its masked luminance. + * + * Every intermediate buffer is a grey image, so we use device buffers rather + * than images and keep the packing conventions of the CPU code (arrays of 2 or + * 4 channel structs) to average several terms in a single pass. + ***/ + +static cl_int _box_mean_cl(const int devid, + const dt_iop_toneequalizer_global_data_t *const gd, + cl_mem dev_buf, + cl_mem dev_tmp, + const int width, + const int height, + const int ch, + const int radius) +{ + // The moving average of the separable box filter reads samples that a + // parallel in-place pass would already have overwritten, so we bounce + // through dev_tmp and land back in dev_buf. + const int kernel_x = (ch == 4) + ? gd->kernel_toneequal_box_mean_x_4c + : gd->kernel_toneequal_box_mean_x_2c; + const int kernel_y = (ch == 4) + ? gd->kernel_toneequal_box_mean_y_4c + : gd->kernel_toneequal_box_mean_y_2c; + + const cl_int err = dt_opencl_enqueue_kernel_1d_args(devid, kernel_x, height, + CLARG(dev_buf), CLARG(dev_tmp), + CLARG(width), CLARG(height), CLARG(radius)); + if(err != CL_SUCCESS) return err; + + return dt_opencl_enqueue_kernel_1d_args(devid, kernel_y, width, + CLARG(dev_tmp), CLARG(dev_buf), + CLARG(width), CLARG(height), CLARG(radius)); +} + + +static cl_int _quantize_cl(const int devid, + const dt_iop_toneequalizer_global_data_t *const gd, + cl_mem dev_in, + cl_mem dev_out, + const int width, + const int height, + const float sampling, + const float clip_min, + const float clip_max) +{ + return dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_toneequal_quantize, + width, height, + CLARG(dev_in), CLARG(dev_out), CLARG(width), CLARG(height), + CLARG(sampling), CLARG(clip_min), CLARG(clip_max)); +} + + +static cl_int _fast_surface_blur_cl(const int devid, + const dt_iop_toneequalizer_global_data_t *const gd, + cl_mem dev_image, + const int width, + const int height, + const int radius, + const float feathering, + const int iterations, + const dt_iop_guided_filter_blending_t filter, + const float quantization, + const float quantize_min, + const float quantize_max) +{ + // Works in-place on a grey image, see fast_surface_blur() + // in common/fast_guided_filter.h + + // A down-scaling of 4 seems empirically safe and consistent no matter the + // image zoom level + const float scaling = 4.0f; + const int ds_radius = (radius < 4) ? 1 : (int)((float)radius / scaling); + + // the downscaled dimensions can round down to zero on thumbnails + const int ds_width = MAX(1, (int)((float)width / scaling)); + const int ds_height = MAX(1, (int)((float)height / scaling)); + + const size_t bsize = (size_t)width * height * sizeof(float); + const size_t ds_bsize = (size_t)ds_width * ds_height * sizeof(float); + + cl_int err = CL_MEM_OBJECT_ALLOCATION_FAILURE; + + cl_mem dev_ab = dt_opencl_alloc_device_buffer(devid, 2 * bsize); + cl_mem dev_ds_image = dt_opencl_alloc_device_buffer(devid, ds_bsize); + cl_mem dev_ds_mask = dt_opencl_alloc_device_buffer(devid, ds_bsize); + cl_mem dev_ds_ab = dt_opencl_alloc_device_buffer(devid, 2 * ds_bsize); + // array of struct : { { guide, mask, guide * guide, guide * mask } } + cl_mem dev_packed = dt_opencl_alloc_device_buffer(devid, 4 * ds_bsize); + // scratch space of the box average, large enough for both channel counts + cl_mem dev_tmp = dt_opencl_alloc_device_buffer(devid, 4 * ds_bsize); + + if(!dev_ab || !dev_ds_image || !dev_ds_mask || !dev_ds_ab || !dev_packed || !dev_tmp) + goto error; + + // Downsample the image for speed-up + err = dt_interpolate_bilinear_cl(devid, dev_image, width, height, + dev_ds_image, ds_width, ds_height, 1); + if(err != CL_SUCCESS) goto error; + + // Iterations of filter models the diffusion, sort of + for(int i = 0; i < iterations; ++i) + { + // (Re)build the mask from the quantized image to help guiding + err = _quantize_cl(devid, gd, dev_ds_image, dev_ds_mask, ds_width, ds_height, + quantization, quantize_min, quantize_max); + if(err != CL_SUCCESS) goto error; + + // Perform the patch-wise variance analyse to get the a and b parameters + // for the linear blending s.t. mask = a * I + b + err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_toneequal_gf_pack, + ds_width, ds_height, + CLARG(dev_ds_mask), CLARG(dev_ds_image), CLARG(dev_packed), + CLARG(ds_width), CLARG(ds_height)); + if(err != CL_SUCCESS) goto error; + + err = _box_mean_cl(devid, gd, dev_packed, dev_tmp, + ds_width, ds_height, 4, ds_radius); + if(err != CL_SUCCESS) goto error; + + err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_toneequal_gf_ab, + ds_width, ds_height, + CLARG(dev_packed), CLARG(dev_ds_ab), + CLARG(ds_width), CLARG(ds_height), CLARG(feathering)); + if(err != CL_SUCCESS) goto error; + + // Compute the patch-wise average of parameters a and b + err = _box_mean_cl(devid, gd, dev_ds_ab, dev_tmp, + ds_width, ds_height, 2, ds_radius); + if(err != CL_SUCCESS) goto error; + + if(i != iterations - 1) + { + // Process the intermediate filtered image + const int blending = DT_GF_BLENDING_LINEAR; + err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_toneequal_gf_blend, + ds_width, ds_height, + CLARG(dev_ds_image), CLARG(dev_ds_ab), + CLARG(ds_width), CLARG(ds_height), CLARG(blending)); + if(err != CL_SUCCESS) goto error; + } + } + + // Upsample the blending parameters a and b + err = dt_interpolate_bilinear_cl(devid, dev_ds_ab, ds_width, ds_height, + dev_ab, width, height, 2); + if(err != CL_SUCCESS) goto error; + + // Finally, blend the guided image + { + const int blending = filter; + err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_toneequal_gf_blend, + width, height, + CLARG(dev_image), CLARG(dev_ab), + CLARG(width), CLARG(height), CLARG(blending)); + } + +error: + dt_opencl_release_mem_object(dev_tmp); + dt_opencl_release_mem_object(dev_packed); + dt_opencl_release_mem_object(dev_ds_ab); + dt_opencl_release_mem_object(dev_ds_mask); + dt_opencl_release_mem_object(dev_ds_image); + dt_opencl_release_mem_object(dev_ab); + return err; +} + + +static cl_int _fast_eigf_surface_blur_cl(const int devid, + const dt_iop_toneequalizer_global_data_t *const gd, + cl_mem dev_image, + const int width, + const int height, + const float sigma, + const float feathering, + const int iterations, + const dt_iop_guided_filter_blending_t filter, + const float quantization, + const float quantize_min, + const float quantize_max) +{ + // Works in-place on a grey image, see fast_eigf_surface_blur() + // in common/eigf.h + const float scaling = fmaxf(fminf(sigma, 4.0f), 1.0f); + const float ds_sigma = fmaxf(sigma / scaling, 1.0f); + + // the downscaled dimensions can round down to zero on thumbnails + const int ds_width = MAX(1, (int)((float)width / scaling)); + const int ds_height = MAX(1, (int)((float)height / scaling)); + + const size_t bsize = (size_t)width * height * sizeof(float); + const size_t ds_bsize = (size_t)ds_width * ds_height * sizeof(float); + + // without quantization the guide is the image itself, which halves the + // number of terms to blur + const gboolean use_mask = (quantization != 0.0f); + const int ch = use_mask ? 4 : 2; + + cl_int err = CL_MEM_OBJECT_ALLOCATION_FAILURE; + + cl_mem dev_mask = use_mask + ? dt_opencl_alloc_device_buffer(devid, bsize) + : NULL; + cl_mem dev_ds_mask = use_mask + ? dt_opencl_alloc_device_buffer(devid, ds_bsize) + : NULL; + cl_mem dev_ds_image = dt_opencl_alloc_device_buffer(devid, ds_bsize); + // average - variance arrays: store the guide and mask averages and variances + cl_mem dev_ds_av = dt_opencl_alloc_device_buffer(devid, ch * ds_bsize); + cl_mem dev_av = dt_opencl_alloc_device_buffer(devid, ch * bsize); + + if(!dev_ds_image || !dev_ds_av || !dev_av + || (use_mask && (!dev_mask || !dev_ds_mask))) + goto error; + + // Iterations of filter models the diffusion, sort of + for(int i = 0; i < iterations; i++) + { + // blend linear for all intermediate images, use filter for last iteration + const int blending = (i == iterations - 1) ? (int)filter : DT_GF_BLENDING_LINEAR; + + err = dt_interpolate_bilinear_cl(devid, dev_image, width, height, + dev_ds_image, ds_width, ds_height, 1); + if(err != CL_SUCCESS) goto error; + + if(use_mask) + { + // (Re)build the mask from the quantized image to help guiding + err = _quantize_cl(devid, gd, dev_image, dev_mask, width, height, + quantization, quantize_min, quantize_max); + if(err != CL_SUCCESS) goto error; + + // Downsample the mask for speed-up + err = dt_interpolate_bilinear_cl(devid, dev_mask, width, height, + dev_ds_mask, ds_width, ds_height, 1); + if(err != CL_SUCCESS) goto error; + + err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_toneequal_eigf_pack_4c, + ds_width, ds_height, + CLARG(dev_ds_mask), CLARG(dev_ds_image), CLARG(dev_ds_av), + CLARG(ds_width), CLARG(ds_height)); + if(err != CL_SUCCESS) goto error; + + err = dt_gaussian_mean_blur_cl(devid, dev_ds_av, + ds_width, ds_height, 4, ds_sigma); + if(err != CL_SUCCESS) goto error; + + err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_toneequal_eigf_finish_4c, + ds_width, ds_height, + CLARG(dev_ds_av), CLARG(ds_width), CLARG(ds_height)); + if(err != CL_SUCCESS) goto error; + + // Upsample the variances and averages + err = dt_interpolate_bilinear_cl(devid, dev_ds_av, ds_width, ds_height, + dev_av, width, height, 4); + if(err != CL_SUCCESS) goto error; + + // Blend the guided image + err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_toneequal_eigf_blend, + width, height, + CLARG(dev_image), CLARG(dev_mask), CLARG(dev_av), + CLARG(width), CLARG(height), CLARG(blending), CLARG(feathering)); + if(err != CL_SUCCESS) goto error; + } + else + { + // no need to build a mask + err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_toneequal_eigf_pack_2c, + ds_width, ds_height, + CLARG(dev_ds_image), CLARG(dev_ds_av), + CLARG(ds_width), CLARG(ds_height)); + if(err != CL_SUCCESS) goto error; + + err = dt_gaussian_mean_blur_cl(devid, dev_ds_av, + ds_width, ds_height, 2, ds_sigma); + if(err != CL_SUCCESS) goto error; + + err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_toneequal_eigf_finish_2c, + ds_width, ds_height, + CLARG(dev_ds_av), CLARG(ds_width), CLARG(ds_height)); + if(err != CL_SUCCESS) goto error; + + // Upsample the variances and averages + err = dt_interpolate_bilinear_cl(devid, dev_ds_av, ds_width, ds_height, + dev_av, width, height, 2); + if(err != CL_SUCCESS) goto error; + + // Blend the guided image + err = dt_opencl_enqueue_kernel_2d_args + (devid, gd->kernel_toneequal_eigf_blend_no_mask, width, height, + CLARG(dev_image), CLARG(dev_av), + CLARG(width), CLARG(height), CLARG(blending), CLARG(feathering)); + if(err != CL_SUCCESS) goto error; + } + } + +error: + dt_opencl_release_mem_object(dev_av); + dt_opencl_release_mem_object(dev_ds_av); + dt_opencl_release_mem_object(dev_ds_image); + dt_opencl_release_mem_object(dev_ds_mask); + dt_opencl_release_mem_object(dev_mask); + return err; +} + + +static cl_int _compute_luminance_mask_cl(const int devid, + const dt_iop_toneequalizer_global_data_t *const gd, + cl_mem dev_in, + cl_mem dev_luminance, + const int width, + const int height, + const dt_iop_toneequalizer_data_t *const d) +{ + // Contrast boosting is done around the average luminance of the mask for the + // plain filters only, see compute_luminance_mask() above + const gboolean boost_contrast = (d->details == DT_TONEEQ_GUIDED) + || (d->details == DT_TONEEQ_EIGF); + const int method = d->method; + const float exposure_boost = d->exposure_boost; + const float fulcrum = boost_contrast ? CONTRAST_FULCRUM : 0.0f; + const float contrast_boost = boost_contrast ? d->contrast_boost : 1.0f; + + const cl_int err = + dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_toneequal_luminance_mask, + width, height, + CLARG(dev_in), CLARG(dev_luminance), CLARG(width), CLARG(height), + CLARG(method), CLARG(exposure_boost), + CLARG(fulcrum), CLARG(contrast_boost)); + if(err != CL_SUCCESS) return err; + + switch(d->details) + { + case DT_TONEEQ_AVG_GUIDED: + return _fast_surface_blur_cl(devid, gd, dev_luminance, width, height, + d->radius, d->feathering, d->iterations, + DT_GF_BLENDING_GEOMEAN, d->quantization, + exp2f(-14.0f), 4.0f); + + case DT_TONEEQ_GUIDED: + return _fast_surface_blur_cl(devid, gd, dev_luminance, width, height, + d->radius, d->feathering, d->iterations, + DT_GF_BLENDING_LINEAR, d->quantization, + exp2f(-14.0f), 4.0f); + + case DT_TONEEQ_AVG_EIGF: + return _fast_eigf_surface_blur_cl(devid, gd, dev_luminance, width, height, + d->radius, d->feathering, d->iterations, + DT_GF_BLENDING_GEOMEAN, d->quantization, + exp2f(-14.0f), 4.0f); + + case DT_TONEEQ_EIGF: + return _fast_eigf_surface_blur_cl(devid, gd, dev_luminance, width, height, + d->radius, d->feathering, d->iterations, + DT_GF_BLENDING_LINEAR, d->quantization, + exp2f(-14.0f), 4.0f); + + case DT_TONEEQ_NONE: + default: + return CL_SUCCESS; + } +} + + +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) +{ + const dt_iop_toneequalizer_data_t *const d = piece->data; + const dt_iop_toneequalizer_global_data_t *const gd = self->global_data; + dt_iop_toneequalizer_gui_data_t *const g = self->gui_data; + + const int devid = piece->pipe->devid; + const int width = roi_in->width; + const int height = roi_in->height; + const size_t num_elem = (size_t)width * height; + + // Get the hash of the upstream pipe to track changes + const dt_hash_t hash = dt_dev_pixelpipe_piece_hash(piece, roi_out, TRUE); + + // Sanity checks + if(width < 1 || height < 1) return DT_OPENCL_PROCESS_CL; + if(roi_in->width < roi_out->width || roi_in->height < roi_out->height) + return DT_OPENCL_PROCESS_CL; // input should be at least as large as output + if(piece->colors != 4) return DT_OPENCL_PROCESS_CL; // we need RGB signal + + // The GUI reads the luminance mask from host memory, so we mirror the + // caching logic of toneeq_process() and copy the mask back whenever the + // upstream pipe state has changed. + gboolean cached = FALSE; + float *luminance = NULL; + + if(self->dev->gui_attached && g) + { + // If the module instance has changed order in the pipe, invalidate the caches + if(g->pipe_order != piece->module->iop_order) + { + dt_iop_gui_enter_critical_section(self); + g->ui_preview_hash = DT_INVALID_HASH; + g->pipe_order = piece->module->iop_order; + g->luminance_valid = FALSE; + g->histogram_valid = FALSE; + dt_iop_gui_leave_critical_section(self); + dt_preview_data_invalidate(&g->pd); + } + + if(dt_pipe_is_full(piece->pipe)) + { + // Re-allocate a new buffer if the full preview size has changed + if(g->full_preview_buf_width != width || g->full_preview_buf_height != height) + { + dt_free_align(g->full_preview_buf); + g->full_preview_buf = dt_alloc_align_float(num_elem); + g->full_preview_buf_width = width; + g->full_preview_buf_height = height; + } + + luminance = g->full_preview_buf; + cached = TRUE; + } + else if(dt_pipe_is_preview(piece->pipe)) + { + // The shared under-cursor service owns the buffer and its locks. + // The resize and the luminance_valid invalidation happen under one + // GUI lock so the GUI never reads a resized, not-yet-recomputed buffer. + luminance = dt_preview_data_resize(&g->pd, width, height, + _toneeq_preview_resized, self); + cached = TRUE; + } + + if(cached && !luminance) + { + dt_control_log(_("tone equalizer failed to allocate memory, check your RAM settings")); + return DT_OPENCL_PROCESS_CL; + } + } + + cl_int err = CL_MEM_OBJECT_ALLOCATION_FAILURE; + + cl_mem dev_luminance = dt_opencl_alloc_device_buffer(devid, num_elem * sizeof(float)); + if(!dev_luminance) goto error; + + // Compute the luminance mask + err = _compute_luminance_mask_cl(devid, gd, dev_in, dev_luminance, + width, height, d); + if(err != CL_SUCCESS) goto error; + + // Keep the host-side cache in sync so that the GUI can compute the histogram + // and read the luminance under the cursor + if(cached) + { + const size_t bsize = num_elem * sizeof(float); + + if(dt_pipe_is_full(piece->pipe)) + { + dt_hash_t saved_hash; + hash_set_get(&g->ui_preview_hash, &saved_hash, &self->gui_lock); + + dt_iop_gui_enter_critical_section(self); + const gboolean luminance_valid = g->luminance_valid; + dt_iop_gui_leave_critical_section(self); + + if(hash != saved_hash || !luminance_valid) + { + /* copy back only if upstream pipe state has changed */ + err = dt_opencl_read_buffer_from_device(devid, luminance, dev_luminance, + 0, bsize, TRUE); + if(err != CL_SUCCESS) goto error; + hash_set_get(&hash, &g->ui_preview_hash, &self->gui_lock); + } + } + else if(dt_pipe_is_preview(piece->pipe)) + { + const dt_hash_t saved_hash = dt_preview_data_get_hash(&g->pd); + + dt_iop_gui_enter_critical_section(self); + const gboolean luminance_valid = g->luminance_valid; + dt_iop_gui_leave_critical_section(self); + + if(saved_hash != hash || !luminance_valid) + { + /* copy back only if upstream pipe state has changed */ + // Flag the cache as being recomputed so the GUI threads never + // read a partially filled buffer, then commit hash + validity + // once the data is ready. + dt_iop_gui_enter_critical_section(self); + g->histogram_valid = FALSE; + g->luminance_valid = FALSE; + dt_iop_gui_leave_critical_section(self); + + err = dt_opencl_read_buffer_from_device(devid, luminance, dev_luminance, + 0, bsize, TRUE); + if(err != CL_SUCCESS) goto error; + dt_preview_data_set_hash(&g->pd, piece); + + dt_iop_gui_enter_critical_section(self); + g->luminance_valid = TRUE; + dt_iop_gui_leave_critical_section(self); + dt_dev_pixelpipe_cache_invalidate_later(piece->pipe, self->iop_order, "toneequal: "); + } + } + } + + // The output dimensions need to be smaller or equal to the input ones + const int out_width = MIN(width, roi_out->width); + const int out_height = MIN(height, roi_out->height); + const int offset_x = (roi_in->x < roi_out->x) ? roi_out->x - roi_in->x : 0; + const int offset_y = (roi_in->y < roi_out->y) ? roi_out->y - roi_in->y : 0; + + // Display output + if(self->dev->gui_attached && g && dt_pipe_is_full(piece->pipe) && g->mask_display) + { + err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_toneequal_display_mask, + out_width, out_height, + CLARG(dev_in), CLARG(dev_luminance), CLARG(dev_out), + CLARG(out_width), CLARG(out_height), CLARG(width), + CLARG(offset_x), CLARG(offset_y)); + if(err != CL_SUCCESS) goto error; + + piece->pipe->mask_display = DT_DEV_PIXELPIPE_DISPLAY_PASSTHRU; + } + else + { + // the correction is interpolated by a series of gaussians, which is + // cheaper on a GPU than uploading the correction LUT used by the CPU code + const float gauss_denom = gaussian_denom(d->smoothing); + + // the 8 octave factors are passed as two float4, so keep the halves in + // their own pointers : CLFLARRAY() casts before it adds an offset + const float *const restrict factors_low = d->factors; + const float *const restrict factors_high = d->factors + PIXEL_CHAN / 2; + + err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_toneequal_apply, + out_width, out_height, + CLARG(dev_in), CLARG(dev_luminance), CLARG(dev_out), + CLARG(out_width), CLARG(out_height), CLARG(width), + CLARG(offset_x), CLARG(offset_y), + CLFLARRAY(4, factors_low), CLFLARRAY(4, factors_high), + CLARG(gauss_denom)); + } + +error: + dt_opencl_release_mem_object(dev_luminance); + return err; +} + +#endif // HAVE_OPENCL + + +void tiling_callback(dt_iop_module_t *self, + dt_dev_pixelpipe_iop_t *piece, + const dt_iop_roi_t *roi_in, + const dt_iop_roi_t *roi_out, + dt_develop_tiling_t *tiling) +{ + const dt_iop_toneequalizer_data_t *const d = piece->data; + + tiling->maxbuf = 1.0f; + tiling->maxbuf_cl = 1.0f; + tiling->overhead = 0; + tiling->overlap = 0; + tiling->align = 1; + + // in and out buffers plus the full size luminance mask, expressed as a + // multiple of a 4 channel image buffer + float factor = 2.25f; + + switch(d->details) + { + case DT_TONEEQ_AVG_GUIDED: + case DT_TONEEQ_GUIDED: + // full size a and b parameters plus the sixteenth-sized working set + factor += 0.5f + 0.25f; + break; + + case DT_TONEEQ_AVG_EIGF: + case DT_TONEEQ_EIGF: + { + // the downscaling factor of the EIGF depends on the filter radius + const float scaling = fmaxf(fminf((float)d->radius, 4.0f), 1.0f); + const float ds = 1.0f / (scaling * scaling); + factor += (d->quantization != 0.0f) + ? 1.25f + 3.5f * ds // mask and averages, plus the gaussian buffers + : 0.5f + 1.75f * ds; + break; + } + + case DT_TONEEQ_NONE: + default: + break; + } + + tiling->factor = factor; + tiling->factor_cl = factor; +} + void modify_roi_in(dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, const dt_iop_roi_t *roi_out, @@ -1566,14 +2192,71 @@ static inline gboolean update_curve_lut(dt_iop_module_t *self) void init_global(dt_iop_module_so_t *self) { + const int program = 45; // toneequal.cl, from programs.conf + dt_iop_toneequalizer_global_data_t *gd = malloc(sizeof(dt_iop_toneequalizer_global_data_t)); self->data = gd; + + gd->kernel_toneequal_luminance_mask = + dt_opencl_create_kernel(program, "toneequal_luminance_mask"); + gd->kernel_toneequal_apply = + dt_opencl_create_kernel(program, "toneequal_apply"); + gd->kernel_toneequal_display_mask = + dt_opencl_create_kernel(program, "toneequal_display_mask"); + gd->kernel_toneequal_quantize = + dt_opencl_create_kernel(program, "toneequal_quantize"); + gd->kernel_toneequal_box_mean_x_2c = + dt_opencl_create_kernel(program, "toneequal_box_mean_x_2c"); + gd->kernel_toneequal_box_mean_y_2c = + dt_opencl_create_kernel(program, "toneequal_box_mean_y_2c"); + gd->kernel_toneequal_box_mean_x_4c = + dt_opencl_create_kernel(program, "toneequal_box_mean_x_4c"); + gd->kernel_toneequal_box_mean_y_4c = + dt_opencl_create_kernel(program, "toneequal_box_mean_y_4c"); + gd->kernel_toneequal_gf_pack = + dt_opencl_create_kernel(program, "toneequal_gf_pack"); + gd->kernel_toneequal_gf_ab = + dt_opencl_create_kernel(program, "toneequal_gf_ab"); + gd->kernel_toneequal_gf_blend = + dt_opencl_create_kernel(program, "toneequal_gf_blend"); + gd->kernel_toneequal_eigf_pack_4c = + dt_opencl_create_kernel(program, "toneequal_eigf_pack_4c"); + gd->kernel_toneequal_eigf_finish_4c = + dt_opencl_create_kernel(program, "toneequal_eigf_finish_4c"); + gd->kernel_toneequal_eigf_pack_2c = + dt_opencl_create_kernel(program, "toneequal_eigf_pack_2c"); + gd->kernel_toneequal_eigf_finish_2c = + dt_opencl_create_kernel(program, "toneequal_eigf_finish_2c"); + gd->kernel_toneequal_eigf_blend = + dt_opencl_create_kernel(program, "toneequal_eigf_blend"); + gd->kernel_toneequal_eigf_blend_no_mask = + dt_opencl_create_kernel(program, "toneequal_eigf_blend_no_mask"); } void cleanup_global(dt_iop_module_so_t *self) { + const dt_iop_toneequalizer_global_data_t *gd = self->data; + + dt_opencl_free_kernel(gd->kernel_toneequal_luminance_mask); + dt_opencl_free_kernel(gd->kernel_toneequal_apply); + dt_opencl_free_kernel(gd->kernel_toneequal_display_mask); + dt_opencl_free_kernel(gd->kernel_toneequal_quantize); + dt_opencl_free_kernel(gd->kernel_toneequal_box_mean_x_2c); + dt_opencl_free_kernel(gd->kernel_toneequal_box_mean_y_2c); + dt_opencl_free_kernel(gd->kernel_toneequal_box_mean_x_4c); + dt_opencl_free_kernel(gd->kernel_toneequal_box_mean_y_4c); + dt_opencl_free_kernel(gd->kernel_toneequal_gf_pack); + dt_opencl_free_kernel(gd->kernel_toneequal_gf_ab); + dt_opencl_free_kernel(gd->kernel_toneequal_gf_blend); + dt_opencl_free_kernel(gd->kernel_toneequal_eigf_pack_4c); + dt_opencl_free_kernel(gd->kernel_toneequal_eigf_finish_4c); + dt_opencl_free_kernel(gd->kernel_toneequal_eigf_pack_2c); + dt_opencl_free_kernel(gd->kernel_toneequal_eigf_finish_2c); + dt_opencl_free_kernel(gd->kernel_toneequal_eigf_blend); + dt_opencl_free_kernel(gd->kernel_toneequal_eigf_blend_no_mask); + free(self->data); self->data = NULL; } From cacba783b9ac3fe9c734615c3678344ef1cb0524 Mon Sep 17 00:00:00 2001 From: Philipp Lutz Date: Sat, 29 Aug 2026 14:11:01 +0000 Subject: [PATCH 2/8] toneequal: keep the full pipe luminance mask on the device Giving tone equalizer an OpenCL path moves its luminance mask into device memory. That mask is not a private intermediate: the CPU code publishes it to the GUI, which reads it from another thread. The GUI cannot read device memory, so a device to host copy has to appear somewhere, and such a copy is a synchronisation point in an otherwise asynchronous command queue. This seam, interleaving the pixelpipe and GUI threads with a GPU queue now in between, is the long standing objection to giving this module a GPU code path at all, so it is worth keeping as narrow as it can be. process_cl() mirrored toneeq_process() and kept both host caches, but only the preview pipe one is ever read. Both GUI consumers take their data from it: update_histogram() through g->pd.buf, and _luminance_from_module_buffer() for the exposure under the cursor. The full pipe copy landed in g->full_preview_buf, which no GUI code reads, and on the OpenCL path the correction is applied on the device from dev_luminance, so no later GPU run read it back either. It was a blocking multi-megabyte transfer into a buffer nobody consumes, issued whenever the roi or the upstream pipe state changed: every pan and zoom step, and every motion event of a feathering or iterations drag, since those invalidate the luminance cache. Drop it, and let the full pipe mask stay on the device. The copy could not simply be deleted. toneeq_process() skips compute_luminance_mask() while g->ui_preview_hash still matches, so a GPU run that advanced that hash without filling the host buffer would let a later CPU run apply a stale mask. That is reachable by turning OpenCL off mid session, by cldevice_..._nocl=toneequal, and by any late fallback to the CPU path. The hash is therefore invalidated rather than advanced, so such a transition recomputes the mask once. That costs about one CPU process() run. The preview pipe copy stays as it is. The GUI genuinely needs that buffer in host memory, it goes through the shared dt_preview_data_t service, that pipe is small, and the blocking read is issued outside the module GUI lock so it cannot stall a GUI thread. Pixel output is unchanged: both kernels read dev_luminance either way, and gui_attached is FALSE on export, so this block is darkroom only. --- data/kernels/toneequal.cl | 2 +- src/iop/toneequal.c | 88 +++++++++++++++------------------------ 2 files changed, 34 insertions(+), 56 deletions(-) diff --git a/data/kernels/toneequal.cl b/data/kernels/toneequal.cl index 833534c3de1..6132ed9ea3e 100644 --- a/data/kernels/toneequal.cl +++ b/data/kernels/toneequal.cl @@ -1,6 +1,6 @@ /* This file is part of darktable, - Copyright (C) 2025 darktable developers. + Copyright (C) 2026 darktable developers. darktable is free software: you can redistribute it and/or modify it under the terms of the GNU General Public License as published by diff --git a/src/iop/toneequal.c b/src/iop/toneequal.c index 42ccfaeb8cb..2c86e642c79 100644 --- a/src/iop/toneequal.c +++ b/src/iop/toneequal.c @@ -1596,9 +1596,8 @@ int process_cl(dt_iop_module_t *self, return DT_OPENCL_PROCESS_CL; // input should be at least as large as output if(piece->colors != 4) return DT_OPENCL_PROCESS_CL; // we need RGB signal - // The GUI reads the luminance mask from host memory, so we mirror the - // caching logic of toneeq_process() and copy the mask back whenever the - // upstream pipe state has changed. + // Only the preview pipe publishes its luminance mask to the GUI, so only + // that one is copied back to host memory; see below. gboolean cached = FALSE; float *luminance = NULL; @@ -1618,17 +1617,19 @@ int process_cl(dt_iop_module_t *self, if(dt_pipe_is_full(piece->pipe)) { - // Re-allocate a new buffer if the full preview size has changed - if(g->full_preview_buf_width != width || g->full_preview_buf_height != height) - { - dt_free_align(g->full_preview_buf); - g->full_preview_buf = dt_alloc_align_float(num_elem); - g->full_preview_buf_width = width; - g->full_preview_buf_height = height; - } - - luminance = g->full_preview_buf; - cached = TRUE; + // The mask stays on the device : the correction is applied there and + // no GUI code reads g->full_preview_buf. Both GUI consumers, the + // histogram and the exposure under the cursor, are fed by the preview + // pipe buffer instead. Copying the full pipe mask back would be a + // blocking multi-megabyte transfer into a buffer nobody reads, and it + // would happen on every roi or upstream change. + // + // toneeq_process() skips compute_luminance_mask() while + // g->ui_preview_hash still matches, so invalidate it here: should the + // pipe fall back to the CPU, it has to recompute the mask rather than + // reuse a host buffer this path never filled. + const dt_hash_t invalid = DT_INVALID_HASH; + hash_set_get(&invalid, &g->ui_preview_hash, &self->gui_lock); } else if(dt_pipe_is_preview(piece->pipe)) { @@ -1661,55 +1662,32 @@ int process_cl(dt_iop_module_t *self, // and read the luminance under the cursor if(cached) { - const size_t bsize = num_elem * sizeof(float); + const dt_hash_t saved_hash = dt_preview_data_get_hash(&g->pd); - if(dt_pipe_is_full(piece->pipe)) - { - dt_hash_t saved_hash; - hash_set_get(&g->ui_preview_hash, &saved_hash, &self->gui_lock); + dt_iop_gui_enter_critical_section(self); + const gboolean luminance_valid = g->luminance_valid; + dt_iop_gui_leave_critical_section(self); + if(saved_hash != hash || !luminance_valid) + { + /* copy back only if upstream pipe state has changed */ + // Flag the cache as being recomputed so the GUI threads never + // read a partially filled buffer, then commit hash + validity + // once the data is ready. dt_iop_gui_enter_critical_section(self); - const gboolean luminance_valid = g->luminance_valid; + g->histogram_valid = FALSE; + g->luminance_valid = FALSE; dt_iop_gui_leave_critical_section(self); - if(hash != saved_hash || !luminance_valid) - { - /* copy back only if upstream pipe state has changed */ - err = dt_opencl_read_buffer_from_device(devid, luminance, dev_luminance, - 0, bsize, TRUE); - if(err != CL_SUCCESS) goto error; - hash_set_get(&hash, &g->ui_preview_hash, &self->gui_lock); - } - } - else if(dt_pipe_is_preview(piece->pipe)) - { - const dt_hash_t saved_hash = dt_preview_data_get_hash(&g->pd); + err = dt_opencl_read_buffer_from_device(devid, luminance, dev_luminance, + 0, num_elem * sizeof(float), TRUE); + if(err != CL_SUCCESS) goto error; + dt_preview_data_set_hash(&g->pd, piece); dt_iop_gui_enter_critical_section(self); - const gboolean luminance_valid = g->luminance_valid; + g->luminance_valid = TRUE; dt_iop_gui_leave_critical_section(self); - - if(saved_hash != hash || !luminance_valid) - { - /* copy back only if upstream pipe state has changed */ - // Flag the cache as being recomputed so the GUI threads never - // read a partially filled buffer, then commit hash + validity - // once the data is ready. - dt_iop_gui_enter_critical_section(self); - g->histogram_valid = FALSE; - g->luminance_valid = FALSE; - dt_iop_gui_leave_critical_section(self); - - err = dt_opencl_read_buffer_from_device(devid, luminance, dev_luminance, - 0, bsize, TRUE); - if(err != CL_SUCCESS) goto error; - dt_preview_data_set_hash(&g->pd, piece); - - dt_iop_gui_enter_critical_section(self); - g->luminance_valid = TRUE; - dt_iop_gui_leave_critical_section(self); - dt_dev_pixelpipe_cache_invalidate_later(piece->pipe, self->iop_order, "toneequal: "); - } + dt_dev_pixelpipe_cache_invalidate_later(piece->pipe, self->iop_order, "toneequal: "); } } From 5a8ca79a95724fd866eda2f7a8434518614c1f3b Mon Sep 17 00:00:00 2001 From: Philipp Lutz Date: Wed, 2 Sep 2026 07:08:03 +0000 Subject: [PATCH 3/8] preview data: allow committing an explicit hash dt_preview_data_set_hash() records dt_dev_pixelpipe_piece_hash(piece, roi, TRUE), which covers the owning module's own params. That is the right key for a buffer that is a copy of the module output, but not for one that only depends on a part of the module state. Add dt_preview_data_set_hash_value() so a module can commit a freshness key it computed itself, and document that dt_preview_data_is_fresh(), which recomputes the cumulative piece hash, does not apply to buffers stored that way. No functional change: no caller uses it yet. --- src/develop/preview_data.c | 10 ++++++++++ src/develop/preview_data.h | 9 +++++++++ 2 files changed, 19 insertions(+) diff --git a/src/develop/preview_data.c b/src/develop/preview_data.c index b26710cce78..26fa146702e 100644 --- a/src/develop/preview_data.c +++ b/src/develop/preview_data.c @@ -139,6 +139,16 @@ void dt_preview_data_set_hash(dt_preview_data_t *pd, dt_iop_gui_leave_critical_section((dt_iop_module_t *)pd->module); } +void dt_preview_data_set_hash_value(dt_preview_data_t *pd, + const dt_hash_t hash) +{ + if(!pd || !pd->module) return; + + dt_iop_gui_enter_critical_section((dt_iop_module_t *)pd->module); + pd->hash = hash; + dt_iop_gui_leave_critical_section((dt_iop_module_t *)pd->module); +} + gboolean dt_preview_data_get(dt_preview_data_t *pd, const size_t x, const size_t y, diff --git a/src/develop/preview_data.h b/src/develop/preview_data.h index 65c12db9f81..1a80026bab6 100644 --- a/src/develop/preview_data.h +++ b/src/develop/preview_data.h @@ -123,6 +123,15 @@ float *dt_preview_data_resize(dt_preview_data_t *pd, void dt_preview_data_set_hash(dt_preview_data_t *pd, const dt_dev_pixelpipe_iop_t *piece); +/** + * Record a freshness hash computed by the caller, for data that does not + * depend on the whole piece hash (e.g. the tone equalizer luminance mask, + * which ignores the module's own band factors). Such data must be checked + * with dt_preview_data_get_hash(), not with is_fresh(). Thread-safe. + */ +void dt_preview_data_set_hash_value(dt_preview_data_t *pd, + const dt_hash_t hash); + /** * Read a component of the per-pixel value at buffer pixel (x, y). * From 753b072de860e59e38b92cc79d2f5c364ad000d7 Mon Sep 17 00:00:00 2001 From: Philipp Lutz Date: Wed, 2 Sep 2026 07:08:03 +0000 Subject: [PATCH 4/8] toneequal: key the luminance mask on its own inputs toneeq_process() and process_cl() both gated mask recomputation on dt_dev_pixelpipe_piece_hash(piece, roi_out, TRUE) With include = TRUE that hashes nodes[0 .. position-1], and module->position is the module's own 1-based index in the pipe, so the module's own params were part of the key. The luminance mask does not depend on the nine band factors nor on `smoothing`; those only feed d->factors and the correction LUT. But moving any of them changed the key, so every frame of a band slider drag recomputed the whole mask on the CPU and on the GPU alike, and on the preview pipe copied a bit-identical mask device to host again. The mask dominates process(): a full CPU run measures ~60 ms/Mpix (-d perf, four cores, mire1.cr2 at 3908x2602 -> 0.575-0.694 s). Key it instead on the upstream pipe only (include = FALSE, which still folds in the roi) plus the params compute_luminance_mask() and _compute_luminance_mask_cl() actually read: method, details, radius (and the blending it derives from), feathering, iterations, quantization, scale, exposure_boost and contrast_boost. That is the same set gui_changed() calls invalidate_luminance_cache() for. --- src/iop/toneequal.c | 36 ++++++++++++++++++++++++++++++------ 1 file changed, 30 insertions(+), 6 deletions(-) diff --git a/src/iop/toneequal.c b/src/iop/toneequal.c index 2c86e642c79..940fad03cf8 100644 --- a/src/iop/toneequal.c +++ b/src/iop/toneequal.c @@ -639,6 +639,30 @@ static void hash_set_get(const dt_hash_t *hash_in, } +static dt_hash_t _luminance_mask_hash(dt_dev_pixelpipe_iop_t *piece, + const dt_iop_roi_t *const roi_out) +{ + // Freshness key of the cached luminance mask. + // + // include = TRUE hashes nodes[0 .. position-1], i.e. our own params too, so + // a band slider drag recomputed the whole mask (~60 ms/Mpix) and re-copied + // it from the device. The mask ignores the band factors and `smoothing`: + // hash the upstream pipe (include = FALSE, roi folded in) plus the params + // compute_luminance_mask() reads, the invalidate list of gui_changed(). + const dt_iop_toneequalizer_data_t *const d = piece->data; + + const float mask_floats[] = { d->blending, d->feathering, d->contrast_boost, + d->exposure_boost, d->quantization, d->scale }; + const int mask_ints[] = { d->radius, d->iterations, + (int)d->method, (int)d->details }; + + dt_hash_t hash = dt_dev_pixelpipe_piece_hash(piece, roi_out, FALSE); + hash = dt_hash(hash, mask_floats, sizeof(mask_floats)); + hash = dt_hash(hash, mask_ints, sizeof(mask_ints)); + return hash; +} + + static void invalidate_luminance_cache(dt_iop_module_t *const self) { // Invalidate the private luminance cache and histogram when @@ -1037,8 +1061,8 @@ void toneeq_process(dt_iop_module_t *self, const size_t height = roi_in->height; const size_t num_elem = width * height; - // Get the hash of the upstream pipe to track changes - const dt_hash_t hash = dt_dev_pixelpipe_piece_hash(piece, roi_out, TRUE); + // Freshness key of the luminance mask cache + const dt_hash_t hash = _luminance_mask_hash(piece, roi_out); // Sanity checks if(width < 1 || height < 1) return; @@ -1151,7 +1175,7 @@ void toneeq_process(dt_iop_module_t *self, dt_iop_gui_leave_critical_section(self); compute_luminance_mask(in, luminance, width, height, d); - dt_preview_data_set_hash(&g->pd, piece); + dt_preview_data_set_hash_value(&g->pd, hash); dt_iop_gui_enter_critical_section(self); g->luminance_valid = TRUE; @@ -1587,8 +1611,8 @@ int process_cl(dt_iop_module_t *self, const int height = roi_in->height; const size_t num_elem = (size_t)width * height; - // Get the hash of the upstream pipe to track changes - const dt_hash_t hash = dt_dev_pixelpipe_piece_hash(piece, roi_out, TRUE); + // Freshness key of the luminance mask cache + const dt_hash_t hash = _luminance_mask_hash(piece, roi_out); // Sanity checks if(width < 1 || height < 1) return DT_OPENCL_PROCESS_CL; @@ -1682,7 +1706,7 @@ int process_cl(dt_iop_module_t *self, err = dt_opencl_read_buffer_from_device(devid, luminance, dev_luminance, 0, num_elem * sizeof(float), TRUE); if(err != CL_SUCCESS) goto error; - dt_preview_data_set_hash(&g->pd, piece); + dt_preview_data_set_hash_value(&g->pd, hash); dt_iop_gui_enter_critical_section(self); g->luminance_valid = TRUE; From 432a770fdc5cb11f9b95c3c583b1b208962edb9f Mon Sep 17 00:00:00 2001 From: Philipp Lutz Date: Wed, 2 Sep 2026 07:38:48 +0000 Subject: [PATCH 5/8] Add entry to RELEASE_NOTES.md --- RELEASE_NOTES.md | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/RELEASE_NOTES.md b/RELEASE_NOTES.md index ec4af100c25..342d5d6afeb 100644 --- a/RELEASE_NOTES.md +++ b/RELEASE_NOTES.md @@ -153,6 +153,10 @@ changes (where available). - Separated the demosaicer's crop-and-scale step into a new hidden module to improve cache efficiency and UI responsiveness. +- Add a OpenCL code path to tone equalizer which gives around + 6-20x faster processing times in comparison to CPU code path, + dependening on CPU & GPU. + ## Other Changes - Added a new collection filter for the original image dimensions. From 2c535781679c4b7b56358debde9e1a4b00eb403a Mon Sep 17 00:00:00 2001 From: Philipp Lutz Date: Thu, 24 Sep 2026 23:47:17 +0200 Subject: [PATCH 6/8] toneequal: take gui_lock the same way in process_cl() process_cl() has its own copy of the luminance mask cache handling of toneeq_process(), so give it the same critical sections: - read g->mask_display in the section that sets up the cache, and display the mask from that snapshot rather than from gui_data, - test the preview mask and clear its flags in one section, and commit its hash and validity flag together in another. Unlike the CPU fill, the blocking device to host read of the preview mask stays outside the lock. On an in-order queue it waits for every kernel queued before it on the preview pipe, and a GUI thread blocked on gui_lock would wait for all of that GPU work too. --- src/iop/toneequal.c | 36 ++++++++++++++++++++---------------- 1 file changed, 20 insertions(+), 16 deletions(-) diff --git a/src/iop/toneequal.c b/src/iop/toneequal.c index 940fad03cf8..5e6c9b9826a 100644 --- a/src/iop/toneequal.c +++ b/src/iop/toneequal.c @@ -1625,17 +1625,20 @@ int process_cl(dt_iop_module_t *self, gboolean cached = FALSE; float *luminance = NULL; + gboolean mask_display = FALSE; if(self->dev->gui_attached && g) { + // gui_lock is recursive, so the dt_preview_data_*() calls below may + // take it again from inside this section + dt_iop_gui_enter_critical_section(self); + mask_display = g->mask_display; // If the module instance has changed order in the pipe, invalidate the caches if(g->pipe_order != piece->module->iop_order) { - dt_iop_gui_enter_critical_section(self); g->ui_preview_hash = DT_INVALID_HASH; g->pipe_order = piece->module->iop_order; g->luminance_valid = FALSE; g->histogram_valid = FALSE; - dt_iop_gui_leave_critical_section(self); dt_preview_data_invalidate(&g->pd); } @@ -1652,8 +1655,7 @@ int process_cl(dt_iop_module_t *self, // g->ui_preview_hash still matches, so invalidate it here: should the // pipe fall back to the CPU, it has to recompute the mask rather than // reuse a host buffer this path never filled. - const dt_hash_t invalid = DT_INVALID_HASH; - hash_set_get(&invalid, &g->ui_preview_hash, &self->gui_lock); + g->ui_preview_hash = DT_INVALID_HASH; } else if(dt_pipe_is_preview(piece->pipe)) { @@ -1664,6 +1666,7 @@ int process_cl(dt_iop_module_t *self, _toneeq_preview_resized, self); cached = TRUE; } + dt_iop_gui_leave_critical_section(self); if(cached && !luminance) { @@ -1689,26 +1692,27 @@ int process_cl(dt_iop_module_t *self, const dt_hash_t saved_hash = dt_preview_data_get_hash(&g->pd); dt_iop_gui_enter_critical_section(self); - const gboolean luminance_valid = g->luminance_valid; - dt_iop_gui_leave_critical_section(self); - - if(saved_hash != hash || !luminance_valid) + const gboolean stale = (saved_hash != hash) || !g->luminance_valid; + if(stale) { - /* copy back only if upstream pipe state has changed */ - // Flag the cache as being recomputed so the GUI threads never - // read a partially filled buffer, then commit hash + validity - // once the data is ready. - dt_iop_gui_enter_critical_section(self); + // Flag the cache as being recomputed so the GUI threads never read a partially filled buffer, + // then commit hash + validity once the data is ready. g->histogram_valid = FALSE; g->luminance_valid = FALSE; - dt_iop_gui_leave_critical_section(self); + } + dt_iop_gui_leave_critical_section(self); + if(stale) + { + // Copy back only if upstream pipe state has changed, unlike the CPU fill, keep this + // outside the critical section: the read waits for every kernel queued before it, and a + // GUI thread blocked on gui_lock would wait for all of that GPU work too err = dt_opencl_read_buffer_from_device(devid, luminance, dev_luminance, 0, num_elem * sizeof(float), TRUE); if(err != CL_SUCCESS) goto error; - dt_preview_data_set_hash_value(&g->pd, hash); dt_iop_gui_enter_critical_section(self); + dt_preview_data_set_hash_value(&g->pd, hash); g->luminance_valid = TRUE; dt_iop_gui_leave_critical_section(self); dt_dev_pixelpipe_cache_invalidate_later(piece->pipe, self->iop_order, "toneequal: "); @@ -1722,7 +1726,7 @@ int process_cl(dt_iop_module_t *self, const int offset_y = (roi_in->y < roi_out->y) ? roi_out->y - roi_in->y : 0; // Display output - if(self->dev->gui_attached && g && dt_pipe_is_full(piece->pipe) && g->mask_display) + if(self->dev->gui_attached && g && dt_pipe_is_full(piece->pipe) && mask_display) { err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_toneequal_display_mask, out_width, out_height, From 52064d673c928df01a9e8415fbafc2fd69d1f870 Mon Sep 17 00:00:00 2001 From: Philipp Lutz Date: Thu, 10 Sep 2026 16:56:59 +0000 Subject: [PATCH 7/8] toneequal: do not self-invalidate the preview cache on the OpenCL path process_cl() called dt_dev_pixelpipe_cache_invalidate_later() for self->iop_order after reading the luminance mask back from the device, mirroring a line toneeq_process() carried at the time the OpenCL path was written. 50e0964 has since removed that line from the CPU path: the call discards the module's own freshly written preview cacheline and every downstream one, while downstream modules consume the RGBA output, not g->pd.buf. Upstream and parameter changes are already tracked by the pixelpipe hashing. Drop it from the OpenCL path too, so the preview cachelines survive a mask readback there as well. --- src/iop/toneequal.c | 1 - 1 file changed, 1 deletion(-) diff --git a/src/iop/toneequal.c b/src/iop/toneequal.c index 5e6c9b9826a..b313aa18f5a 100644 --- a/src/iop/toneequal.c +++ b/src/iop/toneequal.c @@ -1715,7 +1715,6 @@ int process_cl(dt_iop_module_t *self, dt_preview_data_set_hash_value(&g->pd, hash); g->luminance_valid = TRUE; dt_iop_gui_leave_critical_section(self); - dt_dev_pixelpipe_cache_invalidate_later(piece->pipe, self->iop_order, "toneequal: "); } } From 6f3cd270266e47bdf7c196f384ec7f122a55fa33 Mon Sep 17 00:00:00 2001 From: Philipp Lutz Date: Thu, 10 Sep 2026 16:57:17 +0000 Subject: [PATCH 8/8] toneequal: drop the unset scale field from the luminance mask hash _luminance_mask_hash() folded d->scale into the key, but nothing ever assigns that field, so it is the zero left by init_pipe(), and the `scale` parameter it is forwarded to in fast_surface_blur() and fast_eigf_surface_blur() is unused anyway. Hash only what the mask builders read, and say so in the comment instead of retelling the history that lives in the commit log. --- src/iop/toneequal.c | 14 ++++++-------- 1 file changed, 6 insertions(+), 8 deletions(-) diff --git a/src/iop/toneequal.c b/src/iop/toneequal.c index b313aa18f5a..6a37288a1d0 100644 --- a/src/iop/toneequal.c +++ b/src/iop/toneequal.c @@ -642,17 +642,15 @@ static void hash_set_get(const dt_hash_t *hash_in, static dt_hash_t _luminance_mask_hash(dt_dev_pixelpipe_iop_t *piece, const dt_iop_roi_t *const roi_out) { - // Freshness key of the cached luminance mask. - // - // include = TRUE hashes nodes[0 .. position-1], i.e. our own params too, so - // a band slider drag recomputed the whole mask (~60 ms/Mpix) and re-copied - // it from the device. The mask ignores the band factors and `smoothing`: - // hash the upstream pipe (include = FALSE, roi folded in) plus the params - // compute_luminance_mask() reads, the invalidate list of gui_changed(). + // freshness key of the cached luminance mask: the upstream pipe and roi + // (include = FALSE leaves our own params out, the mask ignores the band + // factors and `smoothing`) plus the params the mask builders read. + // keep in sync with compute_luminance_mask(), _compute_luminance_mask_cl() + // and the invalidate list of gui_changed() const dt_iop_toneequalizer_data_t *const d = piece->data; const float mask_floats[] = { d->blending, d->feathering, d->contrast_boost, - d->exposure_boost, d->quantization, d->scale }; + d->exposure_boost, d->quantization }; const int mask_ints[] = { d->radius, d->iterations, (int)d->method, (int)d->details };