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

[Marvin Scholz]

On 16 Dec 2021, at 23:53, Aman Karmani wrote:

> On Thu, Dec 16, 2021 at 2:45 PM Marvin Scholz <epirat07 at gmail.com> 
> wrote:
>
>> On 16 Dec 2021, at 21:28, Aman Karmani wrote:
>>
>>> From: Aman Karmani <aman at tmm1.net>
>>>
>>
>> Thanks for your work on this! Some comments inline:
>>
>>> Signed-off-by: Aman Karmani <aman at 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 at ffmpeg.org
>>> https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
>>>
>>> To unsubscribe, visit link above, or email
>>> ffmpeg-devel-request at ffmpeg.org with subject "unsubscribe".
>>