2 * This file is part of FFmpeg.
4 * FFmpeg is free software; you can redistribute it and/or
5 * modify it under the terms of the GNU Lesser General Public
6 * License as published by the Free Software Foundation; either
7 * version 2.1 of the License, or (at your option) any later version.
9 * FFmpeg is distributed in the hope that it will be useful,
10 * but WITHOUT ANY WARRANTY; without even the implied warranty of
11 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
12 * Lesser General Public License for more details.
14 * You should have received a copy of the GNU Lesser General Public
15 * License along with FFmpeg; if not, write to the Free Software
16 * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
20 #include "libavutil/avassert.h"
21 #include "libavutil/common.h"
22 #include "libavutil/imgutils.h"
23 #include "libavutil/mem.h"
24 #include "libavutil/opt.h"
25 #include "libavutil/pixdesc.h"
30 #include "opencl_source.h"
34 // the integral image may overflow 32bit, consider using 64bit
36 static const enum AVPixelFormat supported_formats[] = {
42 static int is_format_supported(enum AVPixelFormat fmt)
46 for (i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++)
47 if (supported_formats[i] == fmt)
52 typedef struct NLMeansOpenCLContext {
53 OpenCLFilterContext ocf;
55 cl_kernel vert_kernel;
56 cl_kernel horiz_kernel;
57 cl_kernel accum_kernel;
58 cl_kernel average_kernel;
62 cl_mem overflow; // overflow in integral image?
71 cl_command_queue command_queue;
72 } NLMeansOpenCLContext;
74 static int nlmeans_opencl_init(AVFilterContext *avctx, int width, int height)
76 NLMeansOpenCLContext *ctx = avctx->priv;
79 int weight_buf_size = width * height * sizeof(float);
81 ctx->h = ctx->sigma * 10;
82 if (!(ctx->research_size & 1)) {
83 ctx->research_size |= 1;
84 av_log(avctx, AV_LOG_WARNING,
85 "research_size should be odd, set to %d",
89 if (!(ctx->patch_size & 1)) {
91 av_log(avctx, AV_LOG_WARNING,
92 "patch_size should be odd, set to %d",
96 if (!ctx->research_size_uv)
97 ctx->research_size_uv = ctx->research_size;
98 if (!ctx->patch_size_uv)
99 ctx->patch_size_uv = ctx->patch_size;
101 err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_nlmeans, 1);
105 ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
106 ctx->ocf.hwctx->device_id,
108 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
109 "command queue %d.\n", cle);
111 ctx->vert_kernel = clCreateKernel(ctx->ocf.program,
113 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
114 "vert_sum kernel %d.\n", cle);
116 ctx->horiz_kernel = clCreateKernel(ctx->ocf.program,
118 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
119 "horiz_sum kernel %d.\n", cle);
121 ctx->accum_kernel = clCreateKernel(ctx->ocf.program,
122 "weight_accum", &cle);
123 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
124 "accum kernel %d.\n", cle);
126 ctx->average_kernel = clCreateKernel(ctx->ocf.program,
128 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
129 "average kernel %d.\n", cle);
131 ctx->integral_img = clCreateBuffer(ctx->ocf.hwctx->context, 0,
132 4 * width * height * sizeof(cl_int),
134 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
135 "integral image %d.\n", cle);
137 ctx->weight = clCreateBuffer(ctx->ocf.hwctx->context, 0,
138 weight_buf_size, NULL, &cle);
139 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
140 "weight buffer %d.\n", cle);
142 ctx->sum = clCreateBuffer(ctx->ocf.hwctx->context, 0,
143 weight_buf_size, NULL, &cle);
144 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
145 "sum buffer %d.\n", cle);
147 ctx->overflow = clCreateBuffer(ctx->ocf.hwctx->context, 0,
148 sizeof(cl_int), NULL, &cle);
149 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
150 "overflow buffer %d.\n", cle);
152 ctx->initialised = 1;
156 CL_RELEASE_KERNEL(ctx->vert_kernel);
157 CL_RELEASE_KERNEL(ctx->horiz_kernel);
158 CL_RELEASE_KERNEL(ctx->accum_kernel);
159 CL_RELEASE_KERNEL(ctx->average_kernel);
161 CL_RELEASE_MEMORY(ctx->integral_img);
162 CL_RELEASE_MEMORY(ctx->weight);
163 CL_RELEASE_MEMORY(ctx->sum);
164 CL_RELEASE_MEMORY(ctx->overflow);
166 CL_RELEASE_QUEUE(ctx->command_queue);
170 static int nlmeans_plane(AVFilterContext *avctx, cl_mem dst, cl_mem src,
171 cl_int width, cl_int height, cl_int p, cl_int r)
173 NLMeansOpenCLContext *ctx = avctx->priv;
174 const float zero = 0.0f;
175 const size_t worksize1[] = {height};
176 const size_t worksize2[] = {width};
177 const size_t worksize3[2] = {width, height};
178 int i, dx, dy, err = 0, weight_buf_size;
180 int nb_pixel, *tmp = NULL, idx = 0;
183 weight_buf_size = width * height * sizeof(float);
184 cle = clEnqueueFillBuffer(ctx->command_queue, ctx->weight,
185 &zero, sizeof(float), 0, weight_buf_size,
187 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to fill weight buffer: %d.\n",
189 cle = clEnqueueFillBuffer(ctx->command_queue, ctx->sum,
190 &zero, sizeof(float), 0, weight_buf_size,
192 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to fill sum buffer: %d.\n",
195 nb_pixel = (2 * r + 1) * (2 * r + 1) - 1;
196 dxdy = av_malloc(nb_pixel * 2 * sizeof(cl_int));
197 tmp = av_malloc(nb_pixel * 2 * sizeof(int));
202 for (dx = -r; dx <= r; dx++) {
203 for (dy = -r; dy <= r; dy++) {
210 // repack dx/dy seperately, as we want to do four pairs of dx/dy in a batch
211 for (i = 0; i < nb_pixel / 4; i++) {
212 dxdy[i * 8] = tmp[i * 8]; // dx0
213 dxdy[i * 8 + 1] = tmp[i * 8 + 2]; // dx1
214 dxdy[i * 8 + 2] = tmp[i * 8 + 4]; // dx2
215 dxdy[i * 8 + 3] = tmp[i * 8 + 6]; // dx3
216 dxdy[i * 8 + 4] = tmp[i * 8 + 1]; // dy0
217 dxdy[i * 8 + 5] = tmp[i * 8 + 3]; // dy1
218 dxdy[i * 8 + 6] = tmp[i * 8 + 5]; // dy2
219 dxdy[i * 8 + 7] = tmp[i * 8 + 7]; // dy3
223 for (i = 0; i < nb_pixel / 4; i++) {
224 cl_int *dx_cur = dxdy + 8 * i;
225 cl_int *dy_cur = dxdy + 8 * i + 4;
228 // integral(x,y) = sum([u(v,y) - u(v+dx,y+dy)]^2) for v in [0, x]
229 CL_SET_KERNEL_ARG(ctx->horiz_kernel, 0, cl_mem, &ctx->integral_img);
230 CL_SET_KERNEL_ARG(ctx->horiz_kernel, 1, cl_mem, &src);
231 CL_SET_KERNEL_ARG(ctx->horiz_kernel, 2, cl_int, &width);
232 CL_SET_KERNEL_ARG(ctx->horiz_kernel, 3, cl_int, &height);
233 CL_SET_KERNEL_ARG(ctx->horiz_kernel, 4, cl_int4, dx_cur);
234 CL_SET_KERNEL_ARG(ctx->horiz_kernel, 5, cl_int4, dy_cur);
235 cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->horiz_kernel, 1,
236 NULL, worksize1, NULL, 0, NULL, NULL);
237 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue horiz_kernel: %d.\n",
240 // integral(x, y) = sum(integral(x, v)) for v in [0, y]
241 CL_SET_KERNEL_ARG(ctx->vert_kernel, 0, cl_mem, &ctx->integral_img);
242 CL_SET_KERNEL_ARG(ctx->vert_kernel, 1, cl_mem, &ctx->overflow);
243 CL_SET_KERNEL_ARG(ctx->vert_kernel, 2, cl_int, &width);
244 CL_SET_KERNEL_ARG(ctx->vert_kernel, 3, cl_int, &height);
245 cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->vert_kernel,
246 1, NULL, worksize2, NULL, 0, NULL, NULL);
247 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue vert_kernel: %d.\n",
250 // accumulate weights
251 CL_SET_KERNEL_ARG(ctx->accum_kernel, 0, cl_mem, &ctx->sum);
252 CL_SET_KERNEL_ARG(ctx->accum_kernel, 1, cl_mem, &ctx->weight);
253 CL_SET_KERNEL_ARG(ctx->accum_kernel, 2, cl_mem, &ctx->integral_img);
254 CL_SET_KERNEL_ARG(ctx->accum_kernel, 3, cl_mem, &src);
255 CL_SET_KERNEL_ARG(ctx->accum_kernel, 4, cl_int, &width);
256 CL_SET_KERNEL_ARG(ctx->accum_kernel, 5, cl_int, &height);
257 CL_SET_KERNEL_ARG(ctx->accum_kernel, 6, cl_int, &p);
258 CL_SET_KERNEL_ARG(ctx->accum_kernel, 7, cl_float, &ctx->h);
259 CL_SET_KERNEL_ARG(ctx->accum_kernel, 8, cl_int4, dx_cur);
260 CL_SET_KERNEL_ARG(ctx->accum_kernel, 9, cl_int4, dy_cur);
261 cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->accum_kernel,
262 2, NULL, worksize3, NULL, 0, NULL, NULL);
263 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
268 CL_SET_KERNEL_ARG(ctx->average_kernel, 0, cl_mem, &dst);
269 CL_SET_KERNEL_ARG(ctx->average_kernel, 1, cl_mem, &src);
270 CL_SET_KERNEL_ARG(ctx->average_kernel, 2, cl_mem, &ctx->sum);
271 CL_SET_KERNEL_ARG(ctx->average_kernel, 3, cl_mem, &ctx->weight);
272 cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->average_kernel, 2,
273 NULL, worksize3, NULL, 0, NULL, NULL);
274 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue average kernel: %d.\n",
276 cle = clFlush(ctx->command_queue);
277 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to flush command queue: %d.\n", cle);
286 static int nlmeans_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
288 AVFilterContext *avctx = inlink->dst;
289 AVFilterLink *outlink = avctx->outputs[0];
290 NLMeansOpenCLContext *ctx = avctx->priv;
291 AVFrame *output = NULL;
292 AVHWFramesContext *input_frames_ctx;
293 const AVPixFmtDescriptor *desc;
294 enum AVPixelFormat in_format;
296 const cl_int zero = 0;
297 int w, h, err, cle, overflow, p, patch, research;
299 av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n",
300 av_get_pix_fmt_name(input->format),
301 input->width, input->height, input->pts);
303 if (!input->hw_frames_ctx)
304 return AVERROR(EINVAL);
305 input_frames_ctx = (AVHWFramesContext*)input->hw_frames_ctx->data;
306 in_format = input_frames_ctx->sw_format;
308 output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
310 err = AVERROR(ENOMEM);
314 err = av_frame_copy_props(output, input);
318 if (!ctx->initialised) {
319 desc = av_pix_fmt_desc_get(in_format);
320 if (!is_format_supported(in_format)) {
321 err = AVERROR(EINVAL);
322 av_log(avctx, AV_LOG_ERROR, "input format %s not supported\n",
323 av_get_pix_fmt_name(in_format));
326 ctx->chroma_w = AV_CEIL_RSHIFT(inlink->w, desc->log2_chroma_w);
327 ctx->chroma_h = AV_CEIL_RSHIFT(inlink->h, desc->log2_chroma_h);
329 err = nlmeans_opencl_init(avctx, inlink->w, inlink->h);
334 cle = clEnqueueWriteBuffer(ctx->command_queue, ctx->overflow, CL_FALSE,
335 0, sizeof(cl_int), &zero, 0, NULL, NULL);
336 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to initialize overflow"
337 "detection buffer %d.\n", cle);
339 for (p = 0; p < FF_ARRAY_ELEMS(output->data); p++) {
340 src = (cl_mem) input->data[p];
341 dst = (cl_mem) output->data[p];
346 w = p ? ctx->chroma_w : inlink->w;
347 h = p ? ctx->chroma_h : inlink->h;
348 patch = (p ? ctx->patch_size_uv : ctx->patch_size) / 2;
349 research = (p ? ctx->research_size_uv : ctx->research_size) / 2;
350 err = nlmeans_plane(avctx, dst, src, w, h, patch, research);
354 // overflow occurred?
355 cle = clEnqueueReadBuffer(ctx->command_queue, ctx->overflow, CL_FALSE,
356 0, sizeof(cl_int), &overflow, 0, NULL, NULL);
357 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to read overflow: %d.\n", cle);
359 cle = clFinish(ctx->command_queue);
360 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish kernel: %d.\n", cle);
363 av_log(avctx, AV_LOG_ERROR, "integral image overflow %d\n", overflow);
365 av_frame_free(&input);
367 av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
368 av_get_pix_fmt_name(output->format),
369 output->width, output->height, output->pts);
371 return ff_filter_frame(outlink, output);
374 clFinish(ctx->command_queue);
375 av_frame_free(&input);
376 av_frame_free(&output);
380 static av_cold void nlmeans_opencl_uninit(AVFilterContext *avctx)
382 NLMeansOpenCLContext *ctx = avctx->priv;
385 CL_RELEASE_KERNEL(ctx->vert_kernel);
386 CL_RELEASE_KERNEL(ctx->horiz_kernel);
387 CL_RELEASE_KERNEL(ctx->accum_kernel);
388 CL_RELEASE_KERNEL(ctx->average_kernel);
390 CL_RELEASE_MEMORY(ctx->integral_img);
391 CL_RELEASE_MEMORY(ctx->weight);
392 CL_RELEASE_MEMORY(ctx->sum);
393 CL_RELEASE_MEMORY(ctx->overflow);
395 CL_RELEASE_QUEUE(ctx->command_queue);
397 ff_opencl_filter_uninit(avctx);
400 #define OFFSET(x) offsetof(NLMeansOpenCLContext, x)
401 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
402 static const AVOption nlmeans_opencl_options[] = {
403 { "s", "denoising strength", OFFSET(sigma), AV_OPT_TYPE_DOUBLE, { .dbl = 1.0 }, 1.0, 30.0, FLAGS },
404 { "p", "patch size", OFFSET(patch_size), AV_OPT_TYPE_INT, { .i64 = 2*3+1 }, 0, 99, FLAGS },
405 { "pc", "patch size for chroma planes", OFFSET(patch_size_uv), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, 99, FLAGS },
406 { "r", "research window", OFFSET(research_size), AV_OPT_TYPE_INT, { .i64 = 7*2+1 }, 0, 99, FLAGS },
407 { "rc", "research window for chroma planes", OFFSET(research_size_uv), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, 99, FLAGS },
411 AVFILTER_DEFINE_CLASS(nlmeans_opencl);
413 static const AVFilterPad nlmeans_opencl_inputs[] = {
416 .type = AVMEDIA_TYPE_VIDEO,
417 .filter_frame = &nlmeans_opencl_filter_frame,
418 .config_props = &ff_opencl_filter_config_input,
423 static const AVFilterPad nlmeans_opencl_outputs[] = {
426 .type = AVMEDIA_TYPE_VIDEO,
427 .config_props = &ff_opencl_filter_config_output,
432 const AVFilter ff_vf_nlmeans_opencl = {
433 .name = "nlmeans_opencl",
434 .description = NULL_IF_CONFIG_SMALL("Non-local means denoiser through OpenCL"),
435 .priv_size = sizeof(NLMeansOpenCLContext),
436 .priv_class = &nlmeans_opencl_class,
437 .init = &ff_opencl_filter_init,
438 .uninit = &nlmeans_opencl_uninit,
439 .query_formats = &ff_opencl_filter_query_formats,
440 .inputs = nlmeans_opencl_inputs,
441 .outputs = nlmeans_opencl_outputs,
442 .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,