天天看點

ffmpeg開發播放器學習筆記 - Metal 渲染YUV

作者:音視訊流媒體技術

該節是ffmpeg開發播放器學習筆記的第五節《Metal 渲染YUV》

Metal是Apple開發的适用于iOS/macOS/iPadOS平台的圖形渲染與硬體加速庫。Metal 提供對圖形處理器 (GPU) 的接近直接通路,使您能最大程度地發揮 iOS、macOS 和 Apple tvOS app 中的圖形和計算潛能。Metal 建構于易用的低開銷架構之上,而且提供預編譯的 GPU 着色器和精細的資源控制,并支援多線程處理。相對于OpenGL,它是采用了面向對象的設計更易于使用,在Apple的系統平台上可發揮更大的性能優勢。

目标

  • 了解Metal基本使用流程
  • 初始化Metal
  • 了解metal小程式
  • 了解Metal計算線程分布
  • 利用Metal渲染YUV

了解Metal基本使用流程

下面這張流程圖大緻展示了Metal的工程原理:

ffmpeg開發播放器學習筆記 - Metal 渲染YUV

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.将指令緩沖對象送出到緩沖隊列計算

一切資料準備好之後就可以将指令緩沖對象送出到指令隊列,等待指令隊列的排程并完成計算。

相關學習資料推薦,點選下方連結免費報名,先碼住不迷路~】

【免費分享】音視訊學習資料包、大廠面試題、技術視訊和學習路線圖,資料包括(C/C++,Linux,FFmpeg webRTC rtmp hls rtsp ffplay srs 等等)有需要的可以點選加群免費領取~

ffmpeg開發播放器學習筆記 - Metal 渲染YUV

初始化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多很多,這也是硬體加速的本質。 網格看上去類似這樣的:

ffmpeg開發播放器學習筆記 - Metal 渲染YUV

這裡繪制的示意圖為了更貼近本節的内容,是以繪制成了二維網格。需要注意的是,Metal還可以執行一維網格與三維網格,原理都是一樣的。

圖中藍色的部分是實際需要計算的資料,橙色的部分則是無資料的網格。在設定Metal最終計算的風格大小的時候往往會多設定一部分資料,防止某些邊緣界線的資料被遺漏了。這些無資料的網格需要在小程式裡進行過濾。

Metal在計算的時候不是以單一的網格作為計算基礎,而是将一組網格作為一個計算基礎,一次性一組網絡。它看上去是這樣的:

ffmpeg開發播放器學習筆記 - Metal 渲染YUV

這裡模拟一組網格是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的渲染

更多内容請關注微信公衆号<<程式猿搬磚>>

原文 https://juejin.cn/post/6919035400707112967

繼續閱讀