63static void _heal_sub(
const float *
const top_buffer,
const float *
const bottom_buffer,
64 float *
const restrict red_buffer,
float *
const restrict black_buffer,
68 const size_t res_stride = 4 * ((
width + 1) / 2);
72 const int parity =
row & 1;
73 const size_t row_start = (
row+1) * res_stride;
74 float *
const buf1 = parity ? red_buffer + row_start : black_buffer + row_start;
75 float *
const buf2 = parity ? black_buffer + row_start : red_buffer + row_start;
77 for(
size_t col = 0; col <
width/2; col++)
79 const size_t idx = 4 * (
row *
width + 2*col);
82 buf1[4*col + c] = top_buffer[idx + c] - bottom_buffer[idx + c];
83 buf2[4*col + c] = top_buffer[idx+4 + c] - bottom_buffer[idx+4 + c];
90 const size_t res_idx = (
width-1)/2;
94 buf1[4*res_idx + c] = top_buffer[idx + c] - bottom_buffer[idx + c];
95 buf2[4*res_idx + c] = 0.0f;
100 memset(red_buffer, 0, res_stride *
sizeof(
float));
101 memset(red_buffer + (
height+1)*res_stride, 0, res_stride *
sizeof(
float));
102 memset(black_buffer, 0, res_stride *
sizeof(
float));
103 memset(black_buffer + (
height+1)*res_stride, 0, res_stride *
sizeof(
float));
107static void _heal_add(
const float *
const restrict red_buffer,
const float *
const black_buffer,
108 const float *
const restrict second_buffer,
float *
const restrict result_buffer,
113 const size_t res_stride = 4 * ((
width + 1) / 2);
117 const int parity =
row & 1;
118 const size_t row_start = (
row+1) * res_stride;
119 const float *
const restrict buf1 = parity ? red_buffer + row_start : black_buffer + row_start;
120 const float *
const restrict buf2 = parity ? black_buffer + row_start : red_buffer + row_start;
122 for(
size_t col = 0; col <
width/2; col++)
124 const size_t idx = 4 * (
row *
width + 2*col);
127 result_buffer[idx + c] = buf1[4*col + c] + second_buffer[idx + c];
128 result_buffer[idx + 4 + c] = buf2[4*col + c] + second_buffer[idx + 4 + c];
134 const size_t res_idx = (
width-1)/2;
137 result_buffer[idx + c] = buf1[4*res_idx + c] + second_buffer[idx + c];
151#pragma omp declare reduction(vsum:_aligned_pixel:omp_out=_add_float4(omp_out,omp_in)) \
152 initializer(omp_priv = { { 0.0f, 0.0f, 0.0f, 0.0f } })
157 const float *
const restrict neighbor_pixels,
158 const size_t height,
const size_t width,
const unsigned *
const restrict runs,
159 const size_t num_runs,
const size_t start_parity,
const float w)
180#if !(defined(__apple_build_version__) && __apple_build_version__ < 11030000)
183 for(
size_t i = 0;
i < num_runs;
i++)
185 const size_t idx = runs[2*
i];
186 const unsigned count = runs[2*
i+1];
187 const size_t index = (size_t)4 * idx;
190 if(
row == 1) a -= 1.0f;
192 const size_t vert_offset = 4 *
width;
193 const size_t lroffset = 4 * (start_parity ^ (
row & 1));
196 const size_t col = idx %
width;
200 if(col > 0 || lroffset)
204 if(col + 1 <
width || lroffset == 0)
212 diff[c] = w * ((aa * active_pixels[index+c])
213 - (neighbor_pixels[index - vert_offset + c] + neighbor_pixels[index + vert_offset + c]
214 + left[c] + right[c]));
215 active_pixels[index + c] -= diff[c];
216 err.
v[c] += (diff[c] * diff[c]);
221 copy_pixel(left, neighbor_pixels + index - 4 + lroffset);
222 for(
size_t j = 0; j < count; j++)
224 const size_t pixidx = index + 4*j;
229 right[c] = neighbor_pixels[pixidx + lroffset + c];
230 diff[c] = w * (a * active_pixels[pixidx+c]
231 - (neighbor_pixels[pixidx - vert_offset + c] + neighbor_pixels[pixidx + vert_offset + c]
232 + left[c] + right[c]));
233 active_pixels[pixidx + c] -= diff[c];
234 err.
v[c] += (diff[c] * diff[c]);
239 return err.
v[0] + err.
v[1] + err.
v[2];
245 size_t start,
const size_t width,
246 unsigned *
const restrict runs,
size_t count,
size_t *nmask)
252 if(start == 0 && mask[start])
254 runs[2*count] = start_index;
260 gboolean in_run =
FALSE;
261 unsigned run_start = 0;
263 for(col = start; col <
width; col += 2)
276 runs[2*count] = start_index + run_start / 2;
277 runs[2*count + 1] = (col - run_start) / 2;
284 runs[2*count] = start_index + run_start / 2;
285 const unsigned runlen = (col - run_start) / 2;
286 runs[2*count + 1] = runlen;
287 if(runlen > 1 && col >
width)
291 runs[2*count + 2] = runs[2*count] + runs[2*count+1];
292 runs[2*count + 3] = 1;
303 const size_t subwidth,
unsigned *
const restrict runs,
size_t *count,
size_t *nmask)
307 const int parity = start ^ (
row & 1);
308 const size_t index = (
row + 1) * subwidth;
309 const size_t mask_index =
row *
width;
317 const float *
const restrict mask,
const int max_iter)
326 const size_t subwidth = (
width+1)/2;
328 sizeof(
unsigned) * subwidth * (
height + 2),
331 sizeof(
unsigned) * subwidth * (
height + 2),
335 fprintf(stderr,
"_heal_laplace_loop: error allocating memory for healing\n");
340 size_t num_black = 0;
341 size_t nmask_red = 0;
342 size_t nmask_black = 0;
345#pragma omp parallel sections
354 const size_t nmask = nmask_red + nmask_black;
360 const float w = ((2.0f - 1.0f / (0.1575f * sqrtf(nmask) + 0.8f)) * .25f);
362 const float epsilon = (0.1 / 255);
363 const float err_exit = epsilon * epsilon * w * w;
366 for(
int iter = 0; iter < max_iter; iter++)
372 if(err < err_exit)
break;
386void dt_heal(
const float *
const src_buffer,
float *dest_buffer,
const float *
const mask_buffer,
const int width,
387 const int height,
const int ch,
const int max_iter)
391 fprintf(stderr,
"dt_heal: full-color image required\n");
394 const size_t subwidth = 4 * ((
width+1)/2);
399 fprintf(stderr,
"dt_heal: error allocating memory for healing\n");
460 const int width,
const int height,
const int max_iter)
462 cl_int err = CL_SUCCESS;
466 float *src_buffer = NULL;
467 float *dest_buffer = NULL;
472 fprintf(stderr,
"dt_heal_cl: error allocating memory for healing\n");
480 fprintf(stderr,
"dt_heal_cl: error allocating memory for healing\n");
487 if(err != CL_SUCCESS)
494 if(err != CL_SUCCESS)
503 if(err != CL_SUCCESS)
void cleanup(dt_imageio_module_format_t *self)
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)
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 int max_iter)
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)
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 int max_iter)
heal_params_cl_t * dt_heal_init_cl(const int devid)
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)
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 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)
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)
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,...)
dt_heal_cl_global_t * global