Skip to content

Commit b4f1a62

Browse files
jenshannoschwalmTurboGit
jenshannoschwalm
authored and
TurboGit
committed
Improved & fixed highlights DT_DEV_PIXELPIPE_DISPLAY_PASSTHRU mode
If we are in passthru mode we should not clip/modify any data and use the fastest way to do so. As the module roi_in/out handling is somewhat tricky due to it's raw/non-raw support we do this by using clip mode (with a non-clipping value). While doing this the OpenCl code needed some care.
1 parent be97851 commit b4f1a62

File tree

2 files changed

+84
-57
lines changed

2 files changed

+84
-57
lines changed

‎data/kernels/basic.cl‎

Lines changed: 15 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -273,18 +273,26 @@ highlights_4f_clip (read_only image2d_t in, write_only image2d_t out, const int
273273
}
274274

275275
kernel void
276-
highlights_1f_clip (read_only image2d_t in, write_only image2d_t out, const int width, const int height,
277-
global float *clips, const int rx, const int ry, const int filters, global const unsigned char (*const xtrans)[6])
276+
highlights_1f_clip (read_only image2d_t in, write_only image2d_t out,
277+
const int iwidth, const int iheight,
278+
const int owidth, const int oheight,
279+
global float *clips, const int dx, const int dy,
280+
const int filters, global const unsigned char (*const xtrans)[6])
278281
{
279282
const int x = get_global_id(0);
280283
const int y = get_global_id(1);
281284

282-
if(x >= width || y >= height) return;
283-
const int color = (filters == 9u) ? FCxtrans(y, x, xtrans) : FC(y, x, filters);
284-
285-
float pixel = read_imagef(in, sampleri, (int2)(x, y)).x;
285+
if(x >= owidth || y >= oheight) return;
286286

287-
pixel = fmin(clips[color], pixel);
287+
const int irow = y + dy;
288+
const int icol = x + dx;
289+
float pixel = 0.0f;
290+
if((icol >= 0) && (irow >= 0) && (irow < iheight) && (icol < iwidth))
291+
{
292+
const int color = (filters == 9u) ? FCxtrans(irow, icol, xtrans) : FC(irow, icol, filters);
293+
pixel = read_imagef(in, sampleri, (int2)(icol, irow)).x;
294+
pixel = fmin(clips[color], pixel);
295+
}
288296
write_imagef (out, (int2)(x, y), pixel);
289297
}
290298

‎src/iop/highlights.c‎

