On 27 March 2018 at 16:09, Gabriel Machado <gabriel_mach...@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.
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
http://ffmpeg.org/mailman/listinfo/ffmpeg-devel
