FFmpeg
utils.m
Go to the documentation of this file.
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  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 
50 CVMetalTextureRef 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 }
w
uint8_t w
Definition: llviddspenc.c:38
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:180
width
#define width
format
Filter the word “frame” indicates either a video frame or a group of audio as stored in an AVFrame structure Format for each input and each output the list of supported formats For video that means pixel format For audio that means channel sample format(the sample packing is implied by the sample format) and sample rate. The lists are not just lists
ctx
AVFormatContext * ctx
Definition: movenc.c:48
utils.h
NULL
#define NULL
Definition: coverity.c:32
ff_metal_texture_from_pixbuf
CVMetalTextureRef ff_metal_texture_from_pixbuf(void *ctx, CVMetalTextureCacheRef textureCache, CVPixelBufferRef pixbuf, int plane, MTLPixelFormat format)
Definition: utils.m:50
height
#define height
log.h
available
if no frame is available
Definition: filter_design.txt:166
ret
ret
Definition: filter_design.txt:187
ff_metal_compute_encoder_dispatch
void ff_metal_compute_encoder_dispatch(id< MTLDevice > device, id< MTLComputePipelineState > pipeline, id< MTLComputeCommandEncoder > encoder, NSUInteger width, NSUInteger height)
Definition: utils.m:22
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
h
h
Definition: vp9dsp_template.c:2038