On 17 Dec 2021, at 2:14, Aman Karmani wrote:
> On Thu, Dec 16, 2021 at 3:13 PM Marvin Scholz <epira...@gmail.com> wrote:
>
>>
>>
>> On 16 Dec 2021, at 23:53, Aman Karmani wrote:
>>
>>> On Thu, Dec 16, 2021 at 2:45 PM Marvin Scholz <epira...@gmail.com>
>>> wrote:
>>>
>>>> On 16 Dec 2021, at 21:28, Aman Karmani wrote:
>>>>
>>>>> From: Aman Karmani <a...@tmm1.net>
>>>>>
>>>>
>>>> Thanks for your work on this! Some comments inline:
>>>>
>>>>> Signed-off-by: Aman Karmani <a...@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-requ...@ffmpeg.org with subject "unsubscribe".
>>>>
>>
>>
>>
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-requ...@ffmpeg.org with subject "unsubscribe".
