On 16 Dec 2021, at 21:28, Aman Karmani wrote:
From: Aman Karmani <a...@tmm1.net>
Thanks for your work on this! Some comments inline:
Signed-off-by: Aman Karmani <a...@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-requ...@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-requ...@ffmpeg.org with subject "unsubscribe".
