[FFmpeg-devel] [PATCH v4 4/5] avfilter: add metal utilities

[Aman Karmani]

From: Aman Karmani <aman at tmm1.net>

Reviewed-by: Ridley Combs <rcombs at rcombs.me>
Signed-off-by: Aman Karmani <aman at 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..759ebedfba
--- /dev/null
+++ b/libavfilter/metal/utils.m
@@ -0,0 +1,73 @@
+/*
+ * This file is part of FFmpeg.
+ *
+ * FFmpeg is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2.1 of the License, or (at your option) any later version.
+ *
+ * FFmpeg is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with FFmpeg; if not, write to the Free Software
+ * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
+ */
+
+#include "libavutil/log.h"
+#include <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;
+}
-- 
2.33.0