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. 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..6132ed9ea3e --- /dev/null +++ b/data/kernels/toneequal.cl @@ -0,0 +1,549 @@ +/* + This file is part of darktable, + 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 + 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/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). * diff --git a/src/iop/toneequal.c b/src/iop/toneequal.c index 0f9081b7bb1..6a37288a1d0 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; @@ -621,6 +639,28 @@ 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: 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 }; + 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 @@ -813,8 +853,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, @@ -1018,8 +1059,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; @@ -1132,7 +1173,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; @@ -1180,6 +1221,594 @@ 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; + + // 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; + 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 + + // 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; + + 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) + { + g->ui_preview_hash = DT_INVALID_HASH; + g->pipe_order = piece->module->iop_order; + g->luminance_valid = FALSE; + g->histogram_valid = FALSE; + dt_preview_data_invalidate(&g->pd); + } + + if(dt_pipe_is_full(piece->pipe)) + { + // 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. + g->ui_preview_hash = DT_INVALID_HASH; + } + 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; + } + dt_iop_gui_leave_critical_section(self); + + 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 dt_hash_t saved_hash = dt_preview_data_get_hash(&g->pd); + + dt_iop_gui_enter_critical_section(self); + const gboolean stale = (saved_hash != hash) || !g->luminance_valid; + if(stale) + { + // 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); + + 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_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); + } + } + + // 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) && 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 +2195,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; }