ffmpeg开发播放器学习笔记 - Metal 渲染YUV

程序猿搬砖
2,775 阅读10分钟

该节是ffmpeg开发播放器学习笔记的第五节《Metal 渲染YUV》

Metal是Apple开发的适用于iOS/macOS/iPadOS平台的图形渲染与硬件加速库。Metal 提供对图形处理器 (GPU) 的接近直接访问，使您能最大程度地发挥 iOS、macOS 和 Apple tvOS app 中的图形和计算潜能。Metal 构建于易用的低开销架构之上，而且提供预编译的 GPU 着色器和精细的资源控制，并支持多线程处理。相对于OpenGL,它是采用了面向对象的设计更易于使用,在Apple的系统平台上可发挥更大的性能优势。

image

✅ 第一节 - Hello FFmpeg
✅ 第二节 - 软解视频流,渲染 RGB24
✅ 第三节 - 认识YUV
✅ 第四节 - 硬解码,OpenGL渲染YUV
🔔 第五节 - Metal 渲染YUV
📗 第六节 - 解码音频,使用AudioQueue 播放
📗 第七节 - 音视频同步
📗 第八节 - 完善播放控制
📗 第九节 - 倍速播放
📗 第十节 - 增加视频过滤效果
📗 第十一节 - 音频变声

该节 Demo 地址: github.com/czqasngit/f… 实例代码提供了Objective-C与Swift两种实现,为了方便说明,文章引用的是Objective-C代码,因为Swift代码指针看着不简洁。
该节最终效果如下图:

image image

目标

  • 了解Metal基本使用流程
  • 初始化Metal
  • 了解metal小程序
  • 了解Metal计算线程分布
  • 利用Metal渲染YUV

了解Metal基本使用流程

下面这张流程图大致展示了Metal的工程原理: image

1.获取GPU设备实例

要使用Metal进行计算或渲染,首先需要获取到当前系统支持的GPU实例,后续所有的操作都必须建立在这个GPU实例的计算上进行。值得注意的是macOS平台可能会有多个GPU实例。

2.初始化计算管线

Metal使用.metal文件来编写小程序,它的风格类似C++。编写好的.metal文件会在编译时统一生成default.metallib资源文件。通过字符串查找到需要使用的小程序并最终生成计算管理实例。

3.创建指令队列

Metal的计算是通过计算队列实例来管理的,它的目标就是合理的调度GPU计算资源按提交的计算指令一个一个的计算。

4.创建指令缓冲对象

Metal的每一次计算都需要创建一个新的指令缓冲对象,Metal的计算目的是GPU,GPU不能直接访问内存数据,缓冲对象分配好当前计算所需要的显存,在计算时Metal直接读取显存数据进行计算。

5.将内存数据发送到显存

Metal计算所需要的数据需要在提交计算前从内存发送到显存,并在缓冲区中保存。

6.将指令缓冲对象提交到缓冲队列计算

一切数据准备好之后就可以将指令缓冲对象提交到指令队列,等待指令队列的调度并完成计算。

初始化Metal

1.获取MTLDevice

id<MTLDevice> device = MTLCreateSystemDefaultDevice();

通过上面的代码可以获取到一个最优的GPU实例,如果你需要获取所有的GPU实例你可以使用如下代码:

 NSArray<id<MTLDevice>> devices = MTLCopyAllDevices();

2.创建MTLComputePipelineState

创建MTLLibrary

id<MTLLibrary> library = [device newDefaultLibrary];

MTLLibrary实例是所有的metal小程序的集合,它可以理解成metal小程序库。

从MTLLibrary中获取MTLFunction

id<MTLFunction> function = [library newFunctionWithName:@"yuv420ToRGB"];

MTLFunction是一个具体的小程序

生成MTLComputePipelineState实例

id<MTLComputePipelineState> computePipline = [device newComputePipelineStateWithFunction:function error:&error];

