RFC: Gamma corrected pipe input scaling for rgb images

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,