From 714da1fd898f83c7bef38fe427af3692917cbcb2 Mon Sep 17 00:00:00 2001 From: Danil Iashchenko Date: Wed, 20 Jun 2018 04:33:02 +0300 Subject: lavfi: Add boxblur_opencl filter Behaves like existing boxblur filter. --- libavfilter/Makefile | 6 +- libavfilter/allfilters.c | 1 + libavfilter/boxblur.c | 127 ++++++++++++++++++ libavfilter/boxblur.h | 48 +++++++ libavfilter/vf_avgblur_opencl.c | 283 ++++++++++++++++++++++++++-------------- libavfilter/vf_boxblur.c | 116 ++-------------- 6 files changed, 372 insertions(+), 209 deletions(-) create mode 100644 libavfilter/boxblur.c create mode 100644 libavfilter/boxblur.h (limited to 'libavfilter') diff --git a/libavfilter/Makefile b/libavfilter/Makefile index 5b4be4966c..552499558d 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -149,14 +149,16 @@ OBJS-$(CONFIG_ASS_FILTER) += vf_subtitles.o OBJS-$(CONFIG_ATADENOISE_FILTER) += vf_atadenoise.o OBJS-$(CONFIG_AVGBLUR_FILTER) += vf_avgblur.o OBJS-$(CONFIG_AVGBLUR_OPENCL_FILTER) += vf_avgblur_opencl.o opencl.o \ - opencl/avgblur.o + opencl/avgblur.o boxblur.o OBJS-$(CONFIG_BBOX_FILTER) += bbox.o vf_bbox.o OBJS-$(CONFIG_BENCH_FILTER) += f_bench.o 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 f2d27d2424..2b44626028 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -148,6 +148,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 0000000000..4534b456d9 --- /dev/null +++ b/libavfilter/boxblur.c @@ -0,0 +1,127 @@ +/* + * Copyright (c) 2002 Michael Niedermayer + * Copyright (c) 2011 Stefano Sabatini + * 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" + +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 +}; + + +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 0000000000..5694722a9d --- /dev/null +++ b/libavfilter/boxblur.h @@ -0,0 +1,48 @@ +/* + * Copyright (c) 2002 Michael Niedermayer + * Copyright (c) 2011 Stefano Sabatini + * 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 AVFILTER_BOXBLUR_H +#define AVFILTER_BOXBLUR_H + +#include "libavutil/eval.h" +#include "libavutil/pixdesc.h" +#include "libavutil/mem.h" + +#include "avfilter.h" + +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 // AVFILTER_BOXBLUR_H diff --git a/libavfilter/vf_avgblur_opencl.c b/libavfilter/vf_avgblur_opencl.c index 48cebb5887..d1d3eb1983 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,57 @@ 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++) { + CL_SET_KERNEL_ARG(ctx->kernel_horiz, 0, cl_mem, &inter); + CL_SET_KERNEL_ARG(ctx->kernel_horiz, 1, cl_mem, i == 0 ? &src : &dst); + if (!strcmp(avctx->filter->name, "avgblur_opencl")) { + CL_SET_KERNEL_ARG(ctx->kernel_horiz, 2, cl_int, &radius_x); + } else if (!strcmp(avctx->filter->name, "boxblur_opencl")) { + CL_SET_KERNEL_ARG(ctx->kernel_horiz, 2, cl_int, &ctx->radius[p]); + } + + 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); + + CL_SET_KERNEL_ARG(ctx->kernel_vert, 0, cl_mem, &dst); + CL_SET_KERNEL_ARG(ctx->kernel_vert, 1, cl_mem, &inter); + + if (!strcmp(avctx->filter->name, "avgblur_opencl")) { + CL_SET_KERNEL_ARG(ctx->kernel_vert, 2, cl_int, &radius_y); + } else if (!strcmp(avctx->filter->name, "boxblur_opencl")) { + CL_SET_KERNEL_ARG(ctx->kernel_vert, 2, cl_int, &ctx->radius[p]); + } + + 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 +297,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 +327,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 +338,7 @@ static const AVFilterPad avgblur_opencl_inputs[] = { { NULL } }; + static const AVFilterPad avgblur_opencl_outputs[] = { { .name = "default", @@ -324,6 +348,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 +376,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 8e43986846..bde74084f0 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; @@ -138,9 +78,6 @@ static int config_input(AVFilterLink *inlink) AVFilterContext *ctx = inlink->dst; BoxBlurContext *s = ctx->priv; int w = inlink->w, h = inlink->h; - int cw, ch; - double var_values[VARS_NB], res; - char *expr; int ret; if (!(s->temp[0] = av_malloc(2*FFMAX(w, h))) || @@ -150,48 +87,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 +309,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, -- cgit v1.2.3