From 799e270a35d5e1ba70f52e3cf34b2643862597b8 Mon Sep 17 00:00:00 2001 From: Hanno Schwalm Date: Sat, 20 Jun 2026 13:37:33 +0200 Subject: [PATCH] Implement and use OpenCL buffer and image fill functions OpenCL 1.2 offers clEnqueueFillBuffer(), clEnqueueFillImage() and clEnqueueMigrateMemObjects(), they now are linked and made available as - dt_clEnqueueFillBuffer - dt_clEnqueueFillImage - dt_clEnqueueMigrateMemObjects Having the fill functions instead of specific kernels - simplifies code - results in some small perf gains as both are faster than the specific kernels. --- data/kernels/atrous.cl | 11 ------ data/kernels/basecurve.cl | 12 ------- data/kernels/capture.cl | 12 ------- data/kernels/demosaic_markesteijn.cl | 14 -------- data/kernels/denoiseprofile.cl | 14 -------- data/kernels/dwt.cl | 13 ------- data/kernels/nlmeans.cl | 12 ------- src/common/dlopencl.c | 6 ++++ src/common/dlopencl.h | 7 ++++ src/common/dwt.c | 8 ++--- src/common/dwt.h | 1 - src/common/nlmeans_core.c | 16 ++------- src/common/nlmeans_core.h | 1 - src/common/opencl.c | 52 ++++++++++++++++++++++++++++ src/common/opencl.h | 10 ++++++ src/iop/atrous.c | 7 ++-- src/iop/basecurve.c | 7 ++-- src/iop/demosaic.c | 6 ---- src/iop/demosaicing/capture.c | 3 +- src/iop/demosaicing/xtrans.c | 7 ++-- src/iop/denoiseprofile.c | 8 +---- src/iop/nlmeans.c | 7 +--- 22 files changed, 90 insertions(+), 144 deletions(-) diff --git a/data/kernels/atrous.cl b/data/kernels/atrous.cl index fe7baf144905..5a6d97cab48f 100644 --- a/data/kernels/atrous.cl +++ b/data/kernels/atrous.cl @@ -87,17 +87,6 @@ eaw_synthesize(__write_only image2d_t out, write_imagef (out, (int2)(x, y), sum); } -__kernel void -eaw_zero(__write_only image2d_t out, - const int width, - const int height) -{ - const int x = get_global_id(0); - const int y = get_global_id(1); - if(x >= width || y >= height) return; - write_imagef(out, (int2)(x, y), (float4)0.0f); -} - __kernel void eaw_addbuffers(__write_only image2d_t out_out, __read_only image2d_t out_in, diff --git a/data/kernels/basecurve.cl b/data/kernels/basecurve.cl index 24a675fe73ab..2e8caf144727 100644 --- a/data/kernels/basecurve.cl +++ b/data/kernels/basecurve.cl @@ -55,18 +55,6 @@ basecurve_lut(read_only image2d_t in, write_only image2d_t out, const int width, write_imagef (out, (int2)(x, y), pixel); } - -kernel void -basecurve_zero(write_only image2d_t out, const int width, const int height) -{ - const int x = get_global_id(0); - const int y = get_global_id(1); - - if(x >= width || y >= height) return; - - write_imagef (out, (int2)(x, y), (float4)0.0f); -} - /* Original basecurve implementation. Applies a LUT on a per-channel basis which can cause color shifts. diff --git a/data/kernels/capture.cl b/data/kernels/capture.cl index 8bbf13851110..9c73be04321f 100644 --- a/data/kernels/capture.cl +++ b/data/kernels/capture.cl @@ -167,18 +167,6 @@ __kernel void kernel_9x9_div(global float *in, out[i] = luminance[i] / fmax(val, CAPTURE_YMIN); } -__kernel void prefill_clip_mask(global float *mask, - const int width, - const int height) -{ - const int col = get_global_id(0); - const int row = get_global_id(1); - if(col >= width || row >= height) return; - - const int i = mad24(row, width, col); - mask[i] = 1.0f; -} - __kernel void prepare_blend(__read_only image2d_t cfa, __read_only image2d_t dev_out, const int filters, diff --git a/data/kernels/demosaic_markesteijn.cl b/data/kernels/demosaic_markesteijn.cl index fe793736b0f0..58a6ea7c4848 100644 --- a/data/kernels/demosaic_markesteijn.cl +++ b/data/kernels/demosaic_markesteijn.cl @@ -869,20 +869,6 @@ markesteijn_homo_quench(global uchar *homosum1, global uchar *homosum2, const in homosum2[glidx] = hmo2; } -// Initialize output image to zero -kernel void -markesteijn_zero(write_only image2d_t out, const int width, const int height, const int border) -{ - const int x = get_global_id(0); - const int y = get_global_id(1); - - // take sufficient border into account - if(x < border || x >= width-border || y < border || y >= height-border) return; - - write_imagef(out, (int2)(x, y), (float4)0.0f); -} - - // accumulate contributions of all directions into output image kernel void markesteijn_accu(read_only image2d_t in, write_only image2d_t out, global float *rgb, diff --git a/data/kernels/denoiseprofile.cl b/data/kernels/denoiseprofile.cl index c4219cb481c6..f54d076948d6 100644 --- a/data/kernels/denoiseprofile.cl +++ b/data/kernels/denoiseprofile.cl @@ -113,20 +113,6 @@ kernel void denoiseprofile_precondition_Y0U0V0(read_only image2d_t in, write_imagef (out, (int2)(x, y), outpx); } - -kernel void -denoiseprofile_init(global float4* out, const int width, const int height) -{ - const int x = get_global_id(0); - const int y = get_global_id(1); - const int gidx = mad24(y, width, x); - - if(x >= width || y >= height) return; - - out[gidx] = (float4)0.0f; -} - - kernel void denoiseprofile_dist(read_only image2d_t in, global float* U4, const int width, const int height, const int2 q) diff --git a/data/kernels/dwt.cl b/data/kernels/dwt.cl index 642383016405..198361c09ee3 100644 --- a/data/kernels/dwt.cl +++ b/data/kernels/dwt.cl @@ -114,16 +114,3 @@ dwt_hat_transform_col(global float4 *lpass, int width, int height, const int sc, * lpass_mult; } } - -kernel void -dwt_init_buffer(global float4 *buffer, int width, int height) -{ - const int x = get_global_id(0); - const int y = get_global_id(1); - - if(x >= width || y >= height) return; - - const int idx = mad24(y, width, x); - - buffer[idx] = 0.f; -} diff --git a/data/kernels/nlmeans.cl b/data/kernels/nlmeans.cl index ee45c1db3587..5789ce2c1756 100644 --- a/data/kernels/nlmeans.cl +++ b/data/kernels/nlmeans.cl @@ -34,18 +34,6 @@ static inline float ddirac(const int2 q) return ((q.x || q.y) ? 1.0f : 0.0f); } -kernel void nlmeans_init(global float4* out, - const int width, - const int height) -{ - const int x = get_global_id(0); - const int y = get_global_id(1); - if(x >= width || y >= height) return; - - out[mad24(y, width, x)] = (float4)0.0f; -} - - kernel void nlmeans_dist(read_only image2d_t in, global float *U4, const int width, diff --git a/src/common/dlopencl.c b/src/common/dlopencl.c index effa7d727d65..f1e2c1002c62 100644 --- a/src/common/dlopencl.c +++ b/src/common/dlopencl.c @@ -225,6 +225,12 @@ dt_dlopencl_t *dt_dlopencl_init(const char *name) (void (**)(void)) & ocl->symbols->dt_clGetMemObjectInfo); success = success && dt_gmodule_symbol(module, "clGetImageInfo", ((void (**)(void)) & ocl->symbols->dt_clGetImageInfo)); + success = success && dt_gmodule_symbol(module, "clEnqueueMigrateMemObjects", + ((void (**)(void)) & ocl->symbols->dt_clEnqueueMigrateMemObjects)); + success = success && dt_gmodule_symbol(module, "clEnqueueFillBuffer", + ((void (**)(void)) & ocl->symbols->dt_clEnqueueFillBuffer)); + success = success && dt_gmodule_symbol(module, "clEnqueueFillImage", + ((void (**)(void)) & ocl->symbols->dt_clEnqueueFillImage)); } ocl->have_opencl = success; diff --git a/src/common/dlopencl.h b/src/common/dlopencl.h index 5f93c1451b4c..42d3fe7313a5 100644 --- a/src/common/dlopencl.h +++ b/src/common/dlopencl.h @@ -133,6 +133,10 @@ typedef cl_int (*dt_clEnqueueNativeKernel_t)(cl_command_queue, void (*user_func) typedef cl_int (*dt_clEnqueueMarker_t)(cl_command_queue, cl_event *); typedef cl_int (*dt_clEnqueueWaitForEvents_t)(cl_command_queue, cl_uint, const cl_event *); +typedef cl_int (*dt_clEnqueueMigrateMemObjects_t)(cl_command_queue, cl_uint, const cl_mem*, cl_mem_migration_flags, cl_uint, const cl_event*, cl_event*); +typedef cl_int (*dt_clEnqueueFillBuffer_t)(cl_command_queue, cl_mem, const void *, size_t, size_t, size_t, cl_uint, const cl_event *, cl_event *); +typedef cl_int (*dt_clEnqueueFillImage_t)(cl_command_queue, cl_mem, const void *, const size_t *, const size_t *, cl_uint, const cl_event *, cl_event *); + typedef struct dt_dlopencl_symbols_t { dt_clGetPlatformIDs_t dt_clGetPlatformIDs; @@ -206,6 +210,9 @@ typedef struct dt_dlopencl_symbols_t dt_clEnqueueNativeKernel_t dt_clEnqueueNativeKernel; dt_clEnqueueMarker_t dt_clEnqueueMarker; dt_clEnqueueWaitForEvents_t dt_clEnqueueWaitForEvents; + dt_clEnqueueMigrateMemObjects_t dt_clEnqueueMigrateMemObjects; + dt_clEnqueueFillBuffer_t dt_clEnqueueFillBuffer; + dt_clEnqueueFillImage_t dt_clEnqueueFillImage; } dt_dlopencl_symbols_t; diff --git a/src/common/dwt.c b/src/common/dwt.c index b13d249ef325..5b487246bec1 100644 --- a/src/common/dwt.c +++ b/src/common/dwt.c @@ -549,7 +549,6 @@ dt_dwt_cl_global_t *dt_dwt_init_cl_global() g->kernel_dwt_subtract_layer = dt_opencl_create_kernel(program, "dwt_subtract_layer"); g->kernel_dwt_hat_transform_col = dt_opencl_create_kernel(program, "dwt_hat_transform_col"); g->kernel_dwt_hat_transform_row = dt_opencl_create_kernel(program, "dwt_hat_transform_row"); - g->kernel_dwt_init_buffer = dt_opencl_create_kernel(program, "dwt_init_buffer"); return g; } @@ -562,7 +561,6 @@ void dt_dwt_free_cl_global(dt_dwt_cl_global_t *g) dt_opencl_free_kernel(g->kernel_dwt_subtract_layer); dt_opencl_free_kernel(g->kernel_dwt_hat_transform_col); dt_opencl_free_kernel(g->kernel_dwt_hat_transform_row); - dt_opencl_free_kernel(g->kernel_dwt_init_buffer); free(g); } @@ -672,8 +670,7 @@ static cl_int dwt_wavelet_decompose_cl(cl_mem img, dwt_params_cl_t *const p, _dw if(layers == NULL) goto cleanup; // init layer buffer - err = dt_opencl_enqueue_kernel_2d_args(devid, p->global->kernel_dwt_init_buffer, p->width, p->height, - CLARG(layers), CLARG(p->width), CLARG(p->height)); + err = dt_opencl_fill_buffer(devid, layers, (size_t)p->width * p->height, p->ch, 0.0f); if(err != CL_SUCCESS) goto cleanup; if(p->merge_from_scale > 0) @@ -682,8 +679,7 @@ static cl_int dwt_wavelet_decompose_cl(cl_mem img, dwt_params_cl_t *const p, _dw if(merged_layers == NULL) goto cleanup; // init reconstruct buffer - err = dt_opencl_enqueue_kernel_2d_args(devid, p->global->kernel_dwt_init_buffer, p->width, p->height, - CLARG(merged_layers), CLARG(p->width), CLARG(p->height)); + err = dt_opencl_fill_buffer(devid, merged_layers, (size_t)p->width * p->height, p->ch, 0.0f); if(err != CL_SUCCESS) goto cleanup; } diff --git a/src/common/dwt.h b/src/common/dwt.h index c8c88b45ed3c..67704f3664df 100644 --- a/src/common/dwt.h +++ b/src/common/dwt.h @@ -106,7 +106,6 @@ typedef struct dt_dwt_cl_global_t int kernel_dwt_subtract_layer; int kernel_dwt_hat_transform_col; int kernel_dwt_hat_transform_row; - int kernel_dwt_init_buffer; } dt_dwt_cl_global_t; typedef struct dwt_params_cl_t diff --git a/src/common/nlmeans_core.c b/src/common/nlmeans_core.c index 01a7d9e2c131..11663b2570e4 100644 --- a/src/common/nlmeans_core.c +++ b/src/common/nlmeans_core.c @@ -583,18 +583,6 @@ static void get_blocksizes( return; } -// zero output pixels, as we will be accumulating them one patch at a time -static inline cl_int nlmeans_cl_init( - const int devid, - const int kernel, - cl_mem dev_out, - const int height, - const int width) -{ - return dt_opencl_enqueue_kernel_2d_args(devid, kernel, width, height, - CLARG(dev_out), CLARG(width), CLARG(height)); -} - // horizontal pass, add together columns of each patch static inline cl_int nlmeans_cl_horiz( const int devid, @@ -663,7 +651,7 @@ int nlmeans_denoise_cl( get_blocksizes(&hblocksize, &vblocksize, P, devid, params->kernel_horiz, params->kernel_vert); // zero the output buffer into which we will be accumulating results - err = nlmeans_cl_init(devid,params->kernel_init,dev_out,height,width); + err = dt_opencl_fill_buffer(devid, dev_out, (size_t)width * height, 4, 0.0f); if(err != CL_SUCCESS) goto error; const size_t bwidth = ROUNDUP(width, hblocksize); @@ -748,7 +736,7 @@ int nlmeans_denoiseprofile_cl( get_blocksizes(&hblocksize, &vblocksize, P, devid, params->kernel_horiz, params->kernel_vert); // zero the output buffer into which we will be accumulating results - err = nlmeans_cl_init(devid,params->kernel_init,dev_out,height,width); + err = dt_opencl_fill_buffer(devid, dev_out, (size_t)width * height, 4, 0.0f); if(err != CL_SUCCESS) goto error; const size_t bwidth = ROUNDUP(width, hblocksize); diff --git a/src/common/nlmeans_core.h b/src/common/nlmeans_core.h index 13037c102dd7..766c89c18b91 100644 --- a/src/common/nlmeans_core.h +++ b/src/common/nlmeans_core.h @@ -32,7 +32,6 @@ struct dt_nlmeans_param_t int decimate; // set to 1 to search only half the patches in the neighborhood (default = 0) const float* const norm; // array of four per-channel weight factors dt_dev_pixelpipe_type_t pipetype; - int kernel_init; // CL: initialization (runs once) int kernel_dist; // CL: compute channel-normed squared pixel differences (runs for each patch) int kernel_horiz; // CL: horizontal sum (runs for each patch) int kernel_vert; // CL: vertical sum (runs for each patch) diff --git a/src/common/opencl.c b/src/common/opencl.c index d93dc43b04a1..e59bdda0da26 100644 --- a/src/common/opencl.c +++ b/src/common/opencl.c @@ -3359,6 +3359,58 @@ int dt_opencl_unmap_mem_object(const int devid, return err; } +int dt_opencl_fill_buffer(const int devid, + cl_mem buffer, + const size_t pts, + const size_t ch, + const float val) +{ + if(!_cldev_running(devid)) + return DT_OPENCL_NODEVICE; + if(ch < 1 || ch > 4) + return DT_OPENCL_PROCESS_CL; + + cl_event *eventp = _opencl_events_get_slot(devid, "[Fill Float Buffer]"); + const float filler[4] = { val, val, val, val }; + const size_t psize = sizeof(float) * ch; + + const cl_int err = (darktable.opencl->dlocl->symbols->dt_clEnqueueFillBuffer) + (darktable.opencl->dev[devid].cmd_queue, buffer, + &filler, psize, 0, pts * psize, 0, NULL, eventp); + + if(err != CL_SUCCESS) + dt_print(DT_DEBUG_OPENCL, + "[dt_opencl_fill_float%i_buffer] could not fill buffer on device '%s' id=%d: %s", + (int)ch, darktable.opencl->dev[devid].fullname, devid, cl_errstr(err)); + return err; +} + +int dt_opencl_fill_image(const int devid, + cl_mem image, + const size_t *origin, + const size_t *region, + const float val) +{ + if(!_cldev_running(devid)) + return DT_OPENCL_NODEVICE; + + cl_event *eventp = _opencl_events_get_slot(devid, "[Fill Image]"); + const float filler[4] = { val, val, val, val }; + + const size_t org[3] = { origin ? origin[0] : 0, origin ? origin[1] : 0, 0 }; + const size_t reg[3] = { region[0], region[1], 1 }; + + const cl_int err = (darktable.opencl->dlocl->symbols->dt_clEnqueueFillImage) + (darktable.opencl->dev[devid].cmd_queue, image, filler, + org, reg, 0, NULL, eventp); + + if(err != CL_SUCCESS) + dt_print(DT_DEBUG_OPENCL, + "[dt_opencl_fill_image] could not fill image on device '%s' id=%d: %s", + darktable.opencl->dev[devid].fullname, devid, cl_errstr(err)); + return err; +} + void *dt_opencl_alloc_device(const int devid, const int width, const int height, diff --git a/src/common/opencl.h b/src/common/opencl.h index 34a188d8abaa..2f5443bfdcd5 100644 --- a/src/common/opencl.h +++ b/src/common/opencl.h @@ -465,6 +465,16 @@ int dt_opencl_write_host_to_image_raw(const int devid, const int rowpitch, const gboolean blocking); +int dt_opencl_fill_buffer(const int devid, + cl_mem buffer, + const size_t pts, + const size_t ch, + const float val); +int dt_opencl_fill_image(const int devid, + cl_mem image, + const size_t *orig, + const size_t *area, + const float val); void *dt_opencl_copy_host_to_image(const int devid, void *host, const int width, diff --git a/src/iop/atrous.c b/src/iop/atrous.c index a953d74f85d5..c0092b2eb7b4 100644 --- a/src/iop/atrous.c +++ b/src/iop/atrous.c @@ -103,7 +103,6 @@ typedef struct dt_iop_atrous_gui_data_t typedef struct dt_iop_atrous_global_data_t { - int kernel_zero; int kernel_decompose; int kernel_synthesize; int kernel_addbuffers; @@ -412,8 +411,8 @@ int process_cl(dt_iop_module_t *self, if(!dev_detail || !dev_tmp || !dev_tmp2 || !dev_filter) goto error; // clear dev_out to zeros, as we will be incrementally accumulating results there - err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_zero, width, height, - CLARG(dev_out), CLARG(width), CLARG(height)); + const size_t area[2] = { width, height }; + err = dt_opencl_fill_image(devid, dev_out, CLIMG_ORIGIN, area, 0.0f); if(err != CL_SUCCESS) goto error; // the buffers for the buffer ping-pong. We start with dev_in as @@ -627,7 +626,6 @@ void init_global(dt_iop_module_so_t *self) gd->kernel_decompose = dt_opencl_create_kernel(program, "eaw_decompose"); gd->kernel_synthesize = dt_opencl_create_kernel(program, "eaw_synthesize"); #ifdef USE_NEW_CL - gd->kernel_zero = dt_opencl_create_kernel(program, "eaw_zero"); gd->kernel_addbuffers = dt_opencl_create_kernel(program, "eaw_addbuffers"); #endif } @@ -638,7 +636,6 @@ void cleanup_global(dt_iop_module_so_t *self) dt_opencl_free_kernel(gd->kernel_decompose); dt_opencl_free_kernel(gd->kernel_synthesize); #ifdef USE_NEW_CL - dt_opencl_free_kernel(gd->kernel_zero); dt_opencl_free_kernel(gd->kernel_addbuffers); #endif free(self->data); diff --git a/src/iop/basecurve.c b/src/iop/basecurve.c index dc84532889a4..30f3fdbb7a5b 100644 --- a/src/iop/basecurve.c +++ b/src/iop/basecurve.c @@ -340,7 +340,6 @@ typedef struct dt_iop_basecurve_data_t typedef struct dt_iop_basecurve_global_data_t { int kernel_basecurve_lut; - int kernel_basecurve_zero; int kernel_basecurve_legacy_lut; int kernel_basecurve_compute_features; int kernel_basecurve_blur_h; @@ -712,6 +711,7 @@ int process_cl_fusion(dt_iop_module_t *self, // allocate buffers for wavelet transform and blending for(int k = 0, step = 1, w = width, h = height; k < num_levels; k++) { + const size_t area[2] = { w, h }; // coarsest step is some % of image width. dev_col[k] = dt_opencl_alloc_device(devid, w, h, sizeof(float) * 4); if(dev_col[k] == NULL) goto error; @@ -719,8 +719,7 @@ int process_cl_fusion(dt_iop_module_t *self, dev_comb[k] = dt_opencl_alloc_device(devid, w, h, sizeof(float) * 4); if(dev_comb[k] == NULL) goto error; - err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_basecurve_zero, w, h, - CLARG(dev_comb[k]), CLARG(w), CLARG(h)); + err = dt_opencl_fill_image(devid, dev_comb[k], CLIMG_ORIGIN, area, 0.0f); if(err != CL_SUCCESS) goto error; w = (w - 1) / 2 + 1; @@ -1494,7 +1493,6 @@ void init_global(dt_iop_module_so_t *self) dt_iop_basecurve_global_data_t *gd = malloc(sizeof(dt_iop_basecurve_global_data_t)); self->data = gd; gd->kernel_basecurve_lut = dt_opencl_create_kernel(program, "basecurve_lut"); - gd->kernel_basecurve_zero = dt_opencl_create_kernel(program, "basecurve_zero"); gd->kernel_basecurve_legacy_lut = dt_opencl_create_kernel(program, "basecurve_legacy_lut"); gd->kernel_basecurve_compute_features = dt_opencl_create_kernel(program, "basecurve_compute_features"); gd->kernel_basecurve_blur_h = dt_opencl_create_kernel(program, "basecurve_blur_h"); @@ -1514,7 +1512,6 @@ void cleanup_global(dt_iop_module_so_t *self) { dt_iop_basecurve_global_data_t *gd = self->data; dt_opencl_free_kernel(gd->kernel_basecurve_lut); - dt_opencl_free_kernel(gd->kernel_basecurve_zero); dt_opencl_free_kernel(gd->kernel_basecurve_legacy_lut); dt_opencl_free_kernel(gd->kernel_basecurve_compute_features); dt_opencl_free_kernel(gd->kernel_basecurve_blur_h); diff --git a/src/iop/demosaic.c b/src/iop/demosaic.c index 0844af982591..3cb563d55fa2 100644 --- a/src/iop/demosaic.c +++ b/src/iop/demosaic.c @@ -217,7 +217,6 @@ typedef struct dt_iop_demosaic_global_data_t int kernel_markesteijn_homo_max; int kernel_markesteijn_homo_max_corr; int kernel_markesteijn_homo_quench; - int kernel_markesteijn_zero; int kernel_markesteijn_accu; int kernel_markesteijn_final; int kernel_rcd_populate; @@ -233,7 +232,6 @@ typedef struct dt_iop_demosaic_global_data_t int kernel_write_blended_dual; int gaussian_9x9_mul; int gaussian_9x9_div; - int prefill_clip_mask; int prepare_blend; int modify_blend; int show_blend_mask; @@ -1260,7 +1258,6 @@ void init_global(dt_iop_module_so_t *self) gd->kernel_markesteijn_homo_max = dt_opencl_create_kernel(markesteijn, "markesteijn_homo_max"); gd->kernel_markesteijn_homo_max_corr = dt_opencl_create_kernel(markesteijn, "markesteijn_homo_max_corr"); gd->kernel_markesteijn_homo_quench = dt_opencl_create_kernel(markesteijn, "markesteijn_homo_quench"); - gd->kernel_markesteijn_zero = dt_opencl_create_kernel(markesteijn, "markesteijn_zero"); gd->kernel_markesteijn_accu = dt_opencl_create_kernel(markesteijn, "markesteijn_accu"); gd->kernel_markesteijn_final = dt_opencl_create_kernel(markesteijn, "markesteijn_final"); @@ -1280,7 +1277,6 @@ void init_global(dt_iop_module_so_t *self) const int capt = 38; // capture.cl, from programs.conf gd->gaussian_9x9_mul = dt_opencl_create_kernel(capt, "kernel_9x9_mul"); gd->gaussian_9x9_div = dt_opencl_create_kernel(capt, "kernel_9x9_div"); - gd->prefill_clip_mask = dt_opencl_create_kernel(capt, "prefill_clip_mask"); gd->prepare_blend = dt_opencl_create_kernel(capt, "prepare_blend"); gd->modify_blend = dt_opencl_create_kernel(capt, "modify_blend"); gd->show_blend_mask = dt_opencl_create_kernel(capt, "show_blend_mask"); @@ -1326,7 +1322,6 @@ void cleanup_global(dt_iop_module_so_t *self) dt_opencl_free_kernel(gd->kernel_markesteijn_homo_max); dt_opencl_free_kernel(gd->kernel_markesteijn_homo_max_corr); dt_opencl_free_kernel(gd->kernel_markesteijn_homo_quench); - dt_opencl_free_kernel(gd->kernel_markesteijn_zero); dt_opencl_free_kernel(gd->kernel_markesteijn_accu); dt_opencl_free_kernel(gd->kernel_markesteijn_final); dt_opencl_free_kernel(gd->kernel_rcd_populate); @@ -1342,7 +1337,6 @@ void cleanup_global(dt_iop_module_so_t *self) dt_opencl_free_kernel(gd->kernel_write_blended_dual); dt_opencl_free_kernel(gd->gaussian_9x9_mul); dt_opencl_free_kernel(gd->gaussian_9x9_div); - dt_opencl_free_kernel(gd->prefill_clip_mask); dt_opencl_free_kernel(gd->prepare_blend); dt_opencl_free_kernel(gd->modify_blend); dt_opencl_free_kernel(gd->show_blend_mask); diff --git a/src/iop/demosaicing/capture.c b/src/iop/demosaicing/capture.c index 0908ade0b1a7..5edb8aa2bb29 100644 --- a/src/iop/demosaicing/capture.c +++ b/src/iop/demosaicing/capture.c @@ -1062,8 +1062,7 @@ static int _capture_sharpen_cl(dt_iop_module_t *self, if(!blendmask || !luminance || !tmp2 || !tmp1 || !dev_rgb) goto finish; - err = dt_opencl_enqueue_kernel_2d_args(devid, gd->prefill_clip_mask, width, height, - CLARG(tmp2), CLARG(width), CLARG(height)); + err = dt_opencl_fill_buffer(devid, tmp2, pixels, 1, 1.0f); if(err != CL_SUCCESS) goto finish; err = dt_opencl_enqueue_kernel_2d_args(devid, gd->prepare_blend, width, height, diff --git a/src/iop/demosaicing/xtrans.c b/src/iop/demosaicing/xtrans.c index 874ee6464079..da7b90f60f90 100644 --- a/src/iop/demosaicing/xtrans.c +++ b/src/iop/demosaicing/xtrans.c @@ -1991,9 +1991,10 @@ static cl_int process_markesteijn_cl(const dt_iop_module_t *self, if(err != CL_SUCCESS) goto error; } - // initialize output buffer to zero - err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_markesteijn_zero, width, height, - CLARG(dev_out), CLARG(width), CLARG(height), CLARG(pad_tile)); + // initialize output image inner part to zero + const size_t corigin[2] = { pad_tile, pad_tile }; + const size_t carea[2] = { width - 2*pad_tile, height - 2*pad_tile}; + err = dt_opencl_fill_image(devid, dev_out, corigin, carea, 0.0f); if(err != CL_SUCCESS) goto error; // need to get another temp image for the output image diff --git a/src/iop/denoiseprofile.c b/src/iop/denoiseprofile.c index 95a3f951b36a..138f880f9570 100644 --- a/src/iop/denoiseprofile.c +++ b/src/iop/denoiseprofile.c @@ -198,7 +198,6 @@ typedef struct dt_iop_denoiseprofile_global_data_t int kernel_denoiseprofile_precondition; int kernel_denoiseprofile_precondition_v2; int kernel_denoiseprofile_precondition_Y0U0V0; - int kernel_denoiseprofile_init; int kernel_denoiseprofile_dist; int kernel_denoiseprofile_horiz; int kernel_denoiseprofile_vert; @@ -2028,7 +2027,6 @@ static int process_nlmeans_cl(dt_iop_module_t *self, .decimate = 0, .norm = norm2, .pipetype = piece->pipe->type, - .kernel_init = gd->kernel_denoiseprofile_init, .kernel_dist = gd->kernel_denoiseprofile_dist, .kernel_horiz = gd->kernel_denoiseprofile_horiz, .kernel_vert = gd->kernel_denoiseprofile_vert, @@ -2080,8 +2078,7 @@ static int process_nlmeans_cl(dt_iop_module_t *self, else vblocksize = 1; - err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_denoiseprofile_init, width, height, - CLARG(dev_U2), CLARG(width), CLARG(height)); + err = dt_opencl_fill_buffer(devid, dev_U2, (size_t)width * height, 4, 0.0f); if(err != CL_SUCCESS) goto error; const size_t bwidth = ROUNDUP(width, hblocksize); @@ -2748,8 +2745,6 @@ void init_global(dt_iop_module_so_t *self) dt_opencl_create_kernel(program, "denoiseprofile_precondition_v2"); gd->kernel_denoiseprofile_precondition_Y0U0V0 = dt_opencl_create_kernel(program, "denoiseprofile_precondition_Y0U0V0"); - gd->kernel_denoiseprofile_init = - dt_opencl_create_kernel(program, "denoiseprofile_init"); gd->kernel_denoiseprofile_dist = dt_opencl_create_kernel(program, "denoiseprofile_dist"); gd->kernel_denoiseprofile_horiz = @@ -2783,7 +2778,6 @@ void cleanup_global(dt_iop_module_so_t *self) dt_iop_denoiseprofile_global_data_t *gd = self->data; dt_opencl_free_kernel(gd->kernel_denoiseprofile_precondition); dt_opencl_free_kernel(gd->kernel_denoiseprofile_precondition_v2); - dt_opencl_free_kernel(gd->kernel_denoiseprofile_init); dt_opencl_free_kernel(gd->kernel_denoiseprofile_dist); dt_opencl_free_kernel(gd->kernel_denoiseprofile_horiz); dt_opencl_free_kernel(gd->kernel_denoiseprofile_vert); diff --git a/src/iop/nlmeans.c b/src/iop/nlmeans.c index f72f6e55d342..74ba9ca7a7f3 100644 --- a/src/iop/nlmeans.c +++ b/src/iop/nlmeans.c @@ -63,7 +63,6 @@ typedef dt_iop_nlmeans_params_t dt_iop_nlmeans_data_t; typedef struct dt_iop_nlmeans_global_data_t { - int kernel_nlmeans_init; int kernel_nlmeans_dist; int kernel_nlmeans_horiz; int kernel_nlmeans_vert; @@ -201,7 +200,6 @@ int process_cl(dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, cl_mem dev_ .decimate = 0, .norm = norm2, .pipetype = piece->pipe->type, - .kernel_init = gd->kernel_nlmeans_init, .kernel_dist = gd->kernel_nlmeans_dist, .kernel_horiz = gd->kernel_nlmeans_horiz, .kernel_vert = gd->kernel_nlmeans_vert, @@ -261,8 +259,7 @@ int process_cl(dt_iop_module_t *self, dt_dev_pixelpipe_iop_t *piece, cl_mem dev_ size_t sizesl[2]; size_t local[2]; - err = dt_opencl_enqueue_kernel_2d_args(devid, gd->kernel_nlmeans_init, width, height, - CLARG(dev_U2), CLARG(width), CLARG(height)); + err = dt_opencl_fill_buffer(devid, dev_U2, (size_t)width * height, 4, 0.0f); if(err != CL_SUCCESS) goto error; const size_t bwidth = ROUNDUP(width, hblocksize); @@ -389,7 +386,6 @@ void init_global(dt_iop_module_so_t *self) const int program = 5; // nlmeans.cl, from programs.conf dt_iop_nlmeans_global_data_t *gd = malloc(sizeof(dt_iop_nlmeans_global_data_t)); self->data = gd; - gd->kernel_nlmeans_init = dt_opencl_create_kernel(program, "nlmeans_init"); gd->kernel_nlmeans_dist = dt_opencl_create_kernel(program, "nlmeans_dist"); gd->kernel_nlmeans_horiz = dt_opencl_create_kernel(program, "nlmeans_horiz"); gd->kernel_nlmeans_vert = dt_opencl_create_kernel(program, "nlmeans_vert"); @@ -400,7 +396,6 @@ void init_global(dt_iop_module_so_t *self) void cleanup_global(dt_iop_module_so_t *self) { dt_iop_nlmeans_global_data_t *gd = self->data; - dt_opencl_free_kernel(gd->kernel_nlmeans_init); dt_opencl_free_kernel(gd->kernel_nlmeans_dist); dt_opencl_free_kernel(gd->kernel_nlmeans_horiz); dt_opencl_free_kernel(gd->kernel_nlmeans_vert);