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

Song, Ruiling ruiling.song at intel.com
Thu May 31 11:57:50 EEST 2018



> -----Original Message-----
> From: ffmpeg-devel [mailto:ffmpeg-devel-bounces at ffmpeg.org] On Behalf
> Of Danil Iashchenko
> Sent: Tuesday, May 29, 2018 4:44 AM
> To: ffmpeg-devel at ffmpeg.org
> Cc: Danil Iashchenko <danyaschenko at gmail.com>
> Subject: [FFmpeg-devel] [PATCH] libavfilter/boxblur_opencl filter.
> 
> Add opencl version of boxblur filter.
> Behaves like existing libavfilter/boxblur filter.
> 
> Resubmitting patch because I used unnecessary variable inside kernels. Now
> it's fixed.
> 
> ---
>  configure                       |   1 +
>  libavfilter/Makefile            |   2 +
>  libavfilter/allfilters.c        |   1 +
>  libavfilter/opencl/boxblur.cl   |  53 +++++
>  libavfilter/opencl_source.h     |   1 +
>  libavfilter/vf_boxblur_opencl.c | 468
> ++++++++++++++++++++++++++++++++++++++++
>  6 files changed, 526 insertions(+)
>  create mode 100644 libavfilter/opencl/boxblur.cl
>  create mode 100644 libavfilter/vf_boxblur_opencl.c
> 
> diff --git a/configure b/configure
> index e52f8f8..f8a2487 100755
> --- a/configure
> +++ b/configure
> @@ -3302,6 +3302,7 @@ avgblur_opencl_filter_deps="opencl"
>  azmq_filter_deps="libzmq"
>  blackframe_filter_deps="gpl"
>  boxblur_filter_deps="gpl"
> +boxblur_opencl_filter_deps="opencl"
>  bs2b_filter_deps="libbs2b"
>  colormatrix_filter_deps="gpl"
>  convolution_opencl_filter_deps="opencl"
> diff --git a/libavfilter/Makefile b/libavfilter/Makefile
> index c68ef05..97afdef 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_boxblur_opencl.o
> opencl.o \
> +	                                        opencl/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 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/opencl/boxblur.cl b/libavfilter/opencl/boxblur.cl
> new file mode 100644
> index 0000000..9ecb441
> --- /dev/null
> +++ b/libavfilter/opencl/boxblur.cl
> @@ -0,0 +1,53 @@
> +/*
> + * 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
> + */
> +
> +__kernel void boxblur_horiz(__write_only image2d_t dst,
> +                            __read_only  image2d_t src,
> +                            int rad)
> +{
> +    const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
> +                               CLK_ADDRESS_CLAMP_TO_EDGE   |
> +                               CLK_FILTER_NEAREST);
> +
> +    int2   loc = (int2)(get_global_id(0), get_global_id(1));
> +    float4 acc = (float4)(0,0,0,0);
> +
> +    for (int w = loc.x - rad; w < loc.x + rad + 1; w++) {
> +        acc += read_imagef(src, sampler, (int2)(w, loc.y));
> +    }
> +    write_imagef(dst, loc, acc / (rad * 2 + 1));
> +}
> +
> +__kernel void boxblur_vertic(__write_only image2d_t dst,
> +                             __read_only  image2d_t src,
> +                             int radv)
> +{
> +    const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
> +                               CLK_ADDRESS_CLAMP_TO_EDGE   |
> +                               CLK_FILTER_NEAREST);
> +
> +    int2   loc = (int2)(get_global_id(0), get_global_id(1));
> +    float4 acc = (float4)(0,0,0,0);
> +
> +    for (int h = loc.y - radv; h < loc.y + radv + 1; h++) {
> +        acc += read_imagef(src, sampler, (int2)(loc.x, h));
> +    }
> +    write_imagef(dst, loc, acc / (radv * 2 + 1));
> +}
> diff --git a/libavfilter/opencl_source.h b/libavfilter/opencl_source.h
> index 4bb9969..2e09770 100644
> --- a/libavfilter/opencl_source.h
> +++ b/libavfilter/opencl_source.h
> @@ -20,6 +20,7 @@
>  #define AVFILTER_OPENCL_SOURCE_H
> 
>  extern const char *ff_opencl_source_avgblur;
> +extern const char *ff_opencl_source_boxblur;
>  extern const char *ff_opencl_source_convolution;
>  extern const char *ff_opencl_source_overlay;
>  extern const char *ff_opencl_source_unsharp;
> diff --git a/libavfilter/vf_boxblur_opencl.c b/libavfilter/vf_boxblur_opencl.c
> new file mode 100644
> index 0000000..d46fb0f
> --- /dev/null
> +++ b/libavfilter/vf_boxblur_opencl.c
> @@ -0,0 +1,468 @@
> +/*
> + * 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 "libavutil/common.h"
> +#include "libavutil/imgutils.h"
> +#include "libavutil/mem.h"
> +#include "libavutil/opt.h"
> +#include "libavutil/pixdesc.h"
> +#include "libavutil/avstring.h"
> +#include "libavutil/eval.h"
> +
> +
> +#include "avfilter.h"
> +#include "internal.h"
> +#include "opencl.h"
> +#include "opencl_source.h"
> +#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 BoxBlurOpenCLContext {
> +    OpenCLFilterContext ocf;
> +
> +    int              initialised;
> +    cl_kernel        kernel_horiz;
> +    cl_kernel        kernel_vertic;
> +    cl_command_queue command_queue;
> +
> +    FilterParam luma_param;
> +    FilterParam chroma_param;
> +    FilterParam alpha_param;
> +
> +    int hsub, vsub;
> +    int radius[4];
> +    int power[4];
> +} BoxBlurOpenCLContext;
> +
> +
> +#define Y 0
> +#define U 1
> +#define V 2
> +#define A 3
> +
> +
> +static int boxblur_opencl_init(AVFilterContext *avctx)
> +{
> +    BoxBlurOpenCLContext *ctx = avctx->priv;
> +    cl_int cle;
> +    int err;
> +
> +    err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_boxblur,
> 1);
> +    if (err < 0)
> +        goto fail;
> +
> +    ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx-
> >context,
> +                                              ctx->ocf.hwctx->device_id,
> +                                              0, &cle);
> +    if (!ctx->command_queue) {
> +        av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
> +               "command queue: %d.\n", cle);
> +        err = AVERROR(EIO);
> +        goto fail;
> +    }
> +
> +    ctx->kernel_horiz = clCreateKernel(ctx->ocf.program, "boxblur_horiz",
> &cle);
> +    if (!ctx->kernel_horiz) {
> +        av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
> +        err = AVERROR(EIO);
> +        goto fail;
> +    }
> +    ctx->kernel_vertic = clCreateKernel(ctx->ocf.program, "boxblur_vertic",
> &cle);
> +    if (!ctx->kernel_vertic) {
> +        av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
> +        err = AVERROR(EIO);
> +        goto fail;
> +    }
> +
> +    ctx->initialised = 1;
> +    return 0;
> +
> +fail:
> +    if (ctx->command_queue)
> +        clReleaseCommandQueue(ctx->command_queue);
> +    if (ctx->kernel_horiz)
> +        clReleaseKernel(ctx->kernel_horiz);
> +    if (ctx->kernel_vertic)
> +        clReleaseKernel(ctx->kernel_vertic);
> +    return err;
> +}
> +
> +
> +static int boxblur_opencl_make_filter_params(AVFilterLink *inlink)
> +{
> +    const AVPixFmtDescriptor *desc = av_pix_fmt_desc_get(inlink->format);
> +    AVFilterContext    *ctx = inlink->dst;
> +    BoxBlurOpenCLContext *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->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;
> +
> +    return 0;
> +}
> +
> +
> +static int boxblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
> +{
> +    AVFilterContext *avctx = inlink->dst;
> +    AVFilterLink *outlink = avctx->outputs[0];
> +    BoxBlurOpenCLContext *ctx = avctx->priv;
> +    AVFrame *output = NULL;
> +    AVFrame *intermediate = NULL;
> +    cl_int cle;
> +    size_t global_work[2];
> +    cl_mem src, dst, inter;
> +    int err, p, i;
> +
> +    av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n",
> +           av_get_pix_fmt_name(input->format),
> +           input->width, input->height, input->pts);
> +
> +    if (!input->hw_frames_ctx)
> +        return AVERROR(EINVAL);
> +
> +    if (!ctx->initialised) {
> +        err = boxblur_opencl_init(avctx);
> +        if (err < 0)
> +            goto fail;
> +
> +        err = boxblur_opencl_make_filter_params(inlink);
> +        if (err < 0)
> +            goto fail;
> +    }
> +
> +    output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
> +    if (!output) {
> +        err = AVERROR(ENOMEM);
> +        goto fail;
> +    }
> +    intermediate = ff_get_video_buffer(outlink, outlink->w, outlink->h);
> +    if (!intermediate) {
> +        err = AVERROR(ENOMEM);
> +        goto fail;
> +    }
> +
> +    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];
> +
> +        if (!dst)
> +            break;
> +
> +        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 = 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_horiz "
> +                   "radius argument: %d.\n", cle);
> +            goto fail;
> +        }
> +
> +        for (i = 0; i < ctx->power[p]; i++) {
> +            cle = clSetKernelArg(ctx->kernel_horiz, 0, sizeof(cl_mem), &inter);
> +            if (cle != CL_SUCCESS) {
> +                av_log(avctx, AV_LOG_ERROR, "Failed to set kernel_horiz "
> +                       "destination image argument: %d.\n", cle);
> +                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_horiz "
> +                       "source image argument: %d.\n", cle);
> +                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;
> +            }
> +            src = inter;
So, for the next iteration, you will be reading from and writing to 'inter' at the same time, right?
If this is true, it is wrong, you should not do read and write operation on the same buffer unless you explicitly make sure the read/write operation do not have overlap.
For example, if the iteration count is two, you can do like src --> temp -->dst for the horizontal pass. Then when you do vertical pass later, you can do like dst-->temp-->dst.

Thanks!
Ruiling
> +        }
> +
> +        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 = clSetKernelArg(ctx->kernel_vertic, 2, sizeof(cl_int), &ctx-
> >radius[p]);
> +        if (cle != CL_SUCCESS) {
> +            av_log(avctx, AV_LOG_ERROR, "Failed to set kernel_vertic "
> +                   "radius argument: %d.\n", cle);
> +            goto fail;
> +        }
> +
> +        for (i = 0; i < ctx->power[p]; i++) {
> +            cle = clSetKernelArg(ctx->kernel_vertic, 0, sizeof(cl_mem), &dst);
> +            if (cle != CL_SUCCESS) {
> +                av_log(avctx, AV_LOG_ERROR, "Failed to set kernel_vertic "
> +                       "destination image argument: %d.\n", cle);
> +                goto fail;
> +            }
> +            cle = clSetKernelArg(ctx->kernel_vertic, 1, sizeof(cl_mem), &inter);
> +            if (cle != CL_SUCCESS) {
> +                av_log(avctx, AV_LOG_ERROR, "Failed to set kernel_vertic "
> +                       "source image argument: %d.\n", cle);
> +                goto fail;
> +            }
> +            cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx-
> >kernel_vertic, 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;
> +            }
> +            inter = dst;
> +        }
> +
> +    }
> +
> +    cle = clFinish(ctx->command_queue);
> +    if (cle != CL_SUCCESS) {
> +        av_log(avctx, AV_LOG_ERROR, "Failed to finish command
> queue: %d.\n",
> +               cle);
> +        err = AVERROR(EIO);
> +        goto fail;
> +    }
> +
> +    err = av_frame_copy_props(output, input);
> +    if (err < 0)
> +        goto fail;
> +
> +    av_frame_free(&input);
> +    av_frame_free(&intermediate);
> +
> +    av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
> +           av_get_pix_fmt_name(output->format),
> +           output->width, output->height, output->pts);
> +
> +    return ff_filter_frame(outlink, output);
> +
> +fail:
> +    clFinish(ctx->command_queue);
> +    av_frame_free(&input);
> +    av_frame_free(&output);
> +    av_frame_free(&intermediate);
> +    return err;
> +}
> +
> +static av_cold void boxblur_opencl_uninit(AVFilterContext *avctx)
> +{
> +    BoxBlurOpenCLContext *ctx = avctx->priv;
> +    cl_int cle;
> +
> +    if (ctx->kernel_horiz) {
> +        cle = clReleaseKernel(ctx->kernel_horiz);
> +        if (cle != CL_SUCCESS)
> +            av_log(avctx, AV_LOG_ERROR, "Failed to release "
> +                   "kernel: %d.\n", cle);
> +    }
> +    if (ctx->kernel_vertic) {
> +        cle = clReleaseKernel(ctx->kernel_vertic);
> +        if (cle != CL_SUCCESS)
> +            av_log(avctx, AV_LOG_ERROR, "Failed to release "
> +                   "kernel: %d.\n", cle);
> +    }
> +
> +    if (ctx->command_queue) {
> +        cle = clReleaseCommandQueue(ctx->command_queue);
> +        if (cle != CL_SUCCESS)
> +            av_log(avctx, AV_LOG_ERROR, "Failed to release "
> +                   "command queue: %d.\n", cle);
> +    }
> +
> +    ff_opencl_filter_uninit(avctx);
> +}
> +
> +#define OFFSET(x) offsetof(BoxBlurOpenCLContext, x)
> +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM |
> AV_OPT_FLAG_VIDEO_PARAM)
> +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);
> +
> +static const AVFilterPad boxblur_opencl_inputs[] = {
> +    {
> +        .name         = "default",
> +        .type         = AVMEDIA_TYPE_VIDEO,
> +        .filter_frame = &boxblur_opencl_filter_frame,
> +        .config_props = &ff_opencl_filter_config_input,
> +    },
> +    { NULL }
> +};
> +
> +static const AVFilterPad boxblur_opencl_outputs[] = {
> +    {
> +        .name         = "default",
> +        .type         = AVMEDIA_TYPE_VIDEO,
> +        .config_props = &ff_opencl_filter_config_output,
> +    },
> +    { NULL }
> +};
> +
> +AVFilter ff_vf_boxblur_opencl = {
> +    .name           = "boxblur_opencl",
> +    .description    = NULL_IF_CONFIG_SMALL("Apply boxblur filter to input
> video"),
> +    .priv_size      = sizeof(BoxBlurOpenCLContext),
> +    .priv_class     = &boxblur_opencl_class,
> +    .init           = &ff_opencl_filter_init,
> +    .uninit         = &boxblur_opencl_uninit,
> +    .query_formats  = &ff_opencl_filter_query_formats,
> +    .inputs         = boxblur_opencl_inputs,
> +    .outputs        = boxblur_opencl_outputs,
> +    .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
> +};
> --
> 2.7.4
> 
> _______________________________________________
> ffmpeg-devel mailing list
> ffmpeg-devel at ffmpeg.org
> http://ffmpeg.org/mailman/listinfo/ffmpeg-devel


More information about the ffmpeg-devel mailing list