lavfi: Add boxblur_opencl filter

Behaves like existing boxblur filter.

6 files changed, 372 insertions, 209 deletions

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 <michaelni@gmx.at>
+ * 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 <michaelni@gmx.at>
+ * 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<<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 +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,