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,