avfilter: add xfade opencl filter

Changelog

@@ -33,6 +33,7 @@ version <next>:
 - Argonaut Games ADPCM decoder
 - Argonaut Games ASF demuxer
 - xfade video filter
+- xfade_opencl filter
 
 
 version 4.2:

configure

@@ -3596,6 +3596,7 @@ zscale_filter_deps="libzimg const_nan"
 scale_vaapi_filter_deps="vaapi"
 vpp_qsv_filter_deps="libmfx"
 vpp_qsv_filter_select="qsvvpp"
+xfade_opencl_filter_deps="opencl"
 yadif_cuda_filter_deps="ffnvcodec"
 yadif_cuda_filter_deps_any="cuda_nvcc cuda_llvm"
 

doc/filters.texi

@@ -21448,6 +21448,103 @@ Apply a strong blur of both luma and chroma parameters:
 @end example
 @end itemize
 
+@section xfade_opencl
+
+Cross fade two videos with custom transition effect by using OpenCL.
+
+It accepts the following options:
+
+@table @option
+@item transition
+Set one of possible transition effects.
+
+@table @option
+@item custom
+Select custom transition effect, the actual transition description
+will be picked from source and kernel options.
+
+@item fade
+@item wipeleft
+@item wiperight
+@item wipeup
+@item wipedown
+@item slideleft
+@item slideright
+@item slideup
+@item slidedown
+
+Default transition is fade.
+@end table
+
+@item source
+OpenCL program source file for custom transition.
+
+@item kernel
+Set name of kernel to use for custom transition from program source file.
+
+@item duration
+Set duration of video transition.
+
+@item offset
+Set time of start of transition relative to first video.
+@end table
+
+The program source file must contain a kernel function with the given name,
+which will be run once for each plane of the output.  Each run on a plane
+gets enqueued as a separate 2D global NDRange with one work-item for each
+pixel to be generated.  The global ID offset for each work-item is therefore
+the coordinates of a pixel in the destination image.
+
+The kernel function needs to take the following arguments:
+@itemize
+@item
+Destination image, @var{__write_only image2d_t}.
+
+This image will become the output; the kernel should write all of it.
+
+@item
+First Source image, @var{__read_only image2d_t}.
+Second Source image, @var{__read_only image2d_t}.
+
+These are the most recent images on each input.  The kernel may read from
+them to generate the output, but they can't be written to.
+
+@item
+Transition progress, @var{float}. This value is always between 0 and 1 inclusive.
+@end itemize
+
+Example programs:
+
+@itemize
+@item
+Apply dots curtain transition effect:
+@verbatim
+__kernel void blend_images(__write_only image2d_t dst,
+                           __read_only  image2d_t src1,
+                           __read_only  image2d_t src2,
+                           float progress)
+{
+    const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
+                               CLK_FILTER_LINEAR);
+    int2  p = (int2)(get_global_id(0), get_global_id(1));
+    float2 rp = (float2)(get_global_id(0), get_global_id(1));
+    float2 dim = (float2)(get_image_dim(src1).x, get_image_dim(src1).y);
+    rp = rp / dim;
+
+    float2 dots = (float2)(20.0, 20.0);
+    float2 center = (float2)(0,0);
+    float2 unused;
+
+    float4 val1 = read_imagef(src1, sampler, p);
+    float4 val2 = read_imagef(src2, sampler, p);
+    bool next = distance(fract(rp * dots, &unused), (float2)(0.5, 0.5)) < (progress / distance(rp, center));
+
+    write_imagef(dst, p, next ? val1 : val2);
+}
+@end verbatim
+
+@end itemize
+
 @c man end OPENCL VIDEO FILTERS
 
 @chapter VAAPI Video Filters

libavfilter/Makefile

@@ -442,6 +442,7 @@ OBJS-$(CONFIG_WAVEFORM_FILTER)               += vf_waveform.o
 OBJS-$(CONFIG_WEAVE_FILTER)                  += vf_weave.o
 OBJS-$(CONFIG_XBR_FILTER)                    += vf_xbr.o
 OBJS-$(CONFIG_XFADE_FILTER)                  += vf_xfade.o
+OBJS-$(CONFIG_XFADE_OPENCL_FILTER)           += vf_xfade_opencl.o opencl.o opencl/xfade.o
 OBJS-$(CONFIG_XMEDIAN_FILTER)                += vf_xmedian.o framesync.o
 OBJS-$(CONFIG_XSTACK_FILTER)                 += vf_stack.o framesync.o
 OBJS-$(CONFIG_YADIF_FILTER)                  += vf_yadif.o yadif_common.o

libavfilter/allfilters.c

@@ -421,6 +421,7 @@ extern AVFilter ff_vf_waveform;
 extern AVFilter ff_vf_weave;
 extern AVFilter ff_vf_xbr;
 extern AVFilter ff_vf_xfade;
