From fa0107cc0e45dc448caa75bec2b06ec7a6ee5e10 Mon Sep 17 00:00:00 2001 From: Hanno Schwalm Date: Sun, 28 Dec 2025 13:05:21 +0100 Subject: [PATCH] Gamma corrected pipe input scaling Non-raw images can have gamma corrected data as input possibly leading to bad results while we do the initial scaling. This scaling happens when generating the mipmap via `_init_f` and when inserting data into the pixelpipes. What we do now: - we test for the image having exif data as sRGB or AdobeRGB. - the colorin module does a much better check than exif data and possibly sets the profile type to any RGB mode. We keep track on this when setting the input profile and possibly change the image->colorspace (making sure to leave the exif checked data as they are). - dt_iop_clip_and_zoom() got an extra gboolean parameter `gamma`, if that is true the interpolation is encapsulated in rgb->linear and linear->rgb conversions. - we have a new kernel doing the flog gamma correction allowing OpenCL downscaling when inserting into the pipe. This results in improved data visualizing for all pipes and data taken from the preview pipe. qdefrve --- data/kernels/colorspaces.cl | 12 ++++++ src/common/image.h | 3 +- src/common/iop_profile.c | 25 +++++++++++- src/common/iop_profile.h | 1 + src/common/mipmap_cache.c | 8 ++-- src/develop/imageop_math.c | 28 ++++++++++++-- src/develop/imageop_math.h | 2 +- src/develop/pixelpipe_hb.c | 76 ++++++++++++++++++++++++++++++++++--- src/iop/finalscale.c | 2 +- 9 files changed, 140 insertions(+), 17 deletions(-) diff --git a/data/kernels/colorspaces.cl b/data/kernels/colorspaces.cl index 1fa66f8c3dd2..9c807409d3cf 100644 --- a/data/kernels/colorspaces.cl +++ b/data/kernels/colorspaces.cl @@ -81,3 +81,15 @@ colorspaces_transform_rgb_matrix_to_rgb(read_only image2d_t in, write_only image write_imagef(out, (int2)(x, y), pixel); } + +kernel void +colorspaces_transform_gamma(read_only image2d_t in, write_only image2d_t out, const int width, const int height, const float gamma) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + + if(x >= width || y >= height) return; + + const float4 pixel = fmax(0.0f, read_imagef(in, sampleri, (int2)(x, y)) ); + write_imagef(out, (int2)(x, y), dtcl_pow(pixel, gamma)); +} diff --git a/src/common/image.h b/src/common/image.h index 0ab41c21597e..97631e9e6070 100644 --- a/src/common/image.h +++ b/src/common/image.h @@ -107,7 +107,8 @@ typedef enum dt_image_colorspace_t { DT_IMAGE_COLORSPACE_NONE, DT_IMAGE_COLORSPACE_SRGB, - DT_IMAGE_COLORSPACE_ADOBE_RGB + DT_IMAGE_COLORSPACE_ADOBE_RGB, + DT_IMAGE_COLORSPACE_USER_RGB } dt_image_colorspace_t; typedef struct dt_image_raw_parameters_t diff --git a/src/common/iop_profile.c b/src/common/iop_profile.c index 137f7478ef62..ef61b3cace19 100644 --- a/src/common/iop_profile.c +++ b/src/common/iop_profile.c @@ -22,6 +22,7 @@ #include "common/debug.h" #include "common/imagebuf.h" #include "common/matrices.h" +#include "common/image_cache.h" #include "control/control.h" #include "develop/imageop.h" #include "develop/imageop_math.h" @@ -920,8 +921,25 @@ dt_ioppr_set_pipe_input_profile_info(struct dt_develop_t *dev, profile_info = dt_ioppr_add_profile_info_to_list(dev, DT_COLORSPACE_LIN_REC2020, "", intent); } - if(profile_info->type >= DT_COLORSPACE_EMBEDDED_ICC - && profile_info->type <= DT_COLORSPACE_ALTERNATE_MATRIX) + const dt_imgid_t imgid = dev->image_storage.id; + const dt_image_t *cimg = dt_image_cache_get(imgid, 'r'); + const dt_image_colorspace_t ocsp = cimg->colorspace; + dt_image_cache_read_release(cimg); + + const dt_colorspaces_color_profile_type_t ptype = profile_info->type; + const gboolean exif_rgb = ocsp == DT_IMAGE_COLORSPACE_SRGB || ocsp == DT_IMAGE_COLORSPACE_ADOBE_RGB; + const gboolean new_rgb = ptype == DT_COLORSPACE_SRGB || ptype == DT_COLORSPACE_ADOBERGB; + const gboolean user_rgb = ocsp == DT_IMAGE_COLORSPACE_USER_RGB; + if(!exif_rgb // don't fiddle with existing exif data + && ((new_rgb && !user_rgb) || (!new_rgb && user_rgb))) + { + dt_image_t *wimg = dt_image_cache_get(imgid, 'w'); + wimg->colorspace = new_rgb ? DT_IMAGE_COLORSPACE_USER_RGB : DT_IMAGE_COLORSPACE_NONE; + dt_image_cache_write_release_info(wimg, DT_IMAGE_CACHE_RELAXED, NULL); + } + + if(ptype >= DT_COLORSPACE_EMBEDDED_ICC + && ptype <= DT_COLORSPACE_ALTERNATE_MATRIX) { /* We have a camera input matrix, these are not generated from files but in colorin, * so we need to fetch and replace them from somewhere. @@ -1319,6 +1337,8 @@ dt_colorspaces_cl_global_t *dt_colorspaces_init_cl_global() dt_opencl_create_kernel(program, "colorspaces_transform_rgb_matrix_to_lab"); g->kernel_colorspaces_transform_rgb_matrix_to_rgb = dt_opencl_create_kernel(program, "colorspaces_transform_rgb_matrix_to_rgb"); + g->kernel_colorspaces_gamma = + dt_opencl_create_kernel(program, "colorspaces_transform_gamma"); return g; } @@ -1330,6 +1350,7 @@ void dt_colorspaces_free_cl_global(dt_colorspaces_cl_global_t *g) dt_opencl_free_kernel(g->kernel_colorspaces_transform_lab_to_rgb_matrix); dt_opencl_free_kernel(g->kernel_colorspaces_transform_rgb_matrix_to_lab); dt_opencl_free_kernel(g->kernel_colorspaces_transform_rgb_matrix_to_rgb); + dt_opencl_free_kernel(g->kernel_colorspaces_gamma); free(g); } diff --git a/src/common/iop_profile.h b/src/common/iop_profile.h index e4e52162d618..d94b7045c4ed 100644 --- a/src/common/iop_profile.h +++ b/src/common/iop_profile.h @@ -181,6 +181,7 @@ typedef struct dt_colorspaces_cl_global_t int kernel_colorspaces_transform_lab_to_rgb_matrix; int kernel_colorspaces_transform_rgb_matrix_to_lab; int kernel_colorspaces_transform_rgb_matrix_to_rgb; + int kernel_colorspaces_gamma; } dt_colorspaces_cl_global_t; // must be in synch with colorspaces.cl dt_colorspaces_iccprofile_info_cl_t diff --git a/src/common/mipmap_cache.c b/src/common/mipmap_cache.c index 19cc83122ca4..83f896c83e28 100644 --- a/src/common/mipmap_cache.c +++ b/src/common/mipmap_cache.c @@ -1379,10 +1379,12 @@ static void _init_f(dt_mipmap_buffer_t *mipmap_buf, } else { - // downsample + // scale + const gboolean gamma = image->colorspace != DT_IMAGE_COLORSPACE_NONE; dt_print_pipe(DT_DEBUG_PIPE, - "mipmap clip and zoom", NULL, NULL, DT_DEVICE_CPU, &roi_in, &roi_out); - dt_iop_clip_and_zoom(out, (const float *)buf.buf, &roi_out, &roi_in); + "mipmap clip&zoom", NULL, NULL, DT_DEVICE_NONE, &roi_in, &roi_out, "%s", + gamma ? "gamma corrected" : ""); + dt_iop_clip_and_zoom(out, (const float *)buf.buf, &roi_out, &roi_in, gamma); } dt_mipmap_cache_release(&buf); diff --git a/src/develop/imageop_math.c b/src/develop/imageop_math.c index 748d4725c153..2dcf1818002c 100644 --- a/src/develop/imageop_math.c +++ b/src/develop/imageop_math.c @@ -144,15 +144,35 @@ void dt_iop_clip_and_zoom_8(const uint8_t *i, } } -// apply clip and zoom on parts of a supplied full image. -// roi_in and roi_out define which part to work on. +/* apply clip and zoom on parts of a supplied full image, roi_in and roi_out define which part to work on. + gamma correction around scaling supported. + We don't do full RGB->linear and transformation but only use a gamma as that is fully sufficient + for scaling. +*/ void dt_iop_clip_and_zoom(float *out, const float *const in, const dt_iop_roi_t *const roi_out, - const dt_iop_roi_t *const roi_in) + const dt_iop_roi_t *const roi_in, + const gboolean gamma) { const dt_interpolation_t *itor = dt_interpolation_new(DT_INTERPOLATION_USERPREF); - dt_interpolation_resample(itor, out, roi_out, in, roi_in); + float *linear = gamma ? dt_alloc_align_float((size_t)roi_in->width * roi_in->height * 4) : NULL; + if(!linear) + return dt_interpolation_resample(itor, out, roi_out, in, roi_in); + + static const dt_aligned_pixel_t two_point_four = { 2.4f, 2.4f, 2.4f, 2.4f }; + static const dt_aligned_pixel_t rev_two_point_four = { 1.0f / 2.4f, 1.0f / 2.4f, 1.0f / 2.4f, 1.0f / 2.4f }; + + DT_OMP_SIMD(aligned(in, linear : 16)) + for(size_t k = 0; k < (size_t)roi_in->width * roi_in->height*4; k += 4) + dt_vector_powf(&in[k], two_point_four, &linear[k]); + + dt_interpolation_resample(itor, out, roi_out, linear, roi_in); + dt_free_align(linear); + + DT_OMP_SIMD(aligned(out : 16)) + for(size_t k = 0; k < (size_t)roi_out->width * roi_out->height * 4; k += 4) + dt_vector_powf(&out[k], rev_two_point_four, &out[k]); } // apply clip and zoom on the image region supplied in the input buffer. diff --git a/src/develop/imageop_math.h b/src/develop/imageop_math.h index f0347440eabc..8bc306cc396e 100644 --- a/src/develop/imageop_math.h +++ b/src/develop/imageop_math.h @@ -34,7 +34,7 @@ void dt_iop_flip_and_zoom_8(const uint8_t *in, int32_t iw, int32_t ih, uint8_t * /** for homebrew pixel pipe: zoom pixel array. */ void dt_iop_clip_and_zoom(float *out, const float *const in, const struct dt_iop_roi_t *const roi_out, - const struct dt_iop_roi_t *const roi_in); + const struct dt_iop_roi_t *const roi_in, const gboolean gamma); /** zoom pixel array for roi buffers. */ void dt_iop_clip_and_zoom_roi(float *out, const float *const in, const struct dt_iop_roi_t *const roi_out, diff --git a/src/develop/pixelpipe_hb.c b/src/develop/pixelpipe_hb.c index 784b88fb2487..a6e5ac6cee02 100644 --- a/src/develop/pixelpipe_hb.c +++ b/src/develop/pixelpipe_hb.c @@ -1724,15 +1724,81 @@ static gboolean _dev_pixelpipe_process_rec(dt_dev_pixelpipe_t *pipe, roi_in.height = pipe->iheight; roi_in.scale = 1.0f; const gboolean valid_bpp = (bpp == 4 * sizeof(float)); - + const gboolean gamma = dev->image_storage.colorspace != DT_IMAGE_COLORSPACE_NONE; +#ifdef HAVE_OPENCL + const size_t in_size = bpp * roi_in.width * roi_in.height; + const gboolean cl_scale_possible = ((2 * in_size) < dt_opencl_get_device_available(pipe->devid)) + && roi_out->width < roi_in.width + && roi_out->height < roi_in.height; +#else + const gboolean cl_scale_possible = FALSE: +#endif dt_print_pipe(DT_DEBUG_PIPE, "pipe data: clip&zoom", - pipe, module, DT_DEVICE_CPU, &roi_in, roi_out, "%s%s", - valid_bpp ? "" : "requires 4 floats data", - aligned_input ? "" : "non-aligned input buffer"); + pipe, module, pipe->devid, &roi_in, roi_out, "%s%s%s%s", + valid_bpp ? "" : "requires 4 floats data ", + aligned_input ? "" : "non-aligned input buffer ", + cl_scale_possible ? "OpenCL scaling " : "", + gamma ? "gamma corrected" : ""); if(valid_bpp && aligned_input) - dt_iop_clip_and_zoom(*output, pipe->input, roi_out, &roi_in); + { + gboolean done = FALSE; +#ifdef HAVE_OPENCL + if(cl_scale_possible) + { + cl_mem tmp_input = dt_opencl_alloc_device(pipe->devid, roi_in.width, roi_in.height, bpp); + cl_mem linear_input = dt_opencl_alloc_device(pipe->devid, roi_in.width, roi_in.height, bpp); + cl_mem tmp_output = NULL; + cl_mem linear_output = NULL; + if(tmp_input && linear_input) + { + cl_int err = dt_opencl_write_host_to_device(pipe->devid, pipe->input, tmp_input, roi_in.width, roi_in.height, bpp); + if(err == CL_SUCCESS) + { + const float fgamma = 2.4f; + err = dt_opencl_enqueue_kernel_2d_args(pipe->devid, darktable.opencl->colorspaces->kernel_colorspaces_gamma, roi_in.width, roi_in.height, + CLARG(tmp_input), CLARG(linear_input), CLARG(roi_in.width), CLARG(roi_in.height), CLARG(fgamma)); + dt_opencl_release_mem_object(tmp_input); + tmp_input = NULL; + } + if(err == CL_SUCCESS) + { + linear_output = dt_opencl_alloc_device(pipe->devid, roi_out->width, roi_out->height, bpp); + if(!linear_output) err = CL_MEM_OBJECT_ALLOCATION_FAILURE; + if(err == CL_SUCCESS) + err = dt_iop_clip_and_zoom_cl(pipe->devid, linear_output, linear_input, roi_out, &roi_in); + dt_opencl_release_mem_object(linear_input); + linear_input = NULL; + } + if(err == CL_SUCCESS) + { + tmp_output = dt_opencl_alloc_device(pipe->devid, roi_out->width, roi_out->height, bpp); + if(!tmp_output) err = CL_MEM_OBJECT_ALLOCATION_FAILURE; + const float fgamma = 1.0f / 2.4f; + if(err == CL_SUCCESS) + err = dt_opencl_enqueue_kernel_2d_args(pipe->devid, darktable.opencl->colorspaces->kernel_colorspaces_gamma, roi_out->width, roi_out->height, + CLARG(linear_output), CLARG(tmp_output), CLARG(roi_out->width), CLARG(roi_out->height), CLARG(fgamma)); + } + if(err == CL_SUCCESS) + err = dt_opencl_copy_device_to_host(pipe->devid, *output, tmp_output, roi_out->width, roi_out->height, bpp); + if(err == CL_SUCCESS) + done = TRUE; + else + dt_print(DT_DEBUG_PIPE | DT_DEBUG_OPENCL, "OpenCL pipe data: clip&zoom failed"); + } + dt_opencl_release_mem_object(tmp_input); + dt_opencl_release_mem_object(tmp_output); + dt_opencl_release_mem_object(linear_input); + dt_opencl_release_mem_object(linear_output); + } + + if(!done) + dt_iop_clip_and_zoom(*output, pipe->input, roi_out, &roi_in, gamma); +#else + dt_iop_clip_and_zoom(*output, pipe->input, roi_out, &roi_in, gamma); +#endif + } else { memset(*output, 0, (size_t)roi_out->width * roi_out->height * bpp); diff --git a/src/iop/finalscale.c b/src/iop/finalscale.c index a2163ab0cffb..07c73136ca77 100644 --- a/src/iop/finalscale.c +++ b/src/iop/finalscale.c @@ -186,7 +186,7 @@ void process(dt_iop_module_t *self, if(exporting) dt_iop_clip_and_zoom_roi((float *)ovoid, (float *)ivoid, roi_out, roi_in); else - dt_iop_clip_and_zoom((float *)ovoid, (float *)ivoid, roi_out, roi_in); + dt_iop_clip_and_zoom((float *)ovoid, (float *)ivoid, roi_out, roi_in, FALSE); } void commit_params(dt_iop_module_t *self,