指定计算管线计算时需要调用的GPU小程序函数

3.创建MTLCommandQueue

id<MTLCommandQueue> commandQueue = [device newCommandQueue];

到此,Metal使用的初始化工作就完成了

了解metal小程序

metal小程序采用了类似C++风格的编码方式,函数申明使用kernel关键字。

#include <metal_stdlib>
using namespace metal;

kernel void foo(texture2d<float, access::read> texture [[ texture(0) ]],
                constant uint2 &byteSize [[ buffer(1) ]],
                texture2d<float, access::write> outTexture [[ texture(2) ]],
                uint2 gid [[ thread_position_in_grid ]]) {
        
}

textture变量是一个texture2d二维纹理对象,它的数据格式是float。方括号里表示它在Metal框架中的对应的内存中的数据类型是texture,变量位于位置0,只有读取权限。
byteSize变量是一个uint2类型的数据,uint2即包含了两个uint的结构体类型。它对应内存中的Buffer数据类型(即一维数据),变量位于位置1,只有读取权限。
outTexture变量是一个texture2d二维纹理对象,它的数据格式是float。它在Metal框架中的对应的内存中的数据类型是texture,变量位于位置2,只有写入权限。
gid则是Metal框架计算时带入的变量,它表示当前计算的线程号,通过线程号可以获取到具体要计算的数据。这会在稍后进行更详细的说明。

以上是一个基础的函数申明,申明的只读取变量即是从内存发送到显存的数据,只写变量则是计算完成后输出到内存的数据。
更多详细的Metal Language可参考: developer.apple.com/metal/Metal…

了解Metal计算线程

Metal在执行每一个Command Buffer的时候，都将其模拟成一个网格，每一个网格都有一个单独的线程执行。 GPU执行程序的逻辑与CPU最大的不同就是并行执行大量相互不干扰的逻辑,迸发线程数量比CPU多很多,这也是硬件加速的本质。 网格看上去类似这样的:

image

这里绘制的示意图为了更贴近本节的内容,所以绘制成了二维网格。需要注意的是,Metal还可以执行一维网格与三维网格,原理都是一样的。
图中蓝色的部分是实际需要计算的数据,橙色的部分则是无数据的网格。在设定Metal最终计算的风格大小的时候往往会多设置一部分数据,防止某些边缘界线的数据被遗漏了。这些无数据的网格需要在小程序里进行过滤。

Metal在计算的时候不是以单一的网格作为计算基础,而是将一组网格作为一个计算基础，一次性一组网络。它看上去是这样的:

image

这里模拟一组网格是4x4,表明Metal在调度执行的时候一次性计算16个网格。这一组执行完了再执行下一组,因为GPU的线程并发数量也是有限的,配置一个最优的执行大小对计算速度也是有影响的。

利用Metal渲染YUV

1.初始化CVMetalTextureCacheRef

初始化CVMetalTextureCacheRef只需要执行一次

CVMetalTextureCacheRef metalTextureCache = NULL;
CVReturn ret = CVMetalTextureCacheCreate(kCFAllocatorDefault, NULL, device, NULL, &metalTextureCache);

CVMetalTextureCacheRef用于缓存后期创建CVMetalTextureRef与CVPixelBufferRef之间的映射,如果有同样的CVPixelBufferRef实例被再次创建时,可直接使用缓存的CVMetalTextureRef实例。

2.AVFrame转换成CVPixelBufferRef

需要得到用于表达Metal中纹理资源对象实例,需要将AVFrame转换成CVPixelBufferRef

