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; } }