ffmpeg/libavfilter/vf_yadif_videotoolbox.m

/*
 * Copyright (C) 2018 Philip Langdale <philipl@overt.org>
 *               2020 Aman Karmani <aman@tmm1.net>
 *
 * 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 "internal.h"
#include "metal/utils.h"
#include "yadif.h"
#include "libavutil/avassert.h"
#include "libavutil/hwcontext.h"
#include "libavutil/objc.h"

#include <assert.h>

extern char ff_vf_yadif_videotoolbox_metallib_data[];
extern unsigned int ff_vf_yadif_videotoolbox_metallib_len;

typedef struct API_AVAILABLE(macos(10.11), ios(8.0)) YADIFVTContext {
    YADIFContext yadif;

    AVBufferRef       *device_ref;
    AVBufferRef       *input_frames_ref;
    AVHWFramesContext *input_frames;

    id<MTLDevice> mtlDevice;
    id<MTLLibrary> mtlLibrary;
    id<MTLCommandQueue> mtlQueue;
    id<MTLComputePipelineState> mtlPipeline;
    id<MTLFunction> mtlFunction;
    id<MTLBuffer> mtlParamsBuffer;

    CVMetalTextureCacheRef textureCache;
} YADIFVTContext API_AVAILABLE(macos(10.11), ios(8.0));

// Using sizeof(YADIFVTContext) outside of 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 yadif_videotoolbox_init behind a check)
#define YADIF_VT_CTX_SIZE (sizeof(YADIFContext) + sizeof(void*) * 10)

struct mtlYadifParams {
    uint channels;
    uint parity;
    uint tff;
    bool is_second_field;
    bool skip_spatial_check;
    int field_mode;
};

static void call_kernel(AVFilterContext *ctx,
                        id<MTLTexture> dst,
                        id<MTLTexture> prev,
                        id<MTLTexture> cur,
                        id<MTLTexture> next,
                        int channels,
                        int parity,
                        int tff) API_AVAILABLE(macos(10.11), ios(8.0))
{
    YADIFVTContext *s = ctx->priv;
    id<MTLCommandBuffer> buffer = s->mtlQueue.commandBuffer;
    id<MTLComputeCommandEncoder> encoder = buffer.computeCommandEncoder;
    struct mtlYadifParams *params = (struct mtlYadifParams *)s->mtlParamsBuffer.contents;
    *params = (struct mtlYadifParams){
        .channels = channels,
        .parity = parity,
        .tff = tff,
        .is_second_field = !(parity ^ tff),
        .skip_spatial_check = s->yadif.mode&2,
        .field_mode = s->yadif.current_field
    };

    [encoder setTexture:dst  atIndex:0];
    [encoder setTexture:prev atIndex:1];
    [encoder setTexture:cur  atIndex:2];
    [encoder setTexture:next atIndex:3];
    [encoder setBuffer:s->mtlParamsBuffer offset:0 atIndex:4];
    ff_metal_compute_encoder_dispatch(s->mtlDevice, s->mtlPipeline, encoder, dst.width, dst.height);
    [encoder endEncoding];

    [buffer commit];
    [buffer waitUntilCompleted];

    ff_objc_release(&encoder);
    ff_objc_release(&buffer);
}

static void filter(AVFilterContext *ctx, AVFrame *dst,
                   int parity, int tff) API_AVAILABLE(macos(10.11), ios(8.0))
{
    YADIFVTContext *s = ctx->priv;
    YADIFContext *y = &s->yadif;
    int i;

    for (i = 0; i < y->csp->nb_components; i++) {
        int pixel_size, channels;
        const AVComponentDescriptor *comp = &y->csp->comp[i];
        CVMetalTextureRef prev, cur, next, dest;
        id<MTLTexture> tex_prev, tex_cur, tex_next, tex_dest;
        MTLPixelFormat format;

        if (comp->plane < i) {
            // We process planes as a whole, so don't reprocess
            // them for additional components
            continue;
        }

        pixel_size = (comp->depth + comp->shift) / 8;
        channels = comp->step / pixel_size;
        if (pixel_size > 2 || channels > 2) {
            av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n", y->csp->name);
            goto exit;
        }
        switch (pixel_size) {
        case 1:
            format = channels == 1 ? MTLPixelFormatR8Unorm : MTLPixelFormatRG8Unorm;
            break;
        case 2:
            format = channels == 1 ? MTLPixelFormatR16Unorm : MTLPixelFormatRG16Unorm;
            break;
        default:
            av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n", y->csp->name);
            goto exit;
        }
        av_log(ctx, AV_LOG_TRACE,
               "Deinterlacing plane %d: pixel_size: %d channels: %d\n",
               comp->plane, pixel_size, channels);

        prev = ff_metal_texture_from_pixbuf(ctx, s->textureCache, (CVPixelBufferRef)y->prev->data[3], i, format);
        cur  = ff_metal_texture_from_pixbuf(ctx, s->textureCache, (CVPixelBufferRef)y->cur->data[3], i, format);
        next = ff_metal_texture_from_pixbuf(ctx, s->textureCache, (CVPixelBufferRef)y->next->data[3], i, format);
        dest = ff_metal_texture_from_pixbuf(ctx, s->textureCache, (CVPixelBufferRef)dst->data[3], i, format);

        tex_prev = CVMetalTextureGetTexture(prev);
        tex_cur  = CVMetalTextureGetTexture(cur);
        tex_next = CVMetalTextureGetTexture(next);
        tex_dest = CVMetalTextureGetTexture(dest);

        call_kernel(ctx, tex_dest, tex_prev, tex_cur, tex_next,
                         channels, parity, tff);

        CFRelease(prev);
        CFRelease(cur);
        CFRelease(next);
        CFRelease(dest);
    }

    CVBufferPropagateAttachments((CVPixelBufferRef)y->cur->data[3], (CVPixelBufferRef)dst->data[3]);

    if (y->current_field == YADIF_FIELD_END) {
        y->current_field = YADIF_FIELD_NORMAL;
    }

exit:
    return;
}

static av_cold void do_uninit(AVFilterContext *ctx) API_AVAILABLE(macos(10.11), ios(8.0))
{
    YADIFVTContext *s = ctx->priv;
    YADIFContext *y = &s->yadif;

    ff_yadif_uninit(ctx);

    av_buffer_unref(&s->device_ref);
    av_buffer_unref(&s->input_frames_ref);
    s->input_frames = NULL;

    ff_objc_release(&s->mtlParamsBuffer);
    ff_objc_release(&s->mtlFunction);
    ff_objc_release(&s->mtlPipeline);
    ff_objc_release(&s->mtlQueue);
    ff_objc_release(&s->mtlLibrary);
    ff_objc_release(&s->mtlDevice);

    if (s->textureCache) {
        CFRelease(s->textureCache);
        s->textureCache = NULL;
    }
}


static av_cold void yadif_videotoolbox_uninit(AVFilterContext *ctx)
{
    if (@available(macOS 10.11, iOS 8.0, *)) {
        do_uninit(ctx);
    }
}

static av_cold int do_init(AVFilterContext *ctx) API_AVAILABLE(macos(10.11), ios(8.0))
{
    YADIFVTContext *s = ctx->priv;
    NSError *err = nil;
    CVReturn ret;

    s->mtlDevice = MTLCreateSystemDefaultDevice();
    if (!s->mtlDevice) {
        av_log(ctx, AV_LOG_ERROR, "Unable to find Metal device\n");
        goto fail;
    }

    av_log(ctx, AV_LOG_INFO, "Using Metal device: %s\n", s->mtlDevice.name.UTF8String);

    dispatch_data_t libData = dispatch_data_create(
        ff_vf_yadif_videotoolbox_metallib_data,
        ff_vf_yadif_videotoolbox_metallib_len,
        nil,
        nil);
    s->mtlLibrary = [s->mtlDevice newLibraryWithData:libData error:&err];
    dispatch_release(libData);
    libData = nil;
    if (err) {
        av_log(ctx, AV_LOG_ERROR, "Failed to load Metal library: %s\n", err.description.UTF8String);
        goto fail;
    }

    s->mtlFunction = [s->mtlLibrary newFunctionWithName:@"deint"];
    if (!s->mtlFunction) {
        av_log(ctx, AV_LOG_ERROR, "Failed to create Metal function!\n");
        goto fail;
    }

    s->mtlQueue = s->mtlDevice.newCommandQueue;
    if (!s->mtlQueue) {
        av_log(ctx, AV_LOG_ERROR, "Failed to create Metal command queue!\n");
        goto fail;
    }

    s->mtlPipeline = [s->mtlDevice
        newComputePipelineStateWithFunction:s->mtlFunction
        error:&err];
    if (err) {
        av_log(ctx, AV_LOG_ERROR, "Failed to create Metal compute pipeline: %s\n", err.description.UTF8String);
        goto fail;
    }

    s->mtlParamsBuffer = [s->mtlDevice
        newBufferWithLength:sizeof(struct mtlYadifParams)
        options:MTLResourceStorageModeShared];
    if (!s->mtlParamsBuffer) {
        av_log(ctx, AV_LOG_ERROR, "Failed to create Metal buffer for parameters\n");
        goto fail;
    }

    ret = CVMetalTextureCacheCreate(
        NULL,
        NULL,
        s->mtlDevice,
        NULL,
        &s->textureCache
    );
    if (ret != kCVReturnSuccess) {
        av_log(ctx, AV_LOG_ERROR, "Failed to create CVMetalTextureCache: %d\n", ret);
        goto fail;
    }

    return 0;
fail:
    yadif_videotoolbox_uninit(ctx);
    return AVERROR_EXTERNAL;
}

static av_cold int yadif_videotoolbox_init(AVFilterContext *ctx)
{
    if (@available(macOS 10.11, iOS 8.0, *)) {
        // Ensure we calculated YADIF_VT_CTX_SIZE correctly
        static_assert(YADIF_VT_CTX_SIZE == sizeof(YADIFVTContext), "Incorrect YADIF_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_input(AVFilterLink *inlink) API_AVAILABLE(macos(10.11), ios(8.0))
{
    AVFilterContext *ctx = inlink->dst;
    YADIFVTContext *s = ctx->priv;

    if (!inlink->hw_frames_ctx) {
        av_log(ctx, AV_LOG_ERROR, "A hardware frames reference is "
               "required to associate the processing device.\n");
        return AVERROR(EINVAL);
    }

    s->input_frames_ref = av_buffer_ref(inlink->hw_frames_ctx);
    if (!s->input_frames_ref) {
        av_log(ctx, AV_LOG_ERROR, "A input frames reference create "
               "failed.\n");
        return AVERROR(ENOMEM);
    }
    s->input_frames = (AVHWFramesContext*)s->input_frames_ref->data;

    return 0;
}

static int config_input(AVFilterLink *inlink)
{
    AVFilterContext *ctx = inlink->dst;
    if (@available(macOS 10.11, iOS 8.0, *)) {
        return do_config_input(inlink);
    } 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(8.0))
{
    AVHWFramesContext *output_frames;
    AVFilterContext *ctx = link->src;
    YADIFVTContext *s = ctx->priv;
    YADIFContext *y = &s->yadif;
    int ret = 0;

    av_assert0(s->input_frames);
    s->device_ref = av_buffer_ref(s->input_frames->device_ref);
    if (!s->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(s->device_ref);
    if (!link->hw_frames_ctx) {
        av_log(ctx, AV_LOG_ERROR, "Failed to create HW frame context "
               "for output.\n");
        ret = AVERROR(ENOMEM);
        goto exit;
    }

    output_frames = (AVHWFramesContext*)link->hw_frames_ctx->data;

    output_frames->format    = AV_PIX_FMT_VIDEOTOOLBOX;
    output_frames->sw_format = s->input_frames->sw_format;
    output_frames->width     = ctx->inputs[0]->w;
    output_frames->height    = ctx->inputs[0]->h;

    ret = ff_filter_init_hw_frames(ctx, link, 10);
    if (ret < 0)
        goto exit;

    ret = av_hwframe_ctx_init(link->hw_frames_ctx);
    if (ret < 0) {
        av_log(ctx, AV_LOG_ERROR, "Failed to initialise VideoToolbox frame "
               "context for output: %d\n", ret);
        goto exit;
    }

    ret = ff_yadif_config_output_common(link);
    if (ret < 0)
        goto exit;

    y->csp = av_pix_fmt_desc_get(output_frames->sw_format);
    y->filter = filter;

exit:
    return ret;
}

static int config_output(AVFilterLink *link)
{
    AVFilterContext *ctx = link->src;
    if (@available(macOS 10.11, iOS 8.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);
    }
}

#define FLAGS AV_OPT_FLAG_VIDEO_PARAM|AV_OPT_FLAG_FILTERING_PARAM
#define CONST(name, help, val, unit) { name, help, 0, AV_OPT_TYPE_CONST, {.i64=val}, INT_MIN, INT_MAX, FLAGS, unit }

static const AVOption yadif_videotoolbox_options[] = {
    #define OFFSET(x) offsetof(YADIFContext, x)
    { "mode",   "specify the interlacing mode", OFFSET(mode), AV_OPT_TYPE_INT, {.i64=YADIF_MODE_SEND_FRAME}, 0, 3, FLAGS, .unit = "mode"},
    CONST("send_frame",           "send one frame for each frame",                                     YADIF_MODE_SEND_FRAME,           "mode"),
    CONST("send_field",           "send one frame for each field",                                     YADIF_MODE_SEND_FIELD,           "mode"),
    CONST("send_frame_nospatial", "send one frame for each frame, but skip spatial interlacing check", YADIF_MODE_SEND_FRAME_NOSPATIAL, "mode"),
    CONST("send_field_nospatial", "send one frame for each field, but skip spatial interlacing check", YADIF_MODE_SEND_FIELD_NOSPATIAL, "mode"),

    { "parity", "specify the assumed picture field parity", OFFSET(parity), AV_OPT_TYPE_INT, {.i64=YADIF_PARITY_AUTO}, -1, 1, FLAGS, .unit = "parity" },
    CONST("tff",  "assume top field first",    YADIF_PARITY_TFF,  "parity"),
    CONST("bff",  "assume bottom field first", YADIF_PARITY_BFF,  "parity"),
    CONST("auto", "auto detect parity",        YADIF_PARITY_AUTO, "parity"),

    { "deint", "specify which frames to deinterlace", OFFSET(deint), AV_OPT_TYPE_INT, {.i64=YADIF_DEINT_ALL}, 0, 1, FLAGS, .unit = "deint" },
    CONST("all",        "deinterlace all frames",                       YADIF_DEINT_ALL,        "deint"),
    CONST("interlaced", "only deinterlace frames marked as interlaced", YADIF_DEINT_INTERLACED, "deint"),
    #undef OFFSET

    { NULL }
};

AVFILTER_DEFINE_CLASS(yadif_videotoolbox);

static const AVFilterPad yadif_videotoolbox_inputs[] = {
    {
        .name          = "default",
        .type          = AVMEDIA_TYPE_VIDEO,
        .filter_frame  = ff_yadif_filter_frame,
        .config_props  = config_input,
    },
};

static const AVFilterPad yadif_videotoolbox_outputs[] = {
    {
        .name          = "default",
        .type          = AVMEDIA_TYPE_VIDEO,
        .request_frame = ff_yadif_request_frame,
        .config_props  = config_output,
    },
};

const AVFilter ff_vf_yadif_videotoolbox = {
    .name           = "yadif_videotoolbox",
    .description    = NULL_IF_CONFIG_SMALL("YADIF for VideoToolbox frames using Metal compute"),
    .priv_size      = YADIF_VT_CTX_SIZE,
    .priv_class     = &yadif_videotoolbox_class,
    .init           = yadif_videotoolbox_init,
    .uninit         = yadif_videotoolbox_uninit,
    FILTER_SINGLE_PIXFMT(AV_PIX_FMT_VIDEOTOOLBOX),
    FILTER_INPUTS(yadif_videotoolbox_inputs),
    FILTER_OUTPUTS(yadif_videotoolbox_outputs),
    .flags          = AVFILTER_FLAG_SUPPORT_TIMELINE_INTERNAL,
    .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
};