[FFmpeg-devel] [PATCH] libavfilter/boxblur_opencl filter.

Danil Iashchenko danyaschenko at gmail.com
Wed Jun 6 02:45:29 EEST 2018


Behaves like existing boxblur filter. 

---

Thanks! Fixed.

 libavfilter/Makefile            |   2 +
 libavfilter/allfilters.c        |   1 +
 libavfilter/vf_avgblur_opencl.c | 419 ++++++++++++++++++++++++++++++----------
 3 files changed, 324 insertions(+), 98 deletions(-)

diff --git a/libavfilter/Makefile b/libavfilter/Makefile
index c68ef05..6f00059 100644
--- a/libavfilter/Makefile
+++ b/libavfilter/Makefile
@@ -153,6 +153,8 @@ 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_OPENCL_FILTER)         += vf_avgblur_opencl.o opencl.o \
+	                                        opencl/avgblur.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/vf_avgblur_opencl.c b/libavfilter/vf_avgblur_opencl.c
index 48cebb5..d4759de 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.
  *
@@ -23,6 +24,7 @@
 #include "libavutil/mem.h"
 #include "libavutil/opt.h"
 #include "libavutil/pixdesc.h"
+#include "libavutil/eval.h"
 
 #include "avfilter.h"
 #include "internal.h"
@@ -31,6 +33,35 @@
 #include "video.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 AverageBlurOpenCLContext {
     OpenCLFilterContext ocf;
 
@@ -39,13 +70,26 @@ 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 hsub, vsub;
+    int radius[4];
+    int power[4];
+
 } AverageBlurOpenCLContext;
 
 
+#define Y 0
+#define U 1
+#define V 2
+#define A 3
+
+
 static int avgblur_opencl_init(AVFilterContext *avctx)
 {
     AverageBlurOpenCLContext *ctx = avctx->priv;
@@ -80,10 +124,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 +137,122 @@ 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)
+{
+    const AVPixFmtDescriptor *desc = av_pix_fmt_desc_get(inlink->format);
+    AVFilterContext    *ctx = inlink->dst;
+    AverageBlurOpenCLContext *s = ctx->priv;
+    int w = inlink->w, h = inlink->h;
+    int cw, ch;
+    double var_values[VARS_NB], res;
+    char *expr;
+    int ret, i;
+
+    if (!s->luma_param.radius_expr) {
+        av_log(s, 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;
+
+    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);                                         \
+    }
+    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;
+    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 +263,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 +277,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 +294,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 +302,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 +316,98 @@ 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;
+        for (i = 0; i < ctx->power[p]; i++) {
+            cle = clSetKernelArg(ctx->kernel_horiz, 0, sizeof(cl_mem), i == 0 ? &inter : &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_horiz, 1, sizeof(cl_mem), i == 0 ? &src : &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_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), i == 0 ? &dst : &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_vert, 1, sizeof(cl_mem), i == 0 ? &inter : &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_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);
+            if ((i == 0 && ctx->power[p] > 1) || (i && i == ctx->power[p] - 1)) {
+                FFSWAP(cl_mem, inter, dst);
+            }
         }
-        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;
-        }
-
     }
 
     cle = clFinish(ctx->command_queue);
@@ -264,12 +439,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 +469,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 +480,7 @@ static const AVFilterPad avgblur_opencl_inputs[] = {
     { NULL }
 };
 
+
 static const AVFilterPad avgblur_opencl_outputs[] = {
     {
         .name         = "default",
@@ -324,6 +490,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)
+
+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);
+
+
+#if CONFIG_AVGBLUR_OPENCL_FILTER
+
 AVFilter ff_vf_avgblur_opencl = {
     .name           = "avgblur_opencl",
     .description    = NULL_IF_CONFIG_SMALL("Apply average blur filter"),
@@ -336,3 +518,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 */
-- 
2.7.4



More information about the ffmpeg-devel mailing list