From patchwork Fri Dec 17 20:04:14 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Aman Karmani X-Patchwork-Id: 32691 Delivered-To: ffmpegpatchwork2@gmail.com Received: by 2002:a6b:cd86:0:0:0:0:0 with SMTP id d128csp1913754iog; Fri, 17 Dec 2021 12:04:39 -0800 (PST) X-Google-Smtp-Source: ABdhPJyO1uzYVm333HlCjlrvWTjGpfdnhJlNE6tTRz781xon8gSGh4D0U6TQZzRywpaBotSHXIAq X-Received: by 2002:a17:906:49d6:: with SMTP id w22mr3803109ejv.222.1639771479226; Fri, 17 Dec 2021 12:04:39 -0800 (PST) ARC-Seal: i=1; a=rsa-sha256; t=1639771479; cv=none; d=google.com; s=arc-20160816; b=CdbvhAROikiu9oVoAJrprMOjcikc0v44IyGmkFJdt5Jnthp5XVMagXNX6BUTblVQtZ uUwH+jJU4nv+5ZvopnQA/q/7mBsZCGty9Iy4Mylk9W8x2wRasKsnvhM3wSKf3nlDMWT8 wRTgEWWZugLOYQEDsL90zjgJA1v9+zwFd+CVrpr9LhfOhlGs+NKPO8AE+uvw2fVD96xS phYZaDpq/kS2OPnhCOxkh/1ksBKIYsej24FI7OIZOyQCoA318/fnz2nZ2kIzTepzn7yQ aM5092WcZRW0MfK0yoIg7v4vyMIyxexTO1fyCVfObPrAS6GY/WSKH/MiMhnR14NyTrDU c/2w== 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:message-id:date:to:from :dkim-signature:delivered-to; bh=K8OQbyK5/oKc6qsfXhPrdEDXmqcG6Kd3/ij96ngXgfA=; b=Taf3Eo4pkBbNaglmyvGtclmUF2HE8TUMQi5DX+x40s3TM6k0tU97mNJfgqbZmhX8om JxQtGjbAU2w/LZImcgkX+JmX+l6XQDQZaPdrlZyrjnIgEuJ2eGAxrAoFN+dUyHBXyfbH yUOH2KYs0Y/3klih+flAUB6DT9ihiHbFO98ENvEXFZY50u729P/V1JsH3cW6QDKOcLF5 jWtl0cu8dZbOmeC8NPp82qn0LLZFG60pjw91gE/8VzXgemMY4hMkfymB/N1/uXVTzKFM hVnXt7dMUsWODa+TF1LxakdZ0+uDYzimmX9Xp6fq/4pqZujdTQCtE7zMjWZmqZfcqils q+OQ== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20210112 header.b=gLEBXFMa; 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 gn15si1831993ejc.316.2021.12.17.12.04.38; Fri, 17 Dec 2021 12:04:39 -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=gLEBXFMa; 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 EB29D68AEB2; Fri, 17 Dec 2021 22:04:34 +0200 (EET) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-pj1-f48.google.com (mail-pj1-f48.google.com [209.85.216.48]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id AA9CF68A7C8 for ; Fri, 17 Dec 2021 22:04:27 +0200 (EET) Received: by mail-pj1-f48.google.com with SMTP id co15so3219046pjb.2 for ; Fri, 17 Dec 2021 12:04:27 -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:mime-version :content-transfer-encoding; bh=We6ZSNJMkYdroPNtMACcU4vLizMoQgaoCtHLzZ/MUMQ=; b=gLEBXFMaDM1OnqQTHShaJEUCBzzTuzenp9XTub1IK0iNNh6GK0cappT6aASJnWmLFz 6fbCjIMsnNcK+OMpTUsnndhCSHBZniwLJbz6Bg2XEiutbC3j0vygyQPn8h5SwtD1R0pb ojfmPMUTYIz4nbeQDQbSWM70yHseQ+Is+H7vxf8vHyT+rbN1ElzbE1qJdam5F4TLAL+s uSrIC9Mtc7uQSr3v5VREIRv85tCLYSKAgKHC0ci3PTM77Ov9VuZMHpvIl0+Avj2e8LS9 koZ3Prykv4znRA3WEvIgFLanRONyAuQRVKn0aDrehMUq6Qmx7y30I2Zph3a8RQFkGEOg 9nXA== 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 :mime-version:content-transfer-encoding; bh=We6ZSNJMkYdroPNtMACcU4vLizMoQgaoCtHLzZ/MUMQ=; b=gGd5yKVAwXcCCkc46gUPVTuUzZGBjXm8NV/jhTSXIS60gryIZR9WST6F88Wxo8mtJE oCSWZidAO+Um9WVjC/B5eYoHAUjGaPHVkI7U9mvRZfwWhje8oIIPe1vlOCXL7Mvf58YO el1dgzdDV4sgn/wJ6WMDa3ZoOqhvLTAXZ31Ki56Tn+xhwUP6tk5EEdN6lOSvwVu4Xffg 2I9QxRUrkr2LeMn2Ao9dnzWM9GRNfdMMNCiuYnbTzbHUSs8WwHyNo5jd0lzgFcwpOuMM JaW5YbhAOGdBhEyDXBETI664+zSp3V8t0ab0Nai1fPxhC+M0EHeW3JqD19Z4L50W3tzz jDxQ== X-Gm-Message-State: AOAM531PFTbUlQA8zxDgblptFo0acCH0KVZDkA8U7q2p1owCivYy/Bcf 5ZAEWpWfgt5EMAV5IMZsGFf2EsngJyWxIg== X-Received: by 2002:a17:90a:6045:: with SMTP id h5mr14051848pjm.147.1639771465394; Fri, 17 Dec 2021 12:04:25 -0800 (PST) Received: from tmm1-imac.lan ([2600:8802:5501:308f:68db:dc19:cd7b:1dc7]) by smtp.gmail.com with ESMTPSA id k2sm8767301pgh.11.2021.12.17.12.04.24 (version=TLS1_2 cipher=ECDHE-ECDSA-AES128-GCM-SHA256 bits=128/128); Fri, 17 Dec 2021 12:04:24 -0800 (PST) From: Aman Karmani To: ffmpeg-devel@ffmpeg.org Date: Fri, 17 Dec 2021 12:04:14 -0800 Message-Id: <20211217200418.68942-1-ffmpeg@tmm1.net> X-Mailer: git-send-email 2.33.0 MIME-Version: 1.0 Subject: [FFmpeg-devel] [PATCH v4 1/5] avfilter/vf_yadif_cuda: simplify filter definition 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: aman@tmm1.net, Philip Langdale Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" X-TUID: l+hoxXrb5hNL From: Aman Karmani Signed-off-by: Aman Karmani Signed-off-by: Philip Langdale --- libavfilter/vf_yadif_cuda.c | 19 +------------------ 1 file changed, 1 insertion(+), 18 deletions(-) diff --git a/libavfilter/vf_yadif_cuda.c b/libavfilter/vf_yadif_cuda.c index da1ab5a8ff..685b8a2035 100644 --- a/libavfilter/vf_yadif_cuda.c +++ b/libavfilter/vf_yadif_cuda.c @@ -212,23 +212,6 @@ static av_cold void deint_cuda_uninit(AVFilterContext *ctx) s->input_frames = NULL; } -static int deint_cuda_query_formats(AVFilterContext *ctx) -{ - enum AVPixelFormat pix_fmts[] = { - AV_PIX_FMT_CUDA, AV_PIX_FMT_NONE, - }; - int ret; - - if ((ret = ff_formats_ref(ff_make_format_list(pix_fmts), - &ctx->inputs[0]->outcfg.formats)) < 0) - return ret; - if ((ret = ff_formats_ref(ff_make_format_list(pix_fmts), - &ctx->outputs[0]->incfg.formats)) < 0) - return ret; - - return 0; -} - static int config_input(AVFilterLink *inlink) { AVFilterContext *ctx = inlink->dst; @@ -380,9 +363,9 @@ const AVFilter ff_vf_yadif_cuda = { .priv_size = sizeof(DeintCUDAContext), .priv_class = &yadif_cuda_class, .uninit = deint_cuda_uninit, + FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA), FILTER_INPUTS(deint_cuda_inputs), FILTER_OUTPUTS(deint_cuda_outputs), - FILTER_QUERY_FUNC(deint_cuda_query_formats), .flags = AVFILTER_FLAG_SUPPORT_TIMELINE_INTERNAL, .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, }; From patchwork Fri Dec 17 20:04:15 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Aman Karmani X-Patchwork-Id: 32692 Delivered-To: ffmpegpatchwork2@gmail.com Received: by 2002:a6b:cd86:0:0:0:0:0 with SMTP id d128csp1913929iog; Fri, 17 Dec 2021 12:04:49 -0800 (PST) X-Google-Smtp-Source: ABdhPJyH3phzYF1Sp49GaCGQ9jat8SkpBW0/cua0IEv2lnuc1MH1Ly64H1yiHnfAxxKyl2p4LiCX X-Received: by 2002:a17:907:868e:: with SMTP id qa14mr3782530ejc.564.1639771489381; Fri, 17 Dec 2021 12:04:49 -0800 (PST) ARC-Seal: i=1; a=rsa-sha256; t=1639771489; cv=none; d=google.com; s=arc-20160816; b=LARPl7WHGLPvR1hV/z7tkON371jOjBoT9RJY/kPoO2caPEjuaVcu8njSPjt5+Otkro kXIDzNXEqFl8Hu3I+vXEt9Iwi6SFQ7OhkN+pLBI8lC4/H18iWN7IUFwNSbx/byzUaZ6S yezv/E55pB0r9IRyXK8sAUFPzxyTuCG7G7lg0UxNsWQyE2+2z5IG8ZVt3KW5e+Zd/tiE ZaJRTM0pmHRn6A0RSIbwlRaRvSUJHxBpaBO9SorGB5OgusIyMIs4u9PQTltgjRVIsGJm 4l5nLtJH4ROTSwWW0WopfwxivvEeXF73Gl3NdCdVosR3MrlLgGV6N/c86qCRz03kosGV APOw== 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=Hs5WaVDOKFTsbMEcLTL26YTKw2kXTiXgmPURF4oHIoA=; b=H9HMllPYLH/YB1XCe+Z0WvQWYT++05/a5u7XLtxexXDnYRBqEse1TWolJiHsoJaCfx RBll7PvJmxwHqUajyulr3WcTd2iX3ESarF5SwZeHWlsazyu5rBxkPbIGrT3sPOZ2kXOQ bDsxmMT1Q/r3sGmRfSZjwJcWI2Q3QZVnZs0urtk2C+jUxAePoD5gLfe+ft+1blJobTBH c9yGvlxWJWRdNjifeHuRFlfF+p+eh/TqR2bnG2JiOCH9sGFl40xTNeIUF4IdhjnaZm1q 0F/ACY9gSyOpGmAEx9KHLGqjN2xyC7gv5dnnfBt1gC8x9CowrlqfCDZaznsUXcvbk6wL FD5g== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20210112 header.b=HFIEngts; 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 ji16si7037022ejc.940.2021.12.17.12.04.49; Fri, 17 Dec 2021 12:04:49 -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=HFIEngts; 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 1801168AEE9; Fri, 17 Dec 2021 22:04:37 +0200 (EET) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-pl1-f180.google.com (mail-pl1-f180.google.com [209.85.214.180]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id B9D2E68A8B4 for ; Fri, 17 Dec 2021 22:04:28 +0200 (EET) Received: by mail-pl1-f180.google.com with SMTP id q17so2760099plr.11 for ; Fri, 17 Dec 2021 12:04:28 -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=d4KLCz82ByVviXKQuNRrjr5MvomLF+uM7sONuTAa4aM=; b=HFIEngtsGrAaY12JoSGW0CUkYfJg5qVCmLiGpxToyRbvsn5ULk1MHMEHgGU1vurOVO cavxdeNSQStw1Xwze8ivHW82XDuW1OcEzA4ns4OhYslrBzLvS5m+0ZtJeXNq+bUMWgQd xRiRx4X5yJcX2jZasZTiX3RCBS6bdknPVVnNpo+m+2wOAAv3p736wXmgobcSFhOHPd1h 1cvbrCKyUL/D4gaNdyKxrCRULtzfkTtO0/QjO+IjhGPQ6dIR/t8KNJ2pNSarqtSEDSIj rETdxPu/Y1WCVAtBPz+oipBGu5VbA8nFowVaNsfnOPck3uEIi0ieBtQmmNu5xoZO6SSO y9WQ== 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=d4KLCz82ByVviXKQuNRrjr5MvomLF+uM7sONuTAa4aM=; b=ZBYqVWz7UK4UprgXPWBeId0chr3B91a87nNN7LH2l7pY72liIdOhFT0fFEJSKSlsfa ppD3Njx5dxfiwHQklAM8H8hEc5NLsVU5xjWH0Ow3929QuDs0U+PLn7xsXuugrIYNPvD2 rKoTrK3rD5dEt88EybePWS6xG5Crgrc4j1BEkmftRuZr7sO2k+Td0clTqPqQiPjDucPJ EYQg0uTFz0SpJ8p3PBT639xLJZZFDNxPGAOjEMUvBz2n1zkdhVdit6f4stXc0m3jjRCi bCOGacSXqRN+a9Q1beO/l8/7S8auMJAZDWWxqP4XOyQcy9Tg8SRmCHLC18KmvDul4pKD w6BQ== X-Gm-Message-State: AOAM531RNBtcvpbBeon0MUYk/URLsST4dvB1EjYurfILTwlRcSsatTNb daYSP38BSSqnEVQV4o2K1HMiQg3r0+sSrw== X-Received: by 2002:a17:902:e282:b0:148:ef58:10d8 with SMTP id o2-20020a170902e28200b00148ef5810d8mr543318plc.116.1639771466504; Fri, 17 Dec 2021 12:04:26 -0800 (PST) Received: from tmm1-imac.lan ([2600:8802:5501:308f:68db:dc19:cd7b:1dc7]) by smtp.gmail.com with ESMTPSA id k2sm8767301pgh.11.2021.12.17.12.04.25 (version=TLS1_2 cipher=ECDHE-ECDSA-AES128-GCM-SHA256 bits=128/128); Fri, 17 Dec 2021 12:04:25 -0800 (PST) From: Aman Karmani To: ffmpeg-devel@ffmpeg.org Date: Fri, 17 Dec 2021 12:04:15 -0800 Message-Id: <20211217200418.68942-2-ffmpeg@tmm1.net> X-Mailer: git-send-email 2.33.0 In-Reply-To: <20211217200418.68942-1-ffmpeg@tmm1.net> References: <20211217200418.68942-1-ffmpeg@tmm1.net> MIME-Version: 1.0 Subject: [FFmpeg-devel] [PATCH v4 2/5] build: detect Metal.framework and build .metal files 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: aman@tmm1.net Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" X-TUID: U/l7Hqb6g2Pn From: Aman Karmani Reviewed-by: Ridley Combs Signed-off-by: Aman Karmani --- .gitignore | 3 +++ configure | 12 +++++++++++- ffbuild/common.mak | 9 +++++++++ 3 files changed, 23 insertions(+), 1 deletion(-) diff --git a/.gitignore b/.gitignore index 9ed24b542e..1a5bb29ad5 100644 --- a/.gitignore +++ b/.gitignore @@ -19,6 +19,9 @@ *.swp *.ver *.version +*.metal.air +*.metallib +*.metallib.c *.ptx *.ptx.c *.ptx.gz diff --git a/configure b/configure index 5fffcb8afe..32a39f5f5b 100755 --- a/configure +++ b/configure @@ -309,6 +309,7 @@ External library support: if openssl, gnutls or libtls is not used [no] --enable-mediacodec enable Android MediaCodec support [no] --enable-mediafoundation enable encoding via MediaFoundation [auto] + --disable-metal disable Apple Metal framework [autodetect] --enable-libmysofa enable libmysofa, needed for sofalizer filter [no] --enable-openal enable OpenAL 1.1 capture support [no] --enable-opencl enable OpenCL processing [no] @@ -382,6 +383,8 @@ Toolchain options: --dep-cc=DEPCC use dependency generator DEPCC [$cc_default] --nvcc=NVCC use Nvidia CUDA compiler NVCC or clang [$nvcc_default] --ld=LD use linker LD [$ld_default] + --metalcc=METALCC use metal compiler METALCC [$metalcc_default] + --metallib=METALLIB use metal linker METALLIB [$metallib_default] --pkg-config=PKGCONFIG use pkg-config tool PKGCONFIG [$pkg_config_default] --pkg-config-flags=FLAGS pass additional flags to pkgconf [] --ranlib=RANLIB use ranlib RANLIB [$ranlib_default] @@ -2564,6 +2567,8 @@ CMDLINE_SET=" ln_s logfile malloc_prefix + metalcc + metallib nm optflags nvcc @@ -3835,6 +3840,8 @@ host_cc_default="gcc" doxygen_default="doxygen" install="install" ln_s_default="ln -s -f" +metalcc_default="xcrun metal" +metallib_default="xcrun metallib" nm_default="nm -g" pkg_config_default=pkg-config ranlib_default="ranlib" @@ -4435,7 +4442,7 @@ if enabled cuda_nvcc; then fi set_default arch cc cxx doxygen pkg_config ranlib strip sysinclude \ - target_exec x86asmexe + target_exec x86asmexe metalcc metallib enabled cross_compile || host_cc_default=$cc set_default host_cc @@ -6326,6 +6333,7 @@ check_apple_framework CoreFoundation check_apple_framework CoreMedia check_apple_framework CoreVideo check_apple_framework CoreAudio +check_apple_framework Metal enabled avfoundation && { disable coregraphics applicationservices @@ -7620,6 +7628,8 @@ ARFLAGS=$arflags AR_O=$ar_o AR_CMD=$ar NM_CMD=$nm +METALCC=$metalcc +METALLIB=$metallib RANLIB=$ranlib STRIP=$strip STRIPTYPE=$striptype diff --git a/ffbuild/common.mak b/ffbuild/common.mak index 0eb831d434..e79b509425 100644 --- a/ffbuild/common.mak +++ b/ffbuild/common.mak @@ -112,6 +112,15 @@ COMPILE_LASX = $(call COMPILE,CC,LASXFLAGS) $(BIN2CEXE): ffbuild/bin2c_host.o $(HOSTLD) $(HOSTLDFLAGS) $(HOSTLD_O) $^ $(HOSTEXTRALIBS) +%.metal.air: %.metal + $(METALCC) $(patsubst $(SRC_PATH)/%,$(SRC_LINK)/%,$<) -o $@ + +%.metallib: %.metal.air + $(METALLIB) --split-module-without-linking $(patsubst $(SRC_PATH)/%,$(SRC_LINK)/%,$<) -o $@ + +%.metallib.c: %.metallib $(BIN2CEXE) + $(BIN2C) $(patsubst $(SRC_PATH)/%,$(SRC_LINK)/%,$<) $@ $(subst .,_,$(basename $(notdir $@))) + %.ptx: %.cu $(SRC_PATH)/compat/cuda/cuda_runtime.h $(COMPILE_NVCC) From patchwork Fri Dec 17 20:04:16 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Aman Karmani X-Patchwork-Id: 32693 Delivered-To: ffmpegpatchwork2@gmail.com Received: by 2002:a6b:cd86:0:0:0:0:0 with SMTP id d128csp1914121iog; Fri, 17 Dec 2021 12:05:00 -0800 (PST) X-Google-Smtp-Source: ABdhPJwDrCwVGZ+oy8QveUulWFxt/0CBQu+WGvHssMU4wKBgKLWuR7WKhtuAFALvWnNtjw26ORMx X-Received: by 2002:a17:906:5d01:: with SMTP id g1mr3730158ejt.219.1639771500724; Fri, 17 Dec 2021 12:05:00 -0800 (PST) ARC-Seal: i=1; a=rsa-sha256; t=1639771500; cv=none; d=google.com; s=arc-20160816; b=YLivbEum7QIsMVUqKnSfzPXGobwoIm6hK93cxcg/rLX7+SJjMh+nz0gIC0OFw39FEU t7UdTHIneQwOqxBipUWFtmRIFIYES9aT2opGDuoiWF3Ia0Esd7oeoehkwICTlGLejotE MLd2P2BFpnLohZ9I4wEb2CRNHrceGq1Zoubb8H66zEZ1w2nCmaG63A+OvMAVRdVlQEan ZgCSp4opMVokc0Tr+O+tJ51e6IS0oVAAuBgbTtDJgp7pnPlMayvOpHTA72kk4B7Q+ZfA TalQTzYLarm8kpXAnZA3YM/NmsyaiSZaY/eX3wJUBumUByxzw6lzI9iVa7Kwm2lFk17P +oJw== 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=p1aafFJVAVwpgXBjK/UcGjP5OQ+3N/sLx0DBkraKw0E=; b=RQUuBlpwDX6u8NY5rPzm05A2UexJCTYP2oyKxaXWDKTtFwmWSYOsE/8BB3meWNr4LC loyR8El4e9az8CW9CmUEVd3X5Np/X6ww9bhmJXOX4M3iqD92VXj42NJhEtcMP3adxNkE X0+tfgUmBplCQ6FLHzSVF/WV0YwUcNoq1pgznYu32NfvpPh/eqPSDgTGWf3RgWWwnV7Y wQ1W8PnxKvNQbw1HjYHMGW2mlGM9k48xttpDzRcUKPJCNl85bAziBWVcgpBqRyISIZDL 81rqHDLDHAQPUP0Z14m1Ow0hdwpq1XblgvJL58SaOxPikOWvfSUSu2XM+ujEEQYZKQ6Y r5aQ== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20210112 header.b=V4OeBg1G; 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 jg4si7827296ejc.916.2021.12.17.12.05.00; Fri, 17 Dec 2021 12:05:00 -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=V4OeBg1G; 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 2414168AEF7; Fri, 17 Dec 2021 22:04:38 +0200 (EET) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-pj1-f46.google.com (mail-pj1-f46.google.com [209.85.216.46]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id AC9FA68AEB2 for ; Fri, 17 Dec 2021 22:04:29 +0200 (EET) Received: by mail-pj1-f46.google.com with SMTP id a11-20020a17090a854b00b001b11aae38d6so3843243pjw.2 for ; Fri, 17 Dec 2021 12:04:29 -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=hdKYkrQ0cfZaWLOwj5ixJohCvQx8hEwe7EpjbFZdno0=; b=V4OeBg1GxnxtoO/tffs5w/Bu3qIa56zM3G+j+WwmLLdK3c673+QV4JoXL/qa6A8f+R aqD1YCRDRwjfnEr8653QgwLUlVSvb3Li944T27XxBSAGM4Q909WaesiGeBXKCZCPTzHb rF4lEs04XKEWioRhHo/jkZulOo80Y9xUzgNWNHbOgd1d9HrkDWOS1+wicS7SFPARCNjb 8fAj6dKaB5C1p8xLQS0X9wOjTSdd9UZSyYaclTfrz5lvxVzGK+hwnFeFRWW17YtPKixF KgPBK0am9ZzpZ1ATgXqNj4KhvC1PKqMc32SAkZbIE+KPCTPGdY38TX3Jgo2iFjQIC5fy 16rQ== 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=hdKYkrQ0cfZaWLOwj5ixJohCvQx8hEwe7EpjbFZdno0=; b=mqnephI1eUROEg/7yE+45nMOaY7maaLzLvC+4BFUoSGzHeaLhvTGYjxz8ehKFm/Ibq 2s5MSja1UG0MTBd7l9AHg6Mdf2xq6TiOLJm9cUIF4PPoS1+9niBRnplU3N4UXz5btCIO ruriKkB49egGjLyzKQzXuOojKTQVvza03nQfKH2CgX+nJxMx0dY9GjR65cS1DiXqeyE0 ll/smExzBXeNKukvsKahoI8rjAxx5HcD8RUoL3EqC5vhRyEA4a20pFvZp1ssVi915yCv /mQ2gEtyO67+yYW1f+BQWh1OJ9AaI/dG6afAF0KZpcKDkKBk+7IaYfI+9bY0HfrEarQp VwOQ== X-Gm-Message-State: AOAM531v6NbpLFr1ZzKhJ5s0oYxdlkP8F2AVFl8m5OlTVC7bhca+PGj3 kNO8rGcGX4kFnxDalbYFqorl6dhJMRVTrQ== X-Received: by 2002:a17:90a:2f05:: with SMTP id s5mr5567988pjd.231.1639771467660; Fri, 17 Dec 2021 12:04:27 -0800 (PST) Received: from tmm1-imac.lan ([2600:8802:5501:308f:68db:dc19:cd7b:1dc7]) by smtp.gmail.com with ESMTPSA id k2sm8767301pgh.11.2021.12.17.12.04.26 (version=TLS1_2 cipher=ECDHE-ECDSA-AES128-GCM-SHA256 bits=128/128); Fri, 17 Dec 2021 12:04:27 -0800 (PST) From: Aman Karmani To: ffmpeg-devel@ffmpeg.org Date: Fri, 17 Dec 2021 12:04:16 -0800 Message-Id: <20211217200418.68942-3-ffmpeg@tmm1.net> X-Mailer: git-send-email 2.33.0 In-Reply-To: <20211217200418.68942-1-ffmpeg@tmm1.net> References: <20211217200418.68942-1-ffmpeg@tmm1.net> MIME-Version: 1.0 Subject: [FFmpeg-devel] [PATCH v4 3/5] avutil: add obj-c helpers into header-only include 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: aman@tmm1.net Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" X-TUID: +sGdAOT/XDNR From: Aman Karmani Reviewed-by: Ridley Combs Signed-off-by: Aman Karmani --- libavutil/objc.h | 32 ++++++++++++++++++++++++++++++++ 1 file changed, 32 insertions(+) create mode 100644 libavutil/objc.h diff --git a/libavutil/objc.h b/libavutil/objc.h new file mode 100644 index 0000000000..0db993f716 --- /dev/null +++ b/libavutil/objc.h @@ -0,0 +1,32 @@ +/* + * 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 + */ + +#ifndef AVUTIL_OBJC_H +#define AVUTIL_OBJC_H + +#include + +static inline void ff_objc_release(NSObject **obj) +{ + if (*obj) { + [*obj release]; + *obj = nil; + } +} + +#endif /* AVUTIL_OBJC_H */ From patchwork Fri Dec 17 20:04:17 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Aman Karmani X-Patchwork-Id: 32694 Delivered-To: ffmpegpatchwork2@gmail.com Received: by 2002:a6b:cd86:0:0:0:0:0 with SMTP id d128csp1914289iog; Fri, 17 Dec 2021 12:05:12 -0800 (PST) X-Google-Smtp-Source: ABdhPJy8vjLGAIPIKktLPTn0Wi170j+BTgymFPycNQuspq57cKALzv7jacTn+7hhhG+qkgjqcQOc X-Received: by 2002:a17:906:58d5:: with SMTP id e21mr3950616ejs.540.1639771512330; Fri, 17 Dec 2021 12:05:12 -0800 (PST) ARC-Seal: i=1; a=rsa-sha256; t=1639771512; cv=none; d=google.com; s=arc-20160816; b=Az3PYuIH1IcoK5XJf6TMhFawR9ZEwQA5nHt82ZPMv/0eSuQWKdskYAPe5Q5vmOAzhJ 0ieXSyVjhb5Tf8d4aL5yFyvjgq4L62PRSqNkfCqrP1rpRkLOyeOJQPgV+qfmC3fWt81X FM64lM6ogjKPGD2MW01bNuw5PDA9IaHiP8C5KaAkgURYiiL3aAVp8Pbx1HWNJXfNXxci 6eWbbcy/OYvRgTZOJR5nRcE6Jm4fj8ukOjeSePgwoL+RtdktgDu+j7EclxAlGixriWS4 F1UpOrVI9x5aLKRCf5Vjk9IGMxqJ3kiaEuTRLBhJA03cz2k4p0UdDxHvynSqJ8LHUPBn dZ+Q== 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=8gfnr3SjzwJUsW5xkCdlp4vnMkSiEekyLzqFaOIa554=; b=YMN1ate6sp8HmLfOyRL+hZrfrxDqQzd2Q9vI7ZPYkn1E0Y4wQ0K+6qCLzPV48PIcUW R4cvTf2Zq+hARrXxvdfM7HV6o+ixYfm+vscanMT/nZJWsPJm1uJUza0+fT1W+Suh1N1s k7AD+SBZp0nfJz4+FzfgPeyvjXS0fKX32e8BOdYgslNxL8OnloQXmVPv1FmumLfGQRjl qgsgyt+PYbO4dqVW57y8WbqMKXO89Ru13LHzJZpg3FFOQss2X3KCKk0OrPlZonDtjaz2 /tx4sd+SqY2TZKHuu0hdl5XdMJb60SGcFWFDBTd8AIalZReTEB/8WuLsReUsiTQkRREs JzBw== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20210112 header.b=Nn+MMee+; 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 hr32si6444934ejc.956.2021.12.17.12.05.12; Fri, 17 Dec 2021 12:05:12 -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=Nn+MMee+; 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 1E45168ACAE; Fri, 17 Dec 2021 22:04:41 +0200 (EET) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-pl1-f182.google.com (mail-pl1-f182.google.com [209.85.214.182]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 1125468A8F1 for ; Fri, 17 Dec 2021 22:04:31 +0200 (EET) Received: by mail-pl1-f182.google.com with SMTP id u17so2768940plg.9 for ; Fri, 17 Dec 2021 12:04:30 -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=xiWxaPzyrnydAyOwz+8aknzAfeTSTLX6PY3H1TOJNe4=; b=Nn+MMee+BBG2+CXV6xU5oFoC1IZTAfqPBQ1VpPpZiEJ/I3BEUodQPKsxIkxnrQtwdo 8VqTW4bHPYW7NPfXcUCPRaQUPpIWvJOOexqf7CdccaI45BQpbA+2epfUjyWQ7SPFBrp1 8CB9Ir2/IN6MHzOkTPFFjmbqhjRFITmZ3AcHzPOt4eY6Ds9W7pVcCGL1K59GKARHgaAR 4bTdIhynO9gBf2sJHluX3yorupU9uwkb31Ynk9LxdIRhYexWZ5LvtGKv2O0RNk+6T7Tw nG7jRWlYkF2ep9WsZDWAzgeYhoq+6gdAj605vmNmYxTKj8Su7pHXKvEqow5gcO1wM1tP LihA== 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=xiWxaPzyrnydAyOwz+8aknzAfeTSTLX6PY3H1TOJNe4=; b=B67ZmEWBqtgA3liErvmulQhejdtCBBP3piKhVH8MnBC9uzwwSlTIZ2dc8lUZzbpB7h i3VhOJKs5qeCXLSkS5SC7+o2lv1DkRxNn0R/ehD6EPGttCqFSeQBSKKI87K+mkBu6XpY igWdNt5A9MGyAP7jV8dl1ugpiQB19bo9LuWRH3lXb49na3+Hk60fJjTagDHI7q/if45o rUIRBrdo15cqnM9pEX3rd9mQudWoNSUA2134OH82a7KRCWAISrr28umRs4WTG9B9KK3G Ss7CUHnOmk8mQUgxHKq+mYv8ceHmgxurO9BmL5qW7bm9IywkzsADAG2J1bU5eV+lvMrR tw3A== X-Gm-Message-State: AOAM531TseI2Pw43SfOSq6u2nWjdRboSx8kfzR2OB6yXDdha7ffPs0Pq P9Ji1AocV6QYVSVN94qdcg7j3iejwhO51g== X-Received: by 2002:a17:902:8e87:b0:148:a2e8:2c5f with SMTP id bg7-20020a1709028e8700b00148a2e82c5fmr4873543plb.174.1639771468955; Fri, 17 Dec 2021 12:04:28 -0800 (PST) Received: from tmm1-imac.lan ([2600:8802:5501:308f:68db:dc19:cd7b:1dc7]) by smtp.gmail.com with ESMTPSA id k2sm8767301pgh.11.2021.12.17.12.04.27 (version=TLS1_2 cipher=ECDHE-ECDSA-AES128-GCM-SHA256 bits=128/128); Fri, 17 Dec 2021 12:04:28 -0800 (PST) From: Aman Karmani To: ffmpeg-devel@ffmpeg.org Date: Fri, 17 Dec 2021 12:04:17 -0800 Message-Id: <20211217200418.68942-4-ffmpeg@tmm1.net> X-Mailer: git-send-email 2.33.0 In-Reply-To: <20211217200418.68942-1-ffmpeg@tmm1.net> References: <20211217200418.68942-1-ffmpeg@tmm1.net> MIME-Version: 1.0 Subject: [FFmpeg-devel] [PATCH v4 4/5] avfilter: add metal utilities 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: aman@tmm1.net Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" X-TUID: fPXNseau8o6a From: Aman Karmani Reviewed-by: Ridley Combs Signed-off-by: Aman Karmani --- libavfilter/metal/utils.h | 35 +++++++++++++++++++ libavfilter/metal/utils.m | 73 +++++++++++++++++++++++++++++++++++++++ 2 files changed, 108 insertions(+) create mode 100644 libavfilter/metal/utils.h create mode 100644 libavfilter/metal/utils.m diff --git a/libavfilter/metal/utils.h b/libavfilter/metal/utils.h new file mode 100644 index 0000000000..bd0319f63c --- /dev/null +++ b/libavfilter/metal/utils.h @@ -0,0 +1,35 @@ +/* + * 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 + */ + +#ifndef AVFILTER_METAL_UTILS_H +#define AVFILTER_METAL_UTILS_H + +#include +#include + +void ff_metal_compute_encoder_dispatch(id device, + id pipeline, + id encoder, + NSUInteger width, NSUInteger height); + +CVMetalTextureRef ff_metal_texture_from_pixbuf(void *avclass, + CVMetalTextureCacheRef textureCache, + CVPixelBufferRef pixbuf, + int plane, + MTLPixelFormat format); +#endif /* AVFILTER_METAL_UTILS_H */ diff --git a/libavfilter/metal/utils.m b/libavfilter/metal/utils.m new file mode 100644 index 0000000000..759ebedfba --- /dev/null +++ b/libavfilter/metal/utils.m @@ -0,0 +1,73 @@ +/* + * 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 "libavutil/log.h" +#include + +void ff_metal_compute_encoder_dispatch(id device, + id pipeline, + id encoder, + NSUInteger width, NSUInteger height) +{ + [encoder setComputePipelineState:pipeline]; + NSUInteger w = pipeline.threadExecutionWidth; + NSUInteger h = pipeline.maxTotalThreadsPerThreadgroup / w; + MTLSize threadsPerThreadgroup = MTLSizeMake(w, h, 1); + BOOL fallback = YES; + if (@available(macOS 10.15, iOS 11, tvOS 14.5, *)) { + if ([device supportsFamily:MTLGPUFamilyCommon3]) { + MTLSize threadsPerGrid = MTLSizeMake(width, height, 1); + [encoder dispatchThreads:threadsPerGrid threadsPerThreadgroup:threadsPerThreadgroup]; + fallback = NO; + } + } + if (fallback) { + MTLSize threadgroups = MTLSizeMake((width + w - 1) / w, + (height + h - 1) / h, + 1); + [encoder dispatchThreadgroups:threadgroups threadsPerThreadgroup:threadsPerThreadgroup]; + } +} + +CVMetalTextureRef ff_metal_texture_from_pixbuf(void *ctx, + CVMetalTextureCacheRef textureCache, + CVPixelBufferRef pixbuf, + int plane, + MTLPixelFormat format) +{ + CVMetalTextureRef tex = NULL; + CVReturn ret; + + ret = CVMetalTextureCacheCreateTextureFromImage( + NULL, + textureCache, + pixbuf, + NULL, + format, + CVPixelBufferGetWidthOfPlane(pixbuf, plane), + CVPixelBufferGetHeightOfPlane(pixbuf, plane), + plane, + &tex + ); + if (ret != kCVReturnSuccess) { + av_log(ctx, AV_LOG_ERROR, "Failed to create CVMetalTexture from image: %d\n", ret); + return NULL; + } + + return tex; +} From patchwork Fri Dec 17 20:04:18 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Aman Karmani X-Patchwork-Id: 32695 Delivered-To: ffmpegpatchwork2@gmail.com Received: by 2002:a6b:cd86:0:0:0:0:0 with SMTP id d128csp1914441iog; Fri, 17 Dec 2021 12:05:24 -0800 (PST) X-Google-Smtp-Source: ABdhPJyT/3RIPANh9P+57Hf3IDd5gqtSnlyWT2CypNzstFEfPOGiN4eFvqkveTcsYkrfubPOvmg2 X-Received: by 2002:a50:e102:: with SMTP id h2mr4216565edl.298.1639771524109; Fri, 17 Dec 2021 12:05:24 -0800 (PST) ARC-Seal: i=1; a=rsa-sha256; t=1639771524; cv=none; d=google.com; s=arc-20160816; b=NhqrqVZbC/8SkGVW64R5K48o85a8gjfmfvRtiHNjajCFgqfaIffDrauffhONqQFsCj 82xxrk9CNrZuduxAv7BLUQ25jDkED/J/G/DSMgtlqeZaJF3YMTSexpqECxSlOTAea6og 8v7rB+KwdDsUBRoW7uQAp1m9fkG1tcJxFbKMS6b5RlYhacyehjC1NZFzftTNlTA2VCxA IiQ5XYN9G3eBWxjLt+K3DMLaoRXyPi6HD2R8BV3Pcvp8Hql+lga5Ih/KJ95d8b7HKJyM Ls7afGFaImNMhjVLfTJ9cb/IQkrgndxQbhWbSvxI5/gQCqIkzcLJ2KH9p0jZ4X/rypYA VpQg== 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=lUbyaaGThVBkQUtU6zax3+rR8FSl2GBJJqzJmXdbibU=; b=nPfcHE8zLUfx5KibkX4s5RKaCJxwzh9QCqDztS3JDv+vhSpbu1AZG3MN0d35V2vvRs SDfQ+MSKTffmZ0GQl6ANh+HE/l8la0T+1aN2mLi4tsy99OShOgo4NPV54vY0+fOpN4U4 zIuj7Q/nL/VxqDOQD5Bo11Irj/7RWrMQAR/arlVg8CU5a5AgHksIB6MuVmWpcWyW+wzb W80ol6Ga3mS74Pcm0mcfnzlQVjRA6ClKcj33/UkpzRvkAxgyGpMqvlVKHsCvaLQFqkhw A4lyN+oCCuxkZuTwtFPM0eumfVl6CP8xcReS56o1D6yQcyMSrtSDInR/dIOwgRBePcm2 /Hsg== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20210112 header.b=Hh2raLoJ; 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 d6si4769775edx.609.2021.12.17.12.05.23; Fri, 17 Dec 2021 12:05:24 -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=Hh2raLoJ; 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 0FACC68AF04; Fri, 17 Dec 2021 22:04:42 +0200 (EET) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-pg1-f180.google.com (mail-pg1-f180.google.com [209.85.215.180]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 66A8A68AEEC for ; Fri, 17 Dec 2021 22:04:32 +0200 (EET) Received: by mail-pg1-f180.google.com with SMTP id r5so3137913pgi.6 for ; Fri, 17 Dec 2021 12:04:32 -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=1Egk+EAfK8DVV0BYc4hfUrfB5SWF1dyt0Ft0Fdy/tGo=; b=Hh2raLoJ3CW15teGH6ZUlN8O9GzGCTVZ5YnYvhm83uoI047glXkSIjuOkcw6K+jiaI Cko+yU9l7u5j0fskKF089gH/6ZeL+L/T1qjczMeOteExCiSR3PGvPXjJ8Lc6hiETNJgm LRy7OafuoPHiSVmTJfVtppen3pLS7XPawz2pqoMYWQ2d7K0RIs+mhjOzs4HB3UE34pS9 Glo/vfUwPUEMT4NNqJGZL4nfVJ4Dz+hs56g+FyJSPnsg40zHWKCN7uSHRPBeGMl5rvV/ iuZaUPGZEvQ1pJVmvJmjECeCyPCs7NhmfVLe8zvrJsfbadUzzW/rzVsYj9LhAA4ag1jW N9eg== 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=1Egk+EAfK8DVV0BYc4hfUrfB5SWF1dyt0Ft0Fdy/tGo=; b=0X6GtHolFCSZBU6FyA5Ve1BVw3TX3yrN8N6aj5v8KZXZlrqMnn0Eh4cQfmLS23CQbJ lyR8k8TZJgMuR2/Pi00JOl8Cva0bNr8rWssXnrqSoI3LDx5/MqPxefAOd/+GMt0bL102 jcbSa1G5B0/Df53NF5LZ/nohgOMhfmFb+PJpyOjaBvWnM8fxEPIu2T5CvGhTDvj5FrUe cavLynnKddu/jy7DYiIi8EaEp/qJfydlnglglJhW4vnNtdqWuFJXEKBUpjwM6A+4P57o vdhQU+50EXRJEV/qfam1hkHCOZl2JkgFm5Vt/9erm6QvShQmWWh547Qm5KXeLfsVev8A PPBw== X-Gm-Message-State: AOAM531ukyJPZ2igm/McqBzOavlVZys0NREdbAOlBsjcw0gU04xJtKjY abqZJ2J0wSXF+JM6G/ZQNjezJNQxRnQDkg== X-Received: by 2002:a63:be4e:: with SMTP id g14mr4252852pgo.194.1639771470016; Fri, 17 Dec 2021 12:04:30 -0800 (PST) Received: from tmm1-imac.lan ([2600:8802:5501:308f:68db:dc19:cd7b:1dc7]) by smtp.gmail.com with ESMTPSA id k2sm8767301pgh.11.2021.12.17.12.04.29 (version=TLS1_2 cipher=ECDHE-ECDSA-AES128-GCM-SHA256 bits=128/128); Fri, 17 Dec 2021 12:04:29 -0800 (PST) From: Aman Karmani To: ffmpeg-devel@ffmpeg.org Date: Fri, 17 Dec 2021 12:04:18 -0800 Message-Id: <20211217200418.68942-5-ffmpeg@tmm1.net> X-Mailer: git-send-email 2.33.0 In-Reply-To: <20211217200418.68942-1-ffmpeg@tmm1.net> References: <20211217200418.68942-1-ffmpeg@tmm1.net> MIME-Version: 1.0 Subject: [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 Cc: aman@tmm1.net Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" X-TUID: SUwHVRjMSpSy 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, +};