From patchwork Tue May 7 02:44:42 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Ruiling Song X-Patchwork-Id: 13019 Return-Path: X-Original-To: patchwork@ffaux-bg.ffmpeg.org Delivered-To: patchwork@ffaux-bg.ffmpeg.org Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org [79.124.17.100]) by ffaux.localdomain (Postfix) with ESMTP id BF645449655 for ; Tue, 7 May 2019 04:46:03 +0300 (EEST) Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id A541468ADF3; Tue, 7 May 2019 04:46:03 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mga01.intel.com (mga01.intel.com [192.55.52.88]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 7C77368AD6B for ; Tue, 7 May 2019 04:45:56 +0300 (EEST) X-Amp-Result: SKIPPED(no attachment in message) X-Amp-File-Uploaded: False Received: from fmsmga008.fm.intel.com ([10.253.24.58]) by fmsmga101.fm.intel.com with ESMTP/TLS/DHE-RSA-AES256-GCM-SHA384; 06 May 2019 18:45:54 -0700 X-ExtLoop1: 1 X-IronPort-AV: E=Sophos;i="5.60,440,1549958400"; d="scan'208";a="147089386" Received: from ruiling-nuc.sh.intel.com ([10.239.158.179]) by fmsmga008.fm.intel.com with ESMTP; 06 May 2019 18:45:53 -0700 From: Ruiling Song To: ffmpeg-devel@ffmpeg.org Date: Tue, 7 May 2019 10:44:42 +0800 Message-Id: <20190507024442.12841-1-ruiling.song@intel.com> X-Mailer: git-send-email 2.17.1 Subject: [FFmpeg-devel] [PATCH V3] lavfi/opencl: add nlmeans_opencl filter X-BeenThere: ffmpeg-devel@ffmpeg.org X-Mailman-Version: 2.1.20 Precedence: list List-Id: FFmpeg development discussions and patches List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Reply-To: FFmpeg development discussions and patches Cc: Ruiling Song MIME-Version: 1.0 Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" Signed-off-by: Ruiling Song --- configure | 1 + doc/filters.texi | 4 + libavfilter/Makefile | 1 + libavfilter/allfilters.c | 1 + libavfilter/opencl/nlmeans.cl | 115 +++++++++ libavfilter/opencl_source.h | 1 + libavfilter/vf_nlmeans_opencl.c | 443 ++++++++++++++++++++++++++++++++ 7 files changed, 566 insertions(+) create mode 100644 libavfilter/opencl/nlmeans.cl create mode 100644 libavfilter/vf_nlmeans_opencl.c diff --git a/configure b/configure index d644a5b1d4..ee4041e5e0 100755 --- a/configure +++ b/configure @@ -3464,6 +3464,7 @@ mpdecimate_filter_select="pixelutils" minterpolate_filter_select="scene_sad" mptestsrc_filter_deps="gpl" negate_filter_deps="lut_filter" +nlmeans_opencl_filter_deps="opencl" nnedi_filter_deps="gpl" ocr_filter_deps="libtesseract" ocv_filter_deps="libopencv" diff --git a/doc/filters.texi b/doc/filters.texi index 3c15bb95f4..ee0db1809e 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -19240,6 +19240,10 @@ Make every semi-green pixel in the input transparent with some slight blending: @end example @end itemize +@section nlmeans_opencl + +Non-local Means denoise filter through OpenCL, this filter accepts same options as @ref{nlmeans}. + @section overlay_opencl Overlay one video on top of another. diff --git a/libavfilter/Makefile b/libavfilter/Makefile index 59d12ce069..3e409fc62c 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -295,6 +295,7 @@ OBJS-$(CONFIG_MIX_FILTER) += vf_mix.o OBJS-$(CONFIG_MPDECIMATE_FILTER) += vf_mpdecimate.o OBJS-$(CONFIG_NEGATE_FILTER) += vf_lut.o OBJS-$(CONFIG_NLMEANS_FILTER) += vf_nlmeans.o +OBJS-$(CONFIG_NLMEANS_OPENCL_FILTER) += vf_nlmeans_opencl.o opencl.o opencl/nlmeans.o OBJS-$(CONFIG_NNEDI_FILTER) += vf_nnedi.o OBJS-$(CONFIG_NOFORMAT_FILTER) += vf_format.o OBJS-$(CONFIG_NOISE_FILTER) += vf_noise.o diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index ae725cb0e0..fe0f8d7612 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -280,6 +280,7 @@ extern AVFilter ff_vf_mix; extern AVFilter ff_vf_mpdecimate; extern AVFilter ff_vf_negate; extern AVFilter ff_vf_nlmeans; +extern AVFilter ff_vf_nlmeans_opencl; extern AVFilter ff_vf_nnedi; extern AVFilter ff_vf_noformat; extern AVFilter ff_vf_noise; diff --git a/libavfilter/opencl/nlmeans.cl b/libavfilter/opencl/nlmeans.cl new file mode 100644 index 0000000000..72bd681fd6 --- /dev/null +++ b/libavfilter/opencl/nlmeans.cl @@ -0,0 +1,115 @@ +/* + * This file is part of FFmpeg. + * + * FFmpeg is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2.1 of the License, or (at your option) any later version. + * + * FFmpeg is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with FFmpeg; if not, write to the Free Software + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA + */ + +const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE | + CLK_ADDRESS_CLAMP_TO_EDGE | + CLK_FILTER_NEAREST); + +kernel void horiz_sum(__global uint4 *integral_img, + __read_only image2d_t src, + int width, + int height, + int4 dx, + int4 dy) +{ + + int y = get_global_id(0); + int work_size = get_global_size(0); + + uint4 sum = (uint4)(0); + float4 s2; + for (int i = 0; i < width; i++) { + float s1 = read_imagef(src, sampler, (int2)(i, y)).x; + s2.x = read_imagef(src, sampler, (int2)(i + dx.x, y + dy.x)).x; + s2.y = read_imagef(src, sampler, (int2)(i + dx.y, y + dy.y)).x; + s2.z = read_imagef(src, sampler, (int2)(i + dx.z, y + dy.z)).x; + s2.w = read_imagef(src, sampler, (int2)(i + dx.w, y + dy.w)).x; + sum += convert_uint4((s1 - s2) * (s1 - s2) * 255 * 255); + integral_img[y * width + i] = sum; + } +} + +kernel void vert_sum(__global uint4 *integral_img, + __global int *overflow, + int width, + int height) +{ + int x = get_global_id(0); + uint4 sum = 0; + for (int i = 0; i < height; i++) { + if (any((uint4)UINT_MAX - integral_img[i * width + x] < sum)) + atomic_inc(overflow); + integral_img[i * width + x] += sum; + sum = integral_img[i * width + x]; + } +} + +kernel void weight_accum(global float *sum, global float *weight, + global uint4 *integral_img, __read_only image2d_t src, + int width, int height, int p, float h, + int4 dx, int4 dy) +{ + // w(x) = integral_img(x-p, y-p) + + // integral_img(x+p, y+p) - + // integral_img(x+p, y-p) - + // integral_img(x-p, y+p) + // total_sum[x] += w(x, y) * src(x + dx, y + dy) + // total_weight += w(x, y) + + int x = get_global_id(0); + int y = get_global_id(1); + int4 xoff = x + dx; + int4 yoff = y + dy; + uint4 a = 0, b = 0, c = 0, d = 0; + uint4 src_pix = 0; + + // out-of-bounding-box? + int oobb = (x - p) < 0 || (y - p) < 0 || (y + p) >= height || (x + p) >= width; + + src_pix.x = (int)(255 * read_imagef(src, sampler, (int2)(xoff.x, yoff.x)).x); + src_pix.y = (int)(255 * read_imagef(src, sampler, (int2)(xoff.y, yoff.y)).x); + src_pix.z = (int)(255 * read_imagef(src, sampler, (int2)(xoff.z, yoff.z)).x); + src_pix.w = (int)(255 * read_imagef(src, sampler, (int2)(xoff.w, yoff.w)).x); + if (!oobb) { + a = integral_img[(y - p) * width + x - p]; + b = integral_img[(y + p) * width + x - p]; + c = integral_img[(y - p) * width + x + p]; + d = integral_img[(y + p) * width + x + p]; + } + + float4 patch_diff = convert_float4(d + a - c - b); + float4 w = native_exp(-patch_diff / (h * h)); + float w_sum = w.x + w.y + w.z + w.w; + weight[y * width + x] += w_sum; + sum[y * width + x] += dot(w, convert_float4(src_pix)); +} + +kernel void average(__write_only image2d_t dst, + __read_only image2d_t src, + global float *sum, global float *weight) { + int x = get_global_id(0); + int y = get_global_id(1); + int2 dim = get_image_dim(dst); + + float w = weight[y * dim.x + x]; + float s = sum[y * dim.x + x]; + float src_pix = read_imagef(src, sampler, (int2)(x, y)).x; + float r = (s + src_pix * 255) / (1.0f + w) / 255.0f; + if (x < dim.x && y < dim.y) + write_imagef(dst, (int2)(x, y), (float4)(r, 0.0f, 0.0f, 1.0f)); +} diff --git a/libavfilter/opencl_source.h b/libavfilter/opencl_source.h index 51f7178cf2..1a6cd7ca7a 100644 --- a/libavfilter/opencl_source.h +++ b/libavfilter/opencl_source.h @@ -24,6 +24,7 @@ extern const char *ff_opencl_source_colorkey; extern const char *ff_opencl_source_colorspace_common; extern const char *ff_opencl_source_convolution; extern const char *ff_opencl_source_neighbor; +extern const char *ff_opencl_source_nlmeans; extern const char *ff_opencl_source_overlay; extern const char *ff_opencl_source_tonemap; extern const char *ff_opencl_source_transpose; diff --git a/libavfilter/vf_nlmeans_opencl.c b/libavfilter/vf_nlmeans_opencl.c new file mode 100644 index 0000000000..11847add37 --- /dev/null +++ b/libavfilter/vf_nlmeans_opencl.c @@ -0,0 +1,443 @@ +/* + * This file is part of FFmpeg. + * + * FFmpeg is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2.1 of the License, or (at your option) any later version. + * + * FFmpeg is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with FFmpeg; if not, write to the Free Software + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA + */ +#include + +#include "libavutil/avassert.h" +#include "libavutil/common.h" +#include "libavutil/imgutils.h" +#include "libavutil/mem.h" +#include "libavutil/opt.h" +#include "libavutil/pixdesc.h" + +#include "avfilter.h" +#include "internal.h" +#include "opencl.h" +#include "opencl_source.h" +#include "video.h" + +// TODO: +// the integral image may overflow 32bit, consider using 64bit + +static const enum AVPixelFormat supported_formats[] = { + AV_PIX_FMT_YUV420P, + AV_PIX_FMT_YUV444P, + AV_PIX_FMT_GBRP, +}; + +static int is_format_supported(enum AVPixelFormat fmt) +{ + int i; + + for (i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++) + if (supported_formats[i] == fmt) + return 1; + return 0; +} + +typedef struct NLMeansOpenCLContext { + OpenCLFilterContext ocf; + int initialised; + cl_kernel vert_kernel; + cl_kernel horiz_kernel; + cl_kernel accum_kernel; + cl_kernel average_kernel; + cl_mem integral_img; + cl_mem weight; + cl_mem sum; + cl_mem overflow; // overflow in integral image? + double sigma; + float h; + int chroma_w; + int chroma_h; + int patch_size; + int patch_size_uv; + int research_size; + int research_size_uv; + cl_command_queue command_queue; +} NLMeansOpenCLContext; + +static int nlmeans_opencl_init(AVFilterContext *avctx, int width, int height) +{ + NLMeansOpenCLContext *ctx = avctx->priv; + cl_int cle; + int err; + int weight_buf_size = width * height * sizeof(float); + + ctx->h = ctx->sigma * 10; + if (!(ctx->research_size & 1)) { + ctx->research_size |= 1; + av_log(avctx, AV_LOG_WARNING, + "research_size should be odd, set to %d", + ctx->research_size); + } + + if (!(ctx->patch_size & 1)) { + ctx->patch_size |= 1; + av_log(avctx, AV_LOG_WARNING, + "patch_size should be odd, set to %d", + ctx->patch_size); + } + + if (!ctx->research_size_uv) + ctx->research_size_uv = ctx->research_size; + if (!ctx->patch_size_uv) + ctx->patch_size_uv = ctx->patch_size; + + err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_nlmeans, 1); + if (err < 0) + goto fail; + + ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context, + ctx->ocf.hwctx->device_id, + 0, &cle); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL " + "command queue %d.\n", cle); + + ctx->vert_kernel = clCreateKernel(ctx->ocf.program, + "vert_sum", &cle); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create " + "vert_sum kernel %d.\n", cle); + + ctx->horiz_kernel = clCreateKernel(ctx->ocf.program, + "horiz_sum", &cle); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create " + "horiz_sum kernel %d.\n", cle); + + ctx->accum_kernel = clCreateKernel(ctx->ocf.program, + "weight_accum", &cle); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create " + "accum kernel %d.\n", cle); + + ctx->average_kernel = clCreateKernel(ctx->ocf.program, + "average", &cle); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create " + "average kernel %d.\n", cle); + + ctx->integral_img = clCreateBuffer(ctx->ocf.hwctx->context, 0, + 4 * width * height * sizeof(cl_int), + NULL, &cle); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create " + "integral image %d.\n", cle); + + ctx->weight = clCreateBuffer(ctx->ocf.hwctx->context, 0, + weight_buf_size, NULL, &cle); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create " + "weight buffer %d.\n", cle); + + ctx->sum = clCreateBuffer(ctx->ocf.hwctx->context, 0, + weight_buf_size, NULL, &cle); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create " + "sum buffer %d.\n", cle); + + ctx->overflow = clCreateBuffer(ctx->ocf.hwctx->context, 0, + sizeof(cl_int), NULL, &cle); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create " + "overflow buffer %d.\n", cle); + + ctx->initialised = 1; + return 0; + +fail: + CL_RELEASE_KERNEL(ctx->vert_kernel); + CL_RELEASE_KERNEL(ctx->horiz_kernel); + CL_RELEASE_KERNEL(ctx->accum_kernel); + CL_RELEASE_KERNEL(ctx->average_kernel); + + CL_RELEASE_MEMORY(ctx->integral_img); + CL_RELEASE_MEMORY(ctx->weight); + CL_RELEASE_MEMORY(ctx->sum); + CL_RELEASE_MEMORY(ctx->overflow); + + CL_RELEASE_QUEUE(ctx->command_queue); + return err; +} + +static int nlmeans_plane(AVFilterContext *avctx, cl_mem dst, cl_mem src, + cl_int width, cl_int height, cl_int p, cl_int r) +{ + NLMeansOpenCLContext *ctx = avctx->priv; + const float zero = 0.0f; + const size_t worksize1[] = {height}; + const size_t worksize2[] = {width}; + const size_t worksize3[2] = {width, height}; + int i, dx, dy, err = 0, weight_buf_size; + cl_int cle; + int nb_pixel, *tmp, idx = 0; + cl_int *dxdy; + + weight_buf_size = width * height * sizeof(float); + cle = clEnqueueFillBuffer(ctx->command_queue, ctx->weight, + &zero, sizeof(float), 0, weight_buf_size, + 0, NULL, NULL); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to fill weight buffer: %d.\n", + cle); + cle = clEnqueueFillBuffer(ctx->command_queue, ctx->sum, + &zero, sizeof(float), 0, weight_buf_size, + 0, NULL, NULL); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to fill sum buffer: %d.\n", + cle); + + nb_pixel = (2 * r + 1) * (2 * r + 1) - 1; + dxdy = av_malloc(nb_pixel * 2 * sizeof(cl_int)); + tmp = av_malloc(nb_pixel * 2 * sizeof(int)); + + if (!dxdy || !tmp) + goto fail; + + for (dx = -r; dx <= r; dx++) { + for (dy = -r; dy <= r; dy++) { + if (dx || dy) { + tmp[idx++] = dx; + tmp[idx++] = dy; + } + } + } + // repack dx/dy seperately, as we want to do four pairs of dx/dy in a batch + for (i = 0; i < nb_pixel / 4; i++) { + dxdy[i * 8] = tmp[i * 8]; // dx0 + dxdy[i * 8 + 1] = tmp[i * 8 + 2]; // dx1 + dxdy[i * 8 + 2] = tmp[i * 8 + 4]; // dx2 + dxdy[i * 8 + 3] = tmp[i * 8 + 6]; // dx3 + dxdy[i * 8 + 4] = tmp[i * 8 + 1]; // dy0 + dxdy[i * 8 + 5] = tmp[i * 8 + 3]; // dy1 + dxdy[i * 8 + 6] = tmp[i * 8 + 5]; // dy2 + dxdy[i * 8 + 7] = tmp[i * 8 + 7]; // dy3 + } + av_freep(&tmp); + + for (i = 0; i < nb_pixel / 4; i++) { + cl_int *dx_cur = dxdy + 8 * i; + cl_int *dy_cur = dxdy + 8 * i + 4; + + // horizontal pass + // integral(x,y) = sum([u(v,y) - u(v+dx,y+dy)]^2) for v in [0, x] + CL_SET_KERNEL_ARG(ctx->horiz_kernel, 0, cl_mem, &ctx->integral_img); + CL_SET_KERNEL_ARG(ctx->horiz_kernel, 1, cl_mem, &src); + CL_SET_KERNEL_ARG(ctx->horiz_kernel, 2, cl_int, &width); + CL_SET_KERNEL_ARG(ctx->horiz_kernel, 3, cl_int, &height); + CL_SET_KERNEL_ARG(ctx->horiz_kernel, 4, cl_int4, dx_cur); + CL_SET_KERNEL_ARG(ctx->horiz_kernel, 5, cl_int4, dy_cur); + cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->horiz_kernel, 1, + NULL, worksize1, NULL, 0, NULL, NULL); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue horiz_kernel: %d.\n", + cle); + // vertical pass + // integral(x, y) = sum(integral(x, v)) for v in [0, y] + CL_SET_KERNEL_ARG(ctx->vert_kernel, 0, cl_mem, &ctx->integral_img); + CL_SET_KERNEL_ARG(ctx->vert_kernel, 1, cl_mem, &ctx->overflow); + CL_SET_KERNEL_ARG(ctx->vert_kernel, 2, cl_int, &width); + CL_SET_KERNEL_ARG(ctx->vert_kernel, 3, cl_int, &height); + cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->vert_kernel, + 1, NULL, worksize2, NULL, 0, NULL, NULL); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue vert_kernel: %d.\n", + cle); + + // accumulate weights + CL_SET_KERNEL_ARG(ctx->accum_kernel, 0, cl_mem, &ctx->sum); + CL_SET_KERNEL_ARG(ctx->accum_kernel, 1, cl_mem, &ctx->weight); + CL_SET_KERNEL_ARG(ctx->accum_kernel, 2, cl_mem, &ctx->integral_img); + CL_SET_KERNEL_ARG(ctx->accum_kernel, 3, cl_mem, &src); + CL_SET_KERNEL_ARG(ctx->accum_kernel, 4, cl_int, &width); + CL_SET_KERNEL_ARG(ctx->accum_kernel, 5, cl_int, &height); + CL_SET_KERNEL_ARG(ctx->accum_kernel, 6, cl_int, &p); + CL_SET_KERNEL_ARG(ctx->accum_kernel, 7, cl_float, &ctx->h); + CL_SET_KERNEL_ARG(ctx->accum_kernel, 8, cl_int4, dx_cur); + CL_SET_KERNEL_ARG(ctx->accum_kernel, 9, cl_int4, dy_cur); + cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->accum_kernel, + 2, NULL, worksize3, NULL, 0, NULL, NULL); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle); + } + av_freep(&dxdy); + + // average + CL_SET_KERNEL_ARG(ctx->average_kernel, 0, cl_mem, &dst); + CL_SET_KERNEL_ARG(ctx->average_kernel, 1, cl_mem, &src); + CL_SET_KERNEL_ARG(ctx->average_kernel, 2, cl_mem, &ctx->sum); + CL_SET_KERNEL_ARG(ctx->average_kernel, 3, cl_mem, &ctx->weight); + cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->average_kernel, 2, + NULL, worksize3, NULL, 0, NULL, NULL); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue average kernel: %d.\n", + cle); + cle = clFlush(ctx->command_queue); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to flush command queue: %d.\n", cle); +fail: + if (tmp) + av_freep(&tmp); + if (dxdy) + av_freep(&dxdy); + return err; +} + +static int nlmeans_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) +{ + AVFilterContext *avctx = inlink->dst; + AVFilterLink *outlink = avctx->outputs[0]; + NLMeansOpenCLContext *ctx = avctx->priv; + AVFrame *output = NULL; + AVHWFramesContext *input_frames_ctx; + const AVPixFmtDescriptor *desc; + enum AVPixelFormat in_format; + cl_mem src, dst; + const cl_int zero = 0; + int w, h, err, cle, overflow, p, patch, research; + + av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n", + av_get_pix_fmt_name(input->format), + input->width, input->height, input->pts); + + if (!input->hw_frames_ctx) + return AVERROR(EINVAL); + input_frames_ctx = (AVHWFramesContext*)input->hw_frames_ctx->data; + in_format = input_frames_ctx->sw_format; + + output = ff_get_video_buffer(outlink, outlink->w, outlink->h); + if (!output) { + err = AVERROR(ENOMEM); + goto fail; + } + + err = av_frame_copy_props(output, input); + if (err < 0) + goto fail; + + if (!ctx->initialised) { + desc = av_pix_fmt_desc_get(in_format); + if (!is_format_supported(in_format)) { + err = AVERROR(EINVAL); + av_log(avctx, AV_LOG_ERROR, "input format %s not supported\n", + av_get_pix_fmt_name(in_format)); + goto fail; + } + ctx->chroma_w = AV_CEIL_RSHIFT(inlink->w, desc->log2_chroma_w); + ctx->chroma_h = AV_CEIL_RSHIFT(inlink->h, desc->log2_chroma_h); + + err = nlmeans_opencl_init(avctx, inlink->w, inlink->h); + if (err < 0) + goto fail; + } + + cle = clEnqueueWriteBuffer(ctx->command_queue, ctx->overflow, CL_FALSE, + 0, sizeof(cl_int), &zero, 0, NULL, NULL); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to initialize overflow" + "detection buffer %d.\n", cle); + + for (p = 0; p < FF_ARRAY_ELEMS(output->data); p++) { + src = (cl_mem) input->data[p]; + dst = (cl_mem) output->data[p]; + + if (!dst) + break; + av_assert0(src); + w = p ? ctx->chroma_w : inlink->w; + h = p ? ctx->chroma_h : inlink->h; + patch = (p ? ctx->patch_size_uv : ctx->patch_size) / 2; + research = (p ? ctx->research_size_uv : ctx->research_size) / 2; + err = nlmeans_plane(avctx, dst, src, w, h, patch, research); + if (err < 0) + goto fail; + } + // overflow occurred? + cle = clEnqueueReadBuffer(ctx->command_queue, ctx->overflow, CL_FALSE, + 0, sizeof(cl_int), &overflow, 0, NULL, NULL); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to read overflow: %d.\n", cle); + + cle = clFinish(ctx->command_queue); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish kernel: %d.\n", cle); + + if (overflow > 0) + av_log(avctx, AV_LOG_ERROR, "integral image overflow %d\n", overflow); + + av_frame_free(&input); + + av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n", + av_get_pix_fmt_name(output->format), + output->width, output->height, output->pts); + + return ff_filter_frame(outlink, output); + +fail: + clFinish(ctx->command_queue); + av_frame_free(&input); + av_frame_free(&output); + return err; +} + +static av_cold void nlmeans_opencl_uninit(AVFilterContext *avctx) +{ + NLMeansOpenCLContext *ctx = avctx->priv; + cl_int cle; + + CL_RELEASE_KERNEL(ctx->vert_kernel); + CL_RELEASE_KERNEL(ctx->horiz_kernel); + CL_RELEASE_KERNEL(ctx->accum_kernel); + CL_RELEASE_KERNEL(ctx->average_kernel); + + CL_RELEASE_MEMORY(ctx->integral_img); + CL_RELEASE_MEMORY(ctx->weight); + CL_RELEASE_MEMORY(ctx->sum); + CL_RELEASE_MEMORY(ctx->overflow); + + CL_RELEASE_QUEUE(ctx->command_queue); + + ff_opencl_filter_uninit(avctx); +} + +#define OFFSET(x) offsetof(NLMeansOpenCLContext, x) +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) +static const AVOption nlmeans_opencl_options[] = { + { "s", "denoising strength", OFFSET(sigma), AV_OPT_TYPE_DOUBLE, { .dbl = 1.0 }, 1.0, 30.0, FLAGS }, + { "p", "patch size", OFFSET(patch_size), AV_OPT_TYPE_INT, { .i64 = 2*3+1 }, 0, 99, FLAGS }, + { "pc", "patch size for chroma planes", OFFSET(patch_size_uv), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, 99, FLAGS }, + { "r", "research window", OFFSET(research_size), AV_OPT_TYPE_INT, { .i64 = 7*2+1 }, 0, 99, FLAGS }, + { "rc", "research window for chroma planes", OFFSET(research_size_uv), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, 99, FLAGS }, + { NULL } +}; + +AVFILTER_DEFINE_CLASS(nlmeans_opencl); + +static const AVFilterPad nlmeans_opencl_inputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .filter_frame = &nlmeans_opencl_filter_frame, + .config_props = &ff_opencl_filter_config_input, + }, + { NULL } +}; + +static const AVFilterPad nlmeans_opencl_outputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = &ff_opencl_filter_config_output, + }, + { NULL } +}; + +AVFilter ff_vf_nlmeans_opencl = { + .name = "nlmeans_opencl", + .description = NULL_IF_CONFIG_SMALL("Non-local means denoiser through OpenCL"), + .priv_size = sizeof(NLMeansOpenCLContext), + .priv_class = &nlmeans_opencl_class, + .init = &ff_opencl_filter_init, + .uninit = &nlmeans_opencl_uninit, + .query_formats = &ff_opencl_filter_query_formats, + .inputs = nlmeans_opencl_inputs, + .outputs = nlmeans_opencl_outputs, + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +};