[FFmpeg-devel,v3] avfilter: add vf_overlay_videotoolbox

Overlay filter for VideoToolbox hwframes. Unlike most hardware
overlay filters, this filter does not require the two inputs to
have the same pixel format; instead, it will perform format
conversion automatically with hardware accelerated methods.

Signed-off-by: Gnattu OC <gnattuoc@me.com>
---
 Changelog                                     |   1 +
 configure                                     |   1 +
 doc/filters.texi                              |  52 ++
 libavfilter/Makefile                          |   3 +
 libavfilter/allfilters.c                      |   1 +
 libavfilter/metal/utils.h                     |   1 -
 libavfilter/metal/utils.m                     |   7 +-
 .../metal/vf_overlay_videotoolbox.metal       |  58 ++
 libavfilter/vf_overlay_videotoolbox.m         | 551 ++++++++++++++++++
 9 files changed, 672 insertions(+), 3 deletions(-)
 create mode 100644 libavfilter/metal/vf_overlay_videotoolbox.metal
 create mode 100644 libavfilter/vf_overlay_videotoolbox.m

A ping for this.

> On Feb 21, 2024, at 09:53, gnattu <gnattuoc@me.com> wrote:
> 
> Overlay filter for VideoToolbox hwframes. Unlike most hardware
> overlay filters, this filter does not require the two inputs to
> have the same pixel format; instead, it will perform format
> conversion automatically with hardware accelerated methods.
> 
> Signed-off-by: Gnattu OC <gnattuoc@me.com>
> ---
> Changelog                                     |   1 +
> configure                                     |   1 +
> doc/filters.texi                              |  52 ++
> libavfilter/Makefile                          |   3 +
> libavfilter/allfilters.c                      |   1 +
> libavfilter/metal/utils.h                     |   1 -
> libavfilter/metal/utils.m                     |   7 +-
> .../metal/vf_overlay_videotoolbox.metal       |  58 ++
> libavfilter/vf_overlay_videotoolbox.m         | 551 ++++++++++++++++++
> 9 files changed, 672 insertions(+), 3 deletions(-)
> create mode 100644 libavfilter/metal/vf_overlay_videotoolbox.metal
> create mode 100644 libavfilter/vf_overlay_videotoolbox.m
> 
> diff --git a/Changelog b/Changelog
> index 610ee61dd6..3ecfdab81b 100644
> --- a/Changelog
> +++ b/Changelog
> @@ -27,6 +27,7 @@ version <next>:
> - a C11-compliant compiler is now required; note that this requirement
>   will be bumped to C17 in the near future, so consider updating your
>   build environment if it lacks C17 support
> +- VideoToolbox overlay filter
> 
> version 6.1:
> - libaribcaption decoder
> diff --git a/configure b/configure
> index 23066efa32..a7c349d126 100755
> --- a/configure
> +++ b/configure
> @@ -3807,6 +3807,7 @@ overlay_qsv_filter_deps="libmfx"
> overlay_qsv_filter_select="qsvvpp"
> overlay_vaapi_filter_deps="vaapi VAProcPipelineCaps_blend_flags"
> overlay_vulkan_filter_deps="vulkan spirv_compiler"
> +overlay_videotoolbox_filter_deps="metal corevideo coreimage videotoolbox"
> owdenoise_filter_deps="gpl"
> pad_opencl_filter_deps="opencl"
> pan_filter_deps="swresample"
> diff --git a/doc/filters.texi b/doc/filters.texi
> index e0436a5755..bfb77562cb 100644
> --- a/doc/filters.texi
> +++ b/doc/filters.texi
> @@ -19033,6 +19033,58 @@ See @ref{framesync}.
> 
> This filter also supports the @ref{framesync} options.
> 
> +@section overlay_videotoolbox
> +
> +Overlay one video on top of another.
> +
> +This is the VideoToolbox variant of the @ref{overlay} filter.
> +It takes two inputs and has one output. The first input is the "main" video on which the second input is overlaid.
> +It only accepts VideoToolbox frames. The underlying input pixel formats do not have to match.
> +Different input pixel formats and color spaces will be automatically converted using hardware accelerated methods.
> +The final output will have the same pixel format and color space as the "main" input.
> +
> +The filter accepts the following options:
> +
> +@table @option
> +
> +@item x
> +Set the x coordinate of the overlaid video on the main video.
> +Default value is @code{0}.
> +
> +@item y
> +Set the y coordinate of the overlaid video on the main video.
> +Default value is @code{0}.
> +
> +@item eof_action
> +See @ref{framesync}.
> +
> +@item shortest
> +See @ref{framesync}.
> +
> +@item repeatlast
> +See @ref{framesync}.
> +
> +@end table
> +
> +@subsection Examples
> +
> +@itemize
> +@item
> +Overlay an image LOGO at the top-left corner of the INPUT video.
> +The INPUT video is in nv12 format and the LOGO image is in rgba format.
> +@example
> +-hwaccel videotoolbox -i INPUT -i LOGO -codec:v:0 h264_videotoolbox -filter_complex "[0:v]format=nv12,hwupload[a], [1:v]format=rgba,hwupload[b], [a][b]overlay_videotoolbox" OUTPUT
> +@end example
> +@item
> +Overlay an SDR video OVERLAY at the top-left corner of the HDR video MAIN.
> +The INPUT video is in p010 format and the LOGO image is in nv12 format.
> +The OUTPUT video will also be an HDR video with OVERLAY mapped to HDR.
> +@example
> +-hwaccel videotoolbox -i MAIN -i OVERLAY -codec:v:0 hevc_videotoolbox -tag:v hvc1 -filter_complex "[0:v]format=p010,hwupload[a], [1:v]format=nv12,hwupload[b], [a][b]overlay_videotoolbox" OUTPUT
> +@end example
> +
> +@end itemize
> +
> @section owdenoise
> 
> Apply Overcomplete Wavelet denoiser.
> diff --git a/libavfilter/Makefile b/libavfilter/Makefile
> index f6c1d641d6..ea1389ab57 100644
> --- a/libavfilter/Makefile
> +++ b/libavfilter/Makefile
> @@ -413,6 +413,9 @@ OBJS-$(CONFIG_OVERLAY_OPENCL_FILTER)         += vf_overlay_opencl.o opencl.o \
> OBJS-$(CONFIG_OVERLAY_QSV_FILTER)            += vf_overlay_qsv.o framesync.o
> OBJS-$(CONFIG_OVERLAY_VAAPI_FILTER)          += vf_overlay_vaapi.o framesync.o vaapi_vpp.o
> OBJS-$(CONFIG_OVERLAY_VULKAN_FILTER)         += vf_overlay_vulkan.o vulkan.o vulkan_filter.o
> +OBJS-$(CONFIG_OVERLAY_VIDEOTOOLBOX_FILTER)     += vf_overlay_videotoolbox.o framesync.o \
> +                                                metal/vf_overlay_videotoolbox.metallib.o \
> +                                                metal/utils.o
> OBJS-$(CONFIG_OWDENOISE_FILTER)              += vf_owdenoise.o
> OBJS-$(CONFIG_PAD_FILTER)                    += vf_pad.o
> OBJS-$(CONFIG_PAD_OPENCL_FILTER)             += vf_pad_opencl.o opencl.o opencl/pad.o
> diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c
> index 149bf50997..ec9d975ecb 100644
> --- a/libavfilter/allfilters.c
> +++ b/libavfilter/allfilters.c
> @@ -389,6 +389,7 @@ extern const AVFilter ff_vf_overlay_qsv;
> extern const AVFilter ff_vf_overlay_vaapi;
> extern const AVFilter ff_vf_overlay_vulkan;
> extern const AVFilter ff_vf_overlay_cuda;
> +extern const AVFilter ff_vf_overlay_videotoolbox;
> extern const AVFilter ff_vf_owdenoise;
> extern const AVFilter ff_vf_pad;
> extern const AVFilter ff_vf_pad_opencl;
> diff --git a/libavfilter/metal/utils.h b/libavfilter/metal/utils.h
> index 7350d42a35..d79c79751c 100644
> --- a/libavfilter/metal/utils.h
> +++ b/libavfilter/metal/utils.h
> @@ -55,5 +55,4 @@ CVMetalTextureRef ff_metal_texture_from_pixbuf(void *avclass,
>                                                int plane,
>                                                MTLPixelFormat format)
>                                                API_AVAILABLE(macos(10.11), ios(8.0));
> -
> #endif /* AVFILTER_METAL_UTILS_H */
> diff --git a/libavfilter/metal/utils.m b/libavfilter/metal/utils.m
> index f365d3ceea..db5c5f6f10 100644
> --- a/libavfilter/metal/utils.m
> +++ b/libavfilter/metal/utils.m
> @@ -55,6 +55,9 @@ CVMetalTextureRef ff_metal_texture_from_pixbuf(void *ctx,
> {
>     CVMetalTextureRef tex = NULL;
>     CVReturn ret;
> +    bool is_planer = CVPixelBufferIsPlanar(pixbuf);
> +    size_t width = is_planer ? CVPixelBufferGetWidthOfPlane(pixbuf, plane) : CVPixelBufferGetWidth(pixbuf);
> +    size_t height = is_planer ? CVPixelBufferGetHeightOfPlane(pixbuf, plane) : CVPixelBufferGetHeight(pixbuf);
> 
>     ret = CVMetalTextureCacheCreateTextureFromImage(
>         NULL,
> @@ -62,8 +65,8 @@ CVMetalTextureRef ff_metal_texture_from_pixbuf(void *ctx,
>         pixbuf,
>         NULL,
>         format,
> -        CVPixelBufferGetWidthOfPlane(pixbuf, plane),
> -        CVPixelBufferGetHeightOfPlane(pixbuf, plane),
> +        width,
> +        height,
>         plane,
>         &tex
>     );
> diff --git a/libavfilter/metal/vf_overlay_videotoolbox.metal b/libavfilter/metal/vf_overlay_videotoolbox.metal
> new file mode 100644
> index 0000000000..936e57e03e
> --- /dev/null
> +++ b/libavfilter/metal/vf_overlay_videotoolbox.metal
> @@ -0,0 +1,58 @@
> +/*
> + * Copyright (C) 2024 Gnattu OC <gnattuoc@me.com>
> + *
> + * 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 <metal_stdlib>
> +#include <metal_integer>
> +#include <metal_texture>
> +
> +using namespace metal;
> +
> +struct mtlBlendParams {
> +    uint x_position;
> +    uint y_position;
> +};
> +
> +/*
> + * Blend shader for premultiplied alpha textures
> + */
> +kernel void blend_shader(
> +                         texture2d<float, access::read> source [[ texture(0) ]],
> +                         texture2d<float, access::read> mask [[ texture(1) ]],
> +                         texture2d<float, access::write> dest [[ texture(2) ]],
> +                         constant mtlBlendParams& params [[ buffer(3) ]],
> +                         uint2 gid [[ thread_position_in_grid ]])
> +{
> +    const auto mask_size = uint2(mask.get_width(),
> +                                 mask.get_height());
> +    const auto loc_overlay = uint2(params.x_position, params.y_position);
> +    if (gid.x <  loc_overlay.x ||
> +        gid.y <  loc_overlay.y ||
> +        gid.x >= mask_size.x + loc_overlay.x ||
> +        gid.y >= mask_size.y + loc_overlay.y)
> +    {
> +        float4 source_color = source.read(gid);
> +        dest.write(source_color, gid);
> +    } else {
> +        float4 source_color = source.read(gid);
> +        float4 mask_color = mask.read((gid - loc_overlay));
> +        float4 result_color = source_color * (1.0f - mask_color.w) + (mask_color * mask_color.w);
> +        dest.write(result_color, gid);
> +    }
> +}
> diff --git a/libavfilter/vf_overlay_videotoolbox.m b/libavfilter/vf_overlay_videotoolbox.m
> new file mode 100644
> index 0000000000..e2df688729
> --- /dev/null
> +++ b/libavfilter/vf_overlay_videotoolbox.m
> @@ -0,0 +1,551 @@
> +/*
> + * Copyright (C) 2024 Gnattu OC <gnattuoc@me.com>
> + *
> + * 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 <CoreImage/CoreImage.h>
> +#include <VideoToolbox/VideoToolbox.h>
> +#include "internal.h"
> +#include "metal/utils.h"
> +#include "framesync.h"
> +#include "libavutil/hwcontext.h"
> +#include "libavutil/pixdesc.h"
> +#include "libavutil/opt.h"
> +#include "libavutil/objc.h"
> +#include "video.h"
> +
> +#include <assert.h>
> +
> +extern char ff_vf_overlay_videotoolbox_metallib_data[];
> +extern unsigned int ff_vf_overlay_videotoolbox_metallib_len;
> +
> +// Although iOS 8.0 introduced basic Metal support, its feature set is not complete and does not have CoreImage compatability.
> +// We have to set the minimum iOS version to 9.0.
> +typedef struct API_AVAILABLE(macos(10.11), ios(9.0)) OverlayVideoToolboxContext {
> +    AVBufferRef *device_ref;
> +    FFFrameSync fs;
> +
> +    CVMetalTextureCacheRef texture_cache;
> +    CVPixelBufferRef input_main_pixel_buffer_cache;
> +    CVPixelBufferRef input_overlay_pixel_buffer_cache;
> +    CVPixelBufferRef output_pixel_buffer_cache;
> +    CIContext *ci_ctx;
> +    VTPixelTransferSessionRef vt_session;
> +
> +    id<MTLDevice> mtl_device;
> +    id<MTLLibrary> mtl_library;
> +    id<MTLCommandQueue> mtl_queue;
> +    id<MTLComputePipelineState> mtl_pipeline;
> +    id<MTLFunction> mtl_function;
> +    id<MTLBuffer> mtl_params_buffer;
> +
> +    uint x_position;
> +    uint y_position;
> +    uint hwframe_ctx_allocated;
> +} OverlayVideoToolboxContext API_AVAILABLE(macos(10.11), ios(9.0));
> +
> +typedef struct MtlBlendParams {
> +    uint x_position;
> +    uint y_position;
> +} MtlBlendParams;
> +
> +// Using sizeof(OverlayVideoToolboxContext) without an availability check will error
> +// if we're targeting an older OS version, so we need to calculate the size ourselves
> +// (we'll statically verify it's correct in overlay_videotoolbox_init behind a check)
> +#define OVERLAY_VT_CTX_SIZE (sizeof(FFFrameSync) + sizeof(uint) * 3 + sizeof(void*) * 13 + 4)
> +
> +static void call_kernel(AVFilterContext *avctx,
> +                        id<MTLTexture> dst,
> +                        id<MTLTexture> main,
> +                        id<MTLTexture> overlay,
> +                        uint x_position,
> +                        uint y_position) API_AVAILABLE(macos(10.11), ios(9.0))
> +{
> +    OverlayVideoToolboxContext *ctx = avctx->priv;
> +    // Both the command buffer and encoder are auto-released by objc on default.
> +    // Use CFBridgingRetain to get a more C-like behavior.
> +    id<MTLCommandBuffer> buffer = CFBridgingRetain(ctx->mtl_queue.commandBuffer);
> +    id<MTLComputeCommandEncoder> encoder = CFBridgingRetain((__bridge id<MTLCommandBuffer>)buffer.computeCommandEncoder);
> +
> +    MtlBlendParams *params = (MtlBlendParams *)ctx->mtl_params_buffer.contents;
> +    *params = (MtlBlendParams){
> +        .x_position = x_position,
> +        .y_position = y_position,
> +    };
> +
> +    [(__bridge id<MTLComputeCommandEncoder>)encoder setTexture: main atIndex: 0];
> +    [(__bridge id<MTLComputeCommandEncoder>)encoder setTexture: overlay atIndex: 1];
> +    [(__bridge id<MTLComputeCommandEncoder>)encoder setTexture: dst atIndex: 2];
> +    [(__bridge id<MTLComputeCommandEncoder>)encoder setBuffer: ctx->mtl_params_buffer offset: 0 atIndex: 3];
> +    ff_metal_compute_encoder_dispatch(ctx->mtl_device, ctx->mtl_pipeline, (__bridge id<MTLComputeCommandEncoder>)encoder, dst.width, dst.height);
> +    [(__bridge id<MTLComputeCommandEncoder>)encoder endEncoding];
> +
> +    [(__bridge id<MTLCommandBuffer>)buffer commit];
> +    [(__bridge id<MTLCommandBuffer>)buffer waitUntilCompleted];
> +
> +    ff_objc_release(&encoder);
> +    ff_objc_release(&buffer);
> +}
> +
> +// Copies and/or converts one pixel buffer to another.
> +// This transparently handles pixel format and color spaces, and will do a conversion if needed.
> +static int transfer_pixel_buffer(OverlayVideoToolboxContext *ctx, CVPixelBufferRef source, CVPixelBufferRef destination)
> +{
> +    if (@available(macOS 10.8, iOS 16.0, *)) {
> +        int ret = 0;
> +        ret = VTPixelTransferSessionTransferImage(ctx->vt_session, source, destination);
> +        if (ret < 0)
> +            return ret;
> +    } else {
> +        CIImage *temp_image = NULL;
> +        temp_image = CFBridgingRetain([CIImage imageWithCVPixelBuffer: source]);
> +        [(__bridge CIContext*)ctx->ci_ctx render: (__bridge CIImage*)temp_image toCVPixelBuffer: destination];
> +        CFRelease(temp_image);
> +        CVBufferPropagateAttachments(source, destination);
> +    }
> +    return 0;
> +}
> +
> +static int overlay_vt_blend(FFFrameSync *fs) API_AVAILABLE(macos(10.11), ios(9.0))
> +{
> +    AVFilterContext *avctx = fs->parent;
> +    OverlayVideoToolboxContext *ctx = avctx->priv;
> +    AVFilterLink *outlink = avctx->outputs[0];
> +    AVFilterLink *inlink_main = avctx->inputs[0];
> +    AVFilterLink *inlink_overlay = avctx->inputs[1];
> +    AVFrame *input_main, *input_overlay;
> +    AVFrame *output;
> +    AVHWFramesContext *frames_ctx = (AVHWFramesContext*)inlink_main->hw_frames_ctx->data;
> +    AVHWFramesContext *frames_ctx_overlay = (AVHWFramesContext*)inlink_overlay->hw_frames_ctx->data;
> +    const AVPixFmtDescriptor *in_main_desc, *in_overlay_desc;
> +
> +    CVMetalTextureRef main, dst, overlay;
> +    id<MTLTexture> tex_main, tex_overlay, tex_dst;
> +
> +    MTLPixelFormat mtl_format;
> +    OSType cv_format;
> +    int ret;
> +    int i, overlay_planes = 0;
> +
> +    in_main_desc = av_pix_fmt_desc_get(frames_ctx->sw_format);
> +    in_overlay_desc = av_pix_fmt_desc_get(frames_ctx_overlay->sw_format);
> +    if (@available(macOS 11.3, iOS 14.2, *)) {
> +        mtl_format = MTLPixelFormatRGBA16Unorm;
> +        cv_format = kCVPixelFormatType_64RGBALE;
> +    } else {
> +        // On older OS versions, 64-bit RGBA with 16-bit little-endian full-range samples is not supported.
> +        // To handle inputs with color depth greater than 8, convert colors to float type during filtering on these versions.
> +        mtl_format = MTLPixelFormatRGBA16Float;
> +        cv_format = kCVPixelFormatType_64RGBAHalf;
> +    }
> +
> +    // read main and overlay frames from inputs
> +    ret = ff_framesync_get_frame(fs, 0, &input_main, 0);
> +    if (ret < 0)
> +        return ret;
> +    ret = ff_framesync_get_frame(fs, 1, &input_overlay, 0);
> +    if (ret < 0)
> +        return ret;
> +
> +    if (!input_main)
> +        return AVERROR_BUG;
> +
> +    output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
> +    if (!output)
> +        return AVERROR(ENOMEM);
> +
> +    ret = av_frame_copy_props(output, input_main);
> +    if (ret < 0)
> +        return ret;
> +
> +    if (!input_overlay) {
> +        ret = transfer_pixel_buffer(ctx, (CVPixelBufferRef)input_main->data[3], (CVPixelBufferRef)output->data[3]);
> +        if (ret < 0)
> +            return ret;
> +        return ff_filter_frame(outlink, output);
> +    }
> +
> +    for (i = 0; i < in_overlay_desc->nb_components; i++)
> +        overlay_planes = FFMAX(overlay_planes,
> +                               in_overlay_desc->comp[i].plane + 1);
> +
> +    // We need to convert input overlay when it is planar or the color depth does not match
> +    if (overlay_planes > 1 || in_main_desc->comp[0].depth != in_overlay_desc->comp[0].depth) {
> +        if (!ctx->input_overlay_pixel_buffer_cache) {
> +            ret = CVPixelBufferCreate(kCFAllocatorDefault,
> +                                      CVPixelBufferGetWidthOfPlane((CVPixelBufferRef)input_overlay->data[3], 0),
> +                                      CVPixelBufferGetHeightOfPlane((CVPixelBufferRef)input_overlay->data[3], 0),
> +                                      cv_format,
> +                                      (__bridge CFDictionaryRef)@{
> +                                          (NSString *)kCVPixelBufferCGImageCompatibilityKey: @(YES),
> +                                          (NSString *)kCVPixelBufferMetalCompatibilityKey: @(YES)
> +                                      },
> +                                      &ctx->input_overlay_pixel_buffer_cache);
> +            if (ret < 0)
> +                return ret;
> +        }
> +        ret = transfer_pixel_buffer(ctx, (CVPixelBufferRef)input_overlay->data[3], ctx->input_overlay_pixel_buffer_cache);
> +        if (ret < 0)
> +            return ret;
> +        overlay = ff_metal_texture_from_pixbuf(avctx, ctx->texture_cache, ctx->input_overlay_pixel_buffer_cache, 0, mtl_format);
> +    } else {
> +        overlay = ff_metal_texture_from_pixbuf(avctx, ctx->texture_cache, (CVPixelBufferRef)input_overlay->data[3], 0, mtl_format);
> +    }
> +
> +    if (!ctx->input_main_pixel_buffer_cache) {
> +        ret = CVPixelBufferCreate(kCFAllocatorDefault,
> +                                  CVPixelBufferGetWidthOfPlane((CVPixelBufferRef)input_main->data[3], 0),
> +                                  CVPixelBufferGetHeightOfPlane((CVPixelBufferRef)input_main->data[3], 0),
> +                                  cv_format,
> +                                  (__bridge CFDictionaryRef)@{
> +                                      (NSString *)kCVPixelBufferCGImageCompatibilityKey: @(YES),
> +                                      (NSString *)kCVPixelBufferMetalCompatibilityKey: @(YES)
> +                                  },
> +                                  &ctx->input_main_pixel_buffer_cache);
> +        if (ret < 0)
> +            return ret;
> +    }
> +    if (!ctx->output_pixel_buffer_cache) {
> +        ret = CVPixelBufferCreate(kCFAllocatorDefault,
> +                                  CVPixelBufferGetWidthOfPlane((CVPixelBufferRef)input_main->data[3], 0),
> +                                  CVPixelBufferGetHeightOfPlane((CVPixelBufferRef)input_main->data[3], 0),
> +                                  cv_format,
> +                                  (__bridge CFDictionaryRef)@{
> +                                      (NSString *)kCVPixelBufferCGImageCompatibilityKey: @(YES),
> +                                      (NSString *)kCVPixelBufferMetalCompatibilityKey: @(YES)
> +                                  },
> +                                  &ctx->output_pixel_buffer_cache);
> +        if (ret < 0)
> +            return ret;
> +    }
> +
> +    ret = transfer_pixel_buffer(ctx, (CVPixelBufferRef)input_main->data[3], ctx->input_main_pixel_buffer_cache);
> +    if (ret < 0)
> +        return ret;
> +  
> +    main = ff_metal_texture_from_pixbuf(avctx, ctx->texture_cache, ctx->input_main_pixel_buffer_cache, 0, mtl_format);
> +    dst = ff_metal_texture_from_pixbuf(avctx, ctx->texture_cache, ctx->output_pixel_buffer_cache, 0, mtl_format);
> +
> +    tex_main = CVMetalTextureGetTexture(main);
> +    tex_overlay  = CVMetalTextureGetTexture(overlay);
> +    tex_dst = CVMetalTextureGetTexture(dst);
> +
> +    call_kernel(avctx, tex_dst, tex_main, tex_overlay, ctx->x_position, ctx->y_position);
> +
> +    ret = transfer_pixel_buffer(ctx, ctx->output_pixel_buffer_cache, (CVPixelBufferRef)output->data[3]);
> +    if (ret < 0)
> +        return ret;
> +
> +    CFRelease(main);
> +    CFRelease(overlay);
> +    CFRelease(dst);
> +
> +    return ff_filter_frame(outlink, output);
> +}
> +
> +static av_cold void do_uninit(AVFilterContext *avctx) API_AVAILABLE(macos(10.11), ios(9.0))
> +{
> +    OverlayVideoToolboxContext *ctx = avctx->priv;
> +
> +    if (ctx->hwframe_ctx_allocated) {
> +        av_buffer_unref(&ctx->device_ref);
> +        ctx->hwframe_ctx_allocated = 0;
> +    }
> +    ff_framesync_uninit(&ctx->fs);
> +
> +    if (ctx->ci_ctx) {
> +        CFRelease(ctx->ci_ctx);
> +        ctx->ci_ctx = NULL;
> +    }
> +
> +    ff_objc_release(&ctx->mtl_params_buffer);
> +    ff_objc_release(&ctx->mtl_function);
> +    ff_objc_release(&ctx->mtl_pipeline);
> +    ff_objc_release(&ctx->mtl_queue);
> +    ff_objc_release(&ctx->mtl_library);
> +    ff_objc_release(&ctx->mtl_device);
> +
> +    if (ctx->texture_cache) {
> +        CFRelease(ctx->texture_cache);
> +        ctx->texture_cache = NULL;
> +    }
> +    if (ctx->input_main_pixel_buffer_cache) {
> +        CFRelease(ctx->input_main_pixel_buffer_cache);
> +        ctx->input_main_pixel_buffer_cache = NULL;
> +    }
> +    if (ctx->input_overlay_pixel_buffer_cache) {
> +        CFRelease(ctx->input_overlay_pixel_buffer_cache);
> +        ctx->input_overlay_pixel_buffer_cache = NULL;
> +    }
> +    if (ctx->output_pixel_buffer_cache) {
> +        CFRelease(ctx->output_pixel_buffer_cache);
> +        ctx->output_pixel_buffer_cache = NULL;
> +    }
> +    if (ctx->vt_session) {
> +        VTPixelTransferSessionInvalidate(ctx->vt_session);
> +        CFRelease(ctx->vt_session);
> +        ctx->vt_session = NULL;
> +    }
> +}
> +
> +static av_cold void overlay_videotoolbox_uninit(AVFilterContext *ctx)
> +{
> +    if (@available(macOS 10.11, iOS 9.0, *)) {
> +        do_uninit(ctx);
> +    }
> +}
> +
> +static av_cold int do_init(AVFilterContext *avctx) API_AVAILABLE(macos(10.11), ios(9.0))
> +{
> +    OverlayVideoToolboxContext *ctx = avctx->priv;
> +    NSError *err = nil;
> +    CVReturn ret;
> +    dispatch_data_t libData;
> +
> +    ctx->mtl_device = MTLCreateSystemDefaultDevice();
> +    if (!ctx->mtl_device) {
> +        av_log(avctx, AV_LOG_ERROR, "Unable to find Metal device\n");
> +        goto fail;
> +    }
> +
> +    av_log(ctx, AV_LOG_INFO, "Using Metal device: %s\n", ctx->mtl_device.name.UTF8String);
> +
> +    libData = dispatch_data_create(
> +        ff_vf_overlay_videotoolbox_metallib_data,
> +        ff_vf_overlay_videotoolbox_metallib_len,
> +        nil,
> +        nil);
> +
> +    ctx->mtl_library = [ctx->mtl_device newLibraryWithData: libData error: &err];
> +    dispatch_release(libData);
> +    libData = nil;
> +    ctx->mtl_function = [ctx->mtl_library newFunctionWithName: @"blend_shader"];
> +    if (!ctx->mtl_function) {
> +        av_log(avctx, AV_LOG_ERROR, "Failed to create Metal function!\n");
> +        goto fail;
> +    }
> +
> +    ctx->mtl_queue = ctx->mtl_device.newCommandQueue;
> +    if (!ctx->mtl_queue) {
> +        av_log(avctx, AV_LOG_ERROR, "Failed to create Metal command queue!\n");
> +        goto fail;
> +    }
> +
> +    ctx->mtl_pipeline = [ctx->mtl_device
> +        newComputePipelineStateWithFunction: ctx->mtl_function
> +        error: &err];
> +    if (err) {
> +        av_log(ctx, AV_LOG_ERROR, "Failed to create Metal compute pipeline: %s\n", err.description.UTF8String);
> +        goto fail;
> +    }
> +
> +    ctx->mtl_params_buffer = [ctx->mtl_device
> +        newBufferWithLength: sizeof(MtlBlendParams)
> +        options: MTLResourceStorageModeShared];
> +    if (!ctx->mtl_params_buffer) {
> +        av_log(avctx, AV_LOG_ERROR, "Failed to create Metal buffer for parameters\n");
> +        goto fail;
> +    }
> +
> +    ret = CVMetalTextureCacheCreate(
> +        NULL,
> +        NULL,
> +        ctx->mtl_device,
> +        NULL,
> +        &ctx->texture_cache
> +    );
> +    if (ret != kCVReturnSuccess) {
> +        av_log(avctx, AV_LOG_ERROR, "Failed to create CVMetalTextureCache: %d\n", ret);
> +        goto fail;
> +    }
> +
> +    if (@available(macOS 10.8, iOS 16.0, *)) {
> +        ret = VTPixelTransferSessionCreate(NULL, &ctx->vt_session);
> +        if (ret != kCVReturnSuccess) {
> +            av_log(avctx, AV_LOG_ERROR, "Failed to create VTPixelTransferSession: %d\n", ret);
> +            goto fail;
> +        }
> +    } else {
> +        // Use CoreImage as fallback for old OS.
> +        // CoreImage has comparable performance to VTPixelTransferSession, but it supports less pixel formats than VTPixelTransferSession.
> +        // Warn user about possible incorrect results.
> +        av_log(avctx, AV_LOG_WARNING, "VTPixelTransferSessionTransferImage is not available on this OS version, fallback using CoreImage\n");
> +        av_log(avctx, AV_LOG_WARNING, "Try an overlay with BGRA format if you see no overlay\n");
> +        if (@available(macOS 10.15, iOS 13.0, *)) {
> +            ctx->ci_ctx = CFBridgingRetain([CIContext contextWithMTLCommandQueue: ctx->mtl_queue]);
> +        } else {
> +            ctx->ci_ctx = CFBridgingRetain([CIContext contextWithMTLDevice: ctx->mtl_device]);
> +        }
> +    }
> +
> +    ctx->fs.on_event = &overlay_vt_blend;
> +
> +    return 0;
> +fail:
> +    overlay_videotoolbox_uninit(avctx);
> +    return AVERROR_EXTERNAL;
> +}
> +
> +static av_cold int overlay_videotoolbox_init(AVFilterContext *ctx)
> +{
> +    if (@available(macOS 10.11, iOS 9.0, *)) {
> +        // Ensure we calculated OVERLAY_VT_CTX_SIZE correctly
> +        static_assert(OVERLAY_VT_CTX_SIZE == sizeof(OverlayVideoToolboxContext), "Incorrect OVERLAY_VT_CTX_SIZE value!");
> +        return do_init(ctx);
> +    } else {
> +        av_log(ctx, AV_LOG_ERROR, "Metal is not available on this OS version\n");
> +        return AVERROR(ENOSYS);
> +    }
> +}
> +
> +static int do_config_output(AVFilterLink *link) API_AVAILABLE(macos(10.11), ios(9.0))
> +{
> +    AVFilterContext *avctx = link->src;
> +    AVFilterLink *inlink_main = avctx->inputs[0];
> +    AVFilterLink *inlink_overlay = avctx->inputs[1];
> +    OverlayVideoToolboxContext *ctx = avctx->priv;
> +    AVHWFramesContext *main_frames, *output_frames;
> +    AVBufferRef *input_ref;
> +    int ret = 0;
> +
> +    if (!inlink_main->hw_frames_ctx ||
> +        !inlink_overlay->hw_frames_ctx) {
> +        av_log(avctx, AV_LOG_ERROR, "An input HW frames reference is "
> +               "required to associate the processing device.\n");
> +        return AVERROR(EINVAL);
> +    }
> +
> +    input_ref = inlink_main->hw_frames_ctx;
> +    main_frames = (AVHWFramesContext*)input_ref->data;
> +    av_assert0(main_frames);
> +
> +    ctx->device_ref = av_buffer_ref(main_frames->device_ref);
> +    if (!ctx->device_ref) {
> +        av_log(ctx, AV_LOG_ERROR, "A device reference create failed.\n");
> +        return AVERROR(ENOMEM);
> +    }
> +
> +    link->hw_frames_ctx = av_hwframe_ctx_alloc(ctx->device_ref);
> +    if (!link->hw_frames_ctx) {
> +        av_log(avctx, AV_LOG_ERROR, "Failed to create HW frame context "
> +               "for output.\n");
> +        ret = AVERROR(ENOMEM);
> +        return ret;
> +    }
> +    ctx->hwframe_ctx_allocated = 1;
> +
> +    output_frames = (AVHWFramesContext*)link->hw_frames_ctx->data;
> +
> +    output_frames->format    = AV_PIX_FMT_VIDEOTOOLBOX;
> +    output_frames->sw_format = main_frames->sw_format;
> +    output_frames->width     = inlink_main->w;
> +    output_frames->height    = inlink_main->h;
> +
> +    ret = ff_filter_init_hw_frames(avctx, link, 1);
> +    if (ret < 0)
> +        return ret;
> +
> +    ret = av_hwframe_ctx_init(link->hw_frames_ctx);
> +    if (ret < 0) {
> +        av_log(avctx, AV_LOG_ERROR, "Failed to initialise VideoToolbox frame "
> +               "context for output: %d\n", ret);
> +        return ret;
> +    }
> +
> +    link->time_base = inlink_main->time_base;
> +    ctx->fs.time_base = link->time_base;
> +
> +    ret = ff_framesync_init_dualinput(&ctx->fs, avctx);
> +    if (ret < 0)
> +        return ret;
> +
> +    ret = ff_framesync_configure(&ctx->fs);
> +    return ret;
> +}
> +
> +static int config_output(AVFilterLink *link)
> +{
> +    AVFilterContext *ctx = link->src;
> +    if (@available(macOS 10.11, iOS 9.0, *)) {
> +        return do_config_output(link);
> +    } else {
> +        av_log(ctx, AV_LOG_ERROR, "Metal is not available on this OS version\n");
> +        return AVERROR(ENOSYS);
> +    }
> +}
> +
> +static int overlay_videotoolbox_activate(AVFilterContext *avctx)
> +{
> +    OverlayVideoToolboxContext *ctx = avctx->priv;
> +    return ff_framesync_activate(&ctx->fs);
> +}
> +
> +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
> +#define OFFSET(x) offsetof(OverlayVideoToolboxContext, x)
> +
> +static const AVOption overlay_videotoolbox_options[] = {
> +    { "x", "Overlay x position",
> +        OFFSET(x_position), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, INT_MAX, .flags = FLAGS },
> +    { "y", "Overlay y position",
> +        OFFSET(y_position), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, INT_MAX, .flags = FLAGS },
> +    { "eof_action", "Action to take when encountering EOF from secondary input ",
> +        OFFSET(fs.opt_eof_action), AV_OPT_TYPE_INT, { .i64 = EOF_ACTION_REPEAT },
> +        EOF_ACTION_REPEAT, EOF_ACTION_PASS, .flags = FLAGS, .unit = "eof_action" },
> +        { "repeat", "Repeat the previous frame.",   0, AV_OPT_TYPE_CONST, { .i64 = EOF_ACTION_REPEAT }, .flags = FLAGS, .unit = "eof_action" },
> +        { "endall", "End both streams.",            0, AV_OPT_TYPE_CONST, { .i64 = EOF_ACTION_ENDALL }, .flags = FLAGS, .unit = "eof_action" },
> +        { "pass",   "Pass through the main input.", 0, AV_OPT_TYPE_CONST, { .i64 = EOF_ACTION_PASS },   .flags = FLAGS, .unit = "eof_action" },
> +    { "shortest", "force termination when the shortest input terminates", OFFSET(fs.opt_shortest), AV_OPT_TYPE_BOOL, { .i64 = 0 }, 0, 1, FLAGS },
> +    { "repeatlast", "repeat overlay of the last overlay frame", OFFSET(fs.opt_repeatlast), AV_OPT_TYPE_BOOL, { .i64 = 1 }, 0, 1, FLAGS },
> +    { NULL },
> +};
> +
> +FRAMESYNC_DEFINE_CLASS(overlay_videotoolbox, OverlayVideoToolboxContext, fs);
> +
> +static const AVFilterPad overlay_videotoolbox_inputs[] = {
> +    {
> +        .name         = "main",
> +        .type         = AVMEDIA_TYPE_VIDEO,
> +    },
> +    {
> +        .name         = "overlay",
> +        .type         = AVMEDIA_TYPE_VIDEO,
> +    },
> +};
> +
> +static const AVFilterPad overlay_videotoolbox_outputs[] = {
> +    {
> +        .name          = "default",
> +        .type          = AVMEDIA_TYPE_VIDEO,
> +        .config_props  = config_output,
> +    },
> +};
> +
> +const AVFilter ff_vf_overlay_videotoolbox = {
> +    .name           = "overlay_videotoolbox",
> +    .description    = NULL_IF_CONFIG_SMALL("Overlay filter for VideoToolbox frames using Metal compute"),
> +    .priv_size      = OVERLAY_VT_CTX_SIZE,
> +    .priv_class     = &overlay_videotoolbox_class,
> +    .init           = overlay_videotoolbox_init,
> +    .uninit         = overlay_videotoolbox_uninit,
> +    .activate       = overlay_videotoolbox_activate,
> +    .preinit        = overlay_videotoolbox_framesync_preinit,
> +    FILTER_SINGLE_PIXFMT(AV_PIX_FMT_VIDEOTOOLBOX),
> +    FILTER_INPUTS(overlay_videotoolbox_inputs),
> +    FILTER_OUTPUTS(overlay_videotoolbox_outputs),
> +    .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
> +};
> -- 
> 2.39.3 (Apple Git-145)
>