74
74
return IOP_GROUP_CORRECT;
80
return IOP_FLAGS_SUPPORTS_BLENDING;
79
85
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)
81
87
dt_iop_sharpen_data_t *d = (dt_iop_sharpen_data_t *)piece->data;
82
88
dt_iop_sharpen_global_data_t *gd = (dt_iop_sharpen_global_data_t *)self->data;
85
92
const int devid = piece->pipe->devid;
86
93
const int rad = MIN(MAXR, ceilf(d->radius * roi_in->scale / piece->iscale));
89
97
size_t origin[] = {0, 0, 0};
90
98
size_t region[] = {roi_in->width, roi_in->height, 1};
91
err = clEnqueueCopyImage(darktable.opencl->dev[devid].cmd_queue, dev_in, dev_out, origin, origin, region, 0, NULL, NULL);
99
err = dt_opencl_enqueue_copy_image(darktable.opencl->dev[devid].cmd_queue, dev_in, dev_out, origin, origin, region, 0, NULL, NULL);
100
if (err != CL_SUCCESS) goto error;
103
// init gaussian kernel
94
104
float mat[2*(MAXR+1)];
95
105
const int wd = 2*rad+1;
96
106
float *m = mat + rad;
97
const float sigma2 = (2.5*2.5)*(d->radius*roi_in->scale/piece->iscale)*(d->radius*roi_in->scale/piece->iscale);
107
const float sigma2 = (1.0f/(2.5*2.5))*(d->radius*roi_in->scale/piece->iscale)*(d->radius*roi_in->scale/piece->iscale);
98
108
float weight = 0.0f;
99
// init gaussian kernel
100
109
for(int l=-rad; l<=rad; l++) weight += m[l] = expf(- (l*l)/(2.f*sigma2));
101
110
for(int l=-rad; l<=rad; l++) m[l] /= weight;
102
111
size_t sizes[] = {roi_in->width, roi_in->height, 1};
103
cl_mem dev_m = dt_opencl_copy_host_to_device_constant(sizeof(float)*wd, devid, mat);
112
dev_m = dt_opencl_copy_host_to_device_constant(sizeof(float)*wd, devid, mat);
113
if (dev_m == NULL) goto error;
104
114
dt_opencl_set_kernel_arg(darktable.opencl, devid, gd->kernel_sharpen, 0, sizeof(cl_mem), (void *)&dev_in);
105
115
dt_opencl_set_kernel_arg(darktable.opencl, devid, gd->kernel_sharpen, 1, sizeof(cl_mem), (void *)&dev_out);
106
116
dt_opencl_set_kernel_arg(darktable.opencl, devid, gd->kernel_sharpen, 2, sizeof(cl_mem), (void *)&dev_m);
108
118
dt_opencl_set_kernel_arg(darktable.opencl, devid, gd->kernel_sharpen, 4, sizeof(float), (void *)&d->amount);
109
119
dt_opencl_set_kernel_arg(darktable.opencl, devid, gd->kernel_sharpen, 5, sizeof(float), (void *)&d->threshold);
110
120
err = dt_opencl_enqueue_kernel_2d(darktable.opencl, devid, gd->kernel_sharpen, sizes);
111
if(err != CL_SUCCESS) fprintf(stderr, "couldn't enqueue sharpen kernel! %d\n", err);
112
clReleaseMemObject(dev_m);
121
if(err != CL_SUCCESS) goto error;
122
dt_opencl_release_mem_object(dev_m);
126
if (dev_m != NULL) dt_opencl_release_mem_object(dev_m);
127
dt_print(DT_DEBUG_OPENCL, "[opencl_sharpen] couldn't enqueue kernel! %d\n", err);
143
float *const tmp = dt_alloc_align(16, sizeof(float)*ch*roi_out->width*roi_out->height);
127
145
const int wd = 2*rad+1;
130
const float sigma2 = (2.5*2.5)*(data->radius*roi_in->scale/piece->iscale)*(data->radius*roi_in->scale/piece->iscale);
147
const float sigma2 = (1.0f/(2.5*2.5))*(data->radius*roi_in->scale/piece->iscale)*(data->radius*roi_in->scale/piece->iscale);
131
148
float weight = 0.0f;
132
150
// init gaussian kernel
134
for(int l=-rad; l<=rad; l++) for(int k=-rad; k<=rad; k++,m++)
135
weight += *m = expf(- (l*l + k*k)/(2.f*sigma2));
137
for(int l=-rad; l<=rad; l++) for(int k=-rad; k<=rad; k++,m++)
140
// gauss blur the image
142
#pragma omp parallel for default(none) shared(mat, ivoid, ovoid, roi_out, roi_in) schedule(static)
151
for(int l=-rad; l<=rad; l++)
152
weight += mat[l+rad] = expf(- l*l/(2.f*sigma2));
153
for(int l=0; l<wd; l++)
156
// gauss blur the image horizontally
158
#pragma omp parallel for default(none) shared(mat, ivoid, ovoid, roi_out, roi_in) schedule(static)
160
for(int j=0; j<roi_out->height; j++)
162
const float *in = ((float *)ivoid) + ch*(j*roi_in->width + rad);
163
float *out = tmp + ch*(j*roi_out->width + rad);
164
for(int i=rad; i<roi_out->width-rad; i++)
166
const float *inp = in - ch*rad;
167
const float *m = mat;
169
for(int k=-rad; k<=rad; k++,m++,inp+=ch)
177
// gauss blur the image horizontally
179
#pragma omp parallel for default(none) shared(mat, ivoid, ovoid, roi_out, roi_in) schedule(static)
144
181
for(int j=rad; j<roi_out->height-rad; j++)
146
float *in = ((float *)ivoid) + ch*(j*roi_in->width + rad);
183
const float *in = tmp + ch*(j*roi_in->width + rad);
147
184
float *out = ((float *)ovoid) + ch*(j*roi_out->width + rad);
148
185
for(int i=rad; i<roi_out->width-rad; i++)
187
const int step = ch*roi_in->width;
188
const float *inp = in - step*rad;
189
const float *m = mat;
151
190
float sum = 0.0f;
152
for(int l=-rad; l<=rad; l++)
154
float *inp = in + ch*(l*roi_in->width-rad);
155
for(int k=-rad; k<=rad; k++,m++,inp+=ch)
191
for(int k=-rad; k<=rad; k++,m++,inp+=step)
328
367
dt_iop_sharpen_params_t *p = (dt_iop_sharpen_params_t *)self->params;
330
369
self->widget = GTK_WIDGET(gtk_hbox_new(FALSE, 0));
331
g->vbox1 = GTK_VBOX(gtk_vbox_new(FALSE, DT_GUI_IOP_MODULE_CONTROL_SPACING));
332
g->vbox2 = GTK_VBOX(gtk_vbox_new(FALSE, DT_GUI_IOP_MODULE_CONTROL_SPACING));
333
gtk_box_pack_start(GTK_BOX(self->widget), GTK_WIDGET(g->vbox1), FALSE, FALSE, 5);
334
gtk_box_pack_start(GTK_BOX(self->widget), GTK_WIDGET(g->vbox2), TRUE, TRUE, 5);
338
widget = dtgtk_reset_label_new(_("radius"), self, &p->radius, sizeof(float));
339
gtk_box_pack_start(GTK_BOX(g->vbox1), widget, TRUE, TRUE, 0);
340
widget = dtgtk_reset_label_new(_("amount"), self, &p->amount, sizeof(float));
341
gtk_box_pack_start(GTK_BOX(g->vbox1), widget, TRUE, TRUE, 0);
342
widget = dtgtk_reset_label_new(_("threshold"), self, &p->threshold, sizeof(float));
343
gtk_box_pack_start(GTK_BOX(g->vbox1), widget, TRUE, TRUE, 0);
370
g->vbox = GTK_VBOX(gtk_vbox_new(FALSE, DT_GUI_IOP_MODULE_CONTROL_SPACING));
371
gtk_box_pack_start(GTK_BOX(self->widget), GTK_WIDGET(g->vbox), TRUE, TRUE, 5);
345
373
g->scale1 = DTGTK_SLIDER(dtgtk_slider_new_with_range(DARKTABLE_SLIDER_BAR,0.0, 8.0000, 0.100, p->radius, 3));
374
g_object_set (GTK_OBJECT(g->scale1), "tooltip-text", _("spatial extent of the unblurring"), (char *)NULL);
375
dtgtk_slider_set_label(g->scale1,_("radius"));
346
376
g->scale2 = DTGTK_SLIDER(dtgtk_slider_new_with_range(DARKTABLE_SLIDER_BAR,0.0, 2.0000, 0.010, p->amount, 3));
377
g_object_set (GTK_OBJECT(g->scale2), "tooltip-text", _("strength of the sharpen"), (char *)NULL);
378
dtgtk_slider_set_label(g->scale2,_("amount"));
347
379
g->scale3 = DTGTK_SLIDER(dtgtk_slider_new_with_range(DARKTABLE_SLIDER_BAR,0.0, 1.0000, 0.001, p->threshold, 3));
348
gtk_box_pack_start(GTK_BOX(g->vbox2), GTK_WIDGET(g->scale1), TRUE, TRUE, 0);
349
gtk_box_pack_start(GTK_BOX(g->vbox2), GTK_WIDGET(g->scale2), TRUE, TRUE, 0);
350
gtk_box_pack_start(GTK_BOX(g->vbox2), GTK_WIDGET(g->scale3), TRUE, TRUE, 0);
380
g_object_set (GTK_OBJECT(g->scale3), "tooltip-text", _("threshold to activate sharpen"), (char *)NULL);
381
dtgtk_slider_set_label(g->scale3,_("threshold"));
382
gtk_box_pack_start(GTK_BOX(g->vbox), GTK_WIDGET(g->scale1), TRUE, TRUE, 0);
383
gtk_box_pack_start(GTK_BOX(g->vbox), GTK_WIDGET(g->scale2), TRUE, TRUE, 0);
384
gtk_box_pack_start(GTK_BOX(g->vbox), GTK_WIDGET(g->scale3), TRUE, TRUE, 0);
352
386
g_signal_connect (G_OBJECT (g->scale1), "value-changed",
353
387
G_CALLBACK (radius_callback), self);