- (BOOL)setupCVPixelBufferIfNeed:(AVFrame *)frame {
    if(!pixelBufferPoolRef) {
        NSMutableDictionary *pixelBufferAttributes = [[NSMutableDictionary alloc] init];
        if(frame->color_range == AVCOL_RANGE_MPEG) {
            pixelBufferAttributes[_CFToString(kCVPixelBufferPixelFormatTypeKey)] = @(kCVPixelFormatType_420YpCbCr8BiPlanarVideoRange);
        } else {
            pixelBufferAttributes[_CFToString(kCVPixelBufferPixelFormatTypeKey)] = @(kCVPixelFormatType_420YpCbCr8BiPlanarFullRange);
        }
        pixelBufferAttributes[_CFToString(kCVPixelBufferMetalCompatibilityKey)] = @(TRUE);
        pixelBufferAttributes[_CFToString(kCVPixelBufferWidthKey)] = @(frame->width);
        pixelBufferAttributes[_CFToString(kCVPixelBufferHeightKey)] = @(frame->height);
        /// bytes per row(alignment)
        pixelBufferAttributes[_CFToString(kCVPixelBufferBytesPerRowAlignmentKey)] = @(frame->linesize[0]);
//        pixelBufferAttributes[_CFToString(kCVPixelBufferIOSurfacePropertiesKey)] = @{};
        CVReturn cvRet = CVPixelBufferPoolCreate(kCFAllocatorDefault,
                                NULL,
                                (__bridge  CFDictionaryRef)pixelBufferAttributes,
                                &(self->pixelBufferPoolRef));
        if(cvRet != kCVReturnSuccess) {
            NSLog(@"create cv buffer pool failed: %d", cvRet);
            return NO;
        }
    }
    return YES;
}

if(![self setupCVPixelBufferIfNeed:frame]) return NULL;
CVPixelBufferRef _pixelBufferRef;
CVReturn cvRet = CVPixelBufferPoolCreatePixelBuffer(kCFAllocatorDefault, pixelBufferPoolRef, &_pixelBufferRef);
if(cvRet != kCVReturnSuccess) {
    NSLog(@"create cv buffer failed: %d", cvRet);
    return NULL;
}
CVPixelBufferLockBaseAddress(_pixelBufferRef, 0);
/// copy y
size_t yBytesPerRowSize = CVPixelBufferGetBytesPerRowOfPlane(_pixelBufferRef, 0);
void *yBase = CVPixelBufferGetBaseAddressOfPlane(_pixelBufferRef, 0);
memcpy(yBase, frame->data[0], yBytesPerRowSize * frame->height);
/// copy uv
void *uvBase = CVPixelBufferGetBaseAddressOfPlane(_pixelBufferRef, 1);
size_t uvBytesPerRowSize = CVPixelBufferGetBytesPerRowOfPlane(_pixelBufferRef, 1);
memcpy(uvBase, frame->data[1], uvBytesPerRowSize * frame->height / 2);
CVPixelBufferUnlockBaseAddress(_pixelBufferRef, 0);
return _pixelBufferRef;

3.编写小程序

#include <metal_stdlib>
using namespace metal;
///y_inTexture: Y
///uv_inTexture: UV
///byteSize: Y的宽高
///outTexture: RGBA
///gid: 执行线程所在的Grid位置
kernel void yuv420ToRGB(texture2d<float, access::read> y_inTexture [[ texture(0) ]],
                        texture2d<float, access::read> uv_inTexture [[ texture(1) ]],
                        constant uint2 &byteSize [[ buffer(2) ]],
                        texture2d<float, access::write> outTexture [[ texture(3) ]],
                        uint2 gid [[ thread_position_in_grid ]]) {
    /// 超出实际纹理宽高的网格不计算,直接返回
    if(gid.x > byteSize.x || gid.y > byteSize.y) return;
//    if(gid.x % 2 == 0 || gid.y % 2 == 0 || gid.x % 3 == 0 || gid.y % 3 == 0) {
//        outTexture.write(float4(0, 0, 0, 1.0), gid);
//        return;
//    }
    /// 获取y分量数据,由于在创建MetalTexture的时候在方法CVMetalTextureCacheCreateTextureFromImage
    /// 中指定了归一化的格式,所以这里得到的y值范围是[0, 1]
    float4 yFloat4 = y_inTexture.read(gid);

    /// Y与UV在YUV420P格式下的比例是4:1
    /// YUV420P垂直与水平分别是2:1的比例
    /// gid是包含X，Y坐标,所以这里gid/2实际上是缩小了4倍，符合YUV420P中Y与UV的比例
    /// 每4个Y共享一组UV
    float4 uvFloat4 = uv_inTexture.read(gid/2);
    float y = yFloat4.x;
    float cb = uvFloat4.x;
    float cr = uvFloat4.y;
    
    /// 按YCbCr转RGB的公式进行数据转换
    float r = y + 1.403 * (cr - 0.5);
    float g = y - 0.343 * (cb - 0.5) - 0.714 * (cr - 0.5);
    float b = y + 1.770 * (cb - 0.5);
    outTexture.write(float4(r, g, b, 1.0), gid);
        
}