+extern AVFilter ff_vf_xfade_opencl;
 extern AVFilter ff_vf_xmedian;
 extern AVFilter ff_vf_xstack;
 extern AVFilter ff_vf_yadif;

libavfilter/opencl/xfade.cl

+/*
+ * 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
+ */
+
+const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
+                           CLK_FILTER_NEAREST);
+
+__kernel void fade(__write_only image2d_t dst,
+                   __read_only  image2d_t src1,
+                   __read_only  image2d_t src2,
+                   float progress)
+{
+    int2  p = (int2)(get_global_id(0), get_global_id(1));
+
+    float4 val1 = read_imagef(src1, sampler, p);
+    float4 val2 = read_imagef(src2, sampler, p);
+
+    write_imagef(dst, p, mix(val2, val1, progress));
+}
+
+__kernel void wipeleft(__write_only image2d_t dst,
+                       __read_only  image2d_t src1,
+                       __read_only  image2d_t src2,
+                       float progress)
+{
+    int   s = (int)(get_image_dim(src1).x * progress);
+    int2  p = (int2)(get_global_id(0), get_global_id(1));
+
+    float4 val1 = read_imagef(src1, sampler, p);
+    float4 val2 = read_imagef(src2, sampler, p);
+
+    write_imagef(dst, p, p.x > s ? val2 : val1);
+}
+
+__kernel void wiperight(__write_only image2d_t dst,
+                        __read_only  image2d_t src1,
+                        __read_only  image2d_t src2,
+                        float progress)
+{
+    int   s = (int)(get_image_dim(src1).x * (1.f - progress));
+    int2  p = (int2)(get_global_id(0), get_global_id(1));
+
+    float4 val1 = read_imagef(src1, sampler, p);
+    float4 val2 = read_imagef(src2, sampler, p);
+
+    write_imagef(dst, p, p.x > s ? val1 : val2);
+}
+
+__kernel void wipeup(__write_only image2d_t dst,
+                     __read_only  image2d_t src1,
+                     __read_only  image2d_t src2,
+                     float progress)
+{
+    int   s = (int)(get_image_dim(src1).y * progress);
+    int2  p = (int2)(get_global_id(0), get_global_id(1));
+
+    float4 val1 = read_imagef(src1, sampler, p);
+    float4 val2 = read_imagef(src2, sampler, p);
+
+    write_imagef(dst, p, p.y > s ? val2 : val1);
+}
+
+__kernel void wipedown(__write_only image2d_t dst,
+                       __read_only  image2d_t src1,
+                       __read_only  image2d_t src2,
+                       float progress)
+{
+    int   s = (int)(get_image_dim(src1).y * (1.f - progress));
+    int2  p = (int2)(get_global_id(0), get_global_id(1));
+
+    float4 val1 = read_imagef(src1, sampler, p);
+    float4 val2 = read_imagef(src2, sampler, p);
+
+    write_imagef(dst, p, p.y > s ? val1 : val2);
+}
+
+void slide(__write_only image2d_t dst,
+           __read_only  image2d_t src1,
+           __read_only  image2d_t src2,
+           float progress,
+           int2 direction)
+{
+    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)));
+}
+
+__kernel void slidedown(__write_only image2d_t dst,
+                        __read_only  image2d_t src1,
+                        __read_only  image2d_t src2,
+                        float progress)
+{
+    int2 direction = (int2)(0, 1);
+    slide(dst, src1, src2, progress, direction);
+}
+
+__kernel void slideup(__write_only image2d_t dst,
+                      __read_only  image2d_t src1,
+                      __read_only  image2d_t src2,
+                      float progress)
+{
+    int2 direction = (int2)(0, -1);
+    slide(dst, src1, src2, progress, direction);
+}
+
+__kernel void slideleft(__write_only image2d_t dst,
+                        __read_only  image2d_t src1,
+                        __read_only  image2d_t src2,
+                        float progress)
+{
+    int2 direction = (int2)(-1, 0);
+    slide(dst, src1, src2, progress, direction);
+}
+
+__kernel void slideright(__write_only image2d_t dst,
+                         __read_only  image2d_t src1,
+                         __read_only  image2d_t src2,
+                         float progress)
+{
+    int2 direction = (int2)(1, 0);
+    slide(dst, src1, src2, progress, direction);
+}

libavfilter/opencl_source.h

@@ -30,5 +30,6 @@ extern const char *ff_opencl_source_overlay;
 extern const char *ff_opencl_source_tonemap;
 extern const char *ff_opencl_source_transpose;
 extern const char *ff_opencl_source_unsharp;
+extern const char *ff_opencl_source_xfade;
 
 #endif /* AVFILTER_OPENCL_SOURCE_H */

libavfilter/version.h

@@ -30,7 +30,7 @@
 #include "libavutil/version.h"
 
 #define LIBAVFILTER_VERSION_MAJOR   7
-#define LIBAVFILTER_VERSION_MINOR  72
+#define LIBAVFILTER_VERSION_MINOR  73
 #define LIBAVFILTER_VERSION_MICRO 100