63#define DT_HEAL_SQRT_OFFSET 1e-3f
73static inline float _heal_decode(
const float src,
const float correction,
const size_t c,
78 const float root = fmaxf(
_heal_encode(src, c, domain) + correction, 0.0f);
81 return src + correction;
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,
91 const size_t res_stride = 4 * ((
width + 1) / 2);
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;
100 for(
size_t col = 0; col <
width/2; col++)
102 const size_t idx = 4 * (
row *
width + 2*col);
106 buf2[4*col + c] =
_heal_encode(top_buffer[idx+4 + c], c, domain)
114 const size_t res_idx = (
width-1)/2;
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;
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));
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,
137 const size_t res_stride = 4 * ((
width + 1) / 2);
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;
146 for(
size_t col = 0; col <
width/2; col++)
148 const size_t idx = 4 * (
row *
width + 2*col);
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);
158 const size_t res_idx = (
width-1)/2;
161 result_buffer[idx + c] =
_heal_decode(second_buffer[idx + c], buf1[4*res_idx + c], c, domain);
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 } })
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)
204#if !(defined(__apple_build_version__) && __apple_build_version__ < 11030000)
207 for(
size_t i = 0;
i < num_runs;
i++)
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;
214 if(
row == 1) a -= 1.0f;
216 const size_t vert_offset = 4 *
width;
217 const size_t lroffset = 4 * (start_parity ^ (
row & 1));
220 const size_t col = idx %
width;
224 if(col > 0 || lroffset)
228 if(col + 1 <
width || lroffset == 0)
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]);
245 copy_pixel(left, neighbor_pixels + index - 4 + lroffset);
246 for(
size_t j = 0; j < count; j++)
248 const size_t pixidx = index + 4*j;
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]);
263 return err.
v[0] + err.
v[1] + err.
v[2];
269 size_t start,
const size_t width,
270 unsigned *
const restrict runs,
size_t count,
size_t *nmask)
276 if(start == 0 && mask[start])
278 runs[2*count] = start_index;
284 gboolean in_run =
FALSE;
285 unsigned run_start = 0;
287 for(col = start; col <
width; col += 2)
300 runs[2*count] = start_index + run_start / 2;
301 runs[2*count + 1] = (col - run_start) / 2;
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)
315 runs[2*count + 2] = runs[2*count] + runs[2*count+1];
316 runs[2*count + 3] = 1;
327 const size_t subwidth,
unsigned *
const restrict runs,
size_t *count,
size_t *nmask)
331 const int parity = start ^ (
row & 1);
332 const size_t index = (
row + 1) * subwidth;
333 const size_t mask_index =
row *
width;
341 const float *
const restrict mask,
const int max_iter)
350 const size_t subwidth = (
width+1)/2;
352 sizeof(
unsigned) * subwidth * (
height + 2),
355 sizeof(
unsigned) * subwidth * (
height + 2),
359 fprintf(stderr,
"_heal_laplace_loop: error allocating memory for healing\n");
364 size_t num_black = 0;
365 size_t nmask_red = 0;
366 size_t nmask_black = 0;
369#pragma omp parallel sections
378 const size_t nmask = nmask_red + nmask_black;
384 const float w = ((2.0f - 1.0f / (0.1575f * sqrtf(nmask) + 0.8f)) * .25f);
386 const float epsilon = (0.1 / 255);
387 const float err_exit = epsilon * epsilon * w * w;
390 for(
int iter = 0; iter < max_iter; iter++)
396 if(err < err_exit)
break;
410void dt_heal(
const float *
const src_buffer,
float *dest_buffer,
const float *
const mask_buffer,
const int width,
415 fprintf(stderr,
"dt_heal: full-color image required\n");
418 const size_t subwidth = 4 * ((
width+1)/2);
423 fprintf(stderr,
"dt_heal: error allocating memory for healing\n");
486 cl_int err = CL_SUCCESS;
490 float *src_buffer = NULL;
491 float *dest_buffer = NULL;
496 fprintf(stderr,
"dt_heal_cl: error allocating memory for healing\n");
504 fprintf(stderr,
"dt_heal_cl: error allocating memory for healing\n");
511 if(err != CL_SUCCESS)
518 if(err != CL_SUCCESS)
527 if(err != CL_SUCCESS)
void cleanup(dt_imageio_module_format_t *self)
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)
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)
void dt_heal_free_cl_global(void)
void dt_heal_free_cl(heal_params_cl_t *p)
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)
static float _heal_decode(const float src, const float correction, const size_t c, const dt_heal_domain_t domain)
static dt_heal_cl_global_t * _heal_cl_global
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)
void dt_heal_init_cl_global(void)
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)
heal_params_cl_t * dt_heal_init_cl(const int devid)
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)
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)
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)
static float _heal_encode(const float value, const size_t c, const dt_heal_domain_t domain)
#define DT_HEAL_SQRT_OFFSET
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...
#define dt_free(ptr)
g_free() ptr and set it to NULL, skipping both if it is already NULL.
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)
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)
#define DT_OPENCL_SYSMEM_ALLOCATION
#define __OMP_PARALLEL_FOR__(...)
#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)
DT_ALIGNED_PIXEL float dt_aligned_pixel_t[4]
#define for_each_channel(_var,...)
static const dt_aligned_pixel_simd_t value
dt_heal_cl_global_t * global