Ansel 0.0
A darktable fork - bloat + design vision
Loading...
Searching...
No Matches
heal.c
Go to the documentation of this file.
1/*
2 This file is part of darktable,
3 Copyright (C) 2017 Edgardo Hoszowski.
4 Copyright (C) 2019 Andreas Schneider.
5 Copyright (C) 2020 Hubert Kowalski.
6 Copyright (C) 2020-2022 Pascal Obry.
7 Copyright (C) 2021 parafin.
8 Copyright (C) 2021-2022 Ralf Brown.
9 Copyright (C) 2022 Hanno Schwalm.
10 Copyright (C) 2022 Martin Bařinka.
11 Copyright (C) 2024 Alynx Zhou.
12 Copyright (C) 2026 Aurélien PIERRE.
13
14 darktable is free software: you can redistribute it and/or modify
15 it under the terms of the GNU General Public License as published by
16 the Free Software Foundation, either version 3 of the License, or
17 (at your option) any later version.
18
19 darktable is distributed in the hope that it will be useful,
20 but WITHOUT ANY WARRANTY; without even the implied warranty of
21 MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
22 GNU General Public License for more details.
23
24 You should have received a copy of the GNU General Public License
25 along with darktable. If not, see <http://www.gnu.org/licenses/>.
26*/
27
28#include "system/macros.h"
29#include "system/openmp.h"
30#include "system/mem_alloc.h"
31#include "system/simd.h"
33#include "develop/imageop.h"
34#include "math/openmp_maths.h"
35#include "heal.h"
36
37/* Based on the original source code of GIMP's Healing Tool, by Jean-Yves Couleaud
38 *
39 * http://www.gimp.org/
40 *
41 * */
42
43/* NOTES
44 *
45 * The method used here is similar to the lighting invariant correction
46 * method but slightly different: we do not divide the RGB components,
47 * but subtract them I2 = I0 - I1, where I0 is the sample image to be
48 * corrected, I1 is the reference pattern. Then we solve DeltaI=0
49 * (Laplace) with I2 Dirichlet conditions at the borders of the
50 * mask. The solver is a red/black checker Gauss-Seidel with over-relaxation.
51 * It could benefit from a multi-grid evaluation of an initial solution
52 * before the main iteration loop.
53 *
54 * I reduced the convergence criteria to 0.1% (0.001) as we are
55 * dealing here with RGB integer components, more is overkill.
56 *
57 * Jean-Yves Couleaud cjyves@free.fr
58 */
59
60
61/* Offset under the square root: it keeps the transform differentiable at zero and stands in for
62 * the noise floor, where shot noise stops dominating. Scene-linear, 1.0 being white. */
63#define DT_HEAL_SQRT_OFFSET 1e-3f
64
65static inline float _heal_encode(const float value, const size_t c, const dt_heal_domain_t domain)
66{
67 if(domain == DT_HEAL_DOMAIN_SQRT && c < 3) return sqrtf(fmaxf(value, 0.0f) + DT_HEAL_SQRT_OFFSET);
68 return value;
69}
70
71// Inverse of _heal_encode() applied to src + correction. A negative source value keeps its
72// negative part, which the square root cannot carry.
73static inline float _heal_decode(const float src, const float correction, const size_t c,
74 const dt_heal_domain_t domain)
75{
76 if(domain == DT_HEAL_DOMAIN_SQRT && c < 3)
77 {
78 const float root = fmaxf(_heal_encode(src, c, domain) + correction, 0.0f);
79 return root * root - DT_HEAL_SQRT_OFFSET + fminf(src, 0.0f);
80 }
81 return src + correction;
82}
83
84// Subtract bottom from top and store in result as a float; separate 'red' and 'black' pixels into
85// two contiguous regions
86static void _heal_sub(const float *const top_buffer, const float *const bottom_buffer,
87 float *const restrict red_buffer, float *const restrict black_buffer,
88 const size_t width, const size_t height, const dt_heal_domain_t domain)
89{
90 // how many red or black pixels per line? For consistency, we need the larger of the two, so round up
91 const size_t res_stride = 4 * ((width + 1) / 2);
93 for(size_t row = 0; row < height; row++)
94 {
95 const int parity = row & 1;
96 const size_t row_start = (row+1) * res_stride;
97 float *const buf1 = parity ? red_buffer + row_start : black_buffer + row_start;
98 float *const buf2 = parity ? black_buffer + row_start : red_buffer + row_start;
99 // handle the pixels of the row pairwise, one red and one black at a time
100 for(size_t col = 0; col < width/2; col++)
101 {
102 const size_t idx = 4 * (row * width + 2*col);
104 {
105 buf1[4*col + c] = _heal_encode(top_buffer[idx + c], c, domain) - _heal_encode(bottom_buffer[idx + c], c, domain);
106 buf2[4*col + c] = _heal_encode(top_buffer[idx+4 + c], c, domain)
107 - _heal_encode(bottom_buffer[idx+4 + c], c, domain);
108 }
109 }
110 if(width & 1)
111 {
112 // Handle the left-over pixel when the total width is odd. Its color will always be the same as the
113 // left-most pixel, so it goes into buf1
114 const size_t res_idx = (width-1)/2;
115 const size_t idx = 4 * (row * width + (width-1));
117 {
118 buf1[4*res_idx + c] = _heal_encode(top_buffer[idx + c], c, domain) - _heal_encode(bottom_buffer[idx + c], c, domain);
119 buf2[4*res_idx + c] = 0.0f;
120 }
121 }
122 }
123 // clear the top and bottom rows, used for padding
124 memset(red_buffer, 0, res_stride * sizeof(float));
125 memset(red_buffer + (height+1)*res_stride, 0, res_stride * sizeof(float));
126 memset(black_buffer, 0, res_stride * sizeof(float));
127 memset(black_buffer + (height+1)*res_stride, 0, res_stride * sizeof(float));
128}
129
130// Add first to second and store in result, re-interleaving the 'red' and 'black' pixels
131static void _heal_add(const float *const restrict red_buffer, const float *const black_buffer,
132 const float *const restrict second_buffer, float *const restrict result_buffer,
133 const size_t width, const size_t height, const dt_heal_domain_t domain)
134{
135 // how many red or black pixels per line? For consistency, we need the larger of the two, so round up, then
136 // add one to ensure a padding pixel on the right
137 const size_t res_stride = 4 * ((width + 1) / 2);
139 for(size_t row = 0; row < height; row++)
140 {
141 const int parity = row & 1;
142 const size_t row_start = (row+1) * res_stride;
143 const float *const restrict buf1 = parity ? red_buffer + row_start : black_buffer + row_start;
144 const float *const restrict buf2 = parity ? black_buffer + row_start : red_buffer + row_start;
145 // handle the pixels of the row pairwise, one red and one black at a time
146 for(size_t col = 0; col < width/2; col++)
147 {
148 const size_t idx = 4 * (row * width + 2*col);
150 {
151 result_buffer[idx + c] = _heal_decode(second_buffer[idx + c], buf1[4*col + c], c, domain);
152 result_buffer[idx + 4 + c] = _heal_decode(second_buffer[idx + 4 + c], buf2[4*col + c], c, domain);
153 }
154 }
155 if(width & 1)
156 {
157 // handle the left-over pixel when the total width is odd
158 const size_t res_idx = (width-1)/2;
159 const size_t idx = 4 * (row * width + (width-1));
161 result_buffer[idx + c] = _heal_decode(second_buffer[idx + c], buf1[4*res_idx + c], c, domain);
162 }
163 }
164}
165
166// define a custom reduction operation to handle a 3-vector of floats
167// we can't return an array from a function, so wrap the array type in a struct
169#ifdef _OPENMP
170static inline _aligned_pixel _add_float4(_aligned_pixel acc, _aligned_pixel newval)
171{
172 for_each_channel(c) acc.v[c] += newval.v[c];
173 return acc;
174}
175#pragma omp declare reduction(vsum:_aligned_pixel:omp_out=_add_float4(omp_out,omp_in)) \
176 initializer(omp_priv = { { 0.0f, 0.0f, 0.0f, 0.0f } })
177#endif
178
179// Perform one iteration of Gauss-Seidel, and return the sum squared residual.
180static float _heal_laplace_iteration(float *const restrict active_pixels,
181 const float *const restrict neighbor_pixels,
182 const size_t height, const size_t width, const unsigned *const restrict runs,
183 const size_t num_runs, const size_t start_parity, const float w)
184{
185 _aligned_pixel err = { { 0.f } };
186
187 // on each iteration, we adjust each cell of the current color by a weighted fraction of the difference between
188 // it and the sum of the adjacenct cells of the opposite color. Because we've split the cells by color, this
189 // leads to a somewhat different computation of neighbors.
190 // Rearranged, as provided to this function:
191 // r00 r01 r02 ... b00 b01 b02 ...
192 // r10 r11 r12 b10 b11 b12
193 // r20 r21 r22 b20 b21 b22
194 // ... ...
195 // Original layout, from which we need to get the neighbors:
196 // r00 b00 r01 b01 r02 b02 ...
197 // b10 r10 b11 r11 b12 r12
198 // r20 b20 r21 b21 r22 b22
199 // b30 r30 b31 r31 b32 r32
200 // ...
201 // As can be seen, the 'above' and 'below' neighbors of r(i)(j) are always b(i-1)(j) and b(i+1)(j). The
202 // left and right neighbors depend on which color the row starts with: if red, they are b(i)(j-1) and b(i)(j);
203 // if black, they are b(i)(j) and b(i)(j+1). All of the above holds when colors are swapped.
204#if !(defined(__apple_build_version__) && __apple_build_version__ < 11030000) //makes Xcode 11.3.1 compiler crash
205__OMP_PARALLEL_FOR__(reduction(vsum : err)) /* _OPENMP */
206#endif
207 for(size_t i = 0; i < num_runs; i++)
208 {
209 const size_t idx = runs[2*i];
210 const unsigned count = runs[2*i+1];
211 const size_t index = (size_t)4 * idx;
212 const size_t row = idx / width;
213 float a = 4.0f; // four neighboring pixels except at the edges
214 if(row == 1) a -= 1.0f; // we added a padding row at top and bottom
215 if(row == height) a -= 1.0f;
216 const size_t vert_offset = 4 * width;
217 const size_t lroffset = 4 * (start_parity ^ (row & 1)); // how many floats to offset the left/right neighbors
218 if(count == 1)
219 {
220 const size_t col = idx % width;
221 float aa = a;
222 dt_aligned_pixel_t left = { 0.0f };
223 dt_aligned_pixel_t right = { 0.0f };
224 if(col > 0 || lroffset) // first pixel in original stamp?
225 for_each_channel(c) left[c] = neighbor_pixels[index - 4 + lroffset + c];
226 else
227 aa -= 1.0f;
228 if(col + 1 < width || lroffset == 0) // last pixel in original stamp?
229 for_each_channel(c) right[c] = neighbor_pixels[index + lroffset + c];
230 else
231 aa -= 1.0f;
232
234 for_each_channel(c, aligned(active_pixels, neighbor_pixels))
235 {
236 diff[c] = w * ((aa * active_pixels[index+c])
237 - (neighbor_pixels[index - vert_offset + c] + neighbor_pixels[index + vert_offset + c]
238 + left[c] + right[c]));
239 active_pixels[index + c] -= diff[c];
240 err.v[c] += (diff[c] * diff[c]);
241 }
242 continue;
243 }
245 copy_pixel(left, neighbor_pixels + index - 4 + lroffset);
246 for(size_t j = 0; j < count; j++)
247 {
248 const size_t pixidx = index + 4*j;
250 dt_aligned_pixel_t right;
251 for_each_channel(c, aligned(active_pixels,neighbor_pixels))
252 {
253 right[c] = neighbor_pixels[pixidx + lroffset + c];
254 diff[c] = w * (a * active_pixels[pixidx+c]
255 - (neighbor_pixels[pixidx - vert_offset + c] + neighbor_pixels[pixidx + vert_offset + c]
256 + left[c] + right[c]));
257 active_pixels[pixidx + c] -= diff[c];
258 err.v[c] += (diff[c] * diff[c]);
259 left[c] = right[c];
260 }
261 }
262 }
263 return err.v[0] + err.v[1] + err.v[2];
264}
265
266// convert alternating pixels of one row of the opacity mask into a set of runs of opaque pixels of the
267// form (start_index, count), and return the updated number of runs
268static size_t _collect_color_runs(const float *const restrict mask, const size_t start_index,
269 size_t start, const size_t width,
270 unsigned *const restrict runs, size_t count, size_t *nmask)
271{
272 size_t masked = 0;
273 // handle the first and last pixels of the row specially so that we don't need checks on every pixel in the
274 // adaptation loop (which will be run hundreds of times for any stamp large enough that execution time is
275 // non-negligible)
276 if(start == 0 && mask[start])
277 {
278 runs[2*count] = start_index;
279 runs[2*count+1] = 1;
280 count++;
281 masked++;
282 start += 2; // we've processed the first pixel
283 }
284 gboolean in_run = FALSE;
285 unsigned run_start = 0;
286 size_t col;
287 for(col = start; col < width; col += 2)
288 {
289 if(mask[col])
290 {
291 masked++;
292 if(!in_run)
293 {
294 run_start = col;
295 in_run = TRUE;
296 }
297 }
298 else if(in_run)
299 {
300 runs[2*count] = start_index + run_start / 2;
301 runs[2*count + 1] = (col - run_start) / 2;
302 count++;
303 in_run = FALSE;
304 }
305 }
306 if(in_run) // finish off a run that doesn't have a zero pixel to its right
307 {
308 runs[2*count] = start_index + run_start / 2;
309 const unsigned runlen = (col - run_start) / 2;
310 runs[2*count + 1] = runlen;
311 if(runlen > 1 && col > width)
312 {
313 // split off the final pixel into its own run
314 runs[2*count + 1]--;
315 runs[2*count + 2] = runs[2*count] + runs[2*count+1];
316 runs[2*count + 3] = 1;
317 count++;
318 }
319 count++;
320 }
321 *nmask += masked;
322 return count;
323}
324
325// convert one row of the opacity mask into a set of runs of opaque pixels of the form (start_index, count)
326static void collect_runs(const int start, const float *const restrict mask, const size_t width, const size_t height,
327 const size_t subwidth, unsigned *const restrict runs, size_t *count, size_t *nmask)
328{
329 for(size_t row = 0; row < height; row++)
330 {
331 const int parity = start ^ (row & 1);
332 const size_t index = (row + 1) * subwidth;
333 const size_t mask_index = row * width;
334 *count = _collect_color_runs(mask + mask_index, index, parity, width, runs, *count, nmask);
335 }
336}
337
338// Solve the laplace equation for pixels and store the result in-place.
339static void _heal_laplace_loop(float *const restrict red_pixels, float *const restrict black_pixels,
340 const size_t width, const size_t height,
341 const float *const restrict mask, const int max_iter)
342{
343 // we start by converting the opacity mask into runs of nonzero positions, handling the 'red' and 'black'
344 // checkerboarded pixels separately
345 // the worst case is when consecutive red pixels alternate between being in the mask and out (same for black),
346 // in which case we will need exactly as many values in the run-length encoding as pixels; any other
347 // arrangement will yield fewer runs. For any mask a user would have the patience to draw, there will only be
348 // a handful of runs in each row of pixels.
349 // Note that using `unsigned` instead of size_t, the stamp is limited to ~8 gigapixels (the main image can be larger)
350 const size_t subwidth = (width+1)/2; // round up to be able to handle odd widths
351 unsigned *const restrict red_runs = dt_pixelpipe_cache_alloc_align_cache(
352 sizeof(unsigned) * subwidth * (height + 2),
353 0);
354 unsigned *const restrict black_runs = dt_pixelpipe_cache_alloc_align_cache(
355 sizeof(unsigned) * subwidth * (height + 2),
356 0);
357 if(IS_NULL_PTR(red_runs) || IS_NULL_PTR(black_runs))
358 {
359 fprintf(stderr, "_heal_laplace_loop: error allocating memory for healing\n");
360 goto cleanup;
361 }
362
363 size_t num_red = 0;
364 size_t num_black = 0;
365 size_t nmask_red = 0;
366 size_t nmask_black = 0;
367
368#ifdef _OPENMP
369#pragma omp parallel sections
370#endif
371 {
372 collect_runs(1, mask, width, height, subwidth, red_runs, &num_red, &nmask_red);
373 #ifdef _OPENMP
374 #pragma omp section
375 #endif
376 collect_runs(0, mask, width, height, subwidth, black_runs, &num_black, &nmask_black);
377 }
378 const size_t nmask = nmask_red + nmask_black;
379
380 /* Empirically optimal over-relaxation factor. (Benchmarked on
381 * round brushes, at least. I don't know whether aspect ratio
382 * affects it.)
383 */
384 const float w = ((2.0f - 1.0f / (0.1575f * sqrtf(nmask) + 0.8f)) * .25f);
385
386 const float epsilon = (0.1 / 255);
387 const float err_exit = epsilon * epsilon * w * w;
388
389 /* Gauss-Seidel with successive over-relaxation */
390 for(int iter = 0; iter < max_iter; iter++)
391 {
392 // process red/black cells separately
393 float err = _heal_laplace_iteration(black_pixels, red_pixels, height, subwidth, black_runs, num_black, 1, w);
394 err += _heal_laplace_iteration(red_pixels, black_pixels, height, subwidth, red_runs, num_red, 0, w);
395
396 if(err < err_exit) break;
397 }
398
399cleanup:
402}
403
404
405/* Original Algorithm Design:
406 *
407 * T. Georgiev, "Photoshop Healing Brush: a Tool for Seamless Cloning
408 * http://www.tgeorgiev.net/Photoshop_Healing.pdf
409 */
410void dt_heal(const float *const src_buffer, float *dest_buffer, const float *const mask_buffer, const int width,
411 const int height, const int ch, const dt_heal_solver_t solver)
412{
413 if(ch != 4)
414 {
415 fprintf(stderr,"dt_heal: full-color image required\n");
416 return;
417 }
418 const size_t subwidth = 4 * ((width+1)/2); // round up to be able to handle odd widths
419 float *const restrict red_buffer = dt_pixelpipe_cache_alloc_align_float_cache(subwidth * (height + 2), 0);
420 float *const restrict black_buffer = dt_pixelpipe_cache_alloc_align_float_cache(subwidth * (height + 2), 0);
421 if(IS_NULL_PTR(red_buffer) || IS_NULL_PTR(black_buffer))
422 {
423 fprintf(stderr, "dt_heal: error allocating memory for healing\n");
424 goto cleanup;
425 }
426
427 /* subtract pattern from image and store the result split by 'red' and 'black' positions */
428 _heal_sub(dest_buffer, src_buffer, red_buffer, black_buffer, width, height, solver.domain);
429
430 _heal_laplace_loop(red_buffer, black_buffer, width, height, mask_buffer, solver.max_iter);
431
432 /* add solution to original image and store in dest */
433 _heal_add(red_buffer, black_buffer, src_buffer, dest_buffer, width, height, solver.domain);
434
435cleanup:
437 dt_pixelpipe_cache_free_align(black_buffer);
438}
439
440#ifdef HAVE_OPENCL
441
442/* The kernels this subsystem compiles, owned HERE. They used to be handed to
443 * common/opencl.c, parked on the application-wide dt_opencl_t, and read back from it --
444 * a round trip through a god-struct that added nothing but an ordering. opencl.c still
445 * calls init/free, because the kernels must be built after the devices exist, but the
446 * pointer never leaves this file. */
448
455
457{
459 _heal_cl_global = NULL;
460 if(IS_NULL_PTR(g)) return;
461
462 dt_free(g);
463}
464
466{
467
469 if(IS_NULL_PTR(p)) return NULL;
470
472 p->devid = devid;
473
474 return p;
475}
476
478{
479 if(IS_NULL_PTR(p)) return;
480 dt_free(p);
481}
482
483cl_int dt_heal_cl(heal_params_cl_t *p, cl_mem dev_src, cl_mem dev_dest, const float *const mask_buffer,
484 const int width, const int height, const dt_heal_solver_t solver)
485{
486 cl_int err = CL_SUCCESS;
487
488 const int ch = 4;
489
490 float *src_buffer = NULL;
491 float *dest_buffer = NULL;
492
493 src_buffer = dt_pixelpipe_cache_alloc_align_float_cache((size_t)ch * width * height, 0);
494 if(IS_NULL_PTR(src_buffer))
495 {
496 fprintf(stderr, "dt_heal_cl: error allocating memory for healing\n");
498 goto cleanup;
499 }
500
501 dest_buffer = dt_pixelpipe_cache_alloc_align_float_cache((size_t)ch * width * height, 0);
502 if(IS_NULL_PTR(dest_buffer))
503 {
504 fprintf(stderr, "dt_heal_cl: error allocating memory for healing\n");
506 goto cleanup;
507 }
508
509 err = dt_opencl_read_buffer_from_device(p->devid, (void *)src_buffer, dev_src, 0,
510 (size_t)width * height * ch * sizeof(float), CL_TRUE);
511 if(err != CL_SUCCESS)
512 {
513 goto cleanup;
514 }
515
516 err = dt_opencl_read_buffer_from_device(p->devid, (void *)dest_buffer, dev_dest, 0,
517 (size_t)width * height * ch * sizeof(float), CL_TRUE);
518 if(err != CL_SUCCESS)
519 {
520 goto cleanup;
521 }
522
523 // I couldn't make it run fast on opencl (the reduction takes forever), so just call the cpu version
524 dt_heal(src_buffer, dest_buffer, mask_buffer, width, height, ch, solver);
525
526 err = dt_opencl_write_buffer_to_device(p->devid, dest_buffer, dev_dest, 0, sizeof(float) * width * height * ch, CL_TRUE);
527 if(err != CL_SUCCESS)
528 {
529 goto cleanup;
530 }
531
532cleanup:
535
536 return err;
537}
538
539#endif
540// clang-format off
541// modelines: These editor modelines have been set for all relevant files by tools/update_modelines.py
542// vim: shiftwidth=2 expandtab tabstop=2 cindent
543// kate: tab-indents: off; indent-width 2; replace-tabs on; indent-mode cstyle; remove-trailing-spaces modified;
544// clang-format on
#define TRUE
Definition ashift_lsd.c:162
#define FALSE
Definition ashift_lsd.c:158
void cleanup(dt_imageio_module_format_t *self)
Definition avif.c:170
static const int row
void dt_heal(const float *const src_buffer, float *dest_buffer, const float *const mask_buffer, const int width, const int height, const int ch, const dt_heal_solver_t solver)
Definition heal.c:410
static size_t _collect_color_runs(const float *const restrict mask, const size_t start_index, size_t start, const size_t width, unsigned *const restrict runs, size_t count, size_t *nmask)
Definition heal.c:268
void dt_heal_free_cl_global(void)
Definition heal.c:456
void dt_heal_free_cl(heal_params_cl_t *p)
Definition heal.c:477
static void _heal_add(const float *const restrict red_buffer, const float *const black_buffer, const float *const restrict second_buffer, float *const restrict result_buffer, const size_t width, const size_t height, const dt_heal_domain_t domain)
Definition heal.c:131
static float _heal_decode(const float src, const float correction, const size_t c, const dt_heal_domain_t domain)
Definition heal.c:73
static dt_heal_cl_global_t * _heal_cl_global
Definition heal.c:447
static void collect_runs(const int start, const float *const restrict mask, const size_t width, const size_t height, const size_t subwidth, unsigned *const restrict runs, size_t *count, size_t *nmask)
Definition heal.c:326
void dt_heal_init_cl_global(void)
Definition heal.c:449
static void _heal_sub(const float *const top_buffer, const float *const bottom_buffer, float *const restrict red_buffer, float *const restrict black_buffer, const size_t width, const size_t height, const dt_heal_domain_t domain)
Definition heal.c:86
heal_params_cl_t * dt_heal_init_cl(const int devid)
Definition heal.c:465
static void _heal_laplace_loop(float *const restrict red_pixels, float *const restrict black_pixels, const size_t width, const size_t height, const float *const restrict mask, const int max_iter)
Definition heal.c:339
static float _heal_laplace_iteration(float *const restrict active_pixels, const float *const restrict neighbor_pixels, const size_t height, const size_t width, const unsigned *const restrict runs, const size_t num_runs, const size_t start_parity, const float w)
Definition heal.c:180
cl_int dt_heal_cl(heal_params_cl_t *p, cl_mem dev_src, cl_mem dev_dest, const float *const mask_buffer, const int width, const int height, const dt_heal_solver_t solver)
Definition heal.c:483
static float _heal_encode(const float value, const size_t c, const dt_heal_domain_t domain)
Definition heal.c:65
#define DT_HEAL_SQRT_OFFSET
Definition heal.c:63
dt_heal_domain_t
Definition heal.h:35
@ DT_HEAL_DOMAIN_SQRT
Definition heal.h:37
float *const restrict const size_t const size_t ch
#define IS_NULL_PTR(p)
C is way too permissive with !=, == and if(var) checks, which can mean too many things depending on w...
Definition macros.h:96
#define dt_free(ptr)
g_free() ptr and set it to NULL, skipping both if it is already NULL.
Definition mem_alloc.h:171
uint32_t width
Definition mipmap_cache.c:0
uint32_t height
Definition mipmap_cache.c:1
int dt_opencl_write_buffer_to_device(const int devid, void *host, void *device, const size_t offset, const size_t size, const int blocking)
Definition opencl.c:2796
int dt_opencl_read_buffer_from_device(const int devid, void *host, void *device, const size_t offset, const size_t size, const int blocking)
Definition opencl.c:2785
#define DT_OPENCL_SYSMEM_ALLOCATION
Definition opencl.h:62
#define __OMP_PARALLEL_FOR__(...)
Definition openmp.h:95
#define dt_pixelpipe_cache_alloc_align_float_cache(pixels, id)
#define dt_pixelpipe_cache_alloc_align_cache(size, id)
#define dt_pixelpipe_cache_free_align(mem)
static void copy_pixel(float *const __restrict__ out, const float *const __restrict__ in)
Definition simd.h:217
DT_ALIGNED_PIXEL float dt_aligned_pixel_t[4]
Definition simd.h:53
#define for_each_channel(_var,...)
Definition simd.h:87
static const dt_aligned_pixel_simd_t value
Definition simd.h:144
dt_aligned_pixel_t v
Definition eaw.c:199
dt_heal_domain_t domain
Definition heal.h:44
dt_heal_cl_global_t * global
Definition heal.h:61