[FFmpeg-devel] [PATCH v2 RESEND] avfilter: add OpenCL scale filter

[Rostislav Pehlivanov]

On 27 March 2018 at 16:09, Gabriel Machado <gabriel_machado at live.com> wrote:

> Sorry, my mail client corrupted the patch.
>
> ---
>  configure                     |   1 +
>  libavfilter/Makefile          |   1 +
>  libavfilter/allfilters.c      |   1 +
>  libavfilter/opencl/scale.cl   | 111 +++++++++++++++++
>  libavfilter/opencl_source.h   |   1 +
>  libavfilter/vf_scale_opencl.c | 284 ++++++++++++++++++++++++++++++
> ++++++++++++
>  6 files changed, 399 insertions(+)
>  create mode 100644 libavfilter/opencl/scale.cl
>  create mode 100644 libavfilter/vf_scale_opencl.c
>
> diff --git a/configure b/configure
> index 5ccf3ce..4007ee8 100755
> --- a/configure
> +++ b/configure
> @@ -2821,6 +2821,7 @@ v4l2_m2m_deps_any="linux_videodev2_h"
>
>  hwupload_cuda_filter_deps="ffnvcodec"
>  scale_npp_filter_deps="ffnvcodec libnpp"
> +scale_opencl_filter_deps="opencl"
>  scale_cuda_filter_deps="cuda_sdk"
>  thumbnail_cuda_filter_deps="cuda_sdk"
>
> diff --git a/libavfilter/Makefile b/libavfilter/Makefile
> index a90ca30..6303cbd 100644
> --- a/libavfilter/Makefile
> +++ b/libavfilter/Makefile
> @@ -302,6 +302,7 @@ OBJS-$(CONFIG_SAB_FILTER)                    +=
> vf_sab.o
>  OBJS-$(CONFIG_SCALE_FILTER)                  += vf_scale.o scale.o
>  OBJS-$(CONFIG_SCALE_CUDA_FILTER)             += vf_scale_cuda.o
> vf_scale_cuda.ptx.o
>  OBJS-$(CONFIG_SCALE_NPP_FILTER)              += vf_scale_npp.o scale.o
> +OBJS-$(CONFIG_SCALE_OPENCL_FILTER)           += vf_scale_opencl.o
> opencl.o opencl/scale.o
>  OBJS-$(CONFIG_SCALE_QSV_FILTER)              += vf_scale_qsv.o
>  OBJS-$(CONFIG_SCALE_VAAPI_FILTER)            += vf_scale_vaapi.o scale.o
> vaapi_vpp.o
>  OBJS-$(CONFIG_SCALE2REF_FILTER)              += vf_scale.o scale.o
> diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c
> index 1cf1340..3185b17 100644
> --- a/libavfilter/allfilters.c
> +++ b/libavfilter/allfilters.c
> @@ -309,6 +309,7 @@ static void register_all(void)
>      REGISTER_FILTER(SCALE,          scale,          vf);
>      REGISTER_FILTER(SCALE_CUDA,     scale_cuda,     vf);
>      REGISTER_FILTER(SCALE_NPP,      scale_npp,      vf);
> +    REGISTER_FILTER(SCALE_OPENCL,   scale_opencl,   vf);
>      REGISTER_FILTER(SCALE_QSV,      scale_qsv,      vf);
>      REGISTER_FILTER(SCALE_VAAPI,    scale_vaapi,    vf);
>      REGISTER_FILTER(SCALE2REF,      scale2ref,      vf);
> diff --git a/libavfilter/opencl/scale.cl b/libavfilter/opencl/scale.cl
> new file mode 100644
> index 0000000..3436694
> --- /dev/null
> +++ b/libavfilter/opencl/scale.cl
> @@ -0,0 +1,111 @@
> +/*
> + * Copyright (c) 2018 Gabriel Machado
> + *
> + * 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 neighbor(__write_only image2d_t dst,
> +                       __read_only  image2d_t src)
> +{
> +    const sampler_t sampler = (CLK_NORMALIZED_COORDS_TRUE |
> +                               CLK_ADDRESS_CLAMP_TO_EDGE |
> +                               CLK_FILTER_NEAREST);
> +
> +    int2 coord = {get_global_id(0), get_global_id(1)};
> +    int2 size = {get_global_size(0), get_global_size(1)};
> +
> +    float2 pos = (convert_float2(coord) + 0.5) / convert_float2(size);
> +
> +    float4 c = read_imagef(src, sampler, pos);
> +    write_imagef(dst, coord, c);
> +}
> +
> +__kernel void bilinear(__write_only image2d_t dst,
> +                       __read_only  image2d_t src)
> +{
> +    const sampler_t sampler = (CLK_NORMALIZED_COORDS_TRUE |
> +                               CLK_ADDRESS_CLAMP_TO_EDGE |
> +                               CLK_FILTER_LINEAR);
> +
> +    int2 coord = {get_global_id(0), get_global_id(1)};
> +    int2 size = {get_global_size(0), get_global_size(1)};
> +
> +    float2 pos = (convert_float2(coord) + 0.5) / convert_float2(size);
>

Doesn't opencl have an option to use a sampler with non-normalized
addressing mode?


> +
> +    float4 c = read_imagef(src, sampler, pos);
> +    write_imagef(dst, coord, c);
> +}
> +
> +// https://developer.nvidia.com/gpugems/GPUGems/gpugems_ch24.html
> +float MitchellNetravali(float x, float B, float C)
>

You completely ignored what I said. This function doesn't have a license,
you can't use it in a modified form. Either rewrite it or remove it.