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

Mark Thompson sw at jkqxz.net
Fri Jun 8 01:56:07 EEST 2018


On 06/06/18 00:45, Danil Iashchenko wrote:
> 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
     ^
There's a tab here.

>  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
> ...
> +
> +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;

Most of this function is duplicating code from vf_boxblur.c.  Can you move it into another file (boxblur.c?) and then both filters can call it from there?

> +}
> +
> +
>  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);

I don't think you should need to clFinish() after each step?

Also the return value of clFinish() should be checked.

> +            if ((i == 0 && ctx->power[p] > 1) || (i && i == ctx->power[p] - 1)) {
> +                FFSWAP(cl_mem, inter, dst);
> +            }

So the first step does

src   -- horizontal -> inter
inter --  vertical  -> dst

and every step thereafter does:

inter -- horizontal -> dst
dst   --  vertical  -> inter

but dst and inter got swapped after the first step?

After some thought I think that does the right thing, but it could be clearer.  Possibly I am not getting this right, but I think something like:

i == 0 ? src : dst -- horizontal -> inter
inter              --  vertical  -> dst

would do the right thing without any swapping?

>          }
> -        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);

The options and class definition should be inside the #if.

> +
> +
> +#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 */
> 

Doing some testing with this it all looks good.

Thanks,

- Mark


More information about the ffmpeg-devel mailing list