Re: [FFmpeg-devel] [PATCH v3 4/5] avfilter: add metal utilities

["Marvin Scholz" <epirat07@gmail.com>]

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".
>>>>
>>
>>
>>


_______________________________________________
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".