All demosaicer should not write negatives

Demosaicer input has negatives, also it's output even if input didn't.
Make sure we don't push negatives down the pipe.

Can be disabled at compiletime via undefining DT_DEMOSAIC_POSITIVE

Author: jenshannoschwalm

data/kernels/demosaic_rcd.cl

@@ -359,3 +359,15 @@ kernel void demosaic_box3(read_only image2d_t in,
   write_imagef(out, (int2)(col, row), rgb);
 }
 
+kernel void
+image_positive(read_only image2d_t in, write_only image2d_t out, const int width, const int height)
+{
+  const int x = get_global_id(0);
+  const int y = get_global_id(1);
+
+  if(x >= width || y >= height) return;
+
+  float4 pixel = fmax(0.0f, read_imagef(in, sampleri, (int2)(x, y)));
+  write_imagef(out, (int2)(x, y), pixel);
+}
+

src/iop/demosaic.c

@@ -52,6 +52,8 @@ DT_MODULE_INTROSPECTION(6, dt_iop_demosaic_params_t)
 #define DT_DEMOSAIC_XTRANS 1024 // masks for non-Bayer demosaic ops
 #define DT_DEMOSAIC_DUAL 2048   // masks for dual demosaicing methods
 
+#define DT_DEMOSAIC_POSITIVE
+
 typedef enum dt_iop_demosaic_method_t
 {
   // methods for Bayer images
@@ -240,6 +242,7 @@ typedef struct dt_iop_demosaic_global_data_t
   int show_blend_mask;
   int capture_result;
   int final_blend;
+  int image_positive;
   float *gauss_coeffs;
 } dt_iop_demosaic_global_data_t;
 
@@ -899,6 +902,12 @@ void process(dt_iop_module_t *self,
   if(d->color_smoothing != DT_DEMOSAIC_SMOOTH_OFF && no_masking && !run_fast)
     color_smoothing(out, width, height, d->color_smoothing);
 
+#ifdef DT_DEMOSAIC_POSITIVE
+  DT_OMP_FOR_SIMD(aligned(out : 64))
+  for(size_t k = 0; k < (size_t)width * height * 4; k++)
+    out[k] = fmaxf(0.0f, out[k]);
+#endif
+
   if(!direct)
   {
     dt_iop_clip_and_zoom_roi((float *)o, out, roi_out, roi_in);
@@ -1200,8 +1209,13 @@ int process_cl(dt_iop_module_t *self,
     if(err != CL_SUCCESS) goto finish;
   }
 
-  if(!direct)
-     err = dt_iop_clip_and_zoom_roi_cl(devid, dev_out, out_image, roi_out, roi_in);
+#ifdef DT_DEMOSAIC_POSITIVE
+  err = dt_opencl_enqueue_kernel_2d_args(devid, gd->image_positive, iwidth, iheight,
+        CLARG(out_image), CLARG(out_image), CLARG(iwidth), CLARG(iheight));
+#endif
+
+  if(!direct && err == CL_SUCCESS)
+    err = dt_iop_clip_and_zoom_roi_cl(devid, dev_out, out_image, roi_out, roi_in);
 
 finish:
   dt_opencl_release_mem_object(dev_xtrans);
@@ -1284,6 +1298,7 @@ void init_global(dt_iop_module_so_t *self)
   gd->kernel_rcd_step_5_2 = dt_opencl_create_kernel(rcd, "rcd_step_5_2");
   gd->kernel_demosaic_box3 = dt_opencl_create_kernel(rcd, "demosaic_box3");
   gd->kernel_write_blended_dual  = dt_opencl_create_kernel(rcd, "write_blended_dual");
+  gd->image_positive  = dt_opencl_create_kernel(rcd, "image_positive");
 
   const int capt = 38; // capture.cl, from programs.conf
   gd->gaussian_9x9_mul = dt_opencl_create_kernel(capt, "kernel_9x9_mul");
@@ -1357,6 +1372,7 @@ void cleanup_global(dt_iop_module_so_t *self)
   dt_opencl_free_kernel(gd->show_blend_mask);
   dt_opencl_free_kernel(gd->capture_result);
   dt_opencl_free_kernel(gd->final_blend);
+  dt_opencl_free_kernel(gd->image_positive);
   dt_free_align(gd->gauss_coeffs);
   free(self->data);
   self->data = NULL;