| Message ID | 20211216202858.77643-4-ffmpeg@tmm1.net |
|---|---|
| State | New |
| Headers | show |
| Series | [FFmpeg-devel,v3,1/5] avfilter/vf_yadif_cuda: simplify filter definition | expand |
| Context | Check | Description |
|---|---|---|
| andriy/make_x86 | success | Make finished |
| andriy/make_fate_x86 | success | Make fate finished |
| andriy/make_ppc | success | Make finished |
| andriy/make_fate_ppc | success | Make fate finished |
On 16 Dec 2021, at 21:28, Aman Karmani wrote: > From: Aman Karmani <aman@tmm1.net> > Thanks for your work on this! Some comments inline: > Signed-off-by: Aman Karmani <aman@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..5df0ed600e > --- /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; > + } > + } Why not just do an else here instead of the fallback variable? > + 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; > +} > \ No newline at end of file Missing newline at end of file > -- > 2.33.0 > > _______________________________________________ > ffmpeg-devel mailing list > ffmpeg-devel@ffmpeg.org > https://ffmpeg.org/mailman/listinfo/ffmpeg-devel > > To unsubscribe, visit link above, or email > ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
On 16 Dec 2021, at 23:53, Aman Karmani wrote: > On Thu, Dec 16, 2021 at 2:45 PM Marvin Scholz <epirat07@gmail.com> > wrote: > >> On 16 Dec 2021, at 21:28, Aman Karmani wrote: >> >>> From: Aman Karmani <aman@tmm1.net> >>> >> >> Thanks for your work on this! Some comments inline: >> >>> Signed-off-by: Aman Karmani <aman@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..5df0ed600e >>> --- /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; >>> + } >>> + } >> >> Why not just do an else here instead of the fallback variable? >> > > Well there's two if statements, and we need to run the fallback only > when > both fail. So where would the else go? I would need to duplicate it > twice. > Oh, is it not possible to just add the [device supportsFamily:MTLGPUFamilyCommon3] after the @available check? Or does that not work, just like its is not possible to negate a @available for some reason? > >> >>> + 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; >>> +} >>> \ No newline at end of file >> >> Missing newline at end of file >> > > Fixed locally, thanks. > > >> >>> -- >>> 2.33.0 >>> >>> _______________________________________________ >>> ffmpeg-devel mailing list >>> ffmpeg-devel@ffmpeg.org >>> https://ffmpeg.org/mailman/listinfo/ffmpeg-devel >>> >>> To unsubscribe, visit link above, or email >>> ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe". >>
On 17 Dec 2021, at 2:14, Aman Karmani wrote: > On Thu, Dec 16, 2021 at 3:13 PM Marvin Scholz <epirat07@gmail.com> wrote: > >> >> >> On 16 Dec 2021, at 23:53, Aman Karmani wrote: >> >>> On Thu, Dec 16, 2021 at 2:45 PM Marvin Scholz <epirat07@gmail.com> >>> wrote: >>> >>>> On 16 Dec 2021, at 21:28, Aman Karmani wrote: >>>> >>>>> From: Aman Karmani <aman@tmm1.net> >>>>> >>>> >>>> Thanks for your work on this! Some comments inline: >>>> >>>>> Signed-off-by: Aman Karmani <aman@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..5df0ed600e >>>>> --- /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; >>>>> + } >>>>> + } >>>> >>>> Why not just do an else here instead of the fallback variable? >>>> >>> >>> Well there's two if statements, and we need to run the fallback only >>> when >>> both fail. So where would the else go? I would need to duplicate it >>> twice. >>> >> >> Oh, is it not possible to just add the [device >> supportsFamily:MTLGPUFamilyCommon3] >> after the @available check? Or does that not work, just like its is not >> possible to >> negate a @available for some reason? >> > > Indeed, my understanding is that the only way to use it is without any > negation or other conditionals. > > Here's what the compiler says when I try what you suggested: > > OBJCC libavfilter/metal/utils.o > libavfilter/metal/utils.m:31:9: warning: @available does not guard > availability here; use if (@available) instead > [-Wunsupported-availability-guard] > if (@available(macOS 10.15, iOS 11, tvOS 14.5, *) && [device > supportsFamily:MTLGPUFamilyCommon3]) { > ^ > 1 warning generated. > Ah, unfortunate… Nevermind then. > >>> >>>> >>>>> + 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; >>>>> +} >>>>> \ No newline at end of file >>>> >>>> Missing newline at end of file >>>> >>> >>> Fixed locally, thanks. >>> >>> >>>> >>>>> -- >>>>> 2.33.0 >>>>> >>>>> _______________________________________________ >>>>> ffmpeg-devel mailing list >>>>> ffmpeg-devel@ffmpeg.org >>>>> https://ffmpeg.org/mailman/listinfo/ffmpeg-devel >>>>> >>>>> To unsubscribe, visit link above, or email >>>>> ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe". >>>> >> >> >>
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..5df0ed600e --- /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; +} \ No newline at end of file