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