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

Author: jenshannoschwalm

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));
+}

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

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);
 }

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

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);

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.

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,

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);

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,