[FFmpeg-devel] [PATCH] avfilter: add xfade opencl filter

[Mark Thompson]

On 26/01/2020 18:28, Paul B Mahol wrote:
> Signed-off-by: Paul B Mahol <onemda at gmail.com>
> ---
>  configure                     |   1 +
>  doc/filters.texi              |  97 ++++++++
>  libavfilter/Makefile          |   1 +
>  libavfilter/allfilters.c      |   1 +
>  libavfilter/opencl/xfade.cl   | 136 +++++++++++
>  libavfilter/opencl_source.h   |   1 +
>  libavfilter/vf_xfade_opencl.c | 420 ++++++++++++++++++++++++++++++++++
>  7 files changed, 657 insertions(+)
>  create mode 100644 libavfilter/opencl/xfade.cl
>  create mode 100644 libavfilter/vf_xfade_opencl.c
> 
> ...
> +
> +void slide(__write_only image2d_t dst,
> +           __read_only  image2d_t src1,
> +           __read_only  image2d_t src2,
> +           float progress,
> +           int2 direction)
> +{
> +    const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |  \
> +                               CLK_FILTER_NEAREST);

>From the Mali driver:

<source>:87:21: error: non-kernel function variable cannot be declared in constant address space
    const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |  \
                    ^

error: Compiler frontend failed (error code 60)

OpenCL 1.2 §6.9:

"The sampler type (sampler_t) can only be used as the type of a function argument or a
variable declared in the program scope or the outermost scope of a kernel function."

I think just make it global, since you're now using the same sampler in every kernel?

> +    int   w = get_image_dim(src1).x;
> +    int   h = get_image_dim(src1).y;
> +    int2 wh = (int2)(w, h);
> +    int2 uv = (int2)(get_global_id(0), get_global_id(1));
> +    int2 pi = (int2)(progress * w, progress * h);
> +    int2 p = uv + pi * direction;
> +    int2 f = p % wh;
> +
> +    f = f + (int2)(w, h) * (int2)(f.x < 0, f.y < 0);
> +    float4 val1 = read_imagef(src1, sampler, f);
> +    float4 val2 = read_imagef(src2, sampler, f);
> +    write_imagef(dst, uv, mix(val1, val2, (p.y >= 0) * (h > p.y) * (p.x >= 0) * (w > p.x)));
> +}
> +
> ...
> +
> +AVFilter ff_vf_xfade_opencl = {
> +    .name            = "xfade_opencl",
> +    .description     = NULL_IF_CONFIG_SMALL("Cross fade one video with another video."),
> +    .priv_size       = sizeof(XFadeOpenCLContext),
> +    .priv_class      = &xfade_opencl_class,
> +    .init            = &ff_opencl_filter_init,
> +    .uninit          = &xfade_opencl_uninit,
> +    .query_formats   = &ff_opencl_filter_query_formats,
> +    .activate        = &xfade_opencl_activate,
> +    .inputs          = xfade_opencl_inputs,
> +    .outputs         = xfade_opencl_outputs,
> +    .flags_internal  = FF_FILTER_FLAG_HWFRAME_AWARE,
> +};
> 

No other comments from me, so LGTM with that fixed.

Thanks,

- Mark