4.获取Metal纹理资源对象MTLTexture

在Metal中所有的资源对象都是MTLResource,常用的两种类型则是MTLTexture与MTLBuffer,它们都是MTLResource子类型。MTLTexture是纹理(2维或者3维)资源对象,MTLBuffer是连续存储数据的资源对象。以下实例以是yTexture分量的MTLTexture对象:

size_t yWidth = CVPixelBufferGetWidthOfPlane(pixelBuffer, 0);
size_t yHeight = CVPixelBufferGetHeightOfPlane(pixelBuffer, 0);
CVMetalTextureRef yMetalTexture;
CVReturn ret = CVMetalTextureCacheCreateTextureFromImage(kCFAllocatorDefault,
                                                         self->metalTextureCache,
                                                         pixelBuffer,
                                                         NULL,
                                                         MTLPixelFormatR8Unorm,
                                                         yWidth,
                                                         yHeight,
                                                         0,
                                                         &yMetalTexture);

pixelBuffer是CVPixelBufferRef实例,它保存了具体的数据。
MTLPixelFormatR8Unorm指定了数据格式yTexture存储的数据是一个8位无符号整形数据,它的存储范围是[0, 255],归一化后对应的小程序取值范围是[0, 1]。在metal小程序执行时,读取MTLTexture数据一次读取一个8位作为并归一化处理。

得到CVMetalTextureRef即可获取MTLTexture实例

id<MTLTexture> yTexture = CVMetalTextureGetTexture(yMetalTexture);

由于当前CVPixelBufferRef实例中存储的是NV12格式的数据,所以除了Y分量数据,还需要创建UV分量的MTLTexture。代码如下:

size_t uvWidth = CVPixelBufferGetWidthOfPlane(pixelBuffer, 1);
size_t uvHeight = CVPixelBufferGetHeightOfPlane(pixelBuffer, 1);
CVMetalTextureRef uvMetalTexture;
ret = CVMetalTextureCacheCreateTextureFromImage(kCFAllocatorDefault,
                                                self->metalTextureCache,
                                                pixelBuffer,
                                                NULL,
                                                MTLPixelFormatRG8Unorm,
                                                uvWidth,
                                                uvHeight,
                                                1,
                                                &uvMetalTexture);
if(ret != kCVReturnSuccess) return;
id<MTLTexture> uvTexture = CVMetalTextureGetTexture(uvMetalTexture);
if(!uvTexture) return;

MTLPixelFormatRG8Unorm格式指定了由两个8位无符号整形的数据格式。在metal小程序执行时,读取MTLTexture数据一次读取两个8位作为并归一化处理。

5.创建MTLCommandBuffer

每一个指令执行都有单独的运行时数据,这些数据需要开辟GPU显存来存在

id<MTLCommandBuffer> commandBuffer = [self.commandQueue commandBuffer];

6.创建MTLComputeCommandEncoder

内存中的数据需要发送到显存,需要通过MTLComputeCommandEncoder来完成

