From patchwork Thu Dec 16 20:28:54 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Aman Karmani X-Patchwork-Id: 32664 Delivered-To: ffmpegpatchwork2@gmail.com Received: by 2002:a6b:cd86:0:0:0:0:0 with SMTP id d128csp832080iog; Thu, 16 Dec 2021 12:29:19 -0800 (PST) X-Google-Smtp-Source: ABdhPJzdj0sPHLIRbsXzFe6m28assPhj9jiRfAV8DEoyoJp48eizQ5uEyn0I4wbzwVAeaRiZxbWU X-Received: by 2002:a17:906:974c:: with SMTP id o12mr16491134ejy.229.1639686558910; Thu, 16 Dec 2021 12:29:18 -0800 (PST) ARC-Seal: i=1; a=rsa-sha256; t=1639686558; cv=none; d=google.com; s=arc-20160816; b=J9d3c5kqSboMYK+H7adbhVsFQbwRuyC9QH/ttSEpRx1CYhV+bpuLjiMYIGQHOwbts7 omeCT1C1urOvxrsbVzYauNl1Nb6OgyWCPEkKl3tdPK5xFRSm2ngJrO3vvXcG3q24DRLW W+f9NdpjNB0PuAD5N+1ZeytulsqSssr7SdzkvOHXWSWcoiuUeKsa61p45iJUoduCLdLg 1qRYZrj80SdogAVJt/nBUwhoG+FYdjd8FTsldKc0hqzP21cN/4hbe3oEGeZzp86tsGJO agVvyVkzu/st4u3bFiahTXStq/1MF+RWJg4SQyeuasfh/x4ocOBcKSSQdFsQ/mfxpj0j +8gw== 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=LeshDkfhvxzggJvbOZe+NlHu63T0zOqx7sFZ0a6B9jc=; b=jnV+xNr6xop7QbDoK3M7h0iurUMIzORujj2CcimI9cJbiQsvMzZvgRMheo+Oa8/Mf1 ija+Xq9x1vMpS1a4oyS3g4ZdUNKCCdGLl3NFKM9HatgYzOn6RH/HhxmOfYo2cIM4eh9J W/uBJ3NXSXh19BaVIHfUJIZcWZvehxtrMPVEZAQDuUYWu+xWa5VBTDIGaNRbVJ7u84wQ 1uXA8pwiXvzu7CwanO0hfHSmAf8/7hTsaTT7KmyN9HF8L8Vqav8ec+V2ONPOc0vLxXwu NfpMZV4nBKP5op7ZJgXMLbOCDTHTJuGDbrrW1MNZRv7qa8PjG/hoOjtCS0s2FyLQyDp1 Qnwg== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20210112 header.b=h9QBXQRU; 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 f12si4217785edx.390.2021.12.16.12.29.17; Thu, 16 Dec 2021 12:29:18 -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=h9QBXQRU; 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 30A8668AF69; Thu, 16 Dec 2021 22:29:14 +0200 (EET) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-pl1-f176.google.com (mail-pl1-f176.google.com [209.85.214.176]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 353AB68A48B for ; Thu, 16 Dec 2021 22:29:07 +0200 (EET) Received: by mail-pl1-f176.google.com with SMTP id v19so20339878plo.7 for ; Thu, 16 Dec 2021 12:29:07 -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=qElEP9uezkjiQRrVrtzMJc0TH0i6Bm7U9zCep7cv5lc=; b=h9QBXQRUnzMTjABmsmJCTOu3f1CQXnt11FfgXX5QSr/UkyZINhTqGDdd+153a0tgS3 x2HV9uzoXI7hxVdaRpoJtZ1lZB58tS8s4HSVIHXQtDu9RFkXTu8i9YFAve+R2GltNyzJ uVwymveYNl+x9OxpFQiLFA3g/2kwO7x71//IfYDzzczst4UtVQVsEnuZ9JghUL2chVk9 phRBfpzVZadzEecnLlUlhIuiDvC44JAASZg4hmDoagEvU1K7qvHMbtAHMad1aDSjssv1 xEGtKdK7lByacEaWnoHSL9O5KpFa+qWfOHll9klI1l+YJOaR+9XwzCAcH6u1Wp/ScAR8 fGcA== 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=qElEP9uezkjiQRrVrtzMJc0TH0i6Bm7U9zCep7cv5lc=; b=e+0If6sidUm+C5YaBkwTeSbDalAh3pL/X0pArkNk190DRm2zazZzgkNl5RzisXoUSH w0aV+jL9PMvg9ZVeoa1NQpNeLjeLA8biR3VrtU9/FYhleW9dg+JC8ZKuGOs6bUeiV+5k m2rl0kvuTQXgACDqsQn06RK3O6SNP7etNyTrhDdNdHNFKB1yXDh5iMHVro9N0GPntYCu sqDGVWnouGabz1ipBM3VJo2700vVMvOUMYwnO19UWdcYrSZNWG1JGI2sMWTTjz+jfj4x 7t4fXA3wcGrDOnUDjJt3dXJFQQtUu2Yo24+BRu2tqqO8/mhNFuHsc3xh9F7wLz2vqywD SG1A== X-Gm-Message-State: AOAM533XvOnUGCbfkEQgW+ML1QF8o2ERPR/S5DRJ1puGmyBnD2033wDx bTwYFUa76ctWipkoxi9Xx7hEecsygZZaFA== X-Received: by 2002:a17:90a:cb0d:: with SMTP id z13mr8044636pjt.89.1639686544528; Thu, 16 Dec 2021 12:29:04 -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.03 (version=TLS1_2 cipher=ECDHE-ECDSA-AES128-GCM-SHA256 bits=128/128); Thu, 16 Dec 2021 12:29:04 -0800 (PST) From: Aman Karmani To: ffmpeg-devel@ffmpeg.org Date: Thu, 16 Dec 2021 12:28:54 -0800 Message-Id: <20211216202858.77643-1-ffmpeg@tmm1.net> X-Mailer: git-send-email 2.33.0 MIME-Version: 1.0 Subject: [FFmpeg-devel] [PATCH v3 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: philipl@overt.org, Aman Karmani , kernrj@gmail.com Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" X-TUID: 8cpTzquoTWuU From: Aman Karmani Signed-off-by: Aman Karmani --- 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 Thu Dec 16 20:28:55 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Aman Karmani X-Patchwork-Id: 32665 Delivered-To: ffmpegpatchwork2@gmail.com Received: by 2002:a6b:cd86:0:0:0:0:0 with SMTP id d128csp832223iog; Thu, 16 Dec 2021 12:29:29 -0800 (PST) X-Google-Smtp-Source: ABdhPJywhUNIymcmyNRfge3+bNpjZS/imLVvhPfxWKUd48fR+x5jICPFZZ/RrqaVP4MfTgjr+R/3 X-Received: by 2002:a05:6402:13d7:: with SMTP id a23mr405547edx.335.1639686569316; Thu, 16 Dec 2021 12:29:29 -0800 (PST) ARC-Seal: i=1; a=rsa-sha256; t=1639686569; cv=none; d=google.com; s=arc-20160816; b=hX+pwX0scfWXHGBs75NlIX5aKB/NKxNOc64TYrKR6k6I4rZ/ybT0acSVZEGmpNeT39 V2VwKzupvDVV8Og1ok2vByyJd9NJxeKSPvYYW70CgR9Ek3tDnv3NcoE2rYVGgVEsN8lD fH732dI2WDwfpjLIgzLWCsF1Brz495ib/i4NFyjBa9Z8r/3bG54SJmoLQxL6eG33/83u MMqu8YF5F7SvGAD2/1Udsd6PLNe9cz/6P86TjBIfXL3bpboesDdMPsinNGAgwE3g+0q5 WtZ4CQGdza2mhV/ZId2eM7w6jDJOkiniTQ+BKdWFzCBKHMzdqDfouyfK36U9Ptf/IAww CbUQ== 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=Df9vXeLEhiREpdORlGw6s/Czjhh5SfFgvJF49OjpG0A=; b=o5AhZxR3aRzo344HGyealbnSllomhiChvFZy+sH4oYa8JDg+1KBoyBJw7lXWJglpMu U120JgmL67FMYDcFTpLK1u8r4ZzudRcirZQk89ELdQo7iCDXZzSNQ8chw3VxMvLg+Uwu xCMQJ15l/PdPuWLRfiuaxRJVl+B/PYwT2JmFsB31SGRIc4YOmU2U7NcPz93FiNVGLnp2 O3CI/7mHIk7VWSAZgdZE+0f6NvLc8kWR5+1hf4VchINmU/3lLWvTDplkpMr9o9mFVlzE Yx/W4/Bcr0+c2OYN+lRmJxChbJ2d/mpcZXoN1OM2cTMN2XE5FudCAs9eobYZytLpy2gY qweg== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20210112 header.b=TtpjDXB3; 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 gk1si5508210ejc.233.2021.12.16.12.29.28; Thu, 16 Dec 2021 12:29:29 -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=TtpjDXB3; 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 422CE68AF7A; Thu, 16 Dec 2021 22:29:15 +0200 (EET) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org 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 ECF6868ACC3 for ; Thu, 16 Dec 2021 22:29:07 +0200 (EET) Received: by mail-pj1-f43.google.com with SMTP id np6-20020a17090b4c4600b001a90b011e06so453695pjb.5 for ; Thu, 16 Dec 2021 12:29:07 -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=LF+4SgbncZrSlCw5IrzsN1xjU11yS2mEI6IDk07klR0=; b=TtpjDXB3+q+tPi2b1/S6piJKalaxxfKJa5FXyntMro97fonzZ6VH7/9+URUv+BdReR bC84X9jNt3tkmNBs+Nx6adSaSwawoCwH1aTaxhgaX2aaJXBa+7m1774ApZjKoHdPw2Ja LHCOMn5hqHYNezdGXi7hIsya8timO+CIwtNTQe42JslEgj9PzSyaUnnzns6+ou2RAb4D 7EPd80xha2Q5dXr50AyWqhrSXcgTwj/GXUimuhfgb9vT340P26OYzwF62jkOJLy1ADiX xS3DvlUYiRquIfas8kuGLQYajpY8gr62mRyGCFY/r48SqKYM6uncSqsZvJMbmyX7SStg 1x7w== 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=LF+4SgbncZrSlCw5IrzsN1xjU11yS2mEI6IDk07klR0=; b=xuxmHa8hbK1+xMD3XzbmKL9NwJE79he5CwwfiCiRbkQAfwA4/Gisyx6Ey1rY87/Z4b k9fT8qxLDBirUTop8IsvrYMhBsPIup/eeGbV9BCpuzvteGVA8T/2O/XqSOxSzm0UT1RI JpjFKSufeYMi6YTP0xZsspqZReyW/r//RTdGuVe9oxd0UyHmNdLVkZeLWhPoPbAp0Qda 0cik2TfSr77R1gqo5itj82ZtBqYi55K4HI5/UPA1EDN8DhYTZ99gYcw5cbZq0RAgBfCU nGwr23kIgCVHtP+hWeJj3stz+UzA+4o00n5JGKx7UPxQEvT6HtGMkGhkLi3DcGaT2EL3 UhZA== X-Gm-Message-State: AOAM532XTKmmGlc8uWJMCJRBJPnSgK7lKas20PRJe+kvH0wTEB8gPatC XKRhGjVpWRGK0BZyhv0yL6fAKGx6rMCGVg== X-Received: by 2002:a17:902:6a8a:b0:143:905f:aec7 with SMTP id n10-20020a1709026a8a00b00143905faec7mr18087062plk.8.1639686545808; Thu, 16 Dec 2021 12:29:05 -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.04 (version=TLS1_2 cipher=ECDHE-ECDSA-AES128-GCM-SHA256 bits=128/128); Thu, 16 Dec 2021 12:29:05 -0800 (PST) From: Aman Karmani To: ffmpeg-devel@ffmpeg.org Date: Thu, 16 Dec 2021 12:28:55 -0800 Message-Id: <20211216202858.77643-2-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 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: philipl@overt.org, Aman Karmani , kernrj@gmail.com Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" X-TUID: duQU7r2HhtIp From: Aman Karmani Signed-off-by: Aman Karmani --- .gitignore | 3 +++ configure | 8 +++++++- ffbuild/common.mak | 9 +++++++++ 3 files changed, 19 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..ab00b2d7cb 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,7 @@ 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] --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 +2566,7 @@ CMDLINE_SET=" ln_s logfile malloc_prefix + metalcc nm optflags nvcc @@ -3835,6 +3838,7 @@ host_cc_default="gcc" doxygen_default="doxygen" install="install" ln_s_default="ln -s -f" +metalcc_default="xcrun metal" nm_default="nm -g" pkg_config_default=pkg-config ranlib_default="ranlib" @@ -4435,7 +4439,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 enabled cross_compile || host_cc_default=$cc set_default host_cc @@ -6326,6 +6330,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 +7625,7 @@ ARFLAGS=$arflags AR_O=$ar_o AR_CMD=$ar NM_CMD=$nm +METALCC=$metalcc RANLIB=$ranlib STRIP=$strip STRIPTYPE=$striptype diff --git a/ffbuild/common.mak b/ffbuild/common.mak index 0eb831d434..05440911f4 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 + $(METALCC)lib --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 Thu Dec 16 20:28:56 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Aman Karmani X-Patchwork-Id: 32666 Delivered-To: ffmpegpatchwork2@gmail.com Received: by 2002:a6b:cd86:0:0:0:0:0 with SMTP id d128csp832361iog; Thu, 16 Dec 2021 12:29:41 -0800 (PST) X-Google-Smtp-Source: ABdhPJwbJBYb6T222M11/VHaERM9wWuEKzaOCdBKx60sSoARZZ6NQUA2d/PgPqcuhlWYccS7SVjF X-Received: by 2002:a17:906:31c2:: with SMTP id f2mr17486707ejf.341.1639686580898; Thu, 16 Dec 2021 12:29:40 -0800 (PST) ARC-Seal: i=1; a=rsa-sha256; t=1639686580; cv=none; d=google.com; s=arc-20160816; b=fAUFPuU+iIabGT2paYnBiFM41ZhQ7p+Knd7825MCyXERNVXS5MswpJx+Dgk1RjLvw0 gz2wS4BrjKqyFL7TQdmRK52za1j1v0Fpxo0w9o2D5jjNX34Jp8DOG6xwK2b+LxbW48y6 iKw0S5qgXgKxnUdwFk9emIad4VFw/oyJLN66GO9vxxRchoExgWF0FITfot3wd7LSx0Gh goEeMqu8awSsrgx3FLbL2s79Kl8vtVVAtqCHOVcqEi3mdMFckoe8d+hyHKbLVmpXlkQV Z/wIp9bmSww9O2tqnGv7p9YZ5ScKFv+K+vV8qh648oDElMRCKf0zFBmnuJ7QcJqD/yqX GeuQ== 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=eYILe//xkS8hj+kWfduqmg0rGB2DIysTEXeXiZ1vMK0=; b=PN5avAhVBkw9iRB7TUUB7El3jduHyqkOLVXQn/MAYvr+WD/c7gjfn3vaQ3eNb+O81A 5xludceLn/GqvUaqeXKqpz6PR5Nk3Pa0PT1fqbEIbGO4W6ffgHqYHpkqJKT4AzP6qf8+ mv3AJIKtV7lWDMsZhQCN1C3aS/re/AUgXmVWfqQpROiWw4m82TULltMGaBv62VidoL4P 4kOp7qGCpSESnRLg+xhvLhqzh0+FjxqYe3Y9yh0ysVxaoQ6Z279eg+RZHFpii8nJnu2h vNaSqFKHdcOAeupGYXf0scq0BRbfto/7L8MQTcvM6RcaV3KYUhIcTNsR2r8ZMDzzL+s3 SMkQ== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20210112 header.b=ej9UzV3b; 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 dd13si5074322ejc.841.2021.12.16.12.29.40; Thu, 16 Dec 2021 12:29:40 -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=ej9UzV3b; 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 3694268AF85; Thu, 16 Dec 2021 22:29:16 +0200 (EET) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-pj1-f51.google.com (mail-pj1-f51.google.com [209.85.216.51]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 2487768AF56 for ; Thu, 16 Dec 2021 22:29:09 +0200 (EET) Received: by mail-pj1-f51.google.com with SMTP id f18-20020a17090aa79200b001ad9cb23022so459978pjq.4 for ; Thu, 16 Dec 2021 12:29:09 -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=7J5gE8A6u5+d8rZTnUfnUKTgI4tI5h11RBoCpL4Ov9U=; b=ej9UzV3bmb6wA03ZyQk7PgvAHgigZxM4hHFrRP9ev87h7npcNufXxqd8r8lVuJZm/E IRGiDIlYSjuk5IdpSDaR+I7TQXlfZGeOEq8yQ3LQuayXE5pYEvejsrgpl0KnpmCrlYgi nO07vh3lwNf6Ip76F425cqSzuNX9bjZtHjXUMAXtZiK1WYo46Pi99jp22CGWrt6bUgj3 6qrtd2xhqc/tsvWZGqCoAKlEbXCoQDwc4zhs4FVw728ICLMjMre+wGGF2QblaD93WfLu 7CmbU/JVatHkmK2M9l6uJyrA4Btor6MNy6G1CnDksPKOzSYmGIyk0A2NqJ+p3WrJqgYo Rqhg== 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=7J5gE8A6u5+d8rZTnUfnUKTgI4tI5h11RBoCpL4Ov9U=; b=HRRTIHOzDp8oD/DnLY4wn6QI6HAnrpLE8cWgVVuJOOLGwl7zenGG2TVFVvXZh/zWR9 QokPCmSBXSl/FIeIIdcFiZL+Ex6UHdKF73l3qhe1thGelDRK+uGILV/e+yRIYhcc9mKY yTZwjNlGV4bBPfu8bMrcoPWbtMyXjv36u0vgFRxT4TuY96HFSQfa5QRTK9FrMv+3XymF 5XuJQKIm/+iZgtwWxAnbW8dOojwM6kKufHVWmp8iRsVc2/MPSdpC5HtMA1xu2bC1AkRM owjXEHSjy6AcjNvbyDL109LbffL9KuuRmt1IQQlchB1A5wZm1JNKIY1V5zKSr7Zv7Vzc WINA== X-Gm-Message-State: AOAM5314F5uoNE1+nJiPtSGmVjBwM3j+QBNBwvY1HSl0RVovkq2g+Tuw dEGGZHcyC0BHz+rDUUvYdR7/b50wQiHmQg== X-Received: by 2002:a17:902:e892:b0:148:a77a:d312 with SMTP id w18-20020a170902e89200b00148a77ad312mr10250893plg.95.1639686547133; Thu, 16 Dec 2021 12:29:07 -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.05 (version=TLS1_2 cipher=ECDHE-ECDSA-AES128-GCM-SHA256 bits=128/128); Thu, 16 Dec 2021 12:29:06 -0800 (PST) From: Aman Karmani To: ffmpeg-devel@ffmpeg.org Date: Thu, 16 Dec 2021 12:28:56 -0800 Message-Id: <20211216202858.77643-3-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 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: philipl@overt.org, Aman Karmani , kernrj@gmail.com Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" X-TUID: 1RN7KPTZokFm From: Aman Karmani 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..3ca1303394 --- /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 + +inline void ff_objc_release(NSObject **obj) +{ + if (*obj) { + [*obj release]; + *obj = nil; + } +} + +#endif /* AVUTIL_OBJC_H */ \ No newline at end of file From patchwork Thu Dec 16 20:28:57 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Aman Karmani X-Patchwork-Id: 32667 Delivered-To: ffmpegpatchwork2@gmail.com Received: by 2002:a6b:cd86:0:0:0:0:0 with SMTP id d128csp832517iog; Thu, 16 Dec 2021 12:29:53 -0800 (PST) X-Google-Smtp-Source: ABdhPJyQJPzKyYu5JSzCQj39/9W7CnI1ydyMGZxBJt8PGiMkQ/bIK60LLJpHUJN6dsYuXu9x/G8s X-Received: by 2002:a17:907:9152:: with SMTP id l18mr17537067ejs.382.1639686593330; Thu, 16 Dec 2021 12:29:53 -0800 (PST) ARC-Seal: i=1; a=rsa-sha256; t=1639686593; cv=none; d=google.com; s=arc-20160816; b=hpGvPxrIP60CNOdps6sAKx/k+yUs2x/UM0/+VkVbJtLIhaLcAImSJEtZKTZzoORF24 1EVIXliI25fXFXg6iahWkj0N4otJj9yba/q97WJok+H60NNAz9CRApA2shwOjhCNrSTQ bI3F2hpacZmooRKR8YE7NUDsKKoAydYO6qocqpz/3D9/D38UiiAVlvVXskg6w3Ie09EM 1lAgZrQ7lxfZTXkNnByZ1LrCqJj54rJkn9emLwK/xsWdcELWvmh8z5bBA8cw9wdgz1+p 1oNm1m+CuTA15xn5csd47T9Kk6lEt6NrsAdRQvW6KySaYotczKlurCiCgEbIE/Ktrs31 pN+w== 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=S3L6oedYddN+7ZUnsN12449CUJysfZ2qlc2/1pv3a3Y=; b=sXIYoY8CS2ArQNCjKCCMe15z3Ex0CQdwnZvvRuwzIIACMuDxLj8u5aEwsBWhIj2iM6 rdbLkt3wdOLLRXOtXeE9GL6xGIUlMdZRhhUo+gn2xkCnt/K4M3HYGLNzaYZtaVQMMtNv 0w0QxiJc1tCuDugUSjJJR9RaJmJW9h3JXQpifm7+pbXvvuvEOSmvRTg5RQitPI/taXO8 8xztOvstXhMeGwHx7IP3MGWpjmNmkabnMJ3aUysLAn3H8BQJPY1KnXmxMv+ArY1gNeIg TnBUoihLdJ8AKKhOAVF4xwDqDU2aNhPJnRiKymE3oCXS9FqwPfPyq+TZS+agXtjL7bwn 5ojg== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20210112 header.b=lh8AtaCU; 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 sb33si4641287ejc.301.2021.12.16.12.29.53; Thu, 16 Dec 2021 12:29:53 -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=lh8AtaCU; 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 1D9F568AF8A; Thu, 16 Dec 2021 22:29:19 +0200 (EET) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-pj1-f54.google.com (mail-pj1-f54.google.com [209.85.216.54]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 6451768AF7B for ; Thu, 16 Dec 2021 22:29:10 +0200 (EET) Received: by mail-pj1-f54.google.com with SMTP id v16so241292pjn.1 for ; Thu, 16 Dec 2021 12:29:10 -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=rZWMbgT2GNQAa62hMJIeAGRCGooRp8fEShV/SFqtoxE=; b=lh8AtaCUQAh3iLyjKAKI4Z7rcysLtsJG1afjXmliHRd8hjjh0ByHGZXaXvu7F0eCue 5ek50VD/Xx8/AEJtAdWuvhGd/afKT5ouN8KdNlqedBqqaOllWlmjsGA7POCSoyuHd6gt gMVFRWSUw1D6JGDkOTGyr8tpk0CJM+rfhJB4TBmJ83ufL1uj8h/jQw6RFayO8JgNYziI cgx0BcAPVSaFfpV4foaqMFa5sIafZ38Az0X1/7Pl/dqZIeRrTmpDwICWnWADZ5FMTf1C s9hQxe/0IAQ9VLrWhZ6Bi89KwyLJHwZGuKGO57SGA5QZjs3dfLtJDQqD/MjNYPwJzf/I /+hg== 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=rZWMbgT2GNQAa62hMJIeAGRCGooRp8fEShV/SFqtoxE=; b=z7X0resU6stEkDdzKH57HNGHui9BMn+K1arSYrvx7aXfv93WHF18h5fn3QI+/RMFTP DOzyhs8kwXhwYJ5yqmYPJtKDGDom/XgPV9lsAjUgNfd/wrObBr17ydOWvLPAF3ZZXx3o xeK3+eOR9hULEB15sws0pqFOegWE9xOdZuNcay8xmjfTpXY1XkaKXA6ZP0NDKSwTCdhU AXJrEezG77lGfvlYweic4vYXy/vpzNEVXqgRiQtSxNWGno2GAcSZhvUbqETvq91u0qJO fjowdlQyP/+Ctu3E8u/VWrsOKhsgaq4rsnL1Q5P4/ma4icxJKJOgCSrzfBAbOiPUwTbp lnpQ== X-Gm-Message-State: AOAM531OVlpy/qhJPzwx/nCnEjtzeWP5MNUbri30ASL1jAl4joQEINe9 W5W56XPfCbu4UKxwgrIsiFLgYbhuOpkubg== X-Received: by 2002:a17:90b:1d02:: with SMTP id on2mr8041003pjb.226.1639686548374; Thu, 16 Dec 2021 12:29:08 -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.07 (version=TLS1_2 cipher=ECDHE-ECDSA-AES128-GCM-SHA256 bits=128/128); Thu, 16 Dec 2021 12:29:07 -0800 (PST) From: Aman Karmani To: ffmpeg-devel@ffmpeg.org Date: Thu, 16 Dec 2021 12:28:57 -0800 Message-Id: <20211216202858.77643-4-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 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: philipl@overt.org, Aman Karmani , kernrj@gmail.com Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" X-TUID: BxHfITsJS7sD From: Aman Karmani 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..5df0ed600e --- /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; +} \ No newline at end of file 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, +};