[FFmpeg-devel] [PATCH] libavfilter/boxblur_opencl filter.
Danil Iashchenko
danyaschenko at gmail.com
Tue Jun 19 14:43:40 EEST 2018
lavfi: Add boxblur_opencl filter. Behaves like existing boxblur filter.
---
Refactored to use CL_SET_KERNEL_ARG().
configure | 1 +
libavfilter/Makefile | 6 +-
libavfilter/allfilters.c | 1 +
libavfilter/boxblur.c | 105 +++++++++++++++
libavfilter/boxblur.h | 66 ++++++++++
libavfilter/vf_avgblur_opencl.c | 283 ++++++++++++++++++++++++++--------------
libavfilter/vf_boxblur.c | 113 ++--------------
7 files changed, 369 insertions(+), 206 deletions(-)
create mode 100644 libavfilter/boxblur.c
create mode 100644 libavfilter/boxblur.h
diff --git a/configure b/configure
index 333e326..8ca2586 100755
--- a/configure
+++ b/configure
@@ -3311,6 +3311,7 @@ avgblur_opencl_filter_deps="opencl"
azmq_filter_deps="libzmq"
blackframe_filter_deps="gpl"
boxblur_filter_deps="gpl"
+boxblur_opencl_filter_deps="opencl gpl"
bs2b_filter_deps="libbs2b"
colormatrix_filter_deps="gpl"
convolution_opencl_filter_deps="opencl"
diff --git a/libavfilter/Makefile b/libavfilter/Makefile
index 5b4be49..5524995 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 f2d27d2..2b44626 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 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..d1d3eb1 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 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<<s->hsub;
- var_values[VAR_VSUB] = 1<<s->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,
--
2.7.4
More information about the ffmpeg-devel
mailing list