void process(struct dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, void *i, void *o,
const dt_iop_roi_t *roi_in, const dt_iop_roi_t *roi_out)
{
dt_iop_exposure_data_t *d = (dt_iop_exposure_data_t *)piece->data;
commit_params_late(self, piece);
const float black = d->black;
const float white = exposure2white(d->exposure);
const int ch = piece->colors;
const float scale = 1.0 / (white - black);
const __m128 blackv = _mm_set1_ps(black);
const __m128 scalev = _mm_set1_ps(scale);
#ifdef _OPENMP
#pragma omp parallel for default(none) shared(roi_out, i, o) schedule(static)
#endif
for(int k = 0; k < roi_out->height; k++)
{
const float *in = ((float *)i) + (size_t)ch * k * roi_out->width;
float *out = ((float *)o) + (size_t)ch * k * roi_out->width;
for(int j = 0; j < roi_out->width; j++, in += 4, out += 4)
_mm_store_ps(out, (_mm_load_ps(in) - blackv) * scalev);
}
if(piece->pipe->mask_display) dt_iop_alpha_copy(i, o, roi_out->width, roi_out->height);
for(int k = 0; k < 3; k++) piece->pipe->processed_maximum[k] *= scale;
}
int process_cl(struct dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, cl_mem dev_in, cl_mem dev_out,
const dt_iop_roi_t *roi_in, const dt_iop_roi_t *roi_out)
{
dt_iop_exposure_data_t *d = (dt_iop_exposure_data_t *)piece->data;
dt_iop_exposure_global_data_t *gd = (dt_iop_exposure_global_data_t *)self->data;
commit_params_late(self, piece);
cl_int err = -999;
const float black = d->black;
const float white = exposure2white(d->exposure);
const float scale = 1.0 / (white - black);
const int devid = piece->pipe->devid;
const int width = roi_in->width;
const int height = roi_in->height;
size_t sizes[] = { ROUNDUPWD(width), ROUNDUPHT(height), 1 };
dt_opencl_set_kernel_arg(devid, gd->kernel_exposure, 0, sizeof(cl_mem), (void *)&dev_in);
dt_opencl_set_kernel_arg(devid, gd->kernel_exposure, 1, sizeof(cl_mem), (void *)&dev_out);
dt_opencl_set_kernel_arg(devid, gd->kernel_exposure, 2, sizeof(int), (void *)&width);
dt_opencl_set_kernel_arg(devid, gd->kernel_exposure, 3, sizeof(int), (void *)&height);
dt_opencl_set_kernel_arg(devid, gd->kernel_exposure, 4, sizeof(float), (void *)&black);
dt_opencl_set_kernel_arg(devid, gd->kernel_exposure, 5, sizeof(float), (void *)&scale);
err = dt_opencl_enqueue_kernel_2d(devid, gd->kernel_exposure, sizes);
if(err != CL_SUCCESS) goto error;
for(int k = 0; k < 3; k++) piece->pipe->processed_maximum[k] *= scale;
return TRUE;
error:
dt_print(DT_DEBUG_OPENCL, "[opencl_exposure] couldn't enqueue kernel! %d\n", err);
return FALSE;
}
void process(dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, const void *const ivoid, void *const ovoid,
const dt_iop_roi_t *const roi_in, const dt_iop_roi_t *const roi_out)
{
const int ch = piece->colors;
const dt_iop_levels_data_t *const d = (dt_iop_levels_data_t *)piece->data;
if(d->mode == LEVELS_MODE_AUTOMATIC)
{
commit_params_late(self, piece);
}
#ifdef _OPENMP
#pragma omp parallel for default(none) schedule(static)
#endif
for(int k = 0; k < roi_out->height; k++)
{
float *in = (float *)ivoid + (size_t)k * ch * roi_out->width;
float *out = (float *)ovoid + (size_t)k * ch * roi_out->width;
for(int j = 0; j < roi_out->width; j++, in += ch, out += ch)
{
float L_in = in[0] / 100.0f;
if(L_in <= d->levels[0])
{
// Anything below the lower threshold just clips to zero
out[0] = 0.0f;
}
else if(L_in >= d->levels[2])
{
float percentage = (L_in - d->levels[0]) / (d->levels[2] - d->levels[0]);
out[0] = 100.0f * pow(percentage, d->in_inv_gamma);
}
else
{
// Within the expected input range we can use the lookup table
float percentage = (L_in - d->levels[0]) / (d->levels[2] - d->levels[0]);
// out[0] = 100.0 * pow(percentage, d->in_inv_gamma);
out[0] = d->lut[CLAMP((int)(percentage * 0xfffful), 0, 0xffff)];
}
// Preserving contrast
if(in[0] > 0.01f)
{
out[1] = in[1] * out[0] / in[0];
out[2] = in[2] * out[0] / in[0];
}
else
{
out[1] = in[1] * out[0] / 0.01f;
out[2] = in[2] * out[0] / 0.01f;
}
}
}
if(piece->pipe->mask_display) dt_iop_alpha_copy(ivoid, ovoid, roi_out->width, roi_out->height);
}
int process_cl(dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, cl_mem dev_in, cl_mem dev_out,
const dt_iop_roi_t *const roi_in, const dt_iop_roi_t *const roi_out)
{
dt_iop_levels_data_t *d = (dt_iop_levels_data_t *)piece->data;
dt_iop_levels_global_data_t *gd = (dt_iop_levels_global_data_t *)self->data;
if(d->mode == LEVELS_MODE_AUTOMATIC)
{
commit_params_late(self, piece);
}
cl_mem dev_lut = NULL;
cl_int err = -999;
const int devid = piece->pipe->devid;
const int width = roi_out->width;
const int height = roi_out->height;
dev_lut = dt_opencl_copy_host_to_device(devid, d->lut, 256, 256, sizeof(float));
if(dev_lut == NULL) goto error;
size_t sizes[2] = { ROUNDUPWD(width), ROUNDUPHT(height) };
dt_opencl_set_kernel_arg(devid, gd->kernel_levels, 0, sizeof(cl_mem), &dev_in);
dt_opencl_set_kernel_arg(devid, gd->kernel_levels, 1, sizeof(cl_mem), &dev_out);
dt_opencl_set_kernel_arg(devid, gd->kernel_levels, 2, sizeof(int), &width);
dt_opencl_set_kernel_arg(devid, gd->kernel_levels, 3, sizeof(int), &height);
dt_opencl_set_kernel_arg(devid, gd->kernel_levels, 4, sizeof(cl_mem), &dev_lut);
dt_opencl_set_kernel_arg(devid, gd->kernel_levels, 5, sizeof(float), &d->levels[0]);
dt_opencl_set_kernel_arg(devid, gd->kernel_levels, 6, sizeof(float), &d->levels[2]);
dt_opencl_set_kernel_arg(devid, gd->kernel_levels, 7, sizeof(float), &d->in_inv_gamma);
err = dt_opencl_enqueue_kernel_2d(devid, gd->kernel_levels, sizes);
if(err != CL_SUCCESS) goto error;
dt_opencl_release_mem_object(dev_lut);
return TRUE;
error:
if(dev_lut != NULL) dt_opencl_release_mem_object(dev_lut);
dt_print(DT_DEBUG_OPENCL, "[opencl_levels] couldn't enqueue kernel! %d\n", err);
return FALSE;
}
void commit_params (struct dt_iop_module_t *self, dt_iop_params_t *p1, dt_dev_pixelpipe_t *pipe, dt_dev_pixelpipe_iop_t *piece)
{
dt_iop_exposure_params_t *p = (dt_iop_exposure_params_t *)p1;
dt_iop_exposure_data_t *d = (dt_iop_exposure_data_t *)piece->data;
dt_iop_exposure_gui_data_t *g = (dt_iop_exposure_gui_data_t *)self->gui_data;
d->black = p->black;
d->exposure = p->exposure;
self->request_histogram &= ~(DT_REQUEST_ON);
self->request_histogram |= (DT_REQUEST_ONLY_IN_GUI);
self->request_histogram_source = (DT_DEV_PIXELPIPE_PREVIEW);
if(self->dev->gui_attached)
{
g->reprocess_on_next_expose = FALSE;
}
gboolean histogram_is_good = ((self->histogram_stats.bins_count == 16384)
&& (self->histogram != NULL));
d->deflicker_percentile = p->deflicker_percentile;
d->deflicker_target_level = p->deflicker_target_level;
if(p->mode == EXPOSURE_MODE_DEFLICKER && dt_image_is_raw(&self->dev->image_storage))
{
if(p->deflicker_histogram_source == DEFLICKER_HISTOGRAM_SOURCE_SOURCEFILE)
{
if(self->dev->gui_attached)
{
// histogram is precomputed and cached
compute_correction(self, piece, g->deflicker_histogram, &g->deflicker_histogram_stats, &d->exposure);
}
else
{
uint32_t *histogram = NULL;
dt_dev_histogram_stats_t histogram_stats;
deflicker_prepare_histogram(self, &histogram, &histogram_stats);
compute_correction(self, piece, histogram, &histogram_stats, &d->exposure);
free(histogram);
}
d->mode = EXPOSURE_MODE_MANUAL;
}
else
{
if(p->deflicker_histogram_source == DEFLICKER_HISTOGRAM_SOURCE_THUMBNAIL)
{
self->request_histogram |= (DT_REQUEST_ON);
gboolean failed = !histogram_is_good;
if(self->dev->gui_attached && histogram_is_good)
{
/*
* if in GUI, user might zoomed main view => we would get histogram of
* only part of image, so if in GUI we must always use histogram of
* preview pipe, which is always full-size and have biggest size
*/
d->mode = EXPOSURE_MODE_DEFLICKER;
commit_params_late(self, piece);
d->mode = EXPOSURE_MODE_MANUAL;
if(isnan(d->exposure))
failed = TRUE;
}
else if(failed || !(self->dev->gui_attached && histogram_is_good))
{
d->mode = EXPOSURE_MODE_DEFLICKER;
//commit_params_late() will compute correct d->exposure later
self->request_histogram &= ~(DT_REQUEST_ONLY_IN_GUI);
self->request_histogram_source = (DT_DEV_PIXELPIPE_ANY);
if(failed && self->dev->gui_attached)
{
/*
* but sadly we do not yet have a histogram to do so, so this time
* we process asif not in gui, and in expose() immediately
* reprocess, thus everything works as expected
*/
g->reprocess_on_next_expose = TRUE;
}
}
}
}
}
else
{
d->mode = EXPOSURE_MODE_MANUAL;
}
}
