static float dt_iop_exposure_get_white(struct dt_iop_module_t *self)
{
dt_iop_exposure_params_t *p = (dt_iop_exposure_params_t *)self->params;
if(p->mode == EXPOSURE_MODE_DEFLICKER)
{
return p->deflicker_target_level;
}
else
{
return exposure2white(p->exposure);
}
}
static void exposure_set_black(struct dt_iop_module_t *self, const float black)
{
dt_iop_exposure_params_t *p = (dt_iop_exposure_params_t *)self->params;
float b = black;
if(p->black == b) return;
p->black = b;
if(p->black >= exposure2white(p->exposure))
{
exposure_set_white(self, p->black + 0.01);
}
dt_iop_exposure_gui_data_t *g = (dt_iop_exposure_gui_data_t *)self->gui_data;
darktable.gui->reset = 1;
dt_bauhaus_slider_set_soft(g->black, p->black);
darktable.gui->reset = 0;
dt_dev_add_history_item(darktable.develop, self, TRUE);
}
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;
dt_iop_exposure_gui_data_t *g = (dt_iop_exposure_gui_data_t *)self->gui_data;
if(d->mode == EXPOSURE_MODE_DEFLICKER)
{
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;
if(g != NULL && self->dev->gui_attached && piece->pipe->type == DT_DEV_PIXELPIPE_PREVIEW)
{
g->deflicker_computed_exposure = d->exposure;
}
return TRUE;
error:
dt_print(DT_DEBUG_OPENCL, "[opencl_exposure] couldn't enqueue kernel! %d\n", err);
return FALSE;
}
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;
dt_iop_exposure_gui_data_t *g = (dt_iop_exposure_gui_data_t *)self->gui_data;
if(d->mode == EXPOSURE_MODE_DEFLICKER)
{
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;
if(g != NULL && self->dev->gui_attached && piece->pipe->type == DT_DEV_PIXELPIPE_PREVIEW)
{
g->deflicker_computed_exposure = d->exposure;
}
}
static float dt_iop_exposure_get_white(struct dt_iop_module_t *self)
{
dt_iop_exposure_params_t *p = (dt_iop_exposure_params_t *)self->params;
return exposure2white(p->exposure);
}