Lines changed: 69 additions & 50 deletions
Original file line numberDiff line numberDiff line change
@@ -534,29 +534,33 @@ int process_cl(dt_iop_module_t *self,
534534
const dt_iop_roi_t *const roi_in,
535535
const dt_iop_roi_t *const roi_out)
536536
{
537+
dt_dev_pixelpipe_t *pipe = piece->pipe;
537538
dt_iop_highlights_data_t *d = piece->data;
538539
dt_iop_highlights_gui_data_t *g = self->gui_data;
539540
dt_iop_highlights_global_data_t *gd = self->global_data;
540541

541-
const uint32_t filters = piece->pipe->dsc.filters;
542-
const int devid = piece->pipe->devid;
542+
const uint32_t filters = pipe->dsc.filters;
543+
const int devid = pipe->devid;
543544

544-
const gboolean fullpipe = piece->pipe->type & DT_DEV_PIXELPIPE_FULL;
545-
gboolean announce = dt_iop_piece_is_raster_mask_used(piece, BLEND_RASTER_ID);
545+
const gboolean fullpipe = pipe->type & DT_DEV_PIXELPIPE_FULL;
546+
const gboolean passthru = pipe->mask_display == DT_DEV_PIXELPIPE_DISPLAY_PASSTHRU;
547+
const dt_iop_highlights_mode_t dmode = passthru && fullpipe ? DT_IOP_HIGHLIGHTS_CLIP : d->mode;
548+
549+
gboolean announce = passthru ? FALSE : dt_iop_piece_is_raster_mask_used(piece, BLEND_RASTER_ID);
546550

547551
cl_int err = DT_OPENCL_DEFAULT_ERROR;
548552
cl_mem dev_xtrans = NULL;
549553
cl_mem dev_clips = NULL;
550554

551-
if(g && fullpipe)
555+
if(g && fullpipe && !passthru)
552556
{
553557
if(g->hlr_mask_mode != DT_HIGHLIGHTS_MASK_OFF)
554558
{
555-
piece->pipe->mask_display = DT_DEV_PIXELPIPE_DISPLAY_PASSTHRU;
559+
pipe->mask_display = DT_DEV_PIXELPIPE_DISPLAY_PASSTHRU;
556560
if(g->hlr_mask_mode == DT_HIGHLIGHTS_MASK_CLIPPED)
557561
{
558562
const float mclip = d->clip * highlights_clip_magics[d->mode];
559-
const float *c = piece->pipe->dsc.temperature.coeffs;
563+
const float *c = pipe->dsc.temperature.coeffs;
560564
float clips[4] = { mclip * (c[RED] <= 0.0f ? 1.0f : c[RED]),
561565
mclip * (c[GREEN] <= 0.0f ? 1.0f : c[GREEN]),
562566
mclip * (c[BLUE] <= 0.0f ? 1.0f : c[BLUE]),
@@ -565,7 +569,7 @@ int process_cl(dt_iop_module_t *self,
565569
dev_clips = dt_opencl_copy_host_to_device_constant(devid, 4 * sizeof(float), clips);
566570
if(dev_clips == NULL) goto finish;
567571

568-
dev_xtrans = dt_opencl_copy_host_to_device_constant(devid, sizeof(piece->pipe->dsc.xtrans), piece->pipe->dsc.xtrans);
572+
dev_xtrans = dt_opencl_copy_host_to_device_constant(devid, sizeof(pipe->dsc.xtrans), pipe->dsc.xtrans);
569573
if(dev_xtrans == NULL) goto finish;
570574

571575
err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_highlights_false_color, roi_out->width, roi_out->height,
@@ -581,21 +585,21 @@ int process_cl(dt_iop_module_t *self,
581585
}
582586
}
583587

584-
const float clip = d->clip * dt_iop_get_processed_minimum(piece);
588+
const float clip = passthru ? FLT_MAX : (d->clip * dt_iop_get_processed_minimum(piece));
585589

586590
if(!filters)
587591
{
588592
// non-raw images use dedicated kernel which just clips
589593
err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_highlights_4f_clip, roi_in->width, roi_in->height,
590594
CLARG(dev_in), CLARG(dev_out),
591595
CLARG(roi_in->width), CLARG(roi_in->height),
592-
CLARG(d->mode), CLARG(clip));
596+
CLARG(dmode), CLARG(clip));
593597
}
594-
else if(d->mode == DT_IOP_HIGHLIGHTS_OPPOSED)
598+
else if(dmode == DT_IOP_HIGHLIGHTS_OPPOSED)
595599
{
596600
err = process_opposed_cl(self, piece, dev_in, dev_out, roi_in, roi_out);
597601
}
598-
else if(d->mode == DT_IOP_HIGHLIGHTS_LCH && filters != 9u)
602+
else if(dmode == DT_IOP_HIGHLIGHTS_LCH && filters != 9u)
599603
{
600604
// bayer sensor raws with LCH mode
601605
err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_highlights_1f_lch_bayer, roi_in->width, roi_in->height,
@@ -604,7 +608,7 @@ int process_cl(dt_iop_module_t *self,
604608
CLARG(clip), CLARG(roi_out->x), CLARG(roi_out->y),
605609
CLARG(filters));
606610
}
607-
else if(d->mode == DT_IOP_HIGHLIGHTS_LCH && filters == 9u)
611+
else if(dmode == DT_IOP_HIGHLIGHTS_LCH && filters == 9u)
608612
{
609613
// xtrans sensor raws with LCH mode
610614
int blocksizex, blocksizey;
@@ -622,8 +626,7 @@ int process_cl(dt_iop_module_t *self,
622626
else
623627
blocksizex = blocksizey = 1;
624628

625-
dev_xtrans
626-
= dt_opencl_copy_host_to_device_constant(devid, sizeof(piece->pipe->dsc.xtrans), piece->pipe->dsc.xtrans);
629+
dev_xtrans = dt_opencl_copy_host_to_device_constant(devid, sizeof(pipe->dsc.xtrans), pipe->dsc.xtrans);
627630
if(dev_xtrans == NULL) goto finish;
628631

629632
size_t sizes[] = { ROUNDUP(roi_in->width, blocksizex), ROUNDUP(roi_in->height, blocksizey), 1 };
@@ -636,15 +639,15 @@ int process_cl(dt_iop_module_t *self,
636639
CLLOCAL(sizeof(float) * (blocksizex + 4) * (blocksizey + 4)));
637640
err = dt_opencl_enqueue_kernel_2d_with_local(devid, gd->kernel_highlights_1f_lch_xtrans, sizes, local);
638641
}
639-
else if(d->mode == DT_IOP_HIGHLIGHTS_LAPLACIAN)
642+
else if(dmode == DT_IOP_HIGHLIGHTS_LAPLACIAN)
640643
{
641644
const float clipper = d->clip * highlights_clip_magics[DT_IOP_HIGHLIGHTS_LAPLACIAN];
642-
const dt_aligned_pixel_t clips = { clipper * piece->pipe->dsc.processed_maximum[0],
643-
clipper * piece->pipe->dsc.processed_maximum[1],
644-
clipper * piece->pipe->dsc.processed_maximum[2], clip };
645+
const dt_aligned_pixel_t clips = { clipper * pipe->dsc.processed_maximum[0],
646+
clipper * pipe->dsc.processed_maximum[1],
647+
clipper * pipe->dsc.processed_maximum[2], clip };
645648
err = process_laplacian_bayer_cl(self, piece, dev_in, dev_out, roi_in, roi_out, clips);
646649
}
647-
else // (d->mode == DT_IOP_HIGHLIGHTS_CLIP)
650+
else // (dmode == DT_IOP_HIGHLIGHTS_CLIP)
648651
{
649652
const dt_dev_chroma_t *chr = &self->dev->chroma;
650653
dt_aligned_pixel_t clips = { clip, clip, clip, clip};
@@ -654,16 +657,21 @@ int process_cl(dt_iop_module_t *self,
654657
dev_clips = dt_opencl_copy_host_to_device_constant(devid, 4 * sizeof(float), clips);
655658
if(dev_clips == NULL) goto finish;
656659

657-
dev_xtrans = dt_opencl_copy_host_to_device_constant(devid, sizeof(piece->pipe->dsc.xtrans), piece->pipe->dsc.xtrans);
660+
dev_xtrans = dt_opencl_copy_host_to_device_constant(devid, sizeof(pipe->dsc.xtrans), pipe->dsc.xtrans);
658661
if(dev_xtrans == NULL) goto finish;
662+
659663
// raw images with clip mode (both bayer and xtrans)
660-
err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_highlights_1f_clip, roi_in->width, roi_in->height,
664+
const int dy = roi_out->y - roi_in->y;
665+
const int dx = roi_out->x - roi_in->x;
666+
err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_highlights_1f_clip, roi_out->width, roi_out->height,
661667
CLARG(dev_in), CLARG(dev_out),
662668
CLARG(roi_in->width), CLARG(roi_in->height),
663-
CLARG(dev_clips), CLARG(roi_out->x), CLARG(roi_out->y),
669+
CLARG(roi_out->width), CLARG(roi_out->height),
670+
CLARG(dev_clips), CLARG(dx), CLARG(dy),
664671
CLARG(filters), CLARG(dev_xtrans));
665672
}
666673
if(err != CL_SUCCESS) goto finish;
674+
if(passthru) goto finish;
667675

668676
float *mask = NULL;
669677
if(announce)
@@ -686,7 +694,7 @@ int process_cl(dt_iop_module_t *self,
686694
// The guided laplacian and opposed are the modes that keeps signal scene-referred and don't clip highlights to 1
687695
// For the other modes, we need to notify the pipeline that white point has changed
688696
const float m = dt_iop_get_processed_maximum(piece);
689-
for_three_channels(k) piece->pipe->dsc.processed_maximum[k] = m;
697+
for_three_channels(k) pipe->dsc.processed_maximum[k] = m;
690698
}
691699

692700
finish:
@@ -813,32 +821,37 @@ void process(dt_iop_module_t *self,
813821
const dt_iop_roi_t *const roi_in,
814822
const dt_iop_roi_t *const roi_out)
815823
{
816-
const uint32_t filters = piece->pipe->dsc.filters;
824+
dt_dev_pixelpipe_t *pipe = piece->pipe;
825+
const uint32_t filters = pipe->dsc.filters;
817826
dt_iop_highlights_data_t *d = piece->data;
818827
dt_iop_highlights_gui_data_t *g = self->gui_data;
819828

820-
const gboolean fullpipe = piece->pipe->type & DT_DEV_PIXELPIPE_FULL;
821-
const gboolean fastmode = piece->pipe->type & DT_DEV_PIXELPIPE_FAST;
822-
const gboolean scaled = filters == 0 && d->mode != DT_IOP_HIGHLIGHTS_CLIP;
829+
const gboolean fullpipe = pipe->type & DT_DEV_PIXELPIPE_FULL;
830+
const gboolean fastmode = pipe->type & DT_DEV_PIXELPIPE_FAST;
831+
832+
const gboolean passthru = pipe->mask_display == DT_DEV_PIXELPIPE_DISPLAY_PASSTHRU;
833+
dt_iop_highlights_mode_t dmode = passthru ? DT_IOP_HIGHLIGHTS_CLIP : d->mode;
834+
835+
const gboolean scaled = filters == 0 && dmode != DT_IOP_HIGHLIGHTS_CLIP;
823836

824837
float *out = scaled ? dt_alloc_align_float((size_t)roi_in->width * roi_in->height * 4) : NULL;
825-
const gboolean announce = dt_iop_piece_is_raster_mask_used(piece, BLEND_RASTER_ID);
838+
const gboolean announce = passthru ? FALSE : dt_iop_piece_is_raster_mask_used(piece, BLEND_RASTER_ID);
826839

827840
if(!out && scaled)
828841
{
829842
dt_iop_clip_and_zoom_roi((float *)ovoid, (float *)ivoid, roi_out, roi_in);
830843
dt_print_pipe(DT_DEBUG_ALWAYS,
831-
"bypass highlights", piece->pipe, self, DT_DEVICE_CPU, roi_in, roi_out,
844+
"bypass highlights", pipe, self, DT_DEVICE_CPU, roi_in, roi_out,
832845
"can't allocate temp buffer");
833846
dt_iop_piece_clear_raster(piece, NULL);
834847
return;
835848
}
836849

837-
if(g && fullpipe)
850+
if(g && fullpipe && !passthru)
838851
{
839852
if(g->hlr_mask_mode != DT_HIGHLIGHTS_MASK_OFF)
840853
{
841-
piece->pipe->mask_display = DT_DEV_PIXELPIPE_DISPLAY_PASSTHRU;
854+
pipe->mask_display = DT_DEV_PIXELPIPE_DISPLAY_PASSTHRU;
842855
if(g->hlr_mask_mode == DT_HIGHLIGHTS_MASK_CLIPPED)
843856
{
844857
if(scaled)
@@ -858,24 +871,27 @@ void process(dt_iop_module_t *self,
858871

859872
/* While rendering thumnbnails we look for an acceptable lower quality */
860873
gboolean high_quality = TRUE;
861-
if(piece->pipe->type & DT_DEV_PIXELPIPE_THUMBNAIL)
874+
if(pipe->type & DT_DEV_PIXELPIPE_THUMBNAIL)
862875
{
863-
const dt_mipmap_size_t level = dt_mipmap_cache_get_matching_size(piece->pipe->final_width, piece->pipe->final_height);
876+
const dt_mipmap_size_t level = dt_mipmap_cache_get_matching_size(pipe->final_width, pipe->final_height);
864877
const char *min = dt_conf_get_string_const("plugins/lighttable/thumbnail_hq_min_level");
865878
const dt_mipmap_size_t min_s = dt_mipmap_cache_get_min_mip_from_pref(min);
866879
high_quality = (level >= min_s);
867880
}
868881

869-
const float clip = d->clip * dt_iop_get_processed_minimum(piece);
882+
const float clip = passthru ? FLT_MAX : d->clip * dt_iop_get_processed_minimum(piece);
870883

871884
if(filters == 0)
872885
{
873-
if(d->mode == DT_IOP_HIGHLIGHTS_CLIP)
886+
if(dmode == DT_IOP_HIGHLIGHTS_CLIP)
874887
{
875888
process_clip(self, piece, ivoid, ovoid, roi_in, roi_out, clip);
876-
const float m = dt_iop_get_processed_minimum(piece);
877-
for_three_channels(k)
878-
piece->pipe->dsc.processed_maximum[k] = m;
889+
if(!passthru)
890+
{
891+
const float m = dt_iop_get_processed_minimum(piece);
892+
for_three_channels(k)
893+
pipe->dsc.processed_maximum[k] = m;
894+
}
879895
}
880896
else
881897
{
@@ -884,27 +900,28 @@ void process(dt_iop_module_t *self,
884900
dt_free_align(out);
885901
}
886902

903+
if(passthru) return;
904+
887905
float *mask = announce ? _provide_raster_mask(roi_in, roi_out, (float *)ovoid, d->clip, piece) : NULL;
888906
if(mask) dt_iop_piece_set_raster(piece, mask, roi_in, roi_out);
889907
else dt_iop_piece_clear_raster(piece, NULL);
890908

891909
return;
892910
}
893911

894-
const dt_iop_highlights_mode_t dmode = fastmode && (d->mode == DT_IOP_HIGHLIGHTS_SEGMENTS)
895-
? DT_IOP_HIGHLIGHTS_OPPOSED : d->mode;
912+
dmode = fastmode && (dmode == DT_IOP_HIGHLIGHTS_SEGMENTS) ? DT_IOP_HIGHLIGHTS_OPPOSED : dmode;
896913
switch(dmode)
897914
{
898915
case DT_IOP_HIGHLIGHTS_INPAINT: // a1ex's (magiclantern) idea of color inpainting:
899916
{
900917
const float clipper = d->clip * highlights_clip_magics[DT_IOP_HIGHLIGHTS_INPAINT];
901-
const float clips[4] = { clipper * piece->pipe->dsc.processed_maximum[0],
902-
clipper * piece->pipe->dsc.processed_maximum[1],
903-
clipper * piece->pipe->dsc.processed_maximum[2], clip };
918+
const float clips[4] = { clipper * pipe->dsc.processed_maximum[0],
919+
clipper * pipe->dsc.processed_maximum[1],
920+
clipper * pipe->dsc.processed_maximum[2], clip };
904921

905922
if(filters == 9u)
906923
{
907-
const uint8_t(*const xtrans)[6] = (const uint8_t(*const)[6])piece->pipe->dsc.xtrans;
924+
const uint8_t(*const xtrans)[6] = (const uint8_t(*const)[6])pipe->dsc.xtrans;
908925
DT_OMP_FOR()
909926
for(int j = 0; j < roi_out->height; j++)
910927
{
@@ -967,9 +984,9 @@ void process(dt_iop_module_t *self,
967984
case DT_IOP_HIGHLIGHTS_LAPLACIAN:
968985
{
969986
const float clipper = d->clip * highlights_clip_magics[DT_IOP_HIGHLIGHTS_LAPLACIAN];
970-
const dt_aligned_pixel_t clips = { clipper * piece->pipe->dsc.processed_maximum[0],
971-
clipper * piece->pipe->dsc.processed_maximum[1],
972-
clipper * piece->pipe->dsc.processed_maximum[2], clip };
987+
const dt_aligned_pixel_t clips = { clipper * pipe->dsc.processed_maximum[0],
988+
clipper * pipe->dsc.processed_maximum[1],
989+
clipper * pipe->dsc.processed_maximum[2], clip };
973990
process_laplacian_bayer(self, piece, ivoid, ovoid, roi_in, roi_out, clips);
974991
break;
975992
}
@@ -981,17 +998,19 @@ void process(dt_iop_module_t *self,
981998
}
982999
}
9831000

1001+
if(passthru) return;
1002+
9841003
float *mask = announce ? _provide_raster_mask(roi_in, roi_out, (float *)ovoid, d->clip, piece) : NULL;
9851004
if(mask) dt_iop_piece_set_raster(piece, mask, roi_in, roi_out);
9861005
else dt_iop_piece_clear_raster(piece, NULL);
9871006

9881007
// update processed maximum
989-
if((d->mode != DT_IOP_HIGHLIGHTS_LAPLACIAN) && (d->mode != DT_IOP_HIGHLIGHTS_SEGMENTS) && (d->mode != DT_IOP_HIGHLIGHTS_OPPOSED))
1008+
if((dmode != DT_IOP_HIGHLIGHTS_LAPLACIAN) && (dmode != DT_IOP_HIGHLIGHTS_SEGMENTS) && (dmode != DT_IOP_HIGHLIGHTS_OPPOSED) && !passthru)
9901009
{
9911010
// The guided laplacian, inpaint opposed and segmentation modes keep signal scene-referred and don't clip highlights to 1
9921011
// For the other modes, we need to notify the pipeline that white point has changed
9931012
const float m = dt_iop_get_processed_maximum(piece);
994-
for_three_channels(k) piece->pipe->dsc.processed_maximum[k] = m;
1013+
for_three_channels(k) pipe->dsc.processed_maximum[k] = m;
9951014
}
9961015
}
9971016

0 commit comments

Comments
 (0)