From patchwork Thu Dec 16 20:28:58 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Aman Karmani X-Patchwork-Id: 32668 Delivered-To: ffmpegpatchwork2@gmail.com Received: by 2002:a6b:cd86:0:0:0:0:0 with SMTP id d128csp832719iog; Thu, 16 Dec 2021 12:30:07 -0800 (PST) X-Google-Smtp-Source: ABdhPJxVNEjsvp+uQCGGPJ2hfhvE3G77y70Si03AWu32EKaPGuCBnM7V0xvNmoNE548/P0mSPV+0 X-Received: by 2002:a05:6402:1ca2:: with SMTP id cz2mr21682936edb.302.1639686606973; Thu, 16 Dec 2021 12:30:06 -0800 (PST) ARC-Seal: i=1; a=rsa-sha256; t=1639686606; cv=none; d=google.com; s=arc-20160816; b=w1PDzMs9Eoyg4hXL93Y8c1t2wtz8AHuX/Hhxq+rArd5YZ4h7WGOeUcmRJwIDCSKt0C JW6IEZXPRNKhQzDqfdXO9osb1Y8v0QhcOh49PseZD9tgaInH3sISGRs7r/Btg8XoH9pD KPkL+Eno77bxhzXrPFV97MrSavYlYfOo5csnXXMMmP8jPi2oqO7Nm60VSS6vqIGMIbqC G65kwXesckIqeniSVlVM1I245DbyyYiCNa5Yss+fFC4Jv7A05jkZ4bwqOS7fFMbs1ORN q+hJ/CbpL3vjJxT8NaL1JXLSeC8Gf2NFQdy3w8FuRANbO3PyFlzWgY0VhRtMii/eedOc +NpA== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:content-transfer-encoding:cc:reply-to :list-subscribe:list-help:list-post:list-archive:list-unsubscribe :list-id:precedence:subject:mime-version:references:in-reply-to :message-id:date:to:from:dkim-signature:delivered-to; bh=O/2ffRCwUIezXWt0Qxm1GgHRW6HGQbTR1Hl7OqeqhWE=; b=hzLxCwyfeOQYJQfUaQl9Scgc5hAaUvR/OT2xhUFIG4EOZYz0qc7UGSYC+U314CihRa s6PGkcBtmekibOJEWCOfN9qQ/7xewE/FEMZVIIB2Yv/emmN2f8AM3kzsjPcMiV4S+zbz 60jREEaOWqUQu447gzZi7bTeW0vjdKURzRAIsTenxHEQmPTM65ETjl27WiKC9MjDOdhu xjMqtzT+DU5nRnSAzerW5o6VgNtOtqe8MndxXcA2WqA0udVrwjUV2x/exfcKdnb0khDx cM/M7EWQP0Z0XGqffe4/0NGNqs5ZqL83MQFq7tZSHspBBUpIf8uc9xgd9RwPC2+d2XBU o7DA== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20210112 header.b="muscG/6t"; spf=pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) smtp.mailfrom=ffmpeg-devel-bounces@ffmpeg.org Return-Path: Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org. [79.124.17.100]) by mx.google.com with ESMTP id g20si4612930ejt.472.2021.12.16.12.30.06; Thu, 16 Dec 2021 12:30:06 -0800 (PST) Received-SPF: pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) client-ip=79.124.17.100; Authentication-Results: mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20210112 header.b="muscG/6t"; spf=pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) smtp.mailfrom=ffmpeg-devel-bounces@ffmpeg.org Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id 2636068AF8F; Thu, 16 Dec 2021 22:29:20 +0200 (EET) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-pj1-f44.google.com (mail-pj1-f44.google.com [209.85.216.44]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id DE3E968AF51 for ; Thu, 16 Dec 2021 22:29:11 +0200 (EET) Received: by mail-pj1-f44.google.com with SMTP id fv9-20020a17090b0e8900b001a6a5ab1392so489812pjb.1 for ; Thu, 16 Dec 2021 12:29:11 -0800 (PST) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gmail.com; s=20210112; h=sender:from:to:cc:subject:date:message-id:in-reply-to:references :mime-version:content-transfer-encoding; bh=C+GmDgQw3xnt6LeGt/QzU4d4HFhx1CS/hP2L0CR/rLg=; b=muscG/6tote669z4E6JLrMAJmrX9gmDKBuLkAFErzs1fTkgW7zW+Mc+bSwW8kdJJsl nTd9/5mOdOumnSgWq6I3oxSgAGoLrTA+VigtyjvrOPWL4rV2mFbuyjSZtvNK1Gc93t0f J+gZoeg47hQLP4y0gONjnjr6cPgvqnpN3gDbUaNJnkkWxY6sNDG35ei9fyHexCvBoxG0 JtViUkOZFbLDb8oascuXhNZ9bXjvEs0ZNIWYhndoFBtffATuT527olM/lTfmMu0xo1gs qaHWqwJrrkWK6Un6Lmg3uAwGHVaWZC8BHhXqB5/rf9wsRY3yhBaSB2edLULUrdsRH9gz m7Jg== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20210112; h=x-gm-message-state:sender:from:to:cc:subject:date:message-id :in-reply-to:references:mime-version:content-transfer-encoding; bh=C+GmDgQw3xnt6LeGt/QzU4d4HFhx1CS/hP2L0CR/rLg=; b=44Fy65zCEFAw2NbSuuBFmxiJO4H9gyw/C4SamYfYs6d+xTIblm3U2a2H/04+dcVoz1 2inMlmlSjNfV8X2pj+/0mu4iYR3FnyAmMGC6aO6cZmA1WzbLhhd58UzCk0ca/Pr1+Yu2 JkPTIpHeq33hu7jmMyjyLG6MGDlCDmgRl3CiW+ctfXRWiE9XOepeqf+Lip1KR1JQ4pgr bnkkaxT39NsGQ8q890O13g8AUg6YjggN2LH4qKKF6Wft4zE2kiQNBv2EEa27tsISKfmY e6rDeweYHEz5yY+oz2PIR9lvQxFfXqNA9m+huooiA2E309WU5lPtAUudY7kfw4YIgbBv g7SQ== X-Gm-Message-State: AOAM5339CKtJULr02W+yuJgLVhXQb6FwauyIYSIcgHFWE9Vh8vPiJba+ qLXpn1lxyGtxf7Q68PVrMdtgpJBms5Umbg== X-Received: by 2002:a17:903:2448:b0:148:a658:8d31 with SMTP id l8-20020a170903244800b00148a6588d31mr10227741pls.133.1639686549651; Thu, 16 Dec 2021 12:29:09 -0800 (PST) Received: from tmm1-imac.lan ([2600:8802:5501:308f:15c6:3324:e59b:bed5]) by smtp.gmail.com with ESMTPSA id k16sm7534732pfu.183.2021.12.16.12.29.08 (version=TLS1_2 cipher=ECDHE-ECDSA-AES128-GCM-SHA256 bits=128/128); Thu, 16 Dec 2021 12:29:09 -0800 (PST) From: Aman Karmani To: ffmpeg-devel@ffmpeg.org Date: Thu, 16 Dec 2021 12:28:58 -0800 Message-Id: <20211216202858.77643-5-ffmpeg@tmm1.net> X-Mailer: git-send-email 2.33.0 In-Reply-To: <20211216202858.77643-1-ffmpeg@tmm1.net> References: <20211216202858.77643-1-ffmpeg@tmm1.net> MIME-Version: 1.0 Subject: [FFmpeg-devel] [PATCH v3 5/5] avfilter: add vf_deinterlace_videotoolbox X-BeenThere: ffmpeg-devel@ffmpeg.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: FFmpeg development discussions and patches List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Reply-To: FFmpeg development discussions and patches Cc: philipl@overt.org, Aman Karmani , kernrj@gmail.com Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" X-TUID: Q9ZqQFEyKL5X From: Aman Karmani deinterlaces CVPixelBuffers, i.e. AV_PIX_FMT_VIDEOTOOLBOX frames currently implements YADIF, but other algorithms could easily be added to the same filter. for example, an interlaced mpeg2 file can be decoded by avcodec, uploaded into a CVPixelBuffer, deinterlaced by metal, and then encoded to h264 by VideoToolbox as follows: ffmpeg \ -init_hw_device videotoolbox \ -i interlaced.ts \ -vf hwupload,deinterlace_videotoolbox=mode=send_field:deint=interlaced \ -c:v h264_videotoolbox \ -b:v 2000k \ -c:a copy \ -y progressive.ts (note that uploading AVFrame into CVPixelBuffer via hwupload requires 504c60660d3194758823ddd45ceddb86e35d806f) Signed-off-by: Aman Karmani --- configure | 1 + libavfilter/Makefile | 4 + libavfilter/allfilters.c | 1 + .../metal/vf_deinterlace_videotoolbox.metal | 269 ++++++++++++ libavfilter/vf_deinterlace_videotoolbox.m | 400 ++++++++++++++++++ 5 files changed, 675 insertions(+) create mode 100644 libavfilter/metal/vf_deinterlace_videotoolbox.metal create mode 100644 libavfilter/vf_deinterlace_videotoolbox.m diff --git a/configure b/configure index ab00b2d7cb..cbaef21bbf 100755 --- a/configure +++ b/configure @@ -3620,6 +3620,7 @@ cover_rect_filter_deps="avcodec avformat gpl" cropdetect_filter_deps="gpl" deinterlace_qsv_filter_deps="libmfx" deinterlace_vaapi_filter_deps="vaapi" +deinterlace_videotoolbox_filter_deps="metal corevideo videotoolbox" delogo_filter_deps="gpl" denoise_vaapi_filter_deps="vaapi" derain_filter_select="dnn" diff --git a/libavfilter/Makefile b/libavfilter/Makefile index 2fe495df28..4812f88045 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -241,6 +241,10 @@ OBJS-$(CONFIG_DEFLATE_FILTER) += vf_neighbor.o OBJS-$(CONFIG_DEFLICKER_FILTER) += vf_deflicker.o OBJS-$(CONFIG_DEINTERLACE_QSV_FILTER) += vf_deinterlace_qsv.o OBJS-$(CONFIG_DEINTERLACE_VAAPI_FILTER) += vf_deinterlace_vaapi.o vaapi_vpp.o +OBJS-$(CONFIG_DEINTERLACE_VIDEOTOOLBOX_FILTER) += vf_deinterlace_videotoolbox.o \ + metal/vf_deinterlace_videotoolbox.metallib.o \ + metal/utils.o \ + yadif_common.o OBJS-$(CONFIG_DEJUDDER_FILTER) += vf_dejudder.o OBJS-$(CONFIG_DELOGO_FILTER) += vf_delogo.o OBJS-$(CONFIG_DENOISE_VAAPI_FILTER) += vf_misc_vaapi.o vaapi_vpp.o diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index ec57a2c49c..2ed3deb7dd 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -228,6 +228,7 @@ extern const AVFilter ff_vf_deflate; extern const AVFilter ff_vf_deflicker; extern const AVFilter ff_vf_deinterlace_qsv; extern const AVFilter ff_vf_deinterlace_vaapi; +extern const AVFilter ff_vf_deinterlace_videotoolbox; extern const AVFilter ff_vf_dejudder; extern const AVFilter ff_vf_delogo; extern const AVFilter ff_vf_denoise_vaapi; diff --git a/libavfilter/metal/vf_deinterlace_videotoolbox.metal b/libavfilter/metal/vf_deinterlace_videotoolbox.metal new file mode 100644 index 0000000000..50783f2ffe --- /dev/null +++ b/libavfilter/metal/vf_deinterlace_videotoolbox.metal @@ -0,0 +1,269 @@ +/* + * Copyright (C) 2018 Philip Langdale + * 2020 Aman Karmani + * 2020 Stefan Dyulgerov + * + * 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 +#include +#include + +using namespace metal; + +/* + * Parameters + */ + +struct deintParams { + uint channels; + uint parity; + uint tff; + bool is_second_field; + bool skip_spatial_check; + int field_mode; +}; + +/* + * Texture access helpers + */ + +#define accesstype access::sample +const sampler s(coord::pixel); + +template +T tex2D(texture2d tex, uint x, uint y) +{ + return tex.sample(s, float2(x, y)).x; +} + +template <> +float2 tex2D(texture2d tex, uint x, uint y) +{ + return tex.sample(s, float2(x, y)).xy; +} + +template +T tex2D(texture2d tex, uint x, uint y) +{ + return tex.read(uint2(x, y)).x; +} + +template <> +float2 tex2D(texture2d tex, uint x, uint y) +{ + return tex.read(uint2(x, y)).xy; +} + +/* + * YADIF helpers + */ + +template +T spatial_predictor(T a, T b, T c, T d, T e, T f, T g, + T h, T i, T j, T k, T l, T m, T n) +{ + T spatial_pred = (d + k)/2; + T spatial_score = abs(c - j) + abs(d - k) + abs(e - l); + + T score = abs(b - k) + abs(c - l) + abs(d - m); + if (score < spatial_score) { + spatial_pred = (c + l)/2; + spatial_score = score; + score = abs(a - l) + abs(b - m) + abs(c - n); + if (score < spatial_score) { + spatial_pred = (b + m)/2; + spatial_score = score; + } + } + score = abs(d - i) + abs(e - j) + abs(f - k); + if (score < spatial_score) { + spatial_pred = (e + j)/2; + spatial_score = score; + score = abs(e - h) + abs(f - i) + abs(g - j); + if (score < spatial_score) { + spatial_pred = (f + i)/2; + spatial_score = score; + } + } + return spatial_pred; +} + +template +T temporal_predictor(T A, T B, T C, T D, T E, T F, + T G, T H, T I, T J, T K, T L, + T spatial_pred, bool skip_check) +{ + T p0 = (C + H) / 2; + T p1 = F; + T p2 = (D + I) / 2; + T p3 = G; + T p4 = (E + J) / 2; + + T tdiff0 = abs(D - I); + T tdiff1 = (abs(A - F) + abs(B - G)) / 2; + T tdiff2 = (abs(K - F) + abs(G - L)) / 2; + + T diff = max3(tdiff0, tdiff1, tdiff2); + + if (!skip_check) { + T maxi = max3(p2 - p3, p2 - p1, min(p0 - p1, p4 - p3)); + T mini = min3(p2 - p3, p2 - p1, max(p0 - p1, p4 - p3)); + diff = max3(diff, mini, -maxi); + } + + return clamp(spatial_pred, p2 - diff, p2 + diff); +} + +#define T float2 +template <> +T spatial_predictor(T a, T b, T c, T d, T e, T f, T g, + T h, T i, T j, T k, T l, T m, T n) +{ + return T( + spatial_predictor(a.x, b.x, c.x, d.x, e.x, f.x, g.x, + h.x, i.x, j.x, k.x, l.x, m.x, n.x), + spatial_predictor(a.y, b.y, c.y, d.y, e.y, f.y, g.y, + h.y, i.y, j.y, k.y, l.y, m.y, n.y) + ); +} + +template <> +T temporal_predictor(T A, T B, T C, T D, T E, T F, + T G, T H, T I, T J, T K, T L, + T spatial_pred, bool skip_check) +{ + return T( + temporal_predictor(A.x, B.x, C.x, D.x, E.x, F.x, + G.x, H.x, I.x, J.x, K.x, L.x, + spatial_pred.x, skip_check), + temporal_predictor(A.y, B.y, C.y, D.y, E.y, F.y, + G.y, H.y, I.y, J.y, K.y, L.y, + spatial_pred.y, skip_check) + ); +} +#undef T + +/* + * YADIF compute + */ + +template +T yadif_compute_spatial( + texture2d cur, + uint2 pos) +{ + // Calculate spatial prediction + T a = tex2D(cur, pos.x - 3, pos.y - 1); + T b = tex2D(cur, pos.x - 2, pos.y - 1); + T c = tex2D(cur, pos.x - 1, pos.y - 1); + T d = tex2D(cur, pos.x - 0, pos.y - 1); + T e = tex2D(cur, pos.x + 1, pos.y - 1); + T f = tex2D(cur, pos.x + 2, pos.y - 1); + T g = tex2D(cur, pos.x + 3, pos.y - 1); + + T h = tex2D(cur, pos.x - 3, pos.y + 1); + T i = tex2D(cur, pos.x - 2, pos.y + 1); + T j = tex2D(cur, pos.x - 1, pos.y + 1); + T k = tex2D(cur, pos.x - 0, pos.y + 1); + T l = tex2D(cur, pos.x + 1, pos.y + 1); + T m = tex2D(cur, pos.x + 2, pos.y + 1); + T n = tex2D(cur, pos.x + 3, pos.y + 1); + + return spatial_predictor(a, b, c, d, e, f, g, + h, i, j, k, l, m, n); +} + +template +T yadif_compute_temporal( + texture2d cur, + texture2d prev2, + texture2d prev1, + texture2d next1, + texture2d next2, + T spatial_pred, + bool skip_spatial_check, + uint2 pos) +{ + // Calculate temporal prediction + T A = tex2D(prev2, pos.x, pos.y - 1); + T B = tex2D(prev2, pos.x, pos.y + 1); + T C = tex2D(prev1, pos.x, pos.y - 2); + T D = tex2D(prev1, pos.x, pos.y + 0); + T E = tex2D(prev1, pos.x, pos.y + 2); + T F = tex2D(cur, pos.x, pos.y - 1); + T G = tex2D(cur, pos.x, pos.y + 1); + T H = tex2D(next1, pos.x, pos.y - 2); + T I = tex2D(next1, pos.x, pos.y + 0); + T J = tex2D(next1, pos.x, pos.y + 2); + T K = tex2D(next2, pos.x, pos.y - 1); + T L = tex2D(next2, pos.x, pos.y + 1); + + return temporal_predictor(A, B, C, D, E, F, G, H, I, J, K, L, + spatial_pred, skip_spatial_check); +} + +template +T yadif( + texture2d dst, + texture2d prev, + texture2d cur, + texture2d next, + constant deintParams& params, + uint2 pos) +{ + T spatial_pred = yadif_compute_spatial(cur, pos); + + if (params.is_second_field) { + return yadif_compute_temporal(cur, prev, cur, next, next, spatial_pred, params.skip_spatial_check, pos); + } else { + return yadif_compute_temporal(cur, prev, prev, cur, next, spatial_pred, params.skip_spatial_check, pos); + } +} + +/* + * Kernel dispatch + */ + +kernel void deint( + texture2d dst [[texture(0)]], + texture2d prev [[texture(1)]], + texture2d cur [[texture(2)]], + texture2d next [[texture(3)]], + constant deintParams& params [[buffer(4)]], + uint2 pos [[thread_position_in_grid]]) +{ + if ((pos.x >= dst.get_width()) || + (pos.y >= dst.get_height())) { + return; + } + + // Don't modify the primary field + if (pos.y % 2 == params.parity) { + float4 in = cur.read(pos); + dst.write(in, pos); + return; + } + + float2 pred; + if (params.channels == 1) + pred = float2(yadif(dst, prev, cur, next, params, pos)); + else + pred = yadif(dst, prev, cur, next, params, pos); + dst.write(pred.xyyy, pos); +} diff --git a/libavfilter/vf_deinterlace_videotoolbox.m b/libavfilter/vf_deinterlace_videotoolbox.m new file mode 100644 index 0000000000..b18155277b --- /dev/null +++ b/libavfilter/vf_deinterlace_videotoolbox.m @@ -0,0 +1,400 @@ +/* + * Copyright (C) 2018 Philip Langdale + * 2020 Aman Karmani + * + * 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 "yadif.h" +#include +#include +#include +#include + +extern char ff_vf_deinterlace_videotoolbox_metallib_data[]; +extern unsigned int ff_vf_deinterlace_videotoolbox_metallib_len; + +typedef struct DeintVTContext { + YADIFContext yadif; + + AVBufferRef *device_ref; + AVBufferRef *input_frames_ref; + AVHWFramesContext *input_frames; + + id mtlDevice; + id mtlLibrary; + id mtlQueue; + id mtlPipeline; + id mtlFunction; + id mtlParamsBuffer; + + CVMetalTextureCacheRef textureCache; +} DeintVTContext; + +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 dst, + id prev, + id cur, + id next, + int channels, + int parity, + int tff) +{ + DeintVTContext *s = ctx->priv; + id buffer = s->mtlQueue.commandBuffer; + id 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) +{ + DeintVTContext *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 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 int deint_videotoolbox_init(AVFilterContext *ctx) +{ + DeintVTContext *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"); + return AVERROR_EXTERNAL; + } + + av_log(ctx, AV_LOG_INFO, "Using Metal device: %s\n", s->mtlDevice.name.UTF8String); + + dispatch_data_t libData = dispatch_data_create( + ff_vf_deinterlace_videotoolbox_metallib_data, + ff_vf_deinterlace_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); + return AVERROR_EXTERNAL; + } + s->mtlFunction = [s->mtlLibrary newFunctionWithName:@"deint"]; + + s->mtlQueue = s->mtlDevice.newCommandQueue; + if (!s->mtlQueue) { + av_log(ctx, AV_LOG_ERROR, "Failed to create Metal command queue!\n"); + return AVERROR_EXTERNAL; + } + + 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); + return AVERROR_EXTERNAL; + } + + 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"); + return AVERROR_EXTERNAL; + } + + 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); + return AVERROR_EXTERNAL; + } + + return 0; +} + +static av_cold void deint_videotoolbox_uninit(AVFilterContext *ctx) +{ + DeintVTContext *s = ctx->priv; + YADIFContext *y = &s->yadif; + + av_frame_free(&y->prev); + av_frame_free(&y->cur); + av_frame_free(&y->next); + + 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 int config_input(AVFilterLink *inlink) +{ + AVFilterContext *ctx = inlink->dst; + DeintVTContext *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_output(AVFilterLink *link) +{ + AVHWFramesContext *output_frames; + AVFilterContext *ctx = link->src; + DeintVTContext *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; + } + + link->time_base.num = ctx->inputs[0]->time_base.num; + link->time_base.den = ctx->inputs[0]->time_base.den * 2; + link->w = ctx->inputs[0]->w; + link->h = ctx->inputs[0]->h; + + if(y->mode & 1) + link->frame_rate = av_mul_q(ctx->inputs[0]->frame_rate, + (AVRational){2, 1}); + + if (link->w < 3 || link->h < 3) { + av_log(ctx, AV_LOG_ERROR, "Video of less than 3 columns or lines is not supported\n"); + ret = AVERROR(EINVAL); + goto exit; + } + + y->csp = av_pix_fmt_desc_get(output_frames->sw_format); + y->filter = filter; + +exit: + return ret; +} + +#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 deinterlace_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, "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, "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, "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(deinterlace_videotoolbox); + +static const AVFilterPad deint_videotoolbox_inputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .filter_frame = ff_yadif_filter_frame, + .config_props = config_input, + }, +}; + +static const AVFilterPad deint_videotoolbox_outputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .request_frame = ff_yadif_request_frame, + .config_props = config_output, + }, +}; + +AVFilter ff_vf_deinterlace_videotoolbox = { + .name = "deinterlace_videotoolbox", + .description = NULL_IF_CONFIG_SMALL("Deinterlace VideoToolbox frames with Metal compute"), + .priv_size = sizeof(DeintVTContext), + .priv_class = &deinterlace_videotoolbox_class, + .init = deint_videotoolbox_init, + .uninit = deint_videotoolbox_uninit, + FILTER_SINGLE_PIXFMT(AV_PIX_FMT_VIDEOTOOLBOX), + FILTER_INPUTS(deint_videotoolbox_inputs), + FILTER_OUTPUTS(deint_videotoolbox_outputs), + .flags = AVFILTER_FLAG_SUPPORT_TIMELINE_INTERNAL, + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +};