libavfilter/metal/utils.m

   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
  22 void 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     // MAC_OS_X_VERSION_10_15 is only defined on SDKs new enough to include its functionality (including iOS, tvOS, etc)
  32 #ifdef MAC_OS_X_VERSION_10_15
  33     if (@available(macOS 10.15, iOS 11, tvOS 14.5, *)) {
  34         if ([device supportsFamily:MTLGPUFamilyCommon3]) {
  35             MTLSize threadsPerGrid = MTLSizeMake(width, height, 1);
  36             [encoder dispatchThreads:threadsPerGrid threadsPerThreadgroup:threadsPerThreadgroup];
  37             return;
  38         }
  39     }
  40 #endif
  41
  42     // Fallback path, if we took the above one we already returned so none of this is reached
  43     {
  44         MTLSize threadgroups = MTLSizeMake((width + w - 1) / w,
  45                                            (height + h - 1) / h,
  46                                            1);
  47         [encoder dispatchThreadgroups:threadgroups threadsPerThreadgroup:threadsPerThreadgroup];
  48     }
  49 }
  50
  51 CVMetalTextureRef ff_metal_texture_from_pixbuf(void *ctx,
  52                                                CVMetalTextureCacheRef textureCache,
  53                                                CVPixelBufferRef pixbuf,
  54                                                int plane,
  55                                                MTLPixelFormat format)
  56 {
  57     CVMetalTextureRef tex = NULL;
  58     CVReturn ret;
  59
  60     ret = CVMetalTextureCacheCreateTextureFromImage(
  61         NULL,
  62         textureCache,
  63         pixbuf,
  64         NULL,
  65         format,
  66         CVPixelBufferGetWidthOfPlane(pixbuf, plane),
  67         CVPixelBufferGetHeightOfPlane(pixbuf, plane),
  68         plane,
  69         &tex
  70     );
  71     if (ret != kCVReturnSuccess) {
  72         av_log(ctx, AV_LOG_ERROR, "Failed to create CVMetalTexture from image: %d\n", ret);
  73         return NULL;
  74     }
  75
  76     return tex;
  77 }