Skip to content
Browse files

Merge branch 'master' into file_management

  • Loading branch information...
2 parents e565908 + 2ac1cbd commit b41517c7d5a1325fa909d567210bbe72a5ae4c89 @boucman boucman committed May 11, 2012
View
143 data/kernels/extended.cl
@@ -418,5 +418,148 @@ splittoning (read_only image2d_t in, write_only image2d_t out, const int width,
write_imagef (out, (int2)(x, y), pixel);
}
+/* kernels to get the maximum value of an image */
+kernel void
+pixelmax_first (read_only image2d_t in, const int width, const int height, global float *accu, local float *buffer)
+{
+ const int x = get_global_id(0);
+ const int y = get_global_id(1);
+ const int xlsz = get_local_size(0);
+ const int ylsz = get_local_size(1);
+ const int xlid = get_local_id(0);
+ const int ylid = get_local_id(1);
+
+ const int l = ylid * xlsz + xlid;
+
+ buffer[l] = (x < width && y < height) ? read_imagef(in, sampleri, (int2)(x, y)).x : -INFINITY;
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ const int lsz = xlsz * ylsz;
+
+ for(int offset = lsz / 2; offset > 0; offset = offset / 2)
+ {
+ if (l < offset)
+ {
+ float other = buffer[l + offset];
+ float mine = buffer[l];
+ buffer[l] = (mine > other) ? mine : other;
+ }
+ barrier(CLK_LOCAL_MEM_FENCE);
+ }
+
+ const int xgid = get_group_id(0);
+ const int ygid = get_group_id(1);
+ const int xgsz = get_num_groups(0);
+ const int ygsz = get_num_groups(1);
+
+ const int m = ygid * xgsz + xgid;
+ accu[m] = buffer[0];
+}
+
+
+
+__kernel void
+pixelmax_second(global float* input, global float *result, const int length, local float *buffer)
+{
+ int x = get_global_id(0);
+ float accu = -INFINITY;
+
+ while (x < length)
+ {
+ float element = input[x];
+ accu = (accu > element) ? accu : element;
+ x += get_global_size(0);
+ }
+
+ int lid = get_local_id(0);
+ buffer[lid] = accu;
+
+ barrier(CLK_LOCAL_MEM_FENCE);
+
+ for(int offset = get_local_size(0) / 2; offset > 0; offset = offset / 2)
+ {
+ if (lid < offset)
+ {
+ float other = buffer[lid + offset];
+ float mine = buffer[lid];
+ buffer[lid] = (mine > other) ? mine : other;
+ }
+ barrier(CLK_LOCAL_MEM_FENCE);
+ }
+
+ if (lid == 0)
+ {
+ result[get_group_id(0)] = buffer[0];
+ }
+}
+
+
+/* kernel for the global tonemap plugin: reinhard */
+kernel void
+global_tonemap_reinhard (read_only image2d_t in, write_only image2d_t out, const int width, const int height,
+ const float4 parameters)
+{
+ const int x = get_global_id(0);
+ const int y = get_global_id(1);
+
+ if(x >= width || y >= height) return;
+
+ float4 pixel = read_imagef(in, sampleri, (int2)(x, y));
+
+ float l = pixel.x * 0.01f;
+
+ pixel.x = 100.0f * (l/(1.0f + l));
+
+ write_imagef (out, (int2)(x, y), pixel);
+}
+
+
+
+/* kernel for the global tonemap plugin: drago */
+kernel void
+global_tonemap_drago (read_only image2d_t in, write_only image2d_t out, const int width, const int height,
+ const float4 parameters)
+{
+ const int x = get_global_id(0);
+ const int y = get_global_id(1);
+
+ if(x >= width || y >= height) return;
+
+ const float eps = parameters.x;
+ const float ldc = parameters.y;
+ const float bl = parameters.z;
+ const float lwmax = parameters.w;
+
+ float4 pixel = read_imagef(in, sampleri, (int2)(x, y));
+
+ float lw = pixel.x * 0.01f;
+
+ pixel.x = 100.0f * (ldc * log(fmax(eps, lw + 1.0f)) / log(fmax(eps, 2.0f + (pow(lw/lwmax,bl)) * 8.0f)));
+
+ write_imagef (out, (int2)(x, y), pixel);
+}
+
+
+/* kernel for the global tonemap plugin: filmic */
+kernel void
+global_tonemap_filmic (read_only image2d_t in, write_only image2d_t out, const int width, const int height,
+ const float4 parameters)
+{
+ const int x = get_global_id(0);
+ const int y = get_global_id(1);
+
+ if(x >= width || y >= height) return;
+
+ float4 pixel = read_imagef(in, sampleri, (int2)(x, y));
+
+ float l = pixel.x * 0.01f;
+ float m = fmax(0.0f, l - 0.004f);
+
+ pixel.x = 100.0f * ((m*(6.2f*m+.5f))/(m*(6.2f*m+1.7f)+0.06f));
+
+ write_imagef (out, (int2)(x, y), pixel);
+}
+
View
3 src/common/colormatrices.c
@@ -150,6 +150,9 @@ static dt_profiled_colormatrix_t dt_profiled_colormatrices[] =
// Robert Park, ColorChecker Passport, strobe, well lit
{ "NIKON D700", { 789261, 332016, 34149}, { 270386, 985748, -129135}, { 4074, -230209, 999008}, {798172, 826721, 673126}},
+ // Mauro Fuentes, ColorChecker Passport, direct sunlight, well lit
+ { "NIKON D800", { 792038, 268860, 33951}, { 289093, 1169876, -251740}, {-32654, -340393, 1127960}, {782806, 804443, 659058}},
+
// Robert Park, ColorChecker Passport, strobe, well lit
{ "NIKON COOLPIX P7000", { 804947, 229630, 97717}, { 178146, 1138763, -395233}, { 88699, -282013, 1234650}, {809998, 842819, 682144}},
View
128 src/common/curve_tools.c
@@ -33,14 +33,12 @@ float spline_cubic_val( int n, float t[], float tval, float y[],
float ypp[]);
float catmull_rom_val ( int n, float x[], float xval, float y[],
float tangents[]);
-float monotone_hermite_val ( int n, float x[], float xval, float y[],
- float tangents[]);
float *spline_cubic_set( int n, float t[], float y[]);
float *catmull_rom_set ( int n, float x[], float y[]);
float *monotone_hermite_set ( int n, float x[], float y[]);
-float (*spline_val[])(int, float [], float, float [], float []) = {spline_cubic_val, catmull_rom_val, monotone_hermite_val };
+float (*spline_val[])(int, float [], float, float [], float []) = {spline_cubic_val, catmull_rom_val, catmull_rom_val };
float *(*spline_set[])(int, float [], float []) = {spline_cubic_set, catmull_rom_set, monotone_hermite_set};
@@ -419,40 +417,42 @@ float *monotone_hermite_set ( int n, float x[], float y[])
return NULL;
}
}
- delta = (float *)calloc(n-1,sizeof(float));
+
+ delta = (float *)calloc(n,sizeof(float));
//nc_merror(delta, "spline_cubic_set");
- m = (float *)calloc(n-1,sizeof(float));
+ m = (float *)calloc(n,sizeof(float));
//nc_merror(m, "spline_cubic_set");
//calculate the slopes
for (i = 0; i<n-1; i++)
{
delta[i] = (y[i+1]-y[i])/(x[i+1]-x[i]);
}
+ delta[n-1] = delta[n-2];
m[0] = delta[0];
- m[n-2] = delta[n-2];
+ m[n-1] = delta[n-1];
- for (i=1; i<n-2; i++)
+ for (i=1; i<n-1; i++)
{
- m[i] = (delta[i-1]+delta[i])/2.0;
+ m[i] = (delta[i-1]+delta[i])*.5f;
}
- for (i=0; i<n-1; i++)
+ for (i=0; i<n; i++)
{
if (fabs(delta[i]) < EPSILON)
{
- m[i] = 0.0;
- m[i+1] = 0.0;
+ m[i] = 0.0f;
+ m[i+1] = 0.0f;
i++;
}
else
{
alpha = m[i]/delta[i];
beta = m[i+1]/delta[i];
tau = alpha*alpha+beta*beta;
- if (tau > 9.0)
+ if (tau > 9.0f)
{
- m[i] = 3.0*alpha*delta[i]/sqrtf(tau);
- m[i+1] = 3.0*beta*delta[i]/sqrtf(tau);
+ m[i] = 3.0f*alpha*delta[i]/sqrtf(tau);
+ m[i+1] = 3.0f*beta*delta[i]/sqrtf(tau);
i++;
}
}
@@ -496,16 +496,16 @@ float *catmull_rom_set ( int n, float x[], float y[])
}
}
//nc_merror(delta, "spline_cubic_set");
- m = (float *)calloc(n-1,sizeof(float));
+ m = (float *)calloc(n,sizeof(float));
//nc_merror(m, "spline_cubic_set");
//calculate the slopes
- m[0] = (y[1]-y[0]);///(x[1]-x[0]);
+ m[0] = (y[1]-y[0])/(x[1]-x[0]);
for (i=1; i<n-1; i++)
{
- m[i] = (y[i+1]-y[i-1])/2.0;//*(x[i+1]-x[i-1]));
+ m[i] = (y[i+1]-y[i-1])/(x[i+1]-x[i-1]);
}
- m[n-1] = (y[n-1]-y[n-2]);//(x[n-1]-y[n-2]);
+ m[n-1] = (y[n-1]-y[n-2])/(x[n-1]-x[n-2]);
return m;
}
@@ -537,20 +537,13 @@ float interpolate_val( int n, float x[], float xval, float y[], float tangents[]
float catmull_rom_val ( int n, float x[], float xval, float y[],
float tangents[])
{
- float dx;
- float h;
- int i;
- int ival;
- float yval;
- float h00,h01,h10,h11;
- float m0, m1;
//
// Determine the interval [ T(I), T(I+1) ] that contains TVAL.
// Values below T[0] or above T[N-1] use extrapolation.
//
- ival = n - 2;
+ int ival = n - 2;
- for ( i = 0; i < n-2; i++ )
+ for ( int i = 0; i < n-2; i++ )
{
if ( xval < x[i+1] )
{
@@ -559,82 +552,23 @@ float catmull_rom_val ( int n, float x[], float xval, float y[],
}
}
- m0 = tangents[ival];
- m1 = tangents[ival+1];
+ const float m0 = tangents[ival];
+ const float m1 = tangents[ival+1];
//
// In the interval I, the polynomial is in terms of a normalized
// coordinate between 0 and 1.
//
- h = x[ival+1] - x[ival];
- dx = (xval-x[ival])/h;
+ const float h = x[ival+1] - x[ival];
+ const float dx = (xval-x[ival])/h;
+ const float dx2 = dx*dx;
+ const float dx3 = dx*dx2;
- h00 = ( 2.0 * dx*dx*dx) - (3.0 * dx*dx) + 1.0;
- h10 = ( 1.0 * dx*dx*dx) - (2.0 * dx*dx) + dx;
- h01 = (-2.0 * dx*dx*dx) + (3.0 * dx*dx);
- h11 = ( 1.0 * dx*dx*dx) - (1.0 * dx*dx);
+ const float h00 = ( 2.0 * dx3) - (3.0 * dx2) + 1.0;
+ const float h10 = ( 1.0 * dx3) - (2.0 * dx2) + dx;
+ const float h01 = (-2.0 * dx3) + (3.0 * dx2);
+ const float h11 = ( 1.0 * dx3) - (1.0 * dx2);
- h = 1;
- yval = (h00 * y[ival]) + (h10 * h * m0) + (h01 * y[ival+1]) + (h11 * h * m1);
-
- return yval;
-}
-
-/*************************************************************
- * monotone_hermite_val:
- * piecewise monotone hermite spline interpolation
- *
- * n = number of control points
- * x = input x array
- * xval = input value where to interpolate the data
- * y = input y array
- * tangent = input array of tangents
- * output:
- * interpolated value at xval
- *
- *************************************************************/
-float monotone_hermite_val ( int n, float x[], float xval, float y[],
- float tangents[])
-{
- float dx;
- float h;
- int i;
- int ival;
- float yval;
- float h00,h01,h10,h11;
- float m0, m1;
-//
-// Determine the interval [ T(I), T(I+1) ] that contains TVAL.
-// Values below T[0] or above T[N-1] use extrapolation.
-//
- ival = n - 2;
-
- for ( i = 0; i < n-2; i++ )
- {
- if ( xval < x[i+1] )
- {
- ival = i;
- break;
- }
- }
-
- m0 = tangents[ival];
- m1 = tangents[ival+1];
-//
-// In the interval I, the polynomial is in terms of a normalized
-// coordinate between 0 and 1.
-//
- h = x[ival+1] - x[ival];
- dx = (xval-x[ival])/h;
-
- h00 = ( 2.0 * dx*dx*dx) - (3.0 * dx*dx) + 1.0;
- h10 = ( 1.0 * dx*dx*dx) - (2.0 * dx*dx) + dx;
- h01 = (-2.0 * dx*dx*dx) + (3.0 * dx*dx);
- h11 = ( 1.0 * dx*dx*dx) - (1.0 * dx*dx);
-
- //h = 1;
- yval = (h00 * y[ival]) + (h10 * h * m0) + (h01 * y[ival+1]) + (h11 * h * m1);
-
- return yval;
+ return (h00 * y[ival]) + (h10 * h * m0) + (h01 * y[ival+1]) + (h11 * h * m1);
}
View
1 src/develop/pixelpipe_hb.c
@@ -77,6 +77,7 @@ void dt_dev_pixelpipe_set_input(dt_dev_pixelpipe_t *pipe, dt_develop_t *dev, flo
{
pipe->iwidth = width;
pipe->iheight = height;
+ pipe->iflipped = 0;
pipe->iscale = iscale;
pipe->input = input;
pipe->image = dev->image_storage;
View
2 src/develop/pixelpipe_hb.h
@@ -89,6 +89,8 @@ typedef struct dt_dev_pixelpipe_t
float *input;
// width and height of input buffer
int iwidth, iheight;
+ // is image flipped?
+ int iflipped;
// input actually just downscaled buffer? iscale*iwidth = actual width
float iscale;
// dimensions of processed buffer
View
10 src/iop/clipping.c
@@ -369,10 +369,12 @@ void modify_roi_in(struct dt_iop_module_t *self, struct dt_dev_pixelpipe_iop_t *
}
// sanity check.
- roi_in->x = CLAMP(roi_in->x, 0, piece->iwidth);
- roi_in->y = CLAMP(roi_in->y, 0, piece->iheight);
- roi_in->width = CLAMP(roi_in->width, 1, piece->iwidth - roi_in->x);
- roi_in->height = CLAMP(roi_in->height, 1, piece->iheight - roi_in->y);
+ const int scwidth = piece->pipe->iflipped ? piece->pipe->iheight : piece->pipe->iwidth;
+ const int scheight = piece->pipe->iflipped ? piece->pipe->iwidth : piece->pipe->iheight;
+ roi_in->x = CLAMP(roi_in->x, 0, scwidth);
+ roi_in->y = CLAMP(roi_in->y, 0, scheight);
+ roi_in->width = CLAMP(roi_in->width, 1, scwidth - roi_in->x);
+ roi_in->height = CLAMP(roi_in->height, 1, scheight - roi_in->y);
}
// 3rd (final) pass: you get this input region (may be different from what was requested above),
View
8 src/iop/colorcorrection.c
@@ -95,6 +95,14 @@ void init_presets (dt_iop_module_so_t *self)
p.lob = 0.0f;
p.saturation = 1.0f;
dt_gui_presets_add_generic(_("warming filter"), self->op, self->version(), &p, sizeof(p), 1);
+
+ p.hia = 0.95f;
+ p.loa = -3.55f;
+ p.hib = -4.5f;
+ p.lob = -0.0f;
+ p.saturation = 1.0f;
+ dt_gui_presets_add_generic(_("cooling filter"), self->op, self->version(), &p, sizeof(p), 1);
+
}
void init_key_accels(dt_iop_module_so_t *self)
View
9 src/iop/flip.c
@@ -168,10 +168,10 @@ void modify_roi_in(struct dt_iop_module_t *self, struct dt_dev_pixelpipe_iop_t *
roi_in->height = aabb_in[3]-aabb_in[1]+1;
// sanity check.
- roi_in->x = CLAMP(roi_in->x, 0, piece->iwidth);
- roi_in->y = CLAMP(roi_in->y, 0, piece->iheight);
- roi_in->width = CLAMP(roi_in->width, 1, piece->iwidth - roi_in->x);
- roi_in->height = CLAMP(roi_in->height, 1, piece->iheight - roi_in->y);
+ roi_in->x = CLAMP(roi_in->x, 0, piece->pipe->iwidth);
+ roi_in->y = CLAMP(roi_in->y, 0, piece->pipe->iheight);
+ roi_in->width = CLAMP(roi_in->width, 1, piece->pipe->iwidth - roi_in->x);
+ roi_in->height = CLAMP(roi_in->height, 1, piece->pipe->iheight - roi_in->y);
}
// 3rd (final) pass: you get this input region (may be different from what was requested above),
@@ -241,6 +241,7 @@ void commit_params (struct dt_iop_module_t *self, dt_iop_params_t *p1, dt_dev_pi
dt_iop_flip_params_t *p = (dt_iop_flip_params_t *)p1;
dt_iop_flip_data_t *d = (dt_iop_flip_data_t *)piece->data;
d->orientation = p->orientation;
+ piece->pipe->iflipped = d->orientation & 4;
}
void init_pipe (struct dt_iop_module_t *self, dt_dev_pixelpipe_t *pipe, dt_dev_pixelpipe_iop_t *piece)
View
183 src/iop/globaltonemap.c
@@ -29,12 +29,16 @@
#include "develop/develop.h"
#include "develop/imageop.h"
#include "control/control.h"
+#include "common/opencl.h"
#include "gui/accelerators.h"
#include "gui/gtk.h"
#include <gtk/gtk.h>
#include <inttypes.h>
#include <xmmintrin.h>
+#define BLOCKSIZE 2048 /* maximum blocksize. must be a power of 2 and will be automatically reduced if needed */
+#define REDUCESIZE 64
+
// NaN-safe clip: NaN compares false and will result in 0.0
#define CLIP(x) (((x)>=0.0)?((x)<=1.0?(x):1.0):0.0)
DT_MODULE(2)
@@ -78,14 +82,26 @@ typedef struct dt_iop_global_tonemap_data_t
}
dt_iop_global_tonemap_data_t;
+typedef struct dt_iop_global_tonemap_global_data_t
+{
+ int kernel_pixelmax_first;
+ int kernel_pixelmax_second;
+ int kernel_global_tonemap_reinhard;
+ int kernel_global_tonemap_drago;
+ int kernel_global_tonemap_filmic;
+}
+dt_iop_global_tonemap_global_data_t;
+
const char *name()
{
return _("global tonemap");
}
int flags()
{
- return IOP_FLAGS_INCLUDE_IN_STYLES | IOP_FLAGS_SUPPORTS_BLENDING;
+ return IOP_FLAGS_INCLUDE_IN_STYLES | IOP_FLAGS_SUPPORTS_BLENDING; // we can not allow tiling at the moment. drago requires to get the absolute maximum l-value
+ // of the image. tiling conflicts with this need. this implies that opencl will only be used
+ // in darkroom mode, not during export of larger images.
}
int
@@ -199,6 +215,171 @@ void process (struct dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, void
}
}
+#ifdef HAVE_OPENCL
+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_global_tonemap_data_t *d = (dt_iop_global_tonemap_data_t *)piece->data;
+ dt_iop_global_tonemap_global_data_t *gd = (dt_iop_global_tonemap_global_data_t *)self->data;
+
+ cl_int err = -999;
+ cl_mem dev_m = NULL;
+ cl_mem dev_r = NULL;
+ const int devid = piece->pipe->devid;
+ int gtkernel = -1;
+
+ const int width = roi_out->width;
+ const int height = roi_out->height;
+ float parameters[4] = { 0.0f };
+
+ // prepare local work group
+ size_t maxsizes[3] = { 0 }; // the maximum dimensions for a work group
+ size_t workgroupsize = 0; // the maximum number of items in a work group
+ unsigned long localmemsize = 0; // the maximum amount of local memory we can use
+ size_t kernelworkgroupsize = 0; // the maximum amount of items in work group for this kernel
+
+ // make sure blocksize is not too large
+ int blocksize = BLOCKSIZE;
+ if(dt_opencl_get_work_group_limits(devid, maxsizes, &workgroupsize, &localmemsize) == CL_SUCCESS &&
+ dt_opencl_get_kernel_work_group_size(devid, gd->kernel_pixelmax_first, &kernelworkgroupsize) == CL_SUCCESS)
+ {
+ // reduce blocksize step by step until it fits to limits
+ while(blocksize > maxsizes[0] || blocksize > maxsizes[1] || blocksize*blocksize > kernelworkgroupsize
+ || blocksize*blocksize > workgroupsize || blocksize*blocksize*sizeof(float) > localmemsize)
+ {
+ if(blocksize == 1) break;
+ blocksize >>= 1;
+ }
+ }
+ else
+ {
+ blocksize = 1; // slow but safe
+ }
+
+
+
+ // width and height of intermediate buffers. Need to be multiples of BLOCKSIZE
+ const size_t bwidth = width % blocksize == 0 ? width : (width / blocksize + 1)*blocksize;
+ const size_t bheight = height % blocksize == 0 ? height : (height / blocksize + 1)*blocksize;
+
+ switch(d->operator) {
+ case OPERATOR_REINHARD:
+ gtkernel = gd->kernel_global_tonemap_reinhard;
+ break;
+ case OPERATOR_DRAGO:
+ gtkernel = gd->kernel_global_tonemap_drago;
+ break;
+ case OPERATOR_FILMIC:
+ gtkernel = gd->kernel_global_tonemap_filmic;
+ break;
+ }
+
+ if(d->operator == OPERATOR_DRAGO)
+ {
+ size_t sizes[3];
+ size_t local[3];
+
+ int bufsize = (bwidth / blocksize) * (bheight / blocksize);
+ int reducesize = maxsizes[0] < REDUCESIZE ? maxsizes[0] : REDUCESIZE;
+
+ dev_m = dt_opencl_alloc_device_buffer(devid, bufsize*sizeof(float));
+ if(dev_m == NULL) goto error;
+
+ dev_r = dt_opencl_alloc_device_buffer(devid, reducesize*sizeof(float));
+ if(dev_r == NULL) goto error;
+
+ sizes[0] = bwidth;
+ sizes[1] = bheight;
+ sizes[2] = 1;
+ local[0] = blocksize;
+ local[1] = blocksize;
+ local[2] = 1;
+ dt_opencl_set_kernel_arg(devid, gd->kernel_pixelmax_first, 0, sizeof(cl_mem), &dev_in);
+ dt_opencl_set_kernel_arg(devid, gd->kernel_pixelmax_first, 1, sizeof(int), &width);
+ dt_opencl_set_kernel_arg(devid, gd->kernel_pixelmax_first, 2, sizeof(int), &height);
+ dt_opencl_set_kernel_arg(devid, gd->kernel_pixelmax_first, 3, sizeof(cl_mem), &dev_m);
+ dt_opencl_set_kernel_arg(devid, gd->kernel_pixelmax_first, 4, blocksize*blocksize*sizeof(float), NULL);
+ err = dt_opencl_enqueue_kernel_2d_with_local(devid, gd->kernel_pixelmax_first, sizes, local);
+ if(err != CL_SUCCESS) goto error;
+
+
+ sizes[0] = reducesize;
+ sizes[1] = 1;
+ sizes[2] = 1;
+ local[0] = reducesize;
+ local[1] = 1;
+ local[2] = 1;
+ dt_opencl_set_kernel_arg(devid, gd->kernel_pixelmax_second, 0, sizeof(cl_mem), &dev_m);
+ dt_opencl_set_kernel_arg(devid, gd->kernel_pixelmax_second, 1, sizeof(cl_mem), &dev_r);
+ dt_opencl_set_kernel_arg(devid, gd->kernel_pixelmax_second, 2, sizeof(int), &bufsize);
+ dt_opencl_set_kernel_arg(devid, gd->kernel_pixelmax_second, 3, reducesize*sizeof(float), NULL);
+ err = dt_opencl_enqueue_kernel_2d_with_local(devid, gd->kernel_pixelmax_second, sizes, local);
+ if(err != CL_SUCCESS) goto error;
+
+ float maximum[reducesize];
+ err = dt_opencl_read_buffer_from_device(devid, (void*)maximum, dev_r, 0, reducesize*sizeof(float), CL_TRUE);
+ if(err != CL_SUCCESS) goto error;
+
+ dt_opencl_release_mem_object(dev_r);
+ dt_opencl_release_mem_object(dev_m);
+
+ const float eps = 0.0001f;
+ const float lwmax = MAX(eps, (maximum[0]*0.01f));
+ const float ldc = d->drago.max_light * 0.01f / log10f(lwmax + 1.0f);
+ const float bl = logf(MAX(eps, d->drago.bias)) / logf(0.5f);
+
+ parameters[0] = eps;
+ parameters[1] = ldc;
+ parameters[2] = bl;
+ parameters[3] = lwmax;
+ }
+
+ size_t sizes[2] = { ROUNDUPWD(width), ROUNDUPHT(height) };
+ dt_opencl_set_kernel_arg(devid, gtkernel, 0, sizeof(cl_mem), &dev_in);
+ dt_opencl_set_kernel_arg(devid, gtkernel, 1, sizeof(cl_mem), &dev_out);
+ dt_opencl_set_kernel_arg(devid, gtkernel, 2, sizeof(int), &width);
+ dt_opencl_set_kernel_arg(devid, gtkernel, 3, sizeof(int), &height);
+ dt_opencl_set_kernel_arg(devid, gtkernel, 4, 4*sizeof(float), &parameters);
+ err = dt_opencl_enqueue_kernel_2d(devid, gtkernel, sizes);
+ if(err != CL_SUCCESS) goto error;
+ return TRUE;
+
+error:
+ if(dev_m != NULL) dt_opencl_release_mem_object(dev_m);
+ if(dev_r != NULL) dt_opencl_release_mem_object(dev_r);
+ dt_print(DT_DEBUG_OPENCL, "[opencl_global_tonemap] couldn't enqueue kernel! %d\n", err);
+ return FALSE;
+}
+#endif
+
+
+void init_global(dt_iop_module_so_t *module)
+{
+ const int program = 8; // extended.cl from programs.conf
+ dt_iop_global_tonemap_global_data_t *gd = (dt_iop_global_tonemap_global_data_t *)malloc(sizeof(dt_iop_global_tonemap_global_data_t));
+ module->data = gd;
+ gd->kernel_pixelmax_first = dt_opencl_create_kernel(program, "pixelmax_first");
+ gd->kernel_pixelmax_second = dt_opencl_create_kernel(program, "pixelmax_second");
+ gd->kernel_global_tonemap_reinhard = dt_opencl_create_kernel(program, "global_tonemap_reinhard");
+ gd->kernel_global_tonemap_drago = dt_opencl_create_kernel(program, "global_tonemap_drago");
+ gd->kernel_global_tonemap_filmic = dt_opencl_create_kernel(program, "global_tonemap_filmic");
+}
+
+
+void cleanup_global(dt_iop_module_so_t *module)
+{
+ dt_iop_global_tonemap_global_data_t *gd = (dt_iop_global_tonemap_global_data_t *)module->data;
+ dt_opencl_free_kernel(gd->kernel_pixelmax_first);
+ dt_opencl_free_kernel(gd->kernel_pixelmax_second);
+ dt_opencl_free_kernel(gd->kernel_global_tonemap_reinhard);
+ dt_opencl_free_kernel(gd->kernel_global_tonemap_drago);
+ dt_opencl_free_kernel(gd->kernel_global_tonemap_filmic);
+ free(module->data);
+ module->data = NULL;
+}
+
+
+
static void
operator_callback (GtkWidget *combobox, gpointer user_data)
{
View
7 src/iop/tonecurve.c
@@ -334,8 +334,8 @@ void init(dt_iop_module_t *module)
}
},
{ 2, 3, 3 }, // number of nodes per curve
- { CATMULL_ROM, CATMULL_ROM, CATMULL_ROM}, // curve types
- // { MONOTONE_HERMITE, MONOTONE_HERMITE, MONOTONE_HERMITE}, // seems broken (last tangent is off)
+ // { CATMULL_ROM, CATMULL_ROM, CATMULL_ROM}, // curve types
+ { MONOTONE_HERMITE, MONOTONE_HERMITE, MONOTONE_HERMITE},
// { CUBIC_SPLINE, CUBIC_SPLINE, CUBIC_SPLINE},
1, // autoscale_ab
0
@@ -789,12 +789,13 @@ static gboolean dt_iop_tonecurve_motion_notify(GtkWidget *widget, GdkEventMotion
c->selected = 0;
else for(int k=1; k<nodes; k++)
{
- if(tonecurve[k].x > mx || k == nodes-1)
+ if(tonecurve[k].x > mx)
{
c->selected = k;
break;
}
}
+ if(c->selected == -1) c->selected = nodes;
for(int i=nodes; i>c->selected; i--)
{
tonecurve[i].x = tonecurve[i-1].x;

0 comments on commit b41517c

Please sign in to comment.
Something went wrong with that request. Please try again.