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

[Marvin Scholz]

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?

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