From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org [79.124.17.100]) by master.gitmailbox.com (Postfix) with ESMTP id CAB6B4003B for ; Sun, 19 Dec 2021 00:49:11 +0000 (UTC) Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id 386B968AF85; Sun, 19 Dec 2021 02:49:08 +0200 (EET) Received: from mail-pj1-f43.google.com (mail-pj1-f43.google.com [209.85.216.43]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 7B55768AF36 for ; Sun, 19 Dec 2021 02:49:01 +0200 (EET) Received: by mail-pj1-f43.google.com with SMTP id n15-20020a17090a160f00b001a75089daa3so9290352pja.1 for ; Sat, 18 Dec 2021 16:49:01 -0800 (PST) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gmail.com; s=20210112; h=mime-version:references:in-reply-to:from:date:message-id:subject:to; bh=effaXCAL5ZfX3LYzvwoh97aUThh3XSS3aQYigdK4HBI=; b=AgIgtbH9gl/9DfNayCG2WGB4EKhCjJOYMFHjXnn3YJkUjrL+bDIfGpmwHkpeHogWlm qb2ZhbHFUucDRqcaCrFaDD1uueIfWuuFyc/8mZtfWC4XBEBru1bt+tOif+3vwJkRLGxF HbUfBbeKM0F7Sa7JKkTUI7mM/qgtP8e6cKuR1r7QGEQXObSyntX59qDURwILrj3nQ7Bf tPj0IlOA5600urnn2dAcJ3+tWELjC15q0rNwj5yhnbuE9VJXiOCWY1d/IzLtRI6rSGQF BQeImsYJnwIspLt0pwAnqTOLH7PDrxLvoP+DkhWPTGJ1JfMXttJG6EgdzhlMtL4RroCI kVDA== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20210112; h=x-gm-message-state:mime-version:references:in-reply-to:from:date :message-id:subject:to; bh=effaXCAL5ZfX3LYzvwoh97aUThh3XSS3aQYigdK4HBI=; b=V20/nV4mIdCSiMZYeU3U4Ho5lfC+FWq/OVuGiFJOk/LFoG9kvX5wctYn9IsIX+S0Xe nBnSQ+H+O/zvYIS9TZ0moh/f1cZuUzxgbFHVHNnl/LPE/2knXCnjzvht4h6N5m52eiyR 8qQE+Iwe+ulFAxIu8BothFQ2D4CRa5FVPzkY5pbbENSi7bMVbX3xxcwjOTcEhCnNEOdv N3gpn0l/fhteWUSSIkcsg/r+9KINyeKYUuBJ8Ld44G4XiehrbVJ8ssEUzfMAncU7g/yb 4eM+x2LvLIFZHFMYTbqPMjUGhZ9diGZA9++mAHGs7yfTehqVn/U17WlQ7IsUSNVgtlH8 sERQ== X-Gm-Message-State: AOAM533nk5lCvJPfDlD1EpGnQswMyOZVFzVcGpGg4j5djSrwzVbn/93S +dzmY17ZWLgLx2G3HOcwMIS8wXk6eKK2SZ3vtbpTRW9pFCDTJg== X-Google-Smtp-Source: ABdhPJx7R8sRloECGAs2BJABRSDOgQ7oKJOyLmjHaiS4COEEJzPpHoAAdcUy48JsgAX6BTfybkS1RcpUlb+Jda7D7pw= X-Received: by 2002:a17:90b:380c:: with SMTP id mq12mr11995991pjb.32.1639874938917; Sat, 18 Dec 2021 16:48:58 -0800 (PST) MIME-Version: 1.0 References: <20211217200418.68942-1-ffmpeg@tmm1.net> <20211217200418.68942-5-ffmpeg@tmm1.net> <20211217133822.2fa3fa09@fido6> In-Reply-To: From: Pavel Koshevoy Date: Sat, 18 Dec 2021 17:48:23 -0700 Message-ID: To: FFmpeg development discussions and patches X-Content-Filtered-By: Mailman/MimeDel 2.1.29 Subject: Re: [FFmpeg-devel] [PATCH v4 5/5] avfilter: add vf_yadif_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 Content-Type: text/plain; charset="us-ascii" Content-Transfer-Encoding: 7bit Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" Archived-At: List-Archive: List-Post: On Sat, Dec 18, 2021 at 1:02 PM Aman Karmani wrote: > On Fri, Dec 17, 2021 at 1:38 PM Philip Langdale wrote: > > > On Fri, 17 Dec 2021 12:04:18 -0800 > > Aman Karmani wrote: > > > > > From: Aman Karmani > > > > > > deinterlaces CVPixelBuffers, i.e. AV_PIX_FMT_VIDEOTOOLBOX frames > > > > > > for example, an interlaced mpeg2 video 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,yadif_videotoolbox \ > > > -c:v h264_videotoolbox \ > > > -b:v 2000k \ > > > -c:a copy \ > > > -y progressive.ts > > > > > > (note that uploading AVFrame into CVPixelBuffer via hwupload > > > requires 504c60660d3194758823ddd45ceddb86e35d806f) > > > > > > this work is sponsored by Fancy Bits LLC > > > > > > Reviewed-by: Ridley Combs > > > Signed-off-by: Aman Karmani > > > --- > > > configure | 1 + > > > libavfilter/Makefile | 4 + > > > libavfilter/allfilters.c | 1 + > > > libavfilter/metal/vf_yadif_videotoolbox.metal | 269 ++++++++++++ > > > libavfilter/vf_yadif_videotoolbox.m | 406 > > > ++++++++++++++++++ 5 files changed, 681 insertions(+) > > > create mode 100644 libavfilter/metal/vf_yadif_videotoolbox.metal > > > create mode 100644 libavfilter/vf_yadif_videotoolbox.m > > > > > > diff --git a/configure b/configure > > > index 32a39f5f5b..d8b07c8e00 100755 > > > --- a/configure > > > +++ b/configure > > > @@ -3748,6 +3748,7 @@ vpp_qsv_filter_select="qsvvpp" > > > xfade_opencl_filter_deps="opencl" > > > yadif_cuda_filter_deps="ffnvcodec" > > > yadif_cuda_filter_deps_any="cuda_nvcc cuda_llvm" > > > +yadif_videotoolbox_filter_deps="metal corevideo videotoolbox" > > > > > > # examples > > > avio_list_dir_deps="avformat avutil" > > > diff --git a/libavfilter/Makefile b/libavfilter/Makefile > > > index 2fe495df28..9a061ba3c8 100644 > > > --- a/libavfilter/Makefile > > > +++ b/libavfilter/Makefile > > > @@ -519,6 +519,10 @@ OBJS-$(CONFIG_XSTACK_FILTER) += > > > vf_stack.o framesync.o OBJS-$(CONFIG_YADIF_FILTER) > > > += vf_yadif.o yadif_common.o OBJS-$(CONFIG_YADIF_CUDA_FILTER) > > > += vf_yadif_cuda.o vf_yadif_cuda.ptx.o \ yadif_common.o > > > cuda/load_helper.o +OBJS-$(CONFIG_YADIF_VIDEOTOOLBOX_FILTER) += > > > vf_yadif_videotoolbox.o \ > > > + > > > metal/vf_yadif_videotoolbox.metallib.o \ > > > + metal/utils.o \ > > > + yadif_common.o > > > OBJS-$(CONFIG_YAEPBLUR_FILTER) += vf_yaepblur.o > > > OBJS-$(CONFIG_ZMQ_FILTER) += f_zmq.o > > > OBJS-$(CONFIG_ZOOMPAN_FILTER) += vf_zoompan.o > > > diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c > > > index ec57a2c49c..26f1c73505 100644 > > > --- a/libavfilter/allfilters.c > > > +++ b/libavfilter/allfilters.c > > > @@ -496,6 +496,7 @@ extern const AVFilter ff_vf_xmedian; > > > extern const AVFilter ff_vf_xstack; > > > extern const AVFilter ff_vf_yadif; > > > extern const AVFilter ff_vf_yadif_cuda; > > > +extern const AVFilter ff_vf_yadif_videotoolbox; > > > extern const AVFilter ff_vf_yaepblur; > > > extern const AVFilter ff_vf_zmq; > > > extern const AVFilter ff_vf_zoompan; > > > diff --git a/libavfilter/metal/vf_yadif_videotoolbox.metal > > > b/libavfilter/metal/vf_yadif_videotoolbox.metal new file mode 100644 > > > index 0000000000..50783f2ffe > > > --- /dev/null > > > +++ b/libavfilter/metal/vf_yadif_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_yadif_videotoolbox.m > > > b/libavfilter/vf_yadif_videotoolbox.m new file mode 100644 > > > index 0000000000..af83a73e89 > > > --- /dev/null > > > +++ b/libavfilter/vf_yadif_videotoolbox.m > > > @@ -0,0 +1,406 @@ > > > +/* > > > + * 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_yadif_videotoolbox_metallib_data[]; > > > +extern unsigned int ff_vf_yadif_videotoolbox_metallib_len; > > > + > > > +typedef struct YADIFVTContext { > > > + 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; > > > +} YADIFVTContext; > > > + > > > +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) > > > +{ > > > + YADIFVTContext *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) > > > +{ > > > + 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 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 yadif_videotoolbox_uninit(AVFilterContext *ctx) > > > +{ > > > + YADIFVTContext *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 av_cold int yadif_videotoolbox_init(AVFilterContext *ctx) > > > +{ > > > + 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 int config_input(AVFilterLink *inlink) > > > +{ > > > + 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_output(AVFilterLink *link) > > > +{ > > > + 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; > > > + } > > > + > > > + 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 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, "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(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, > > > + }, > > > +}; > > > + > > > +AVFilter ff_vf_yadif_videotoolbox = { > > > + .name = "yadif_videotoolbox", > > > + .description = NULL_IF_CONFIG_SMALL("YADIF for VideoToolbox > > > frames using Metal compute"), > > > + .priv_size = sizeof(YADIFVTContext), > > > + .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, > > > +}; > > > > LGTM for the general part. I'll take your work that the metal specific > > parts work as intended. > > > > Patchset applied. Thanks to everyone who helped review on and off-list. > > > It appears to break the build for me: ``` MAN doc/libavfilter.3 GEN libswresample/libswresample.ver LD libswscale/libswscale.6.dylib LD libpostproc/libpostproc.56.dylib LD libswresample/libswresample.4.dylib STRIP libavcodec/x86/vp9itxfm.o GEN libavcodec/libavcodec.ver LD libavcodec/libavcodec.59.dylib ld: warning: could not create compact unwind for _ff_cfhd_init_vlcs: stack subq instruction is too different from dwarf stack size ld: warning: could not create compact unwind for _ff_rl_init_vlc: stack subq instruction is too different from dwarf stack size LD libavformat/libavformat.59.dylib LD libavfilter/libavfilter.8.dylib clang: error: no such file or directory: 'libavfilter/metal/vf_yadif_videotoolbox.metallib.o' make: *** [libavfilter/libavfilter.8.dylib] Error 1 real 3m22.511s user 16m21.483s sys 1m43.498s ``` I initially tried --disable-metal, but that didn't work: ``` Unknown option "--disable-metal". See /Users/pavel/src/ffmpeg/configure --help for available options. real 0m0.271s user 0m0.164s sys 0m0.105s ``` I was able to work-around the linker error with --disable-filter=yadif_videotoolbox Pavel. _______________________________________________ ffmpeg-devel mailing list ffmpeg-devel@ffmpeg.org https://ffmpeg.org/mailman/listinfo/ffmpeg-devel To unsubscribe, visit link above, or email ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".