17 const uint32_t filters,
const uint8_t (*
const xtrans)[6],
const int only_vng_linear)
19 static const signed char terms[]
20 = { -2, -2, +0, -1, 1, 0x01, -2, -2, +0, +0, 2, 0x01, -2, -1, -1, +0, 1, 0x01, -2, -1, +0, -1, 1, 0x02,
21 -2, -1, +0, +0, 1, 0x03, -2, -1, +0, +1, 2, 0x01, -2, +0, +0, -1, 1, 0x06, -2, +0, +0, +0, 2, 0x02,
22 -2, +0, +0, +1, 1, 0x03, -2, +1, -1, +0, 1, 0x04, -2, +1, +0, -1, 2, 0x04, -2, +1, +0, +0, 1, 0x06,
23 -2, +1, +0, +1, 1, 0x02, -2, +2, +0, +0, 2, 0x04, -2, +2, +0, +1, 1, 0x04, -1, -2, -1, +0, 1, 0x80,
24 -1, -2, +0, -1, 1, 0x01, -1, -2, +1, -1, 1, 0x01, -1, -2, +1, +0, 2, 0x01, -1, -1, -1, +1, 1, 0x88,
25 -1, -1, +1, -2, 1, 0x40, -1, -1, +1, -1, 1, 0x22, -1, -1, +1, +0, 1, 0x33, -1, -1, +1, +1, 2, 0x11,
26 -1, +0, -1, +2, 1, 0x08, -1, +0, +0, -1, 1, 0x44, -1, +0, +0, +1, 1, 0x11, -1, +0, +1, -2, 2, 0x40,
27 -1, +0, +1, -1, 1, 0x66, -1, +0, +1, +0, 2, 0x22, -1, +0, +1, +1, 1, 0x33, -1, +0, +1, +2, 2, 0x10,
28 -1, +1, +1, -1, 2, 0x44, -1, +1, +1, +0, 1, 0x66, -1, +1, +1, +1, 1, 0x22, -1, +1, +1, +2, 1, 0x10,
29 -1, +2, +0, +1, 1, 0x04, -1, +2, +1, +0, 2, 0x04, -1, +2, +1, +1, 1, 0x04, +0, -2, +0, +0, 2, 0x80,
30 +0, -1, +0, +1, 2, 0x88, +0, -1, +1, -2, 1, 0x40, +0, -1, +1, +0, 1, 0x11, +0, -1, +2, -2, 1, 0x40,
31 +0, -1, +2, -1, 1, 0x20, +0, -1, +2, +0, 1, 0x30, +0, -1, +2, +1, 2, 0x10, +0, +0, +0, +2, 2, 0x08,
32 +0, +0, +2, -2, 2, 0x40, +0, +0, +2, -1, 1, 0x60, +0, +0, +2, +0, 2, 0x20, +0, +0, +2, +1, 1, 0x30,
33 +0, +0, +2, +2, 2, 0x10, +0, +1, +1, +0, 1, 0x44, +0, +1, +1, +2, 1, 0x10, +0, +1, +2, -1, 2, 0x40,
34 +0, +1, +2, +0, 1, 0x60, +0, +1, +2, +1, 1, 0x20, +0, +1, +2, +2, 1, 0x10, +1, -2, +1, +0, 1, 0x80,
35 +1, -1, +1, +1, 1, 0x88, +1, +0, +1, +2, 1, 0x08, +1, +0, +2, -1, 1, 0x40, +1, +0, +2, +1, 1, 0x10 };
36 static const signed char chood[]
37 = { -1, -1, -1, 0, -1, +1, 0, +1, +1, +1, +1, 0, +1, -1, 0, -1 };
38 int *ip, *code[16][16];
43 const int prow = (filters == 9) ? 6 : 8;
44 const int pcol = (filters == 9) ? 6 : 2;
45 const int colors = (filters == 9) ? 3 : 4;
48 uint32_t filters4 = filters;
51 else if((filters & 3) == 1)
52 filters4 = filters | 0x03030303u;
54 filters4 = filters | 0x0c0c0c0cu;
59 if(only_vng_linear)
return;
65 fprintf(stderr,
"[demosaic] not able to allocate VNG buffer\n");
68 for(
int row = 0; row < 3; row++) brow[row] = (
float(*)[4])buffer + row *
width;
69 ip = (
int *)(buffer +
sizeof(**brow) *
width * 3);
71 for(
int row = 0; row < prow; row++)
72 for(
int col = 0; col < pcol; col++)
75 const signed char *cp = terms;
76 for(
int t = 0; t < 64; t++)
78 const int y1 = *cp++, x1 = *cp++;
79 const int y2 = *cp++, x2 = *cp++;
81 const int grads = *cp++;
82 const int color =
fcol(row + y1, col + x1, filters4, xtrans);
83 if(
fcol(row + y2, col + x2, filters4, xtrans) != color)
continue;
85 = (
fcol(row, col + 1, filters4, xtrans) == color &&
fcol(row + 1, col, filters4, xtrans) == color)
88 if(abs(y1 - y2) == diag && abs(x1 - x2) == diag)
continue;
89 *ip++ = (y1 *
width + x1) * 4 + color;
90 *ip++ = (y2 *
width + x2) * 4 + color;
92 for(
int g = 0; g < 8; g++)
93 if(grads & 1 << g) *ip++ = g;
98 for(
int g = 0; g < 8; g++)
100 const int y = *cp++, x = *cp++;
101 *ip++ = (y *
width + x) * 4;
102 const int color =
fcol(row, col, filters4, xtrans);
103 if(
fcol(row + y, col + x, filters4, xtrans) != color
104 &&
fcol(row + y * 2, col + x * 2, filters4, xtrans) == color)
105 *ip++ = (y *
width + x) * 8 + color;
111 for(
int row = 2; row <
height - 2; row++)
114#pragma omp parallel for default(none) \
115 dt_omp_firstprivate(colors, pcol, prow, roi_in, width, xtrans) \
116 shared(row, code, brow, out, filters4) \
120 for(
int col = 2; col <
width - 2; col++)
123 float gval[8] = { 0.0f };
124 float *pix = out + 4 * (row *
width + col);
125 ip = code[(row + roi_in->
y) % prow][(col + roi_in->
x) % pcol];
126 while((g = ip[0]) != INT_MAX)
128 float diff = fabsf(pix[g] - pix[ip[1]]) * ip[2];
131 if((g = ip[-1]) == -1)
continue;
133 while((g = *ip++) != -1) gval[g] += diff;
136 float gmin = gval[0], gmax = gval[0];
137 for(g = 1; g < 8; g++)
139 if(gmin > gval[g]) gmin = gval[g];
140 if(gmax < gval[g]) gmax = gval[g];
144 memcpy(brow[2][col], pix,
sizeof(*out) * 4);
147 const float thold = gmin + (gmax * 0.5f);
148 dt_aligned_pixel_t sum = { 0.0f };
149 const int color =
fcol(row + roi_in->
y, col + roi_in->
x, filters4, xtrans);
151 for(g = 0; g < 8; g++, ip += 2)
155 for(
int c = 0; c <
colors; c++)
156 if(c == color && ip[1])
157 sum[c] += (pix[c] + pix[ip[1]]) * 0.5f;
159 sum[c] += pix[ip[0] + c];
163 for(
int c = 0; c <
colors; c++)
165 float tot = pix[color];
166 if(c != color) tot += (sum[c] - sum[color]) / num;
167 brow[2][col][c] = tot;
171 memcpy(out + 4 * ((row - 2) *
width + 2), brow[0] + 2,
sizeof(*out) * 4 * (
width - 4));
173 for(
int g = 0; g < 4; g++) brow[(g - 1) & 3] = brow[g];
176 memcpy(out + (4 * ((
height - 4) *
width + 2)), brow[0] + 2,
sizeof(*out) * 4 * (
width - 4));
177 memcpy(out + (4 * ((
height - 3) *
width + 2)), brow[1] + 2,
sizeof(*out) * 4 * (
width - 4));
183#pragma omp parallel for default(none) \
184 dt_omp_firstprivate(height, width) \
188 for(
int i = 0; i <
height *
width; i++) out[i * 4 + 1] = (out[i * 4 + 1] + out[i * 4 + 3]) / 2.0f;
195 const gboolean smooth,
const int only_vng_linear)
200 const uint8_t(*
const xtrans)[6] = (
const uint8_t(*
const)[6])piece->pipe->dsc.xtrans;
204 if(piece->pipe->dsc.filters == 9u)
205 filters4 = piece->pipe->dsc.filters;
206 else if((piece->pipe->dsc.filters & 3) == 1)
207 filters4 = piece->pipe->dsc.filters | 0x03030303u;
209 filters4 = piece->pipe->dsc.filters | 0x0c0c0c0cu;
211 const int size = (filters4 == 9u) ? 6 : 16;
212 const int colors = (filters4 == 9u) ? 3 : 4;
213 const int prow = (filters4 == 9u) ? 6 : 8;
214 const int pcol = (filters4 == 9u) ? 6 : 2;
215 const int devid = piece->pipe->devid;
217 const float processed_maximum[4]
218 = { piece->pipe->dsc.processed_maximum[0], piece->pipe->dsc.processed_maximum[1],
219 piece->pipe->dsc.processed_maximum[2], 1.0f };
223 cl_mem dev_tmp = NULL;
224 cl_mem dev_aux = NULL;
225 cl_mem dev_xtrans = NULL;
226 cl_mem dev_lookup = NULL;
227 cl_mem dev_code = NULL;
228 cl_mem dev_ips = NULL;
229 cl_mem dev_green_eq = NULL;
232 int32_t(*
lookup)[16][32] = NULL;
234 if(piece->pipe->dsc.filters == 9u)
237 = dt_opencl_copy_host_to_device_constant(devid,
sizeof(piece->pipe->dsc.xtrans), piece->pipe->dsc.xtrans);
238 if(dev_xtrans == NULL)
goto error;
252 const size_t lookup_size = (size_t)16 * 16 * 32 *
sizeof(int32_t);
253 lookup = malloc(lookup_size);
255 for(
int row = 0; row <
size; row++)
256 for(
int col = 0; col <
size; col++)
258 int32_t *ip = &(
lookup[row][col][1]);
260 const int f =
fcol(row + roi_in->y, col + roi_in->x, filters4, xtrans);
262 for(
int y = -1; y <= 1; y++)
263 for(
int x = -1; x <= 1; x++)
265 const int weight = 1 << ((y == 0) + (x == 0));
266 const int color =
fcol(row + y + roi_in->y, col + x + roi_in->x, filters4, xtrans);
267 if(color ==
f)
continue;
268 *ip++ = (y << 16) | (x & 0xffffu);
273 lookup[row][col][0] = (ip - &(
lookup[row][col][0])) / 3;
284 static const signed char terms[]
285 = { -2, -2, +0, -1, 1, 0x01, -2, -2, +0, +0, 2, 0x01, -2, -1, -1, +0, 1, 0x01, -2, -1, +0, -1, 1, 0x02,
286 -2, -1, +0, +0, 1, 0x03, -2, -1, +0, +1, 2, 0x01, -2, +0, +0, -1, 1, 0x06, -2, +0, +0, +0, 2, 0x02,
287 -2, +0, +0, +1, 1, 0x03, -2, +1, -1, +0, 1, 0x04, -2, +1, +0, -1, 2, 0x04, -2, +1, +0, +0, 1, 0x06,
288 -2, +1, +0, +1, 1, 0x02, -2, +2, +0, +0, 2, 0x04, -2, +2, +0, +1, 1, 0x04, -1, -2, -1, +0, 1, 0x80,
289 -1, -2, +0, -1, 1, 0x01, -1, -2, +1, -1, 1, 0x01, -1, -2, +1, +0, 2, 0x01, -1, -1, -1, +1, 1, 0x88,
290 -1, -1, +1, -2, 1, 0x40, -1, -1, +1, -1, 1, 0x22, -1, -1, +1, +0, 1, 0x33, -1, -1, +1, +1, 2, 0x11,
291 -1, +0, -1, +2, 1, 0x08, -1, +0, +0, -1, 1, 0x44, -1, +0, +0, +1, 1, 0x11, -1, +0, +1, -2, 2, 0x40,
292 -1, +0, +1, -1, 1, 0x66, -1, +0, +1, +0, 2, 0x22, -1, +0, +1, +1, 1, 0x33, -1, +0, +1, +2, 2, 0x10,
293 -1, +1, +1, -1, 2, 0x44, -1, +1, +1, +0, 1, 0x66, -1, +1, +1, +1, 1, 0x22, -1, +1, +1, +2, 1, 0x10,
294 -1, +2, +0, +1, 1, 0x04, -1, +2, +1, +0, 2, 0x04, -1, +2, +1, +1, 1, 0x04, +0, -2, +0, +0, 2, 0x80,
295 +0, -1, +0, +1, 2, 0x88, +0, -1, +1, -2, 1, 0x40, +0, -1, +1, +0, 1, 0x11, +0, -1, +2, -2, 1, 0x40,
296 +0, -1, +2, -1, 1, 0x20, +0, -1, +2, +0, 1, 0x30, +0, -1, +2, +1, 2, 0x10, +0, +0, +0, +2, 2, 0x08,
297 +0, +0, +2, -2, 2, 0x40, +0, +0, +2, -1, 1, 0x60, +0, +0, +2, +0, 2, 0x20, +0, +0, +2, +1, 1, 0x30,
298 +0, +0, +2, +2, 2, 0x10, +0, +1, +1, +0, 1, 0x44, +0, +1, +1, +2, 1, 0x10, +0, +1, +2, -1, 2, 0x40,
299 +0, +1, +2, +0, 1, 0x60, +0, +1, +2, +1, 1, 0x20, +0, +1, +2, +2, 1, 0x10, +1, -2, +1, +0, 1, 0x80,
300 +1, -1, +1, +1, 1, 0x88, +1, +0, +1, +2, 1, 0x08, +1, +0, +2, -1, 1, 0x40, +1, +0, +2, +1, 1, 0x10 };
301 static const signed char chood[]
302 = { -1, -1, -1, 0, -1, +1, 0, +1, +1, +1, +1, 0, +1, -1, 0, -1 };
304 const size_t ips_size = (size_t)prow * pcol * 352 *
sizeof(
int);
305 ips = malloc(ips_size);
310 for(
int row = 0; row < prow; row++)
311 for(
int col = 0; col < pcol; col++)
313 code[row][col] = ip - ips;
314 const signed char *cp = terms;
315 for(
int t = 0; t < 64; t++)
317 const int y1 = *cp++, x1 = *cp++;
318 const int y2 = *cp++, x2 = *cp++;
320 const int grads = *cp++;
321 const int color =
fcol(row + y1, col + x1, filters4, xtrans);
322 if(
fcol(row + y2, col + x2, filters4, xtrans) != color)
continue;
324 = (
fcol(row, col + 1, filters4, xtrans) == color &&
fcol(row + 1, col, filters4, xtrans) == color)
327 if(abs(y1 - y2) == diag && abs(x1 - x2) == diag)
continue;
328 *ip++ = (y1 << 16) | (x1 & 0xffffu);
329 *ip++ = (y2 << 16) | (x2 & 0xffffu);
330 *ip++ = (color << 16) | (
weight & 0xffffu);
331 for(
int g = 0;
g < 8;
g++)
332 if(grads & 1 << g) *ip++ =
g;
337 for(
int g = 0;
g < 8;
g++)
339 const int y = *cp++, x = *cp++;
340 *ip++ = (y << 16) | (x & 0xffffu);
341 const int color =
fcol(row, col, filters4, xtrans);
342 if(
fcol(row + y, col + x, filters4, xtrans) != color
343 &&
fcol(row + y * 2, col + x * 2, filters4, xtrans) == color)
345 *ip++ = (2*y << 16) | (2*x & 0xffffu);
357 dev_lookup = dt_opencl_copy_host_to_device_constant(devid, lookup_size,
lookup);
358 if(dev_lookup == NULL)
goto error;
360 dev_code = dt_opencl_copy_host_to_device_constant(devid,
sizeof(code), code);
361 if(dev_code == NULL)
goto error;
363 dev_ips = dt_opencl_copy_host_to_device_constant(devid, ips_size, ips);
364 if(dev_ips == NULL)
goto error;
369 dev_green_eq = dt_opencl_alloc_device(devid, roi_in->width, roi_in->height,
sizeof(
float));
370 if(dev_green_eq == NULL)
goto error;
372 if(!green_equilibration_cl(self, piece, dev_in, dev_green_eq, roi_in))
375 dev_in = dev_green_eq;
378 int width = roi_out->width;
379 int height = roi_out->height;
383 dev_tmp = dt_opencl_alloc_device(devid, roi_in->width, roi_in->height,
sizeof(
float) * 4);
384 if(dev_tmp == NULL)
goto error;
388 const int border = 1;
390 size_t sizes[3] = { ROUNDUPDWD(
width, devid), ROUNDUPDHT(
height, devid), 1 };
401 if(err != CL_SUCCESS)
goto error;
406 dt_opencl_local_buffer_t locopt
407 = (dt_opencl_local_buffer_t){ .xoffset = 2*1, .xfactor = 1, .yoffset = 2*1, .yfactor = 1,
408 .cellsize = 1 *
sizeof(float), .overhead = 0,
409 .sizex = 1 << 8, .sizey = 1 << 8 };
411 if(!dt_opencl_local_buffer_opt(devid, gd->kernel_vng_lin_interpolate, &locopt))
414 size_t sizes[3] = { ROUNDUP(
width, locopt.sizex), ROUNDUP(
height, locopt.sizey), 1 };
415 size_t local[3] = { locopt.sizex, locopt.sizey, 1 };
423 sizeof(
float) * (locopt.sizex + 2) * (locopt.sizey + 2), NULL);
425 if(err != CL_SUCCESS)
goto error;
430 dt_opencl_local_buffer_t locopt
431 = (dt_opencl_local_buffer_t){ .xoffset = 2*2, .xfactor = 1, .yoffset = 2*2, .yfactor = 1,
432 .cellsize = 4 *
sizeof(float), .overhead = 0,
433 .sizex = 1 << 8, .sizey = 1 << 8 };
435 if(!dt_opencl_local_buffer_opt(devid, gd->kernel_vng_interpolate, &locopt))
438 size_t sizes[3] = { ROUNDUP(
width, locopt.sizex), ROUNDUP(
height, locopt.sizey), 1 };
439 size_t local[3] = { locopt.sizex, locopt.sizey, 1 };
451 dt_opencl_set_kernel_arg(devid, gd->kernel_vng_interpolate, 11,
sizeof(
float) * 4 * (locopt.sizex + 4) * (locopt.sizey + 4), NULL);
453 if(err != CL_SUCCESS)
goto error;
458 const int border = 2;
460 size_t sizes[3] = { ROUNDUPDWD(
width, devid), ROUNDUPDHT(
height, devid), 1 };
471 if(err != CL_SUCCESS)
goto error;
477 size_t origin[] = { 0, 0, 0 };
479 err = dt_opencl_enqueue_copy_image(devid, dev_aux, dev_tmp, origin, origin, region);
480 if(err != CL_SUCCESS)
goto error;
482 size_t sizes[3] = { ROUNDUPDWD(
width, devid), ROUNDUPDHT(
height, devid), 1 };
488 if(err != CL_SUCCESS)
goto error;
519 if((data->color_smoothing) && smooth)
521 if(!color_smoothing_cl(self, piece, dev_out, dev_out, roi_out, data->color_smoothing))
static void error(char *msg)
Definition ashift_lsd.c:191
#define TRUE
Definition ashift_lsd.c:151
#define FALSE
Definition ashift_lsd.c:147
static void lin_interpolate(float *out, const float *const in, const dt_iop_roi_t *const roi_out, const dt_iop_roi_t *const roi_in, const uint32_t filters, const uint8_t(*const xtrans)[6])
Definition basic.c:4
int width
Definition bilateral.h:1
int height
Definition bilateral.h:1
static float lookup(read_only image2d_t lut, const float x)
Definition color_conversion.h:70
void dt_print(dt_debug_thread_t thread, const char *msg,...)
Definition darktable.c:1395
@ DT_DEBUG_OPENCL
Definition darktable.h:478
#define dt_free_align(A)
Definition darktable.h:334
@ DT_IOP_GREEN_EQ_NO
Definition demosaic.c:93
@ DT_DEV_DETAIL_MASK_DEMOSAIC
Definition develop.h:110
static void weight(const float *c1, const float *c2, const float sharpen, dt_aligned_pixel_t weight)
Definition eaw.c:29
const dt_collection_filter_flag_t colors[6]
Definition filter.c:274
static float f(const float t, const float c, const float x)
Definition graduatednd.c:173
#define FILTERS_ARE_4BAYER(filters)
Definition imageio.h:40
static int fcol(const int row, const int col, const uint32_t filters, const uint8_t(*const xtrans)[6])
Definition imageop_math.h:222
size_t size
Definition mipmap_cache.c:3
c
Definition derive_filmic_v6_gamut_mapping.py:11
g
Definition derive_filmic_v6_gamut_mapping.py:18
static int dt_opencl_enqueue_kernel_2d(const int dev, const int kernel, const size_t *sizes)
Definition opencl.h:560
static int dt_opencl_set_kernel_arg(const int dev, const int kernel, const size_t size, const void *arg)
Definition opencl.h:556
static void dt_opencl_release_mem_object(void *mem)
Definition opencl.h:601
static int dt_opencl_enqueue_kernel_2d_with_local(const int dev, const int kernel, const size_t *sizes, const size_t *local)
Definition opencl.h:564
Definition pixelpipe_hb.h:46
void * data
Definition pixelpipe_hb.h:49
Definition demosaic.c:178
Definition demosaic.c:120
int x
Definition imageop.h:33
int width
Definition imageop.h:33
int height
Definition imageop.h:33
int y
Definition imageop.h:33
#define dt_alloc_align(B)
Definition tests/cache.c:22
static void vng_interpolate(float *out, const float *const in, const dt_iop_roi_t *const roi_out, const dt_iop_roi_t *const roi_in, const uint32_t filters, const uint8_t(*const xtrans)[6], const int only_vng_linear)
Definition vng.c:15