* [FFmpeg-devel] [PATCH v3 1/5] avfilter/vf_yadif_cuda: simplify filter definition
@ 2021-12-16 20:28 Aman Karmani
2021-12-16 20:28 ` [FFmpeg-devel] [PATCH v3 2/5] build: detect Metal.framework and build .metal files Aman Karmani
` (4 more replies)
0 siblings, 5 replies; 12+ messages in thread
From: Aman Karmani @ 2021-12-16 20:28 UTC (permalink / raw)
To: ffmpeg-devel; +Cc: philipl, Aman Karmani, kernrj
From: Aman Karmani <aman@tmm1.net>
Signed-off-by: Aman Karmani <aman@tmm1.net>
---
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,
};
--
2.33.0
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 12+ messages in thread
* [FFmpeg-devel] [PATCH v3 2/5] build: detect Metal.framework and build .metal files
2021-12-16 20:28 [FFmpeg-devel] [PATCH v3 1/5] avfilter/vf_yadif_cuda: simplify filter definition Aman Karmani
@ 2021-12-16 20:28 ` Aman Karmani
2021-12-17 8:54 ` Martin Storsjö
2021-12-16 20:28 ` [FFmpeg-devel] [PATCH v3 3/5] avutil: add obj-c helpers into header-only include Aman Karmani
` (3 subsequent siblings)
4 siblings, 1 reply; 12+ messages in thread
From: Aman Karmani @ 2021-12-16 20:28 UTC (permalink / raw)
To: ffmpeg-devel; +Cc: philipl, Aman Karmani, kernrj
From: Aman Karmani <aman@tmm1.net>
Signed-off-by: Aman Karmani <aman@tmm1.net>
---
.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)
--
2.33.0
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 12+ messages in thread
* [FFmpeg-devel] [PATCH v3 3/5] avutil: add obj-c helpers into header-only include
2021-12-16 20:28 [FFmpeg-devel] [PATCH v3 1/5] avfilter/vf_yadif_cuda: simplify filter definition Aman Karmani
2021-12-16 20:28 ` [FFmpeg-devel] [PATCH v3 2/5] build: detect Metal.framework and build .metal files Aman Karmani
@ 2021-12-16 20:28 ` Aman Karmani
2021-12-16 20:28 ` [FFmpeg-devel] [PATCH v3 4/5] avfilter: add metal utilities Aman Karmani
` (2 subsequent siblings)
4 siblings, 0 replies; 12+ messages in thread
From: Aman Karmani @ 2021-12-16 20:28 UTC (permalink / raw)
To: ffmpeg-devel; +Cc: philipl, Aman Karmani, kernrj
From: Aman Karmani <aman@tmm1.net>
Signed-off-by: Aman Karmani <aman@tmm1.net>
---
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 <Foundation/Foundation.h>
+
+inline void ff_objc_release(NSObject **obj)
+{
+ if (*obj) {
+ [*obj release];
+ *obj = nil;
+ }
+}
+
+#endif /* AVUTIL_OBJC_H */
\ No newline at end of file
--
2.33.0
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 12+ messages in thread
* [FFmpeg-devel] [PATCH v3 4/5] avfilter: add metal utilities
2021-12-16 20:28 [FFmpeg-devel] [PATCH v3 1/5] avfilter/vf_yadif_cuda: simplify filter definition Aman Karmani
2021-12-16 20:28 ` [FFmpeg-devel] [PATCH v3 2/5] build: detect Metal.framework and build .metal files Aman Karmani
2021-12-16 20:28 ` [FFmpeg-devel] [PATCH v3 3/5] avutil: add obj-c helpers into header-only include Aman Karmani
@ 2021-12-16 20:28 ` Aman Karmani
2021-12-16 22:45 ` Marvin Scholz
2021-12-16 20:28 ` [FFmpeg-devel] [PATCH v3 5/5] avfilter: add vf_deinterlace_videotoolbox Aman Karmani
2021-12-17 3:30 ` [FFmpeg-devel] [PATCH v3 1/5] avfilter/vf_yadif_cuda: simplify filter definition Philip Langdale
4 siblings, 1 reply; 12+ messages in thread
From: Aman Karmani @ 2021-12-16 20:28 UTC (permalink / raw)
To: ffmpeg-devel; +Cc: philipl, Aman Karmani, kernrj
From: Aman Karmani <aman@tmm1.net>
Signed-off-by: Aman Karmani <aman@tmm1.net>
---
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 <Metal/Metal.h>
+#include <CoreVideo/CoreVideo.h>
+
+void ff_metal_compute_encoder_dispatch(id<MTLDevice> device,
+ id<MTLComputePipelineState> pipeline,
+ id<MTLComputeCommandEncoder> 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 <libavfilter/metal/utils.h>
+
+void ff_metal_compute_encoder_dispatch(id<MTLDevice> device,
+ id<MTLComputePipelineState> pipeline,
+ id<MTLComputeCommandEncoder> 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
--
2.33.0
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 12+ messages in thread
* [FFmpeg-devel] [PATCH v3 5/5] avfilter: add vf_deinterlace_videotoolbox
2021-12-16 20:28 [FFmpeg-devel] [PATCH v3 1/5] avfilter/vf_yadif_cuda: simplify filter definition Aman Karmani
` (2 preceding siblings ...)
2021-12-16 20:28 ` [FFmpeg-devel] [PATCH v3 4/5] avfilter: add metal utilities Aman Karmani
@ 2021-12-16 20:28 ` Aman Karmani
2021-12-17 3:39 ` Philip Langdale
2021-12-17 3:30 ` [FFmpeg-devel] [PATCH v3 1/5] avfilter/vf_yadif_cuda: simplify filter definition Philip Langdale
4 siblings, 1 reply; 12+ messages in thread
From: Aman Karmani @ 2021-12-16 20:28 UTC (permalink / raw)
To: ffmpeg-devel; +Cc: philipl, Aman Karmani, kernrj
From: Aman Karmani <aman@tmm1.net>
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 <aman@tmm1.net>
---
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 <philipl@overt.org>
+ * 2020 Aman Karmani <aman@tmm1.net>
+ * 2020 Stefan Dyulgerov <stefan.dyulgerov@gmail.com>
+ *
+ * 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 <metal_stdlib>
+#include <metal_integer>
+#include <metal_texture>
+
+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 <typename T>
+T tex2D(texture2d<float, access::sample> tex, uint x, uint y)
+{
+ return tex.sample(s, float2(x, y)).x;
+}
+
+template <>
+float2 tex2D<float2>(texture2d<float, access::sample> tex, uint x, uint y)
+{
+ return tex.sample(s, float2(x, y)).xy;
+}
+
+template <typename T>
+T tex2D(texture2d<float, access::read> tex, uint x, uint y)
+{
+ return tex.read(uint2(x, y)).x;
+}
+
+template <>
+float2 tex2D<float2>(texture2d<float, access::read> tex, uint x, uint y)
+{
+ return tex.read(uint2(x, y)).xy;
+}
+
+/*
+ * YADIF helpers
+ */
+
+template<typename T>
+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<typename T>
+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>(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>(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 <typename T>
+T yadif_compute_spatial(
+ texture2d<float, accesstype> cur,
+ uint2 pos)
+{
+ // Calculate spatial prediction
+ T a = tex2D<T>(cur, pos.x - 3, pos.y - 1);
+ T b = tex2D<T>(cur, pos.x - 2, pos.y - 1);
+ T c = tex2D<T>(cur, pos.x - 1, pos.y - 1);
+ T d = tex2D<T>(cur, pos.x - 0, pos.y - 1);
+ T e = tex2D<T>(cur, pos.x + 1, pos.y - 1);
+ T f = tex2D<T>(cur, pos.x + 2, pos.y - 1);
+ T g = tex2D<T>(cur, pos.x + 3, pos.y - 1);
+
+ T h = tex2D<T>(cur, pos.x - 3, pos.y + 1);
+ T i = tex2D<T>(cur, pos.x - 2, pos.y + 1);
+ T j = tex2D<T>(cur, pos.x - 1, pos.y + 1);
+ T k = tex2D<T>(cur, pos.x - 0, pos.y + 1);
+ T l = tex2D<T>(cur, pos.x + 1, pos.y + 1);
+ T m = tex2D<T>(cur, pos.x + 2, pos.y + 1);
+ T n = tex2D<T>(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 <typename T>
+T yadif_compute_temporal(
+ texture2d<float, accesstype> cur,
+ texture2d<float, accesstype> prev2,
+ texture2d<float, accesstype> prev1,
+ texture2d<float, accesstype> next1,
+ texture2d<float, accesstype> next2,
+ T spatial_pred,
+ bool skip_spatial_check,
+ uint2 pos)
+{
+ // Calculate temporal prediction
+ T A = tex2D<T>(prev2, pos.x, pos.y - 1);
+ T B = tex2D<T>(prev2, pos.x, pos.y + 1);
+ T C = tex2D<T>(prev1, pos.x, pos.y - 2);
+ T D = tex2D<T>(prev1, pos.x, pos.y + 0);
+ T E = tex2D<T>(prev1, pos.x, pos.y + 2);
+ T F = tex2D<T>(cur, pos.x, pos.y - 1);
+ T G = tex2D<T>(cur, pos.x, pos.y + 1);
+ T H = tex2D<T>(next1, pos.x, pos.y - 2);
+ T I = tex2D<T>(next1, pos.x, pos.y + 0);
+ T J = tex2D<T>(next1, pos.x, pos.y + 2);
+ T K = tex2D<T>(next2, pos.x, pos.y - 1);
+ T L = tex2D<T>(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 <typename T>
+T yadif(
+ texture2d<float, access::write> dst,
+ texture2d<float, accesstype> prev,
+ texture2d<float, accesstype> cur,
+ texture2d<float, accesstype> next,
+ constant deintParams& params,
+ uint2 pos)
+{
+ T spatial_pred = yadif_compute_spatial<T>(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<float, access::write> dst [[texture(0)]],
+ texture2d<float, accesstype> prev [[texture(1)]],
+ texture2d<float, accesstype> cur [[texture(2)]],
+ texture2d<float, accesstype> 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<float>(dst, prev, cur, next, params, pos));
+ else
+ pred = yadif<float2>(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 <philipl@overt.org>
+ * 2020 Aman Karmani <aman@tmm1.net>
+ *
+ * This file is part of FFmpeg.
+ *
+ * FFmpeg is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2.1 of the License, or (at your option) any later version.
+ *
+ * FFmpeg is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with FFmpeg; if not, write to the Free Software
+ * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
+ */
+
+#include "internal.h"
+#include "yadif.h"
+#include <libavutil/avassert.h>
+#include <libavutil/hwcontext.h>
+#include <libavutil/objc.h>
+#include <libavfilter/metal/utils.h>
+
+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> mtlDevice;
+ id<MTLLibrary> mtlLibrary;
+ id<MTLCommandQueue> mtlQueue;
+ id<MTLComputePipelineState> mtlPipeline;
+ id<MTLFunction> mtlFunction;
+ id<MTLBuffer> 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<MTLTexture> dst,
+ id<MTLTexture> prev,
+ id<MTLTexture> cur,
+ id<MTLTexture> next,
+ int channels,
+ int parity,
+ int tff)
+{
+ DeintVTContext *s = ctx->priv;
+ id<MTLCommandBuffer> buffer = s->mtlQueue.commandBuffer;
+ id<MTLComputeCommandEncoder> encoder = buffer.computeCommandEncoder;
+ struct mtlYadifParams *params = (struct mtlYadifParams *)s->mtlParamsBuffer.contents;
+ *params = (struct mtlYadifParams){
+ .channels = channels,
+ .parity = parity,
+ .tff = tff,
+ .is_second_field = !(parity ^ tff),
+ .skip_spatial_check = s->yadif.mode&2,
+ .field_mode = s->yadif.current_field
+ };
+
+ [encoder setTexture:dst atIndex:0];
+ [encoder setTexture:prev atIndex:1];
+ [encoder setTexture:cur atIndex:2];
+ [encoder setTexture:next atIndex:3];
+ [encoder setBuffer:s->mtlParamsBuffer offset:0 atIndex:4];
+ ff_metal_compute_encoder_dispatch(s->mtlDevice, s->mtlPipeline, encoder, dst.width, dst.height);
+ [encoder endEncoding];
+
+ [buffer commit];
+ [buffer waitUntilCompleted];
+
+ ff_objc_release(&encoder);
+ ff_objc_release(&buffer);
+}
+
+static void filter(AVFilterContext *ctx, AVFrame *dst,
+ int parity, int tff)
+{
+ 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<MTLTexture> tex_prev, tex_cur, tex_next, tex_dest;
+ MTLPixelFormat format;
+
+ if (comp->plane < i) {
+ // We process planes as a whole, so don't reprocess
+ // them for additional components
+ continue;
+ }
+
+ pixel_size = (comp->depth + comp->shift) / 8;
+ channels = comp->step / pixel_size;
+ if (pixel_size > 2 || channels > 2) {
+ av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n", y->csp->name);
+ goto exit;
+ }
+ switch (pixel_size) {
+ case 1:
+ format = channels == 1 ? MTLPixelFormatR8Unorm : MTLPixelFormatRG8Unorm;
+ break;
+ case 2:
+ format = channels == 1 ? MTLPixelFormatR16Unorm : MTLPixelFormatRG16Unorm;
+ break;
+ default:
+ av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n", y->csp->name);
+ goto exit;
+ }
+ av_log(ctx, AV_LOG_TRACE,
+ "Deinterlacing plane %d: pixel_size: %d channels: %d\n",
+ comp->plane, pixel_size, channels);
+
+ prev = ff_metal_texture_from_pixbuf(ctx, s->textureCache, (CVPixelBufferRef)y->prev->data[3], i, format);
+ cur = ff_metal_texture_from_pixbuf(ctx, s->textureCache, (CVPixelBufferRef)y->cur->data[3], i, format);
+ next = ff_metal_texture_from_pixbuf(ctx, s->textureCache, (CVPixelBufferRef)y->next->data[3], i, format);
+ dest = ff_metal_texture_from_pixbuf(ctx, s->textureCache, (CVPixelBufferRef)dst->data[3], i, format);
+
+ tex_prev = CVMetalTextureGetTexture(prev);
+ tex_cur = CVMetalTextureGetTexture(cur);
+ tex_next = CVMetalTextureGetTexture(next);
+ tex_dest = CVMetalTextureGetTexture(dest);
+
+ call_kernel(ctx, tex_dest, tex_prev, tex_cur, tex_next,
+ channels, parity, tff);
+
+ CFRelease(prev);
+ CFRelease(cur);
+ CFRelease(next);
+ CFRelease(dest);
+ }
+
+ CVBufferPropagateAttachments((CVPixelBufferRef)y->cur->data[3], (CVPixelBufferRef)dst->data[3]);
+
+ if (y->current_field == YADIF_FIELD_END) {
+ y->current_field = YADIF_FIELD_NORMAL;
+ }
+
+exit:
+ return;
+}
+
+static av_cold 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,
+};
--
2.33.0
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [FFmpeg-devel] [PATCH v3 4/5] avfilter: add metal utilities
2021-12-16 20:28 ` [FFmpeg-devel] [PATCH v3 4/5] avfilter: add metal utilities Aman Karmani
@ 2021-12-16 22:45 ` Marvin Scholz
[not found] ` <CAK=uwuwsNg+Ujk0rx_FXP4KEhEzZ9qXJFKCAJkQqgJ8fkNKOHQ@mail.gmail.com>
0 siblings, 1 reply; 12+ messages in thread
From: Marvin Scholz @ 2021-12-16 22:45 UTC (permalink / raw)
To: FFmpeg development discussions and patches; +Cc: kernrj, Aman Karmani, philipl
On 16 Dec 2021, at 21:28, Aman Karmani wrote:
> From: Aman Karmani <aman@tmm1.net>
>
Thanks for your work on this! Some comments inline:
> Signed-off-by: Aman Karmani <aman@tmm1.net>
> ---
> 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 <Metal/Metal.h>
> +#include <CoreVideo/CoreVideo.h>
> +
> +void ff_metal_compute_encoder_dispatch(id<MTLDevice> device,
> + id<MTLComputePipelineState>
> pipeline,
> + id<MTLComputeCommandEncoder>
> 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 <libavfilter/metal/utils.h>
> +
> +void ff_metal_compute_encoder_dispatch(id<MTLDevice> device,
> + id<MTLComputePipelineState>
> pipeline,
> + id<MTLComputeCommandEncoder>
> 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;
> + }
> + }
Why not just do an else here instead of the fallback variable?
> + 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
Missing newline at end of file
> --
> 2.33.0
>
> _______________________________________________
> ffmpeg-devel mailing list
> ffmpeg-devel@ffmpeg.org
> https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
>
> To unsubscribe, visit link above, or email
> ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [FFmpeg-devel] [PATCH v3 4/5] avfilter: add metal utilities
[not found] ` <CAK=uwuwsNg+Ujk0rx_FXP4KEhEzZ9qXJFKCAJkQqgJ8fkNKOHQ@mail.gmail.com>
@ 2021-12-16 23:13 ` Marvin Scholz
[not found] ` <CAK=uwuxHQ28r6EOQKf7FE-QxRu2XYheLzg54S_YuLdLax=OjEA@mail.gmail.com>
0 siblings, 1 reply; 12+ messages in thread
From: Marvin Scholz @ 2021-12-16 23:13 UTC (permalink / raw)
To: Aman Karmani
Cc: Richard Kern, FFmpeg development discussions and patches, philipl
On 16 Dec 2021, at 23:53, Aman Karmani wrote:
> On Thu, Dec 16, 2021 at 2:45 PM Marvin Scholz <epirat07@gmail.com>
> wrote:
>
>> On 16 Dec 2021, at 21:28, Aman Karmani wrote:
>>
>>> From: Aman Karmani <aman@tmm1.net>
>>>
>>
>> Thanks for your work on this! Some comments inline:
>>
>>> Signed-off-by: Aman Karmani <aman@tmm1.net>
>>> ---
>>> 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 <Metal/Metal.h>
>>> +#include <CoreVideo/CoreVideo.h>
>>> +
>>> +void ff_metal_compute_encoder_dispatch(id<MTLDevice> device,
>>> + id<MTLComputePipelineState>
>>> pipeline,
>>> + id<MTLComputeCommandEncoder>
>>> 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 <libavfilter/metal/utils.h>
>>> +
>>> +void ff_metal_compute_encoder_dispatch(id<MTLDevice> device,
>>> + id<MTLComputePipelineState>
>>> pipeline,
>>> + id<MTLComputeCommandEncoder>
>>> 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;
>>> + }
>>> + }
>>
>> Why not just do an else here instead of the fallback variable?
>>
>
> Well there's two if statements, and we need to run the fallback only
> when
> both fail. So where would the else go? I would need to duplicate it
> twice.
>
Oh, is it not possible to just add the [device
supportsFamily:MTLGPUFamilyCommon3]
after the @available check? Or does that not work, just like its is not
possible to
negate a @available for some reason?
>
>>
>>> + 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
>>
>> Missing newline at end of file
>>
>
> Fixed locally, thanks.
>
>
>>
>>> --
>>> 2.33.0
>>>
>>> _______________________________________________
>>> ffmpeg-devel mailing list
>>> ffmpeg-devel@ffmpeg.org
>>> https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
>>>
>>> To unsubscribe, visit link above, or email
>>> ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
>>
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [FFmpeg-devel] [PATCH v3 4/5] avfilter: add metal utilities
[not found] ` <CAK=uwuxHQ28r6EOQKf7FE-QxRu2XYheLzg54S_YuLdLax=OjEA@mail.gmail.com>
@ 2021-12-17 1:32 ` Marvin Scholz
0 siblings, 0 replies; 12+ messages in thread
From: Marvin Scholz @ 2021-12-17 1:32 UTC (permalink / raw)
To: Aman Karmani
Cc: Richard Kern, FFmpeg development discussions and patches, philipl
On 17 Dec 2021, at 2:14, Aman Karmani wrote:
> On Thu, Dec 16, 2021 at 3:13 PM Marvin Scholz <epirat07@gmail.com> wrote:
>
>>
>>
>> On 16 Dec 2021, at 23:53, Aman Karmani wrote:
>>
>>> On Thu, Dec 16, 2021 at 2:45 PM Marvin Scholz <epirat07@gmail.com>
>>> wrote:
>>>
>>>> On 16 Dec 2021, at 21:28, Aman Karmani wrote:
>>>>
>>>>> From: Aman Karmani <aman@tmm1.net>
>>>>>
>>>>
>>>> Thanks for your work on this! Some comments inline:
>>>>
>>>>> Signed-off-by: Aman Karmani <aman@tmm1.net>
>>>>> ---
>>>>> 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 <Metal/Metal.h>
>>>>> +#include <CoreVideo/CoreVideo.h>
>>>>> +
>>>>> +void ff_metal_compute_encoder_dispatch(id<MTLDevice> device,
>>>>> + id<MTLComputePipelineState>
>>>>> pipeline,
>>>>> + id<MTLComputeCommandEncoder>
>>>>> 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 <libavfilter/metal/utils.h>
>>>>> +
>>>>> +void ff_metal_compute_encoder_dispatch(id<MTLDevice> device,
>>>>> + id<MTLComputePipelineState>
>>>>> pipeline,
>>>>> + id<MTLComputeCommandEncoder>
>>>>> 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;
>>>>> + }
>>>>> + }
>>>>
>>>> Why not just do an else here instead of the fallback variable?
>>>>
>>>
>>> Well there's two if statements, and we need to run the fallback only
>>> when
>>> both fail. So where would the else go? I would need to duplicate it
>>> twice.
>>>
>>
>> Oh, is it not possible to just add the [device
>> supportsFamily:MTLGPUFamilyCommon3]
>> after the @available check? Or does that not work, just like its is not
>> possible to
>> negate a @available for some reason?
>>
>
> Indeed, my understanding is that the only way to use it is without any
> negation or other conditionals.
>
> Here's what the compiler says when I try what you suggested:
>
> OBJCC libavfilter/metal/utils.o
> libavfilter/metal/utils.m:31:9: warning: @available does not guard
> availability here; use if (@available) instead
> [-Wunsupported-availability-guard]
> if (@available(macOS 10.15, iOS 11, tvOS 14.5, *) && [device
> supportsFamily:MTLGPUFamilyCommon3]) {
> ^
> 1 warning generated.
>
Ah, unfortunate… Nevermind then.
>
>>>
>>>>
>>>>> + 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
>>>>
>>>> Missing newline at end of file
>>>>
>>>
>>> Fixed locally, thanks.
>>>
>>>
>>>>
>>>>> --
>>>>> 2.33.0
>>>>>
>>>>> _______________________________________________
>>>>> ffmpeg-devel mailing list
>>>>> ffmpeg-devel@ffmpeg.org
>>>>> https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
>>>>>
>>>>> To unsubscribe, visit link above, or email
>>>>> ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
>>>>
>>
>>
>>
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [FFmpeg-devel] [PATCH v3 1/5] avfilter/vf_yadif_cuda: simplify filter definition
2021-12-16 20:28 [FFmpeg-devel] [PATCH v3 1/5] avfilter/vf_yadif_cuda: simplify filter definition Aman Karmani
` (3 preceding siblings ...)
2021-12-16 20:28 ` [FFmpeg-devel] [PATCH v3 5/5] avfilter: add vf_deinterlace_videotoolbox Aman Karmani
@ 2021-12-17 3:30 ` Philip Langdale
4 siblings, 0 replies; 12+ messages in thread
From: Philip Langdale @ 2021-12-17 3:30 UTC (permalink / raw)
To: ffmpeg-devel
On Thu, 16 Dec 2021 12:28:54 -0800
Aman Karmani <ffmpeg@tmm1.net> wrote:
> From: Aman Karmani <aman@tmm1.net>
>
> Signed-off-by: Aman Karmani <aman@tmm1.net>
> ---
> 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,
> };
LGTM. Thanks.
--phil
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [FFmpeg-devel] [PATCH v3 5/5] avfilter: add vf_deinterlace_videotoolbox
2021-12-16 20:28 ` [FFmpeg-devel] [PATCH v3 5/5] avfilter: add vf_deinterlace_videotoolbox Aman Karmani
@ 2021-12-17 3:39 ` Philip Langdale
0 siblings, 0 replies; 12+ messages in thread
From: Philip Langdale @ 2021-12-17 3:39 UTC (permalink / raw)
To: ffmpeg-devel
On Thu, 16 Dec 2021 12:28:58 -0800
Aman Karmani <ffmpeg@tmm1.net> wrote:
> From: Aman Karmani <aman@tmm1.net>
>
> 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 <aman@tmm1.net>
Ignoring my lack of familiarity with the Apple specific parts, this
looks reasonable, although I have the same question that Paul had about
naming. I know that you're thinking that future metal accelerated
algorithms could go in this filter, but the precedent in the codebase
is that each algorithm is a separate filter, and so if you ever wanted
to add a second one down the line, we'd be expecting it to be separate
anyway.
--phil
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [FFmpeg-devel] [PATCH v3 2/5] build: detect Metal.framework and build .metal files
2021-12-16 20:28 ` [FFmpeg-devel] [PATCH v3 2/5] build: detect Metal.framework and build .metal files Aman Karmani
@ 2021-12-17 8:54 ` Martin Storsjö
2021-12-17 19:52 ` Aman Karmani
0 siblings, 1 reply; 12+ messages in thread
From: Martin Storsjö @ 2021-12-17 8:54 UTC (permalink / raw)
To: FFmpeg development discussions and patches; +Cc: kernrj, Aman Karmani, philipl
On Thu, 16 Dec 2021, Aman Karmani wrote:
> From: Aman Karmani <aman@tmm1.net>
>
> Signed-off-by: Aman Karmani <aman@tmm1.net>
> ---
> .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 $@
Hmm, so does this try to run "xcrun metallib" instead of "xcrun metal"? I
think that can be kinda brittle, e.g. if someone wants to configure a
custom build env, where METALCC expands to e.g.
"my-wrapped-metal-compiler.sh".
I guess it feels a bit boring to need to define two separate variables,
but if it really is two separate tools, then I think that'd be the best
for clarity.
// Martin
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 12+ messages in thread
* Re: [FFmpeg-devel] [PATCH v3 2/5] build: detect Metal.framework and build .metal files
2021-12-17 8:54 ` Martin Storsjö
@ 2021-12-17 19:52 ` Aman Karmani
0 siblings, 0 replies; 12+ messages in thread
From: Aman Karmani @ 2021-12-17 19:52 UTC (permalink / raw)
To: FFmpeg development discussions and patches; +Cc: philipl, Richard Kern
On Fri, Dec 17, 2021 at 12:54 AM Martin Storsjö <martin@martin.st> wrote:
> On Thu, 16 Dec 2021, Aman Karmani wrote:
>
> > From: Aman Karmani <aman@tmm1.net>
> >
> > Signed-off-by: Aman Karmani <aman@tmm1.net>
> > ---
> > .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 $@
>
> Hmm, so does this try to run "xcrun metallib" instead of "xcrun metal"? I
> think that can be kinda brittle, e.g. if someone wants to configure a
> custom build env, where METALCC expands to e.g.
> "my-wrapped-metal-compiler.sh".
> I guess it feels a bit boring to need to define two separate variables,
> but if it really is two separate tools, then I think that'd be the best
> for clarity.
>
Good catch, I forgot about this little hack.
I agree and will split the variables.
>
> // Martin
>
> _______________________________________________
> ffmpeg-devel mailing list
> ffmpeg-devel@ffmpeg.org
> https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
>
> To unsubscribe, visit link above, or email
> ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
>
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
^ permalink raw reply [flat|nested] 12+ messages in thread
end of thread, other threads:[~2021-12-17 19:52 UTC | newest]
Thread overview: 12+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-12-16 20:28 [FFmpeg-devel] [PATCH v3 1/5] avfilter/vf_yadif_cuda: simplify filter definition Aman Karmani
2021-12-16 20:28 ` [FFmpeg-devel] [PATCH v3 2/5] build: detect Metal.framework and build .metal files Aman Karmani
2021-12-17 8:54 ` Martin Storsjö
2021-12-17 19:52 ` Aman Karmani
2021-12-16 20:28 ` [FFmpeg-devel] [PATCH v3 3/5] avutil: add obj-c helpers into header-only include Aman Karmani
2021-12-16 20:28 ` [FFmpeg-devel] [PATCH v3 4/5] avfilter: add metal utilities Aman Karmani
2021-12-16 22:45 ` Marvin Scholz
[not found] ` <CAK=uwuwsNg+Ujk0rx_FXP4KEhEzZ9qXJFKCAJkQqgJ8fkNKOHQ@mail.gmail.com>
2021-12-16 23:13 ` Marvin Scholz
[not found] ` <CAK=uwuxHQ28r6EOQKf7FE-QxRu2XYheLzg54S_YuLdLax=OjEA@mail.gmail.com>
2021-12-17 1:32 ` Marvin Scholz
2021-12-16 20:28 ` [FFmpeg-devel] [PATCH v3 5/5] avfilter: add vf_deinterlace_videotoolbox Aman Karmani
2021-12-17 3:39 ` Philip Langdale
2021-12-17 3:30 ` [FFmpeg-devel] [PATCH v3 1/5] avfilter/vf_yadif_cuda: simplify filter definition Philip Langdale
Git Inbox Mirror of the ffmpeg-devel mailing list - see https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
This inbox may be cloned and mirrored by anyone:
git clone --mirror https://master.gitmailbox.com/ffmpegdev/0 ffmpegdev/git/0.git
# If you have public-inbox 1.1+ installed, you may
# initialize and index your mirror using the following commands:
public-inbox-init -V2 ffmpegdev ffmpegdev/ https://master.gitmailbox.com/ffmpegdev \
ffmpegdev@gitmailbox.com
public-inbox-index ffmpegdev
Example config snippet for mirrors.
AGPL code for this site: git clone https://public-inbox.org/public-inbox.git