From patchwork Fri Jun 8 23:37:51 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Danil Iashchenko X-Patchwork-Id: 9330 Delivered-To: ffmpegpatchwork@gmail.com Received: by 2002:a02:11c:0:0:0:0:0 with SMTP id c28-v6csp1451679jad; Fri, 8 Jun 2018 16:44:00 -0700 (PDT) X-Google-Smtp-Source: ADUXVKJUFziCO9egYAKwyfdku0gj0P4Oq7sX7TgUxQVtb++JOISrlqrB0WJ/BNAMGsLRcJf9nq6O X-Received: by 2002:a1c:b947:: with SMTP id j68-v6mr2560816wmf.144.1528501440150; Fri, 08 Jun 2018 16:44:00 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1528501440; cv=none; d=google.com; s=arc-20160816; b=kNXWt8V/6KO3+k8Q1pQEG9CHrbuezBTjmDHZ0b6koavdAhDrlETPP1Wu4nD06UCiyM /sc5oGs7S3r5AyV1SQyqYxvsJXyEvwUgl+lGJCI5lcNpeLuS4dnw7o4Z2eVOlVrv85Rd J3ehGZ0e8ZcQTVX+eF4PpvIQWbsgTg5cogqLudxu+QxeHl+MdB1k1xEIu8YTVCBc5XiA fL48b4dcFxpyuhBV7l6rQwLaJ2aD9LRa1ViPCSS6d+LbVY6wbti1wfXNIDkvTmCr660/ /w9sv9n55W7J3pgBITikkpXiPyt/HvxrCL42mmImzMEFJXEJqTlwOTgYkHPwzB7JcdW0 bnpA== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:content-transfer-encoding:mime-version:cc:reply-to :list-subscribe:list-help:list-post:list-archive:list-unsubscribe :list-id:precedence:subject:references:in-reply-to:message-id:date :to:from:dkim-signature:delivered-to:arc-authentication-results; bh=SZht3UUUVd9AOBbGX1K/2USAMoVQPmBtm3luEEAJJPk=; b=Y5vQ26fR/6y3kd8Cl2cqIuFIwQfEt3Sv0ghcflx0wNODT56RbeYgP3Oa5DaeFsQga0 aAemd5ss1yrOmfApLtfxzHNiuMyZlyL3oEtc1/iWF/D8uzybdS1g7JKMshkCbapXsglw 1bRQPMiEeVRR5NwVGRWXD+3VlZUHYUkpgC8Jzkizjnp0X+7rM105+1sSq120WthLU5Ld EyLDHsdG6kwwdfSmWupx7nXUlvqdyYtWGiqXaWO0t8rwy66UtL5AIuFUFv/mBZSwgibG pRKbB1U/sWvhKQNygZ4s3F3GXV407Shf8h+lmAvoGrFtCL9ihESqF9wpMqLolNVpqr5G fAcw== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20161025 header.b=rQIiMLoG; spf=pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) smtp.mailfrom=ffmpeg-devel-bounces@ffmpeg.org; dmarc=fail (p=NONE sp=QUARANTINE dis=NONE) header.from=gmail.com Return-Path: Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org. [79.124.17.100]) by mx.google.com with ESMTP id n131-v6si1963302wmb.16.2018.06.08.16.43.59; Fri, 08 Jun 2018 16:44:00 -0700 (PDT) Received-SPF: pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) client-ip=79.124.17.100; Authentication-Results: mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20161025 header.b=rQIiMLoG; spf=pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) smtp.mailfrom=ffmpeg-devel-bounces@ffmpeg.org; dmarc=fail (p=NONE sp=QUARANTINE dis=NONE) header.from=gmail.com Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id 395AB689B08; Sat, 9 Jun 2018 02:43:09 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-lf0-f68.google.com (mail-lf0-f68.google.com [209.85.215.68]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id EE935689723 for ; Sat, 9 Jun 2018 02:43:02 +0300 (EEST) Received: by mail-lf0-f68.google.com with SMTP id u4-v6so22415382lff.3 for ; Fri, 08 Jun 2018 16:43:51 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gmail.com; s=20161025; h=from:to:cc:subject:date:message-id:in-reply-to:references; bh=ehSdRvcfYtgSSYc3s9pmS60tzWHWkTJPrp2xWvCHQJg=; b=rQIiMLoGdDLWuDonVrp77UMcH2bQGsxV7gUZ0s1NoOayvsvC0NUYaBCuPH1Hr4DNni lSvGZU1NuP+a7uSU8ykRplxdfImO1NMiV+koOAL3Wi8Dg+k7ln1fRkKv0HFULoaGcfe+ amdi3OddsIQH1N7SRl258MFLBl1c/5biEFyN0PE7yzVbGuBEQKB3H2Hjf+3fcDg3TeIF vUttaKGsb+QSe29V9RXjTfIinBi83uDwrqefxnkHROzKpQlTbA8aTvnuVoRNvfuCWKnQ ixo61qTgcPIDwQ6Nl3cQHGMJ86GL9ieprUNvkd0Xy8jX/5HFsP0iZCOON2KwrrDhSqZh 7OqQ== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20161025; h=x-gm-message-state:from:to:cc:subject:date:message-id:in-reply-to :references; bh=ehSdRvcfYtgSSYc3s9pmS60tzWHWkTJPrp2xWvCHQJg=; b=iVZXz7lxtIh23OyA7e7ACq9AXPnd6sIrulJ3HpxSl/7oTPZ9O8y7VTsYi5aaeYFI9B AygyrLsPdy9XfbzK4Q8pWT/ikYCgD1nRocSL5aBleUrE3NGw11KrXjbdp3e1mSrJh9fF 3jdsIDaxTRYE6vC/+vMSMi5h9EJVstQh4D6+ualRcaqv/ixYp7L7s4qcsvpPgPn6Fi3K AVYZaO7N5voZt2A1ZntBewjdWLs7MsBXxixSgvPaBv5s3/k5QCndrjiiJ9luEmd4xEod uXG4II0fFEAYKK2TUpaWAMw5Yqa2RGIda9lhcVzlSUv4Z0pxgUC6BqbKxdsYbXA9mKen x2TA== X-Gm-Message-State: APt69E2Nbdwx3lVsBjZ2HFyFPwiN7zeggInYH5q2tffwuP31cfpZeJqI wKpr/kuglTnSStE8oftuoGAk+y0= X-Received: by 2002:a2e:9b91:: with SMTP id z17-v6mr4214825lji.121.1528501075314; Fri, 08 Jun 2018 16:37:55 -0700 (PDT) Received: from dan-acer.lan (campus.ifmo.ru. [194.85.161.2]) by smtp.gmail.com with ESMTPSA id a2-v6sm3997855ljb.80.2018.06.08.16.37.54 (version=TLS1_2 cipher=ECDHE-RSA-AES128-SHA bits=128/128); Fri, 08 Jun 2018 16:37:54 -0700 (PDT) From: Danil Iashchenko To: ffmpeg-devel@ffmpeg.org Date: Sat, 9 Jun 2018 02:37:51 +0300 Message-Id: <1528501071-15950-1-git-send-email-danyaschenko@gmail.com> X-Mailer: git-send-email 2.7.4 In-Reply-To: <4608b3aa-8133-7dc9-bd6d-352684c02c98@jkqxz.net> References: <4608b3aa-8133-7dc9-bd6d-352684c02c98@jkqxz.net> Subject: [FFmpeg-devel] [PATCH] libavfilter/boxblur_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: Danil Iashchenko MIME-Version: 1.0 Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" --- Thanks, fixed. libavfilter/Makefile | 4 +- libavfilter/allfilters.c | 1 + libavfilter/boxblur.c | 105 +++++++++++++ libavfilter/boxblur.h | 66 +++++++++ libavfilter/vf_avgblur_opencl.c | 320 +++++++++++++++++++++++++++------------- libavfilter/vf_boxblur.c | 113 ++------------ 6 files changed, 404 insertions(+), 205 deletions(-) create mode 100644 libavfilter/boxblur.c create mode 100644 libavfilter/boxblur.h diff --git a/libavfilter/Makefile b/libavfilter/Makefile index c68ef05..067210f 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -152,7 +152,9 @@ OBJS-$(CONFIG_BITPLANENOISE_FILTER) += vf_bitplanenoise.o OBJS-$(CONFIG_BLACKDETECT_FILTER) += vf_blackdetect.o OBJS-$(CONFIG_BLACKFRAME_FILTER) += vf_blackframe.o OBJS-$(CONFIG_BLEND_FILTER) += vf_blend.o framesync.o -OBJS-$(CONFIG_BOXBLUR_FILTER) += vf_boxblur.o +OBJS-$(CONFIG_BOXBLUR_FILTER) += vf_boxblur.o boxblur.o +OBJS-$(CONFIG_BOXBLUR_OPENCL_FILTER) += vf_avgblur_opencl.o opencl.o \ + opencl/avgblur.o boxblur.o OBJS-$(CONFIG_BWDIF_FILTER) += vf_bwdif.o OBJS-$(CONFIG_CHROMAKEY_FILTER) += vf_chromakey.o OBJS-$(CONFIG_CIESCOPE_FILTER) += vf_ciescope.o diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index b44093d..97d92a0 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -146,6 +146,7 @@ extern AVFilter ff_vf_blackdetect; extern AVFilter ff_vf_blackframe; extern AVFilter ff_vf_blend; extern AVFilter ff_vf_boxblur; +extern AVFilter ff_vf_boxblur_opencl; extern AVFilter ff_vf_bwdif; extern AVFilter ff_vf_chromakey; extern AVFilter ff_vf_ciescope; diff --git a/libavfilter/boxblur.c b/libavfilter/boxblur.c new file mode 100644 index 0000000..1d4895e --- /dev/null +++ b/libavfilter/boxblur.c @@ -0,0 +1,105 @@ +/* + * Copyright (c) 2018 Danil Iashchenko + * + * 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 "boxblur.h" + + +int ff_boxblur_eval_filter_params(AVFilterLink *inlink, + FilterParam *luma_param, + FilterParam *chroma_param, + FilterParam *alpha_param) +{ + const AVPixFmtDescriptor *desc = av_pix_fmt_desc_get(inlink->format); + AVFilterContext *ctx = inlink->dst; + int w = inlink->w, h = inlink->h; + int cw, ch; + double var_values[VARS_NB], res; + char *expr; + int ret; + + if (!luma_param->radius_expr) { + av_log(ctx, AV_LOG_ERROR, "Luma radius expression is not set.\n"); + return AVERROR(EINVAL); + } + + /* fill missing params */ + if (!chroma_param->radius_expr) { + chroma_param->radius_expr = av_strdup(luma_param->radius_expr); + if (!chroma_param->radius_expr) + return AVERROR(ENOMEM); + } + if (chroma_param->power < 0) + chroma_param->power = luma_param->power; + + if (!alpha_param->radius_expr) { + alpha_param->radius_expr = av_strdup(luma_param->radius_expr); + if (!alpha_param->radius_expr) + return AVERROR(ENOMEM); + } + if (alpha_param->power < 0) + alpha_param->power = luma_param->power; + + var_values[VAR_W] = inlink->w; + var_values[VAR_H] = inlink->h; + var_values[VAR_CW] = cw = w>>(desc->log2_chroma_w); + var_values[VAR_CH] = ch = h>>(desc->log2_chroma_h); + var_values[VAR_HSUB] = 1<<(desc->log2_chroma_w); + var_values[VAR_VSUB] = 1<<(desc->log2_chroma_h); + +#define EVAL_RADIUS_EXPR(comp) \ + expr = comp->radius_expr; \ + ret = av_expr_parse_and_eval(&res, expr, var_names, var_values, \ + NULL, NULL, NULL, NULL, NULL, 0, ctx); \ + comp->radius = res; \ + if (ret < 0) { \ + av_log(NULL, AV_LOG_ERROR, \ + "Error when evaluating " #comp " radius expression '%s'\n", expr); \ + return ret; \ + } + + EVAL_RADIUS_EXPR(luma_param); + EVAL_RADIUS_EXPR(chroma_param); + EVAL_RADIUS_EXPR(alpha_param); + + av_log(ctx, AV_LOG_VERBOSE, + "luma_radius:%d luma_power:%d " + "chroma_radius:%d chroma_power:%d " + "alpha_radius:%d alpha_power:%d " + "w:%d chroma_w:%d h:%d chroma_h:%d\n", + luma_param ->radius, luma_param ->power, + chroma_param->radius, chroma_param->power, + alpha_param ->radius, alpha_param ->power, + w, cw, h, ch); + + +#define CHECK_RADIUS_VAL(w_, h_, comp) \ + if (comp->radius < 0 || \ + 2*comp->radius > FFMIN(w_, h_)) { \ + av_log(ctx, AV_LOG_ERROR, \ + "Invalid " #comp " radius value %d, must be >= 0 and <= %d\n", \ + comp->radius, FFMIN(w_, h_)/2); \ + return AVERROR(EINVAL); \ + } + CHECK_RADIUS_VAL(w, h, luma_param); + CHECK_RADIUS_VAL(cw, ch, chroma_param); + CHECK_RADIUS_VAL(w, h, alpha_param); + + return 0; +} diff --git a/libavfilter/boxblur.h b/libavfilter/boxblur.h new file mode 100644 index 0000000..94b2730 --- /dev/null +++ b/libavfilter/boxblur.h @@ -0,0 +1,66 @@ +/* + * Copyright (c) 2018 Danil Iashchenko + * + * 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 + */ + +#ifndef BOXBLUR_H +#define BOXBLUR_H + +#include "libavutil/eval.h" +#include "libavutil/pixdesc.h" +#include "libavutil/mem.h" + +#include "avfilter.h" + +static const char *const var_names[] = { + "w", + "h", + "cw", + "ch", + "hsub", + "vsub", + NULL +}; + +enum var_name { + VAR_W, + VAR_H, + VAR_CW, + VAR_CH, + VAR_HSUB, + VAR_VSUB, + VARS_NB +}; + +typedef struct FilterParam { + int radius; + int power; + char *radius_expr; +} FilterParam; + +#define Y 0 +#define U 1 +#define V 2 +#define A 3 + +int ff_boxblur_eval_filter_params(AVFilterLink *inlink, + FilterParam *luma_param, + FilterParam *chroma_param, + FilterParam *alpha_param); + +#endif // BOXBLUR_H diff --git a/libavfilter/vf_avgblur_opencl.c b/libavfilter/vf_avgblur_opencl.c index 48cebb5..9854b6b 100644 --- a/libavfilter/vf_avgblur_opencl.c +++ b/libavfilter/vf_avgblur_opencl.c @@ -1,5 +1,6 @@ /* * Copyright (c) 2018 Dylan Fernando + * Copyright (c) 2018 Danil Iashchenko * * This file is part of FFmpeg. * @@ -20,16 +21,14 @@ #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" - +#include "boxblur.h" typedef struct AverageBlurOpenCLContext { OpenCLFilterContext ocf; @@ -39,10 +38,16 @@ typedef struct AverageBlurOpenCLContext { cl_kernel kernel_vert; cl_command_queue command_queue; - int radius; + int radiusH; int radiusV; int planes; + FilterParam luma_param; + FilterParam chroma_param; + FilterParam alpha_param; + int radius[4]; + int power[4]; + } AverageBlurOpenCLContext; @@ -80,10 +85,6 @@ static int avgblur_opencl_init(AVFilterContext *avctx) goto fail; } - if (ctx->radiusV <= 0) { - ctx->radiusV = ctx->radius; - } - ctx->initialised = 1; return 0; @@ -97,6 +98,60 @@ fail: return err; } + +static int avgblur_opencl_make_filter_params(AVFilterLink *inlink) +{ + AVFilterContext *ctx = inlink->dst; + AverageBlurOpenCLContext *s = ctx->priv; + int i; + + if (s->radiusV <= 0) { + s->radiusV = s->radiusH; + } + + for (i = 0; i < 4; i++) { + s->power[i] = 1; + } + return 0; +} + + +static int boxblur_opencl_make_filter_params(AVFilterLink *inlink) +{ + AVFilterContext *ctx = inlink->dst; + AverageBlurOpenCLContext *s = ctx->priv; + int err, i; + + err = ff_boxblur_eval_filter_params(inlink, + &s->luma_param, + &s->chroma_param, + &s->alpha_param); + + if (err != 0) { + av_log(ctx, AV_LOG_ERROR, "Failed to evaluate " + "filter params: %d.\n", err); + return err; + } + + s->radius[Y] = s->luma_param.radius; + s->radius[U] = s->radius[V] = s->chroma_param.radius; + s->radius[A] = s->alpha_param.radius; + + s->power[Y] = s->luma_param.power; + s->power[U] = s->power[V] = s->chroma_param.power; + s->power[A] = s->alpha_param.power; + + for (i = 0; i < 4; i++) { + if (s->power[i] == 0) { + s->power[i] = 1; + s->radius[i] = 0; + } + } + + return 0; +} + + static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) { AVFilterContext *avctx = inlink->dst; @@ -107,7 +162,7 @@ static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) cl_int cle; size_t global_work[2]; cl_mem src, dst, inter; - int err, p, radius_x, radius_y; + int err, p, radius_x, radius_y, i; av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n", av_get_pix_fmt_name(input->format), @@ -121,6 +176,16 @@ static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) if (err < 0) goto fail; + if (!strcmp(avctx->filter->name, "avgblur_opencl")) { + err = avgblur_opencl_make_filter_params(inlink); + if (err < 0) + goto fail; + } else if (!strcmp(avctx->filter->name, "boxblur_opencl")) { + err = boxblur_opencl_make_filter_params(inlink); + if (err < 0) + goto fail; + } + } output = ff_get_video_buffer(outlink, outlink->w, outlink->h); @@ -128,7 +193,6 @@ static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) err = AVERROR(ENOMEM); goto fail; } - intermediate = ff_get_video_buffer(outlink, outlink->w, outlink->h); if (!intermediate) { err = AVERROR(ENOMEM); @@ -137,13 +201,13 @@ static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) for (p = 0; p < FF_ARRAY_ELEMS(output->data); p++) { src = (cl_mem) input->data[p]; - dst = (cl_mem)output->data[p]; - inter = (cl_mem) intermediate->data[p]; + dst = (cl_mem) output->data[p]; + inter = (cl_mem)intermediate->data[p]; if (!dst) break; - radius_x = ctx->radius; + radius_x = ctx->radiusH; radius_y = ctx->radiusV; if (!(ctx->planes & (1 << p))) { @@ -151,88 +215,94 @@ static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) radius_y = 0; } - cle = clSetKernelArg(ctx->kernel_horiz, 0, sizeof(cl_mem), &inter); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "destination image argument: %d.\n", cle); - err = AVERROR_UNKNOWN; - goto fail; - } - cle = clSetKernelArg(ctx->kernel_horiz, 1, sizeof(cl_mem), &src); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "source image argument: %d.\n", cle); - err = AVERROR_UNKNOWN; - goto fail; - } - cle = clSetKernelArg(ctx->kernel_horiz, 2, sizeof(cl_int), &radius_x); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "sizeX argument: %d.\n", cle); - err = AVERROR_UNKNOWN; - goto fail; - } - - err = ff_opencl_filter_work_size_from_image(avctx, global_work, - intermediate, p, 0); - if (err < 0) - goto fail; - - av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d " - "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n", - p, global_work[0], global_work[1]); - - cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel_horiz, 2, NULL, - global_work, NULL, - 0, NULL, NULL); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n", - cle); - err = AVERROR(EIO); - goto fail; - } - - cle = clSetKernelArg(ctx->kernel_vert, 0, sizeof(cl_mem), &dst); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "destination image argument: %d.\n", cle); - err = AVERROR_UNKNOWN; - goto fail; - } - cle = clSetKernelArg(ctx->kernel_vert, 1, sizeof(cl_mem), &inter); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "source image argument: %d.\n", cle); - err = AVERROR_UNKNOWN; - goto fail; - } - cle = clSetKernelArg(ctx->kernel_vert, 2, sizeof(cl_int), &radius_y); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "sizeY argument: %d.\n", cle); - err = AVERROR_UNKNOWN; - goto fail; - } - - err = ff_opencl_filter_work_size_from_image(avctx, global_work, - output, p, 0); - if (err < 0) - goto fail; - av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d " "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n", p, global_work[0], global_work[1]); - cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel_vert, 2, NULL, - global_work, NULL, - 0, NULL, NULL); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n", - cle); - err = AVERROR(EIO); - goto fail; + for (i = 0; i < ctx->power[p]; i++) { + cle = clSetKernelArg(ctx->kernel_horiz, 0, sizeof(cl_mem), &inter); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "destination image argument: %d.\n", cle); + err = AVERROR_UNKNOWN; + goto fail; + } + cle = clSetKernelArg(ctx->kernel_horiz, 1, sizeof(cl_mem), i == 0 ? &src : &dst); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "source image argument: %d.\n", cle); + err = AVERROR_UNKNOWN; + goto fail; + } + + if (!strcmp(avctx->filter->name, "avgblur_opencl")) { + cle = clSetKernelArg(ctx->kernel_horiz, 2, sizeof(cl_int), &radius_x); + } else if (!strcmp(avctx->filter->name, "boxblur_opencl")) { + cle = clSetKernelArg(ctx->kernel_horiz, 2, sizeof(cl_int), &ctx->radius[p]); + } + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "radius argument: %d.\n", cle); + err = AVERROR_UNKNOWN; + goto fail; + } + err = ff_opencl_filter_work_size_from_image(avctx, global_work, + i == 0 ? intermediate : output, p, 0); + if (err < 0) + goto fail; + + cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel_horiz, 2, NULL, + global_work, NULL, + 0, NULL, NULL); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n", + cle); + err = AVERROR(EIO); + goto fail; + } + cle = clFinish(ctx->command_queue); + + err = ff_opencl_filter_work_size_from_image(avctx, global_work, + i == 0 ? output : intermediate, p, 0); + + + cle = clSetKernelArg(ctx->kernel_vert, 0, sizeof(cl_mem), &dst); + + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "destination image argument: %d.\n", cle); + err = AVERROR_UNKNOWN; + goto fail; + } + cle = clSetKernelArg(ctx->kernel_vert, 1, sizeof(cl_mem), &inter); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "source image argument: %d.\n", cle); + err = AVERROR_UNKNOWN; + goto fail; + } + if (!strcmp(avctx->filter->name, "avgblur_opencl")) { + cle = clSetKernelArg(ctx->kernel_vert, 2, sizeof(cl_int), &radius_y); + } else if (!strcmp(avctx->filter->name, "boxblur_opencl")) { + cle = clSetKernelArg(ctx->kernel_vert, 2, sizeof(cl_int), &ctx->radius[p]); + } + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "radius argument: %d.\n", cle); + err = AVERROR_UNKNOWN; + goto fail; + } + + cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel_vert, 2, NULL, + global_work, NULL, + 0, NULL, NULL); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n", + cle); + err = AVERROR(EIO); + goto fail; + } } - } cle = clFinish(ctx->command_queue); @@ -264,12 +334,12 @@ fail: return err; } + static av_cold void avgblur_opencl_uninit(AVFilterContext *avctx) { AverageBlurOpenCLContext *ctx = avctx->priv; cl_int cle; - if (ctx->kernel_horiz) { cle = clReleaseKernel(ctx->kernel_horiz); if (cle != CL_SUCCESS) @@ -294,16 +364,6 @@ static av_cold void avgblur_opencl_uninit(AVFilterContext *avctx) ff_opencl_filter_uninit(avctx); } -#define OFFSET(x) offsetof(AverageBlurOpenCLContext, x) -#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) -static const AVOption avgblur_opencl_options[] = { - { "sizeX", "set horizontal size", OFFSET(radius), AV_OPT_TYPE_INT, {.i64=1}, 1, 1024, FLAGS }, - { "planes", "set planes to filter", OFFSET(planes), AV_OPT_TYPE_INT, {.i64=0xF}, 0, 0xF, FLAGS }, - { "sizeY", "set vertical size", OFFSET(radiusV), AV_OPT_TYPE_INT, {.i64=0}, 0, 1024, FLAGS }, - { NULL } -}; - -AVFILTER_DEFINE_CLASS(avgblur_opencl); static const AVFilterPad avgblur_opencl_inputs[] = { { @@ -315,6 +375,7 @@ static const AVFilterPad avgblur_opencl_inputs[] = { { NULL } }; + static const AVFilterPad avgblur_opencl_outputs[] = { { .name = "default", @@ -324,6 +385,22 @@ static const AVFilterPad avgblur_opencl_outputs[] = { { NULL } }; + +#define OFFSET(x) offsetof(AverageBlurOpenCLContext, x) +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) + +#if CONFIG_AVGBLUR_OPENCL_FILTER + +static const AVOption avgblur_opencl_options[] = { + { "sizeX", "set horizontal size", OFFSET(radiusH), AV_OPT_TYPE_INT, {.i64=1}, 1, 1024, FLAGS }, + { "planes", "set planes to filter", OFFSET(planes), AV_OPT_TYPE_INT, {.i64=0xF}, 0, 0xF, FLAGS }, + { "sizeY", "set vertical size", OFFSET(radiusV), AV_OPT_TYPE_INT, {.i64=0}, 0, 1024, FLAGS }, + { NULL } +}; + +AVFILTER_DEFINE_CLASS(avgblur_opencl); + + AVFilter ff_vf_avgblur_opencl = { .name = "avgblur_opencl", .description = NULL_IF_CONFIG_SMALL("Apply average blur filter"), @@ -336,3 +413,44 @@ AVFilter ff_vf_avgblur_opencl = { .outputs = avgblur_opencl_outputs, .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, }; + +#endif /* CONFIG_AVGBLUR_OPENCL_FILTER */ + + +#if CONFIG_BOXBLUR_OPENCL_FILTER + +static const AVOption boxblur_opencl_options[] = { + { "luma_radius", "Radius of the luma blurring box", OFFSET(luma_param.radius_expr), AV_OPT_TYPE_STRING, {.str="2"}, .flags = FLAGS }, + { "lr", "Radius of the luma blurring box", OFFSET(luma_param.radius_expr), AV_OPT_TYPE_STRING, {.str="2"}, .flags = FLAGS }, + { "luma_power", "How many times should the boxblur be applied to luma", OFFSET(luma_param.power), AV_OPT_TYPE_INT, {.i64=2}, 0, INT_MAX, .flags = FLAGS }, + { "lp", "How many times should the boxblur be applied to luma", OFFSET(luma_param.power), AV_OPT_TYPE_INT, {.i64=2}, 0, INT_MAX, .flags = FLAGS }, + + { "chroma_radius", "Radius of the chroma blurring box", OFFSET(chroma_param.radius_expr), AV_OPT_TYPE_STRING, {.str=NULL}, .flags = FLAGS }, + { "cr", "Radius of the chroma blurring box", OFFSET(chroma_param.radius_expr), AV_OPT_TYPE_STRING, {.str=NULL}, .flags = FLAGS }, + { "chroma_power", "How many times should the boxblur be applied to chroma", OFFSET(chroma_param.power), AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX, .flags = FLAGS }, + { "cp", "How many times should the boxblur be applied to chroma", OFFSET(chroma_param.power), AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX, .flags = FLAGS }, + + { "alpha_radius", "Radius of the alpha blurring box", OFFSET(alpha_param.radius_expr), AV_OPT_TYPE_STRING, {.str=NULL}, .flags = FLAGS }, + { "ar", "Radius of the alpha blurring box", OFFSET(alpha_param.radius_expr), AV_OPT_TYPE_STRING, {.str=NULL}, .flags = FLAGS }, + { "alpha_power", "How many times should the boxblur be applied to alpha", OFFSET(alpha_param.power), AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX, .flags = FLAGS }, + { "ap", "How many times should the boxblur be applied to alpha", OFFSET(alpha_param.power), AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX, .flags = FLAGS }, + + { NULL } +}; + +AVFILTER_DEFINE_CLASS(boxblur_opencl); + +AVFilter ff_vf_boxblur_opencl = { + .name = "boxblur_opencl", + .description = NULL_IF_CONFIG_SMALL("Apply boxblur filter to input video"), + .priv_size = sizeof(AverageBlurOpenCLContext), + .priv_class = &boxblur_opencl_class, + .init = &ff_opencl_filter_init, + .uninit = &avgblur_opencl_uninit, + .query_formats = &ff_opencl_filter_query_formats, + .inputs = avgblur_opencl_inputs, + .outputs = avgblur_opencl_outputs, + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +}; + +#endif /* CONFIG_BOXBLUR_OPENCL_FILTER */ diff --git a/libavfilter/vf_boxblur.c b/libavfilter/vf_boxblur.c index 8e43986..8368c88 100644 --- a/libavfilter/vf_boxblur.c +++ b/libavfilter/vf_boxblur.c @@ -27,39 +27,13 @@ #include "libavutil/avstring.h" #include "libavutil/common.h" -#include "libavutil/eval.h" #include "libavutil/opt.h" -#include "libavutil/pixdesc.h" #include "avfilter.h" #include "formats.h" #include "internal.h" #include "video.h" +#include "boxblur.h" -static const char *const var_names[] = { - "w", - "h", - "cw", - "ch", - "hsub", - "vsub", - NULL -}; - -enum var_name { - VAR_W, - VAR_H, - VAR_CW, - VAR_CH, - VAR_HSUB, - VAR_VSUB, - VARS_NB -}; - -typedef struct FilterParam { - int radius; - int power; - char *radius_expr; -} FilterParam; typedef struct BoxBlurContext { const AVClass *class; @@ -73,40 +47,6 @@ typedef struct BoxBlurContext { uint8_t *temp[2]; ///< temporary buffer used in blur_power() } BoxBlurContext; -#define Y 0 -#define U 1 -#define V 2 -#define A 3 - -static av_cold int init(AVFilterContext *ctx) -{ - BoxBlurContext *s = ctx->priv; - - if (!s->luma_param.radius_expr) { - av_log(ctx, AV_LOG_ERROR, "Luma radius expression is not set.\n"); - return AVERROR(EINVAL); - } - - /* fill missing params */ - if (!s->chroma_param.radius_expr) { - s->chroma_param.radius_expr = av_strdup(s->luma_param.radius_expr); - if (!s->chroma_param.radius_expr) - return AVERROR(ENOMEM); - } - if (s->chroma_param.power < 0) - s->chroma_param.power = s->luma_param.power; - - if (!s->alpha_param.radius_expr) { - s->alpha_param.radius_expr = av_strdup(s->luma_param.radius_expr); - if (!s->alpha_param.radius_expr) - return AVERROR(ENOMEM); - } - if (s->alpha_param.power < 0) - s->alpha_param.power = s->luma_param.power; - - return 0; -} - static av_cold void uninit(AVFilterContext *ctx) { BoxBlurContext *s = ctx->priv; @@ -150,48 +90,16 @@ static int config_input(AVFilterLink *inlink) s->hsub = desc->log2_chroma_w; s->vsub = desc->log2_chroma_h; - var_values[VAR_W] = inlink->w; - var_values[VAR_H] = inlink->h; - var_values[VAR_CW] = cw = w>>s->hsub; - var_values[VAR_CH] = ch = h>>s->vsub; - var_values[VAR_HSUB] = 1<hsub; - var_values[VAR_VSUB] = 1<vsub; - -#define EVAL_RADIUS_EXPR(comp) \ - expr = s->comp##_param.radius_expr; \ - ret = av_expr_parse_and_eval(&res, expr, var_names, var_values, \ - NULL, NULL, NULL, NULL, NULL, 0, ctx); \ - s->comp##_param.radius = res; \ - if (ret < 0) { \ - av_log(NULL, AV_LOG_ERROR, \ - "Error when evaluating " #comp " radius expression '%s'\n", expr); \ - return ret; \ - } - EVAL_RADIUS_EXPR(luma); - EVAL_RADIUS_EXPR(chroma); - EVAL_RADIUS_EXPR(alpha); - - av_log(ctx, AV_LOG_VERBOSE, - "luma_radius:%d luma_power:%d " - "chroma_radius:%d chroma_power:%d " - "alpha_radius:%d alpha_power:%d " - "w:%d chroma_w:%d h:%d chroma_h:%d\n", - s->luma_param .radius, s->luma_param .power, - s->chroma_param.radius, s->chroma_param.power, - s->alpha_param .radius, s->alpha_param .power, - w, cw, h, ch); - -#define CHECK_RADIUS_VAL(w_, h_, comp) \ - if (s->comp##_param.radius < 0 || \ - 2*s->comp##_param.radius > FFMIN(w_, h_)) { \ - av_log(ctx, AV_LOG_ERROR, \ - "Invalid " #comp " radius value %d, must be >= 0 and <= %d\n", \ - s->comp##_param.radius, FFMIN(w_, h_)/2); \ - return AVERROR(EINVAL); \ + ret = ff_boxblur_eval_filter_params(inlink, + &s->luma_param, + &s->chroma_param, + &s->alpha_param); + + if (ret != 0) { + av_log(ctx, AV_LOG_ERROR, "Failed to evaluate " + "filter params: %d.\n", ret); + return ret; } - CHECK_RADIUS_VAL(w, h, luma); - CHECK_RADIUS_VAL(cw, ch, chroma); - CHECK_RADIUS_VAL(w, h, alpha); s->radius[Y] = s->luma_param.radius; s->radius[U] = s->radius[V] = s->chroma_param.radius; @@ -404,7 +312,6 @@ AVFilter ff_vf_boxblur = { .description = NULL_IF_CONFIG_SMALL("Blur the input."), .priv_size = sizeof(BoxBlurContext), .priv_class = &boxblur_class, - .init = init, .uninit = uninit, .query_formats = query_formats, .inputs = avfilter_vf_boxblur_inputs,