[FFmpeg-devel] lavfi: add gblur_opencl filter

---
 configure                     |   1 +
 libavfilter/Makefile          |   2 +
 libavfilter/allfilters.c      |   1 +
 libavfilter/opencl/gblur.cl   |  62 +++++++
 libavfilter/opencl_source.h   |   1 +
 libavfilter/vf_gblur_opencl.c | 370 ++++++++++++++++++++++++++++++++++++++++++
 6 files changed, 437 insertions(+)
 create mode 100644 libavfilter/opencl/gblur.cl
 create mode 100644 libavfilter/vf_gblur_opencl.c

On Thu, Apr 25, 2019 at 22:41:49 +1000, Dylan Fernando wrote:

> +static const AVOption gblur_opencl_options[] = {
> +    { "sigma",  "set horizontal size",  OFFSET(sigma), AV_OPT_TYPE_FLOAT, {.dbl=0.5},   0.0, 1024, FLAGS },
                   ^

Shouldn't this be "set sigma"?

> +    { "planes", "set planes to filter", OFFSET(planes),  AV_OPT_TYPE_INT, {.i64=0xF}, 0,  0xF, FLAGS },
> +    { "sigmaV", "set vertical sigma",   OFFSET(sigmaV), AV_OPT_TYPE_FLOAT, {.dbl=-1},   -1, 1024, FLAGS },
> +    { NULL }

Please also add documentation to doc/filters.texi. You could do this by
reference to the "gblur" filter, but as this filter's options are
slightly different, you will need to copy the relevant sections.

Furthermore, CamelCase variables aren't usually accepted for ffmpeg own
variables, but this mirrors the gblur filter, so - oh well.

> +    matrix_horiz = av_malloc(matrix_bytes_horiz);
> +    if (!matrix_horiz) {
> +        av_freep(&matrix_horiz);

If av_malloc() returned 0/NULL, does it ever need to be freed???

> +    matrix_vert = av_malloc(matrix_bytes_vert);
> +    if (!matrix_vert) {
> +        av_freep(&matrix_vert);

Ditto

I can't judge on the rest.

Cheers,
Moritz

On Thu, 25 Apr 2019 at 10:42 pm, Dylan Fernando <dylanf123@gmail.com> wrote:

> ---
>  configure                     |   1 +
>  libavfilter/Makefile          |   2 +
>  libavfilter/allfilters.c      |   1 +
>  libavfilter/opencl/gblur.cl   |  62 +++++++
>  libavfilter/opencl_source.h   |   1 +
>  libavfilter/vf_gblur_opencl.c | 370
> ++++++++++++++++++++++++++++++++++++++++++
>  6 files changed, 437 insertions(+)
>  create mode 100644 libavfilter/opencl/gblur.cl
>  create mode 100644 libavfilter/vf_gblur_opencl.c
>
> diff --git a/configure b/configure
> index bbeaf2fadc..8c1d3cdf92 100755
> --- a/configure
> +++ b/configure
> @@ -3451,6 +3451,7 @@ freezedetect_filter_select="scene_sad"
>  frei0r_filter_deps="frei0r libdl"
>  frei0r_src_filter_deps="frei0r libdl"
>  fspp_filter_deps="gpl"
> +gblur_opencl_filter_deps="opencl"
>  geq_filter_deps="gpl"
>  histeq_filter_deps="gpl"
>  hqdn3d_filter_deps="gpl"
> diff --git a/libavfilter/Makefile b/libavfilter/Makefile
> index fef6ec5c55..230315ef39 100644
> --- a/libavfilter/Makefile
> +++ b/libavfilter/Makefile
> @@ -243,6 +243,8 @@ OBJS-$(CONFIG_FREEZEDETECT_FILTER)           +=
> vf_freezedetect.o
>  OBJS-$(CONFIG_FREI0R_FILTER)                 += vf_frei0r.o
>  OBJS-$(CONFIG_FSPP_FILTER)                   += vf_fspp.o
>  OBJS-$(CONFIG_GBLUR_FILTER)                  += vf_gblur.o
> +OBJS-$(CONFIG_GBLUR_OPENCL_FILTER)           += vf_gblur_opencl.o
> opencl.o \
> +                                                opencl/gblur.o
>  OBJS-$(CONFIG_GEQ_FILTER)                    += vf_geq.o
>  OBJS-$(CONFIG_GRADFUN_FILTER)                += vf_gradfun.o
>  OBJS-$(CONFIG_GRAPHMONITOR_FILTER)           += f_graphmonitor.o
> diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c
> index c51ae0f3c7..cb0fc051cc 100644
> --- a/libavfilter/allfilters.c
> +++ b/libavfilter/allfilters.c
> @@ -229,6 +229,7 @@ extern AVFilter ff_vf_freezedetect;
>  extern AVFilter ff_vf_frei0r;
>  extern AVFilter ff_vf_fspp;
>  extern AVFilter ff_vf_gblur;
> +extern AVFilter ff_vf_gblur_opencl;
>  extern AVFilter ff_vf_geq;
>  extern AVFilter ff_vf_gradfun;
>  extern AVFilter ff_vf_graphmonitor;
> diff --git a/libavfilter/opencl/gblur.cl b/libavfilter/opencl/gblur.cl
> new file mode 100644
> index 0000000000..4fece30d4a
> --- /dev/null
> +++ b/libavfilter/opencl/gblur.cl
> @@ -0,0 +1,62 @@
> +/*
> + * Copyright (c) 2018 Dylan Fernando
> + *
> + * 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 gblur_conv_horz(__write_only image2d_t dst,
> +                              __read_only  image2d_t src,
> +                              int coef_matrix_dim,
> +                              __constant float *coef_matrix)
> +{
> +    const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
> +                               CLK_ADDRESS_CLAMP_TO_EDGE   |
> +                               CLK_FILTER_NEAREST);
> +
> +    const int half_matrix_dim = (coef_matrix_dim / 2);
> +    int2 loc = (int2)(get_global_id(0), get_global_id(1));
> +    float4 convPix = (float4)(0.0f, 0.0f, 0.0f, 0.0f);
> +
> +    for (int conv_j = -half_matrix_dim; conv_j <= half_matrix_dim;
> conv_j++) {
> +        float4 px = read_imagef(src, sampler, loc + (int2)(conv_j, 0));
> +        convPix += px * coef_matrix[(conv_j + half_matrix_dim)];
> +    }
> +
> +    write_imagef(dst, loc, convPix);
> +}
> +
> +__kernel void gblur_conv_vert(__write_only image2d_t dst,
> +                              __read_only  image2d_t src,
> +                              int coef_matrix_dim,
> +                              __constant float *coef_matrix)
> +{
> +    const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
> +                               CLK_ADDRESS_CLAMP_TO_EDGE   |
> +                               CLK_FILTER_NEAREST);
> +
> +    const int half_matrix_dim = (coef_matrix_dim / 2);
> +    int2 loc = (int2)(get_global_id(0), get_global_id(1));
> +    float4 convPix = (float4)(0.0f, 0.0f, 0.0f, 0.0f);
> +
> +    for (int conv_j = -half_matrix_dim; conv_j <= half_matrix_dim;
> conv_j++) {
> +        float4 px = read_imagef(src, sampler, loc + (int2)(0, conv_j));
> +        convPix += px * coef_matrix[(conv_j + half_matrix_dim)];
> +    }
> +
> +    write_imagef(dst, loc, convPix);
> +}
> diff --git a/libavfilter/opencl_source.h b/libavfilter/opencl_source.h
> index 4118138c30..be7e826c4c 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_gblur;
>  extern const char *ff_opencl_source_colorspace_common;
>  extern const char *ff_opencl_source_convolution;
>  extern const char *ff_opencl_source_neighbor;
> diff --git a/libavfilter/vf_gblur_opencl.c b/libavfilter/vf_gblur_opencl.c
> new file mode 100644
> index 0000000000..86dd10ee74
> --- /dev/null
> +++ b/libavfilter/vf_gblur_opencl.c
> @@ -0,0 +1,370 @@
> +/*
> + * 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/opt.h"
> +
> +#include "avfilter.h"
> +#include "internal.h"
> +#include "opencl.h"
> +#include "opencl_source.h"
> +#include "video.h"
> +
> +typedef struct GBlurOpenCLContext {
> +    OpenCLFilterContext ocf;
> +
> +    int              initialised;
> +    cl_kernel        kernel_horz;
> +    cl_kernel        kernel_vert;
> +
> +    cl_command_queue command_queue;
> +
> +    int planes;
> +    float sigma;
> +    float sigmaV;
> +
> +    cl_mem matrix_horiz;
> +    cl_mem matrix_vert;
> +    int kernel_dimension_horiz;
> +    int kernel_dimension_vert;
> +
> +} GBlurOpenCLContext;
> +
> +
> +static int gblur_opencl_init(AVFilterContext *avctx)
> +{
> +    GBlurOpenCLContext *ctx = avctx->priv;
> +    cl_int cle;
> +    int err;
> +    err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_gblur,
> 1);
> +    if (err < 0)
> +        goto fail;
> +
> +    ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
> +                                              ctx->ocf.hwctx->device_id,
> +                                              0, &cle);
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
> +                     "command queue %d.\n", cle);
> +
> +    ctx->kernel_horz = clCreateKernel(ctx->ocf.program,
> "gblur_conv_horz", &cle);
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
> +                     "kernel %d.\n", cle);
> +    ctx->kernel_vert = clCreateKernel(ctx->ocf.program,
> "gblur_conv_vert", &cle);
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
> +                     "kernel %d.\n", cle);
> +
> +    ctx->initialised = 1;
> +    return 0;
> +
> +fail:
> +    if (ctx->command_queue)
> +        clReleaseCommandQueue(ctx->command_queue);
> +    if (ctx->kernel_horz)
> +        clReleaseKernel(ctx->kernel_horz);
> +    if (ctx->kernel_vert)
> +        clReleaseKernel(ctx->kernel_vert);
> +    return err;
> +}
> +
> +
> +static int gblur_opencl_make_filter_params(AVFilterLink *inlink)
> +{
> +    AVFilterContext    *ctx = inlink->dst;
> +    GBlurOpenCLContext *s = ctx->priv;
> +
> +    float *matrix_horiz = NULL;
> +    float *matrix_vert = NULL;
> +    size_t matrix_bytes_horiz;
> +    size_t matrix_bytes_vert;
> +    cl_int cle;
> +
> +    if (s->sigmaV < 0) {
> +        s->sigmaV = s->sigma;
> +    }
> +
> +
> +    int kernel_dimension_horiz = (int)ceilf(6 * s->sigma);
> +    if (kernel_dimension_horiz % 2 == 0)
> +        kernel_dimension_horiz++;
> +    matrix_bytes_horiz = sizeof(float)*kernel_dimension_horiz;
> +
> +    matrix_horiz = av_malloc(matrix_bytes_horiz);
> +    if (!matrix_horiz) {
> +        av_freep(&matrix_horiz);
> +        return AVERROR(ENOMEM);
> +    }
> +
> +    double acc = 0;
> +
> +    for (int i = 0; i < kernel_dimension_horiz; i++) {
> +        int x = i - (kernel_dimension_horiz / 2);
> +
> +        matrix_horiz[i] = (1 / sqrt(2 * 3.14159*pow(s->sigma,
> 2)))*exp(-(pow(x, 2) / (2 * pow(s->sigma, 2))));
> +        acc += matrix_horiz[i];
> +    }
> +
> +    int kernel_dimension_vert = (int)ceilf(6 * s->sigmaV);
> +    if (kernel_dimension_vert % 2 == 0)
> +        kernel_dimension_vert++;
> +
> +    matrix_bytes_vert = sizeof(float)*kernel_dimension_vert;
> +    matrix_vert = av_malloc(matrix_bytes_vert);
> +    if (!matrix_vert) {
> +        av_freep(&matrix_vert);
> +        return AVERROR(ENOMEM);
> +    }
> +
> +    acc = 0;
> +    for (int i = 0; i < kernel_dimension_vert; i++) {
> +        int x = i - (kernel_dimension_vert / 2);
> +
> +        matrix_vert[i] = (1 / sqrt(2 * 3.14159*pow(s->sigmaV,
> 2)))*exp(-(pow(x, 2) / (2 * pow(s->sigmaV, 2))));
> +        acc += matrix_vert[i];
> +    }
> +
> +    s->kernel_dimension_horiz = kernel_dimension_horiz;
> +    s->matrix_horiz = clCreateBuffer(s->ocf.hwctx->context,
> +                               CL_MEM_READ_ONLY |
> +                               CL_MEM_COPY_HOST_PTR |
> +                               CL_MEM_HOST_NO_ACCESS,
> +                               matrix_bytes_horiz, matrix_horiz, &cle);
> +
> +    s->kernel_dimension_vert = kernel_dimension_vert;
> +    s->matrix_vert = clCreateBuffer(s->ocf.hwctx->context,
> +                               CL_MEM_READ_ONLY |
> +                               CL_MEM_COPY_HOST_PTR |
> +                               CL_MEM_HOST_NO_ACCESS,
> +                               matrix_bytes_vert, matrix_vert, &cle);
> +
> +    av_freep(&matrix_horiz);
> +    av_freep(&matrix_vert);
> +
> +    return 0;
> +}
> +
> +static int gblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
> +{
> +    AVFilterContext    *avctx = inlink->dst;
> +    AVFilterLink     *outlink = avctx->outputs[0];
> +    GBlurOpenCLContext *ctx = avctx->priv;
> +
> +    AVFrame *output = NULL;
> +    AVFrame *intermediate = NULL;
> +    cl_int cle;
> +    size_t global_work[2];
> +    cl_mem src, dst, inter;
> +    size_t origin[3] = {0, 0, 0};
> +    size_t region[3] = {0, 0, 1};
> +    int err, p;
> +
> +    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 = gblur_opencl_init(avctx);
> +        if (err < 0)
> +            goto fail;
> +
> +        err = gblur_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;
> +
> +        if (!(ctx->planes & (1 << p))) {
> +            err = ff_opencl_filter_work_size_from_image(avctx, region,
> output, p, 0);
> +            if (err < 0)
> +                goto fail;
> +
> +            cle = clEnqueueCopyImage(ctx->command_queue, src, dst,
> +                                        origin, origin, region, 0, NULL,
> NULL);
> +            CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to copy plane %d:
> %d.\n",
> +                                p, cle);
> +        }
> +        else {
> +            CL_SET_KERNEL_ARG(ctx->kernel_horz, 0, cl_mem,   &inter);
> +            CL_SET_KERNEL_ARG(ctx->kernel_horz, 1, cl_mem,   &src);
> +            CL_SET_KERNEL_ARG(ctx->kernel_horz, 2, cl_int,
>  &ctx->kernel_dimension_horiz);
> +            CL_SET_KERNEL_ARG(ctx->kernel_horz, 3, cl_mem,
>  &ctx->matrix_horiz);
> +
> +            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_horz, 2, NULL,
> +                                            global_work, NULL,
> +                                            0, NULL, NULL);
> +            CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue horizontal "
> +                                "kernel: %d.\n", cle);
> +
> +
> +            CL_SET_KERNEL_ARG(ctx->kernel_vert, 0, cl_mem,   &dst);
> +            CL_SET_KERNEL_ARG(ctx->kernel_vert, 1, cl_mem,   &inter);
> +            CL_SET_KERNEL_ARG(ctx->kernel_vert, 2, cl_int,
>  &ctx->kernel_dimension_vert);
> +            CL_SET_KERNEL_ARG(ctx->kernel_vert, 3, cl_mem,
>  &ctx->matrix_vert);
> +
> +            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);
> +            CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue vertical "
> +                                "kernel: %d.\n", cle);
> +        }
> +    }
> +
> +    cle = clFinish(ctx->command_queue);
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue:
> %d.\n", cle);
> +
> +    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 gblur_opencl_uninit(AVFilterContext *avctx)
> +{
> +    GBlurOpenCLContext *ctx = avctx->priv;
> +    cl_int cle;
> +
> +    if (ctx->kernel_horz) {
> +        cle = clReleaseKernel(ctx->kernel_horz);
> +        if (cle != CL_SUCCESS)
> +            av_log(avctx, AV_LOG_ERROR, "Failed to release "
> +                   "kernel: %d.\n", cle);
> +    }
> +
> +    if (ctx->kernel_vert) {
> +        cle = clReleaseKernel(ctx->kernel_vert);
> +        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);
> +}
> +
> +
> +static const AVFilterPad gblur_opencl_inputs[] = {
> +    {
> +        .name         = "default",
> +        .type         = AVMEDIA_TYPE_VIDEO,
> +        .filter_frame = &gblur_opencl_filter_frame,
> +        .config_props = &ff_opencl_filter_config_input,
> +    },
> +    { NULL }
> +};
> +
> +
> +static const AVFilterPad gblur_opencl_outputs[] = {
> +    {
> +        .name         = "default",
> +        .type         = AVMEDIA_TYPE_VIDEO,
> +        .config_props = &ff_opencl_filter_config_output,
> +    },
> +    { NULL }
> +};
> +
> +
> +#define OFFSET(x) offsetof(GBlurOpenCLContext, x)
> +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
> +
> +#if CONFIG_GBLUR_OPENCL_FILTER
> +
> +static const AVOption gblur_opencl_options[] = {
> +    { "sigma",  "set horizontal size",  OFFSET(sigma), AV_OPT_TYPE_FLOAT,
> {.dbl=0.5},   0.0, 1024, FLAGS },
> +    { "planes", "set planes to filter", OFFSET(planes),  AV_OPT_TYPE_INT,
> {.i64=0xF}, 0,  0xF, FLAGS },
> +    { "sigmaV", "set vertical sigma",   OFFSET(sigmaV),
> AV_OPT_TYPE_FLOAT, {.dbl=-1},   -1, 1024, FLAGS },
> +    { NULL }
> +};
> +
> +AVFILTER_DEFINE_CLASS(gblur_opencl);
> +
> +
> +AVFilter ff_vf_gblur_opencl = {
> +    .name           = "gblur_opencl",
> +    .description    = NULL_IF_CONFIG_SMALL("Apply gaussian blur filter"),
> +    .priv_size      = sizeof(GBlurOpenCLContext),
> +    .priv_class     = &gblur_opencl_class,
> +    .init           = &ff_opencl_filter_init,
> +    .uninit         = &gblur_opencl_uninit,
> +    .query_formats  = &ff_opencl_filter_query_formats,
> +    .inputs         = gblur_opencl_inputs,
> +    .outputs        = gblur_opencl_outputs,
> +    .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
> +};
> +
> +#endif /* CONFIG_GBLUR_OPENCL_FILTER */
> --
> 2.15.2 (Apple Git-101.1)


> Anyone have any feedback?

On 4/30/19, Dylan Fernando <dylanf123@gmail.com> wrote:
>
> Anyone have any feedback?

If I'm not mistaken there is already one available.

On Tue, Apr 30, 2019 at 15:45:23 +0200, Paul B Mahol wrote:
> On 4/30/19, Dylan Fernando <dylanf123@gmail.com> wrote:
> >
> > Anyone have any feedback?
>
> If I'm not mistaken there is already one available.

Dylan did post a v2 of this patch (with changes made to those things I
commented on), but unfortunately posted this question in reply to v1.

Moritz

On Tue, 30 Apr 2019 at 11:45 pm, Paul B Mahol <onemda@gmail.com> wrote:

> On 4/30/19, Dylan Fernando <dylanf123@gmail.com> wrote:
> >
> > Anyone have any feedback?
>
> If I'm not mistaken there is already one available.
> _______________________________________________
> ffmpeg-devel mailing list
> ffmpeg-devel@ffmpeg.org
> https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
>
> To unsubscribe, visit link above, or email
> ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".


How do I run the available filter?

>

On 5/2/19, Dylan Fernando <dylanf123@gmail.com> wrote:
> On Tue, 30 Apr 2019 at 11:45 pm, Paul B Mahol <onemda@gmail.com> wrote:
>
>> On 4/30/19, Dylan Fernando <dylanf123@gmail.com> wrote:
>> >
>> > Anyone have any feedback?
>>
>> If I'm not mistaken there is already one available.
>> _______________________________________________
>> ffmpeg-devel mailing list
>> ffmpeg-devel@ffmpeg.org
>> https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
>>
>> To unsubscribe, visit link above, or email
>> ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
>
>
> How do I run the available filter?

Do you ask how to use gblur video filter?

On Thu, May 2, 2019 at 11:27 PM Paul B Mahol <onemda@gmail.com> wrote:

> On 5/2/19, Dylan Fernando <dylanf123@gmail.com> wrote:
> > On Tue, 30 Apr 2019 at 11:45 pm, Paul B Mahol <onemda@gmail.com> wrote:
> >
> >> On 4/30/19, Dylan Fernando <dylanf123@gmail.com> wrote:
> >> >
> >> > Anyone have any feedback?
> >>
> >> If I'm not mistaken there is already one available.
> >> _______________________________________________
> >> ffmpeg-devel mailing list
> >> ffmpeg-devel@ffmpeg.org
> >> https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
> >>
> >> To unsubscribe, visit link above, or email
> >> ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
> >
> >
> > How do I run the available filter?
>
> Do you ask how to use gblur video filter?
> _______________________________________________
> ffmpeg-devel mailing list
> ffmpeg-devel@ffmpeg.org
> https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
>
> To unsubscribe, visit link above, or email
> ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".

 Sorry, I mean how do I run the opencl gblur filter, if there is one
available