id<MTLComputeCommandEncoder> commandEncoder = [commandBuffer computeCommandEncoder];

设置指令缓冲对象执行时的计算管线

[commandEncoder setComputePipelineState:self.computePipline];

设置小程序的纹理参数

[commandEncoder setTexture:yTexture atIndex:0];
[commandEncoder setTexture:uvTexture atIndex:1];

设置Metal执行时的实际网格大小

simd_uint2 byteSize = simd_make_uint2((uint32_t)yWidth, (uint32_t)yHeight);
[commandEncoder setBytes:&byteSize length:sizeof(simd_uint2) atIndex:2];

由于Metal只能访问MTLResource资源对象,这里设置的&byteSize相当于最终也会转化成MTLBuffer。也可以先创建MTLBuffer,配置好数据再设置到小程序对应的变量位置。

设置输出纹理对象

CAMetalLayer *layer = (CAMetalLayer *)self.layer;
id<CAMetalDrawable> drawable = [layer nextDrawable];
[commandEncoder setTexture:drawable.texture atIndex:3];

当前绘制的视图继承至MetalKit提供的MTKView,它提供了可绘制RGBA的纹理资源对象。

设置Metal执行的线程组(一次性执行多少个网格-线程)与线程组数量

NSUInteger threadExecutionWidth = self.computePipline.threadExecutionWidth;
NSUInteger maxTotalThreadsPerThreadgroup = self.computePipline.maxTotalThreadsPerThreadgroup;
MTLSize threadgroupSize = MTLSizeMake(threadExecutionWidth,
                                      maxTotalThreadsPerThreadgroup / threadExecutionWidth,
                                      1);
MTLSize threadgroupCount = MTLSizeMake((yWidth  + threadgroupSize.width -  1) / threadgroupSize.width,
                                       (yHeight + threadgroupSize.height - 1) / threadgroupSize.height,
                                       1);
[commandEncoder dispatchThreadgroups:threadgroupCount threadsPerThreadgroup:threadgroupSize];

MTLComputePipelineState提供了当前GPU实例的最大的一次性执行线程的最大宽与一次性可以执行的最大线程数。通过这两个数据可以计算出一次性执行线程组的大小,也可以自行设置大小,但不能超过。
这里将threadgroupCount设置成超出了实际纹理宽高的大小,防止边界遗漏。
纹理宽度使用Y分量的宽高,UV分量的宽高为Y分量的1/4。

提交并显示RGBA纹理

[commandEncoder endEncoding];
[commandBuffer addCompletedHandler:^(id<MTLCommandBuffer> _Nonnull buffer) {
    CVBufferRelease(yMetalTexture);
    CVBufferRelease(uvMetalTexture);
}];
[commandBuffer presentDrawable:drawable];
[commandBuffer commit];



到此,利用Metal渲染YUV就完成。

总结:

  • 了解Metal基本使用流程,Metal基于面向对象,相对于OpenGL更方便使用。Metal还提供了计算管线,可轻松利用GPU能力实现大并发无依赖的数据计算。
  • 初始化Metal使用环境,了解了MTLDevice,MTLComputePipelineState等对象的作用
  • 了解metal小程序,编写了将YUV转换成RGB的metal小程序
  • 了解Metal计算线程的执行流程与逻辑
  • 利用Metal与MetalKit完成了YUV的渲染

更多内容请关注微信公众号<<程序猿搬砖>>

avatar
文章
阅读
粉丝
相关推荐
ffmpeg使用指南
2.4k阅读
 · 
7点赞
FFmpeg技术详解
2.5k阅读
 · 
12点赞
RK3568开发笔记（十一）：开发版buildroot固件移植一个ffmpeg播放rtsp的播放器Demo
3.8k阅读
 · 
1点赞
Nodejs 第十六章（ffmpeg）
3.1k阅读
 · 
49点赞
ffmpeg 入门与疑问
1.9k阅读
 · 
28点赞