Skip to content

Opencl diff part4 #20709

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
TurboGit merged 2 commits into from
Mar 31, 2026
Merged

Opencl diff part4 #20709

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
d9f2d66
Fix highlights clip mode if no temperature data is provided
jenshannoschwalm Mar 29, 2026
8017a87
Graduated ND improvements & maintenance
jenshannoschwalm Mar 30, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 1 addition & 2 deletions data/kernels/basic.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -251,8 +251,7 @@ exposure (read_only image2d_t in, write_only image2d_t out, const int width, con

/* kernel for the highlights plugin. */
kernel void
highlights_4f_clip (read_only image2d_t in, write_only image2d_t out, const int width, const int height,
const int mode, const float clip)
highlights_4f_clip (read_only image2d_t in, write_only image2d_t out, const int width, const int height, const float clip)
{
const int x = get_global_id(0);
const int y = get_global_id(1);
Expand Down
2 changes: 2 additions & 0 deletions data/kernels/common.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,6 +34,8 @@ constant sampler_t samplerA = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE
#define M_PI_F 3.14159265358979323846f // should be defined by the OpenCL compiler acc. to standard
#endif

#define M_LN2f 0.69314718055994530942f

#define LUT_ELEM 512 // gamut LUT number of elements:

#define RED 0
Expand Down
64 changes: 29 additions & 35 deletions data/kernels/extended.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,6 @@
/*
This file is part of darktable,
copyright (c) 2009--2014 Ulrich Pegelow
Copyright (C) 2009-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
Expand All @@ -22,58 +22,52 @@


__kernel void
graduatedndp (read_only image2d_t in, write_only image2d_t out, const int width, const int height, const float4 color,
const float density, const float length_base, const float length_inc_x, const float length_inc_y)
graduatedndp (read_only image2d_t in,
write_only image2d_t out,
const int width,
const int height,
const float4 color,
const float density,
const float length_base,
const float length_inc_x,
const float length_inc_y)
{
const int x = get_global_id(0);
const int y = get_global_id(1);

if(x >= width || y >= height) return;

float4 pixel = read_imagef(in, sampleri, (int2)(x, y));

const float len = length_base + y*length_inc_y + x*length_inc_x;

const float t = 0.693147181f * (density * clipf(0.5f+len)/8.0f);
const float d1 = t * t * 0.5f;
const float d2 = d1 * t * 0.333333333f;
const float d3 = d2 * t * 0.25f;
float dens = 1.0f + t + d1 + d2 + d3;
dens *= dens;
dens *= dens;
dens *= dens;
float4 pixel = readpixel(in, x, y);

pixel.xyz = fmax((float4)0.0f, pixel / (color + ((float4)1.0f - color) * (float4)dens)).xyz;

write_imagef (out, (int2)(x, y), pixel);
float len = length_base + y*length_inc_y + x*length_inc_x;
float dens = dtcl_exp2(density * clipf(0.5f + len));
pixel = pixel / (color + ((float4)1.0f - color) * (float4)dens);
write_imagef (out, (int2)(x, y), fmax(0.0f, pixel));
}


__kernel void
graduatedndm (read_only image2d_t in, write_only image2d_t out, const int width, const int height, const float4 color,
const float density, const float length_base, const float length_inc_x, const float length_inc_y)
graduatedndm (read_only image2d_t in,
write_only image2d_t out,
const int width,
const int height,
const float4 color,
const float density,
const float length_base,
const float length_inc_x,
const float length_inc_y)
{
const int x = get_global_id(0);
const int y = get_global_id(1);

if(x >= width || y >= height) return;

float4 pixel = read_imagef(in, sampleri, (int2)(x, y));

const float len = length_base + y*length_inc_y + x*length_inc_x;

const float t = 0.693147181f * (-density * clipf(0.5f-len)/8.0f);
const float d1 = t * t * 0.5f;
const float d2 = d1 * t * 0.333333333f;
const float d3 = d2 * t * 0.25f;
float dens = 1.0f + t + d1 + d2 + d3;
dens *= dens;
dens *= dens;
dens *= dens;
float4 pixel = readpixel(in, x, y);

pixel.xyz = fmax((float4)0.0f, pixel * (color + ((float4)1.0f - color) * (float4)dens)).xyz;

write_imagef (out, (int2)(x, y), pixel);
float len = length_base + y*length_inc_y + x*length_inc_x;
float dens = dtcl_exp2(-density * clipf(0.5f - len));
pixel = pixel * (color + ((float4)1.0f - color) * (float4)dens);
write_imagef(out, (int2)(x, y), fmax(0.0f, pixel));
}

__kernel void
Expand Down
10 changes: 6 additions & 4 deletions src/common/math.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,7 +49,9 @@
#define M_PI_F 3.14159265358979323846f
#endif /* !M_PI_F */

#define DT_M_LN2f (0.6931471805599453f)
#ifndef M_LN2f
#define M_LN2f 0.69314718055994530942f
#endif

// clip channel value to be between 0 and 1
// NaN-safe: NaN compares false and will result in 0.0
Expand Down Expand Up @@ -152,13 +154,13 @@ static inline float scharr_gradient(const float *p, const int w)

static inline float Log2(const float x)
{
return (x > 0.0f) ? (logf(x) / DT_M_LN2f) : x;
return (x > 0.0f) ? (logf(x) / M_LN2f) : x;
}

static inline float Log2Thres(const float x,
const float Thres)
{
return logf(x > Thres ? x : Thres) / DT_M_LN2f;
return logf(x > Thres ? x : Thres) / M_LN2f;
}

// ensure that any changes here are synchronized with data/kernels/extended.cl
Expand All @@ -179,7 +181,7 @@ static inline float fastlog2(const float x)
// ensure that any changes here are synchronized with data/kernels/extended.cl
static inline float fastlog(const float x)
{
return DT_M_LN2f * fastlog2(x);
return M_LN2f * fastlog2(x);
}

// multiply 3x3 matrix with 3x1 vector
Expand Down
8 changes: 4 additions & 4 deletions src/iop/basicadj.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1058,16 +1058,16 @@ static void _get_auto_exp(const uint32_t *const histogram, const unsigned int hi
// compute exposure compensation as geometric mean of the amount that
// sets the mean or median at middle gray, and the amount that sets the estimated top
// of the histogram at or near clipping.
const float expcomp1 = (logf(midgray * scale / (ave - shc + midgray * shc))) / DT_M_LN2f;
const float expcomp1 = (logf(midgray * scale / (ave - shc + midgray * shc))) / M_LN2f;
float expcomp2;

if(overex == 0) // image is not overexposed
{
expcomp2 = 0.5f * ((15.5f - histcompr - (2.f * oct7 - oct6)) + logf(scale / rawmax) / DT_M_LN2f);
expcomp2 = 0.5f * ((15.5f - histcompr - (2.f * oct7 - oct6)) + logf(scale / rawmax) / M_LN2f);
}
else
{
expcomp2 = 0.5f * ((15.5f - histcompr - (2.f * octile[7] - octile[6])) + logf(scale / rawmax) / DT_M_LN2f);
expcomp2 = 0.5f * ((15.5f - histcompr - (2.f * octile[7] - octile[6])) + logf(scale / rawmax) / M_LN2f);
}

if(fabsf(expcomp1) - fabsf(expcomp2) > 1.f) // for great expcomp
Expand All @@ -1079,7 +1079,7 @@ static void _get_auto_exp(const uint32_t *const histogram, const unsigned int hi
expcomp = 0.5 * (double)expcomp1 + 0.5 * (double)expcomp2; // for small expcomp
}

const float gain = expf(expcomp * DT_M_LN2f);
const float gain = expf(expcomp * M_LN2f);

const float corr = sqrtf(gain * scale / rawmax);
black = shc * corr;
Expand Down
49 changes: 22 additions & 27 deletions src/iop/graduatednd.c
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,6 @@
/*
This file is part of darktable,
Copyright (C) 2010-2025 darktable developers.
Copyright (C) 2010-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
Expand Down Expand Up @@ -127,8 +127,8 @@ typedef struct dt_iop_graduatednd_data_t
float hardness; // Default 0% = soft and 100% = hard
float rotation; // 2*PI -180 - +180
float offset; // Default 50%, centered, can be offsetted...
float color[4]; // RGB color of gradient
float color1[4]; // inverted color (1 - c)
dt_aligned_pixel_t color; // RGB color of gradient
dt_aligned_pixel_t color1; // inverted color (1 - c)
} dt_iop_graduatednd_data_t;


Expand Down Expand Up @@ -719,9 +719,9 @@ int scrolled(
{
float dens;
if(up)
dens = fminf(8.0, p->density + 0.1);
dens = fminf(8.0f, p->density + 0.1f);
else
dens = fmaxf(-8.0, p->density - 0.1);
dens = fmaxf(-8.0f, p->density - 0.1f);
if(dens != p->density)
{
dt_bauhaus_slider_set(g->density, dens);
Expand All @@ -732,9 +732,9 @@ int scrolled(
{
float comp;
if(up)
comp = fminf(100.0, p->hardness + 1.0);
comp = fminf(100.0f, p->hardness + 1.0f);
else
comp = fmaxf(0.0, p->hardness - 1.0);
comp = fmaxf(0.0f, p->hardness - 1.0f);
if(comp != p->hardness)
{
dt_bauhaus_slider_set(g->hardness, comp);
Expand All @@ -747,18 +747,18 @@ int scrolled(
DT_OMP_DECLARE_SIMD(simdlen(4))
static inline float _density_times_length(const float dens, const float length)
{
return (dens * CLAMP(0.5f + length, 0.0f, 1.0f) / 8.0f);
return (dens * CLIP(0.5f + length) / 8.0f);
}

DT_OMP_DECLARE_SIMD(simdlen(4))
static inline float _compute_density(const float dens, const float length)
{
#if 1
#ifdef __FAST_MATH__
// !!! approximation is ok only when highest density is 8
// for input x = (data->density * CLIP( 0.5+length ), calculate 2^x as (e^(ln2*x/8))^8
// use exp2f approximation to calculate e^(ln2*x/8)
// in worst case - density==8,CLIP(0.5-length) == 1.0 it gives 0.6% of error
const float t = DT_M_LN2f * _density_times_length(dens,length);
const float t = M_LN2f * _density_times_length(dens,length);
const float d1 = t * t * 0.5f;
const float d2 = d1 * t * 0.333333333f;
const float d3 = d2 * t * 0.25f;
Expand Down Expand Up @@ -799,8 +799,8 @@ void process(dt_iop_module_t *self,
const float sinv = sinf(v);
const float cosv = cosf(v);
const float cosv_hh_inv = cosv * hh_inv;
const float filter_radie = sqrtf((hh * hh) + (hw * hw)) / hh;
const float offset = data->offset / 100.0f * 2;
const float filter_radie = hypotf(hh, hw) / hh;
const float offset = data->offset / 100.0f * 2.0f;

const float filter_hardness = (1.0f / filter_radie)
/ (1.0f - (0.5f + (data->hardness / 100.0f) * 0.9f / 2.0f)) * 0.5f;
Expand Down Expand Up @@ -927,7 +927,8 @@ void process(dt_iop_module_t *self,
#ifdef HAVE_OPENCL
int process_cl(dt_iop_module_t *self,
dt_dev_pixelpipe_iop_t *piece,
cl_mem dev_in, cl_mem dev_out,
cl_mem dev_in,
cl_mem dev_out,
const dt_iop_roi_t *const roi_in,
const dt_iop_roi_t *const roi_out)
{
Expand All @@ -949,20 +950,14 @@ int process_cl(dt_iop_module_t *self,
const float v = deg2radf(-data->rotation);
const float sinv = sinf(v);
const float cosv = cosf(v);
const float filter_radie = sqrtf((hh * hh) + (hw * hw)) / hh;
const float offset = data->offset / 100.0f * 2;
const float filter_radie = hypotf(hh, hw) / hh;
const float offset = data->offset / 100.0f * 2.0f;
const float density = data->density;

#if 1
const float filter_hardness = 1.0 / filter_radie
/ (1.0 - (0.5 + (data->hardness / 100.0) * 0.9 / 2.0)) * 0.5;
#else
const float hardness = data->hardness / 100.0f;
const float t = 1.0f - .8f / (.8f + hardness);
const float c = 1.0f + 1000.0f * powf(4.0, hardness);
#endif
const float filter_hardness = (1.0f / filter_radie)
/ (1.0f - (0.5f + (data->hardness / 100.0f) * 0.9f / 2.0f)) * 0.5f;

const float length_base = (sinv * (-1.0 + ix * hw_inv) - cosv * (-1.0 + iy * hh_inv) - 1.0 + offset)
const float length_base = (sinv * (-1.0 + ix * hw_inv) - cosv * (-1.0f + iy * hh_inv) - 1.0f + offset)
* filter_hardness;
const float length_inc_y = -cosv * hh_inv * filter_hardness;
const float length_inc_x = sinv * hw_inv * filter_hardness;
Expand All @@ -971,7 +966,7 @@ int process_cl(dt_iop_module_t *self,
const int kernel = density > 0 ? gd->kernel_graduatedndp : gd->kernel_graduatedndm;

return dt_opencl_enqueue_kernel_2d_args(devid, kernel, width, height,
CLARG(dev_in), CLARG(dev_out), CLARG(width), CLARG(height), CLARRAY(4, data->color), CLARG(density),
CLARG(dev_in), CLARG(dev_out), CLARG(width), CLARG(height), CLARG(data->color), CLARG(density),
CLARG(length_base), CLARG(length_inc_x), CLARG(length_inc_y));
}
#endif
Expand Down Expand Up @@ -1026,9 +1021,9 @@ void commit_params(dt_iop_module_t *self,
d->color[3] = 0.0f;

if(d->density < 0)
for(int l = 0; l < 4; l++) d->color[l] = 1.0 - d->color[l];
for(int l = 0; l < 4; l++) d->color[l] = 1.0f - d->color[l];

for(int l = 0; l < 4; l++) d->color1[l] = 1.0 - d->color[l];
for(int l = 0; l < 4; l++) d->color1[l] = 1.0f - d->color[l];
}

void init_pipe(dt_iop_module_t *self,
Expand Down
Loading
Loading