• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
1/*
2 * This file is part of FFmpeg.
3 *
4 * FFmpeg is free software; you can redistribute it and/or
5 * modify it under the terms of the GNU Lesser General Public
6 * License as published by the Free Software Foundation; either
7 * version 2.1 of the License, or (at your option) any later version.
8 *
9 * FFmpeg is distributed in the hope that it will be useful,
10 * but WITHOUT ANY WARRANTY; without even the implied warranty of
11 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
12 * Lesser General Public License for more details.
13 *
14 * You should have received a copy of the GNU Lesser General Public
15 * License along with FFmpeg; if not, write to the Free Software
16 * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
17 */
18
19#include "libavutil/log.h"
20#include "utils.h"
21
22void ff_metal_compute_encoder_dispatch(id<MTLDevice> device,
23                                       id<MTLComputePipelineState> pipeline,
24                                       id<MTLComputeCommandEncoder> encoder,
25                                       NSUInteger width, NSUInteger height)
26{
27    [encoder setComputePipelineState:pipeline];
28    NSUInteger w = pipeline.threadExecutionWidth;
29    NSUInteger h = pipeline.maxTotalThreadsPerThreadgroup / w;
30    MTLSize threadsPerThreadgroup = MTLSizeMake(w, h, 1);
31    BOOL fallback = YES;
32    // MAC_OS_X_VERSION_10_15 is only defined on SDKs new enough to include its functionality (including iOS, tvOS, etc)
33#ifdef MAC_OS_X_VERSION_10_15
34    if (@available(macOS 10.15, iOS 11, tvOS 14.5, *)) {
35        if ([device supportsFamily:MTLGPUFamilyCommon3]) {
36            MTLSize threadsPerGrid = MTLSizeMake(width, height, 1);
37            [encoder dispatchThreads:threadsPerGrid threadsPerThreadgroup:threadsPerThreadgroup];
38            fallback = NO;
39        }
40    }
41#endif
42    if (fallback) {
43        MTLSize threadgroups = MTLSizeMake((width + w - 1) / w,
44                                           (height + h - 1) / h,
45                                           1);
46        [encoder dispatchThreadgroups:threadgroups threadsPerThreadgroup:threadsPerThreadgroup];
47    }
48}
49
50CVMetalTextureRef ff_metal_texture_from_pixbuf(void *ctx,
51                                               CVMetalTextureCacheRef textureCache,
52                                               CVPixelBufferRef pixbuf,
53                                               int plane,
54                                               MTLPixelFormat format)
55{
56    CVMetalTextureRef tex = NULL;
57    CVReturn ret;
58
59    ret = CVMetalTextureCacheCreateTextureFromImage(
60        NULL,
61        textureCache,
62        pixbuf,
63        NULL,
64        format,
65        CVPixelBufferGetWidthOfPlane(pixbuf, plane),
66        CVPixelBufferGetHeightOfPlane(pixbuf, plane),
67        plane,
68        &tex
69    );
70    if (ret != kCVReturnSuccess) {
71        av_log(ctx, AV_LOG_ERROR, "Failed to create CVMetalTexture from image: %d\n", ret);
72        return NULL;
73    }
74
75    return tex;
76}
77