歡迎大家前往騰訊雲+社群，獲取更多騰訊海量技術實踐乾貨哦~

本文由落影發表於雲+社群專欄

正文

本文介紹Metal和Metal Shader Language，以及Metal和OpenGL ES的差異性，也是實現入門教程的心得總結。

一、Metal

Metal 是一個和 OpenGL ES 類似的面向底層的圖形程式設計介面，可以直接操作GPU；支援iOS和OS X，提供圖形渲染和通用計算能力。（不支援模擬器）

圖片來源 https://www.invasivecode.com/weblog/metal-image-processing

MTLDevice 物件代表GPU，通常使用MTLCreateSystemDefaultDevice獲取預設的GPU； MTLCommandQueue由device建立，用於建立和組織MTLCommandBuffer，保證指令（MTLCommandBuffer）有序地傳送到GPU；MTLCommandBuffer會提供一些encoder，包括編碼繪製指令的MTLRenderCommandEncoder、編碼計算指令的MTLComputeCommandEncoder、編碼快取紋理拷貝指令的MTLBlitCommandEncoder。對於一個commandBuffer，只有呼叫encoder的結束操作，才能進行下一個encoder的建立，同時可以設定執行完指令的回撥。 每一幀都會產生一個MTLCommandBuffer物件，用於填放指令； GPUs的型別很多，每一種都有各自的接收和執行指令方式，在MTLCommandEncoder把指令進行封裝後，MTLCommandBuffer再做聚合到一次提交裡。 MTLRenderPassDescriptor 是一個輕量級的臨時物件，裡面存放較多屬性配置，供MTLCommandBuffer建立MTLRenderCommandEncoder物件用。

MTLRenderPassDescriptor 用來更方便建立MTLRenderCommandEncoder，由MetalKit的view設定屬性，並且在每幀重新整理時都會提供新的MTLRenderPassDescriptor；MTLRenderCommandEncoder在建立的時候，會隱式的呼叫一次clear的命令。 最後再呼叫present和commit介面。

Metal的viewport是3D的區域，包括寬高和近/遠平面。

深度緩衝最大值為1，最小值為0，如下面這兩個都不會顯示。

    // clipSpacePosition為深度緩衝
    out.clipSpacePosition = vector_float4(0.0, 0.0, -0.1, 1.0);
    out.clipSpacePosition = vector_float4(0.0, 0.0, 1.1, 1.0);
 

渲染管道

Metal把輸入、處理、輸出的管道看成是對指定資料的渲染指令，比如輸入頂點資料，輸出渲染後紋理。 MTLRenderPipelineState 表示渲染管道，最主要的三個過程：頂點處理、光柵化、片元處理：

轉換幾何形狀資料為幀快取中的顏色畫素，叫做點陣化（rasterizing），也叫光柵化。其實就是根據頂點的資料，檢測畫素中心是否在三角形內，確定具體哪些畫素需要渲染。 對開發者而言，頂點處理和片元處理是可程式設計的，光柵化是固定的（不可見）。 頂點函式在每個頂點被繪製時都會呼叫，比如說繪製一個三角形，會呼叫三次頂點函式。頂點處理函式返回的物件裡，必須有帶[[position]]描述符的屬性，表面這個屬性是用來計算下一步的光柵化；返回值沒有描述符的部分，則會進行插值處理。

插值處理

畫素處理是針對每一個要渲染的畫素進行處理，返回值通常是4個浮點數，表示RGBA的顏色。

在編譯的時候，Xcode會單獨編譯.metal的檔案，但不會進行連結；需要在app執行時，手動進行連結。 在包裡，可以看到default.metallib，這是對metal shader的編譯結果。

MTLFunction可以用來建立MTLRenderPipelineState物件，MTLRenderPipelineState代表的是圖形渲染的管道； 在呼叫device的newRenderPipelineStateWithDescriptor:error介面時，會進行頂點、畫素函式的連結，形成一個影象處理管道； MTLRenderPipelineDescriptor包括名稱、頂點處理函式、片元處理函式、輸出顏色格式。

setVertexBytes:length:atIndex:這介面的長度限制是4k（4096bytes），對於超過的場景應該使用MTLBuffer。MTLBuffer是GPU能夠直接讀取的記憶體，用來儲存大量的資料；（常用於頂點資料） newBufferWithLength:options:方法用來建立MTLBuffer，引數是大小和訪問方式；MTLResourceStorageModeShared是預設的訪問方式。

紋理

Metal要求所有的紋理都要符合MTLPixelFormat上面的某一種格式，每個格式都代表對影象資料的不同描述方式。 例如MTLPixelFormatBGRA8Unorm格式，記憶體佈局如下：

每個畫素有32位，分別代表BRGA。 MTLTextureDescriptor 用來設定紋理屬性，例如紋理大小和畫素格式。 MTLBuffer用於儲存頂點資料，MTLTexture則用於儲存紋理資料；MTLTexture在建立之後，需要呼叫replaceRegion:mipmapLevel:withBytes:bytesPerRow:填充紋理資料；因為影象資料一般按行進行儲存，所以需要每行的畫素大小。

[[texture(index)]] 用來描述紋理引數，比如說 samplingShader(RasterizerData in [[stage_in]], texture2d<half> colorTexture [[ texture(AAPLTextureIndexBaseColor) ]]) 在讀取紋理的時候，需要兩個引數，一個是sampler和texture coordinate，前者是取樣器，後者是紋理座標。 讀取紋理其實就把對應紋理座標的畫素顏色讀取出來。 紋理座標預設是（0，0）到（1，1），如下：

有時候，紋理的座標會超過1，取樣器會根據事前設定的mag_filter::引數進行計算。

通用計算

通用圖形計算是general-purpose GPU，簡稱GPGPU。 GPU可以用於加密、機器學習、金融等，圖形繪製和圖形計算並不是互斥的，Metal可以同時使用計算管道進行圖形計算，並且用渲染管道進行渲染。

計算管道只有一個步驟，就是kernel function（核心函式），核心函式直接讀取並寫入資源，不像渲染管道需要經過多個步驟； MTLComputePipelineState 代表一個計算處理管道，只需要一個核心函式就可以建立，相比之下，渲染管道需要頂點和片元兩個處理函式；

每次核心函式執行，都會有一個唯一的gid值； 核心函式的執行次數需要事先指定，這個次數由格子大小決定。

threadgroup 指的是設定的處理單元，這個值要根據具體的裝置進行區別，但必須是足夠小的，能讓GPU執行； threadgroupCount 是需要處理的次數，一般來說threadgroupCount*threadgroup=需要處理的大小。

效能相關

臨時物件（建立和銷燬是廉價的，它們的建立方法都返回 autoreleased物件） 1.Command Buffers 2.Command Encoders 程式碼中不需要持有。

高消耗物件（在效能相關的程式碼裡應該儘量重用它,避免反覆建立） 1.Command Queues 2.Buffers 3.Textures 5.Compute States 6.Render Pipeline States 程式碼中需長期持有。

Metal常用的四種資料型別：half、float、short(ushort)、int(uint)。 GPU的暫存器是16位，half是效能消耗最低的資料型別；float需要兩次讀取、消耗兩倍的暫存器空間、兩倍的頻寬、兩倍的電量。 為了提升效能，half和float之間的轉換由硬體來完成，不佔用任何開銷。 同時，Metal自帶的函式都是經過優化的。 在float和half資料型別混合的計算中，為了保持精度會自動將half轉成float來處理，所以如果想用half節省開銷的話，要避免和float混用。 Metal同樣不擅長處理control flow，應該儘可能使用使用三元表示式，取代簡單的if判斷。

此部分參考自WWDC

常見的圖形渲染管道

二、Metal Shader Language

Metal Shader Language的使用場景有兩個，分別是圖形渲染和通用計算；基於C++ 14，執行在GPU上，GPU的特點：頻寬大，並行處理，記憶體小，對條件語句處理較慢（等待時間長）。 Metal著色語言使用clang和 LLVM，支援過載函式，但不支援圖形渲染和通用計算入口函式的過載、遞迴函式呼叫、new和delete操作符、虛擬函式、異常處理、函式指標等，也不能用C++ 11的標準庫。

基本函式

shader有三個基本函式：

• 頂點函式（vertex），對每個頂點進行處理，生成資料並輸出到繪製管線；
• 畫素函式（fragment），對光柵化後的每個畫素點進行處理，生成資料並輸出到繪製管線；
• 通用計算函式（kernel），是平行計算的函式，其返回值型別必須為void；

頂點函式相關的修飾符：

• [[vertex_id]] vertex_id是頂點shader每次處理的index，用於定位當前的頂點
• [[instance_id]] instance_id是單個例項多次渲染時，用於表明當前索引；
• [[clip_distance]]，float 或者 float[n]， n必須是編譯時常量；
• [[point_size]]，float；
• [[position]]，float4；

如果一個頂點函式的返回值不是void，那麼返回值必須包含頂點位置； 如果返回值是float4，預設表示位置，可以不帶[[ position ]]修飾符； 如果一個頂點函式的返回值是結構體，那麼結構體必須包含“[[ position ]]”修飾的變數。

畫素函式相關的修飾符：

• [[color(m)]] float或half等，m必須是編譯時常量，表示輸入值從一個顏色attachment中讀取，m用於指定從哪個顏色attachment中讀取；
• [[front_facing]] bool，如果畫素所屬片元是正面則為true；
• [[point_coord]] float2，表示點圖元的位置，取值範圍是0.0到1.0；
• [[position]] float4，表示畫素對應的視窗相對座標(x, y, z, 1/w)；
• [[sample_id]] uint，The sample number of the sample currently being processed.
• [[sample_mask]] uint，The set of samples covered by the primitive generating the fragmentduring multisample rasterization.

以上都是輸入相關的描述符。畫素函式的返回值是單個畫素的輸出，包括一個或是多個渲染結果顏色值，一個深度值，還有一個sample遮罩，對應的輸出描述符是[[color(m)]] floatn、[[depth(depth_qualifier)]] float、[[sample_mask]] uint。

struct LYFragmentOutput {
    // color attachment 0
    float4 color_float [[color(0)]];// color attachment 1
    int4 color_int4 [[color(1)]];// color attachment 2
    uint4 color_uint4 [[color(2)]];};
fragment LYFragmentOutput fragment_shader( ... ) { ... };

需要注意，顏色attachment的引數設定要和畫素函式的輸入和輸出的資料型別匹配。

Metal支援一個功能，叫做前置深度測試（early depth testing），允許在畫素著色器執行之前執行深度測試。如果一個畫素被覆蓋，則會放棄渲染。使用方式是在fragment關鍵字前面加上[[early_fragment_tests]]： [[early_fragment_tests]] fragment float4 samplingShader(..) 使用前置深度測試的要求是不能在fragment shader對深度進行寫操作。 深度測試還不熟悉的，可以看LearnOpenGL關於深度測試的介紹。

引數的地址空間選擇

Metal種的記憶體訪問主要有兩種方式：Device模式和Constant模式，由程式碼中顯式指定。 Device模式是比較通用的訪問模式，使用限制比較少，而Constant模式是為了多次讀取而設計的快速訪問只讀模式，通過Constant記憶體模式訪問的引數的資料的位元組數量是固定的，特點總結為：

• Device支援讀寫，並且沒有size的限制；
• Constant是隻讀，並且限定大小；

如何選擇Device和Constant模式？ 先看資料size是否會變化，再看訪問的頻率高低，只有那些固定size且經常訪問的部分適合使用constant模式，其他的均用Device。

// Metal關鍵函式用到的指標引數要用地址空間修飾符（device, threadgroup, or constant） 如下
vertex RasterizerData // 返回給片元著色器的結構體
vertexShader(uint vertexID [[ vertex_id ]], // vertex_id是頂點shader每次處理的index，用於定位當前的頂點
             constant LYVertex *vertexArray [[ buffer(0) ]]); // buffer表明是快取資料，0是索引

地址空間的修飾符共有四個，device、threadgroup、constant、thread。 頂點函式（vertex）、畫素函式（fragment）、通用計算函式（kernel）的指標或引用引數，都必須帶有地址空間修飾符號。 對於頂點函式（vertex）和畫素函式（fragment），其指標或引用引數必須定義在device或是constant地址空間； 對於通用計算函式（kernel），其指標或引用引數必須定義在device或是threadgroup或是constant地址空間； void tranforms(device int *source_data, threadgroup int *dest_data, constant float *param_data) {/*...*/}; 如上使用了三種地址空間修飾符，因為有threadgroup修飾符，tranforms函式只能被通用計算函式呼叫。

constant地址空間用於從裝置記憶體池分配儲存的快取物件，是隻讀的。constant地址空間的指標或引用可以做函式的引數，向宣告為常量的變數賦值會產生編譯錯誤，宣告常量但是沒有賦予初始值也會產生編譯錯誤。 在shader中，函式之外的變數（相當於全域性變數），其地址空間必須是constant。

device地址空間用於從裝置記憶體池分配出來的快取物件，可讀也可寫。一個快取物件可以被宣告成一個標量、向量或是使用者自定義結構體的指標或是引用。快取物件使用的記憶體實際大小，應該在CPU側呼叫時就確定。 紋理物件總是在device地址空間分配記憶體，所以紋理型別可以省略修飾符。

threadgroup地址空間用於通用計算函式變數的記憶體分配，變數被一個執行緒組的所有的執行緒共享，threadgroup地址空間分配的變數不能用於圖形繪製函式。

thread地址空間用於每個執行緒內部的記憶體分配，被thread修飾的變數在其他執行緒無法訪問，在圖形繪製或是通用計算函式內宣告的變數是thread地址空間分配。 如下一段程式碼，包括device、threadgroup、thread的使用：

typedef struct
{
    half3 kRec709Luma; // position的修飾符表示這個是頂點
    
} TransParam;

kernel void
sobelKernel(texture2d<half, access::read>  sourceTexture  [[texture(LYFragmentTextureIndexTextureSource)]],
                texture2d<half, access::write> destTexture [[texture(LYFragmentTextureIndexTextureDest)]],
                uint2                          grid         [[thread_position_in_grid]],
            device TransParam *param [[buffer(0)]], // param.kRec709Luma = half3(0.2126, 0.7152, 0.0722); // 把rgba轉成亮度值
            threadgroup float3 *localBuffer [[threadgroup(0)]]) // threadgroup地址空間，這裡並沒有使用到；
{
    // 邊界保護
    if(grid.x <= destTexture.get_width() && grid.y <= destTexture.get_height())
    {
        thread half4 color  = sourceTexture.read(grid); // 初始顏色
        thread half gray   = dot(color.rgb, half3(param->kRec709Luma)); // 轉換成亮度
        destTexture.write(half4(gray, gray, gray, 1.0), grid); // 寫回對應紋理
    }
}

資料結構

Metal中常用的資料結構有向量、矩陣、原子資料型別、快取、紋理、取樣器、陣列、使用者自定義結構體。

half 是16bit是浮點數 0.5h float 是32bit的浮點數 0.5f size_t 是64bit的無符號整數 通常用於sizeof的返回值 ptrdiff_t 是64bit的有符號整數 通常用於指標的差值 half2、half3、half4、float2、float3、float4等，是向量型別，表達方式為基礎型別+向量維數。矩陣類似half4x4、half3x3、float4x4、float3x3。 double、long、long long不支援。

對於向量的訪問，比如說vec=float4(1.0f, 1.0f, 1.0f, 1.0f)，其訪問方式可以是vec[0]、vec[1]，也可以是vec.x、vec.y，也可以是vec.r、vec.g。（.xyzw和.rgba，前者對應三維座標，後者對應RGB顏色空間） 只取部分、亂序取均可，比如說我們常用到的color=texture.bgra。

資料對齊 char3、uchar3的size是4Bytes，而不是3Bytes； 類似的，int是4Bytes，但int3是16而不是12Bytes； 矩陣是由一組向量構成，按照向量的維度對齊；float3x3由3個float3向量構成，那麼每個float3的size是16Bytes； 隱式型別轉換（Implicit Type Conversions） 向量到向量或是標量的隱式轉換會導致編譯錯誤，比如int4 i; float4 f = i; // compile error，無法將一個4維的整形向量轉換為4維的浮點向量。 標量到向量的隱式轉換，是標量被賦值給向量的每一個分量。 float4 f = 2.0f; // f = (2.0f, 2.0f, 2.0f, 2.0f) 標量到矩陣、向量到矩陣的隱式轉換，矩陣到矩陣和向量及標量的隱式轉換會導致編譯錯誤。

紋理資料結構不支援指標和引用，紋理資料結構包括精度和access描述符，access修飾符描述紋理如何被訪問，有三種描述符：sample、read、write，如下：

kernel void
sobelKernel(texture2d<half, access::read>  sourceTexture  [[texture(LYFragmentTextureIndexTextureSource)]],
                texture2d<half, access::write> destTexture [[texture(LYFragmentTextureIndexTextureDest)]],
                uint2                          grid         [[thread_position_in_grid]])

Sampler是取樣器，決定如何對一個紋理進行取樣操作。定址模式，過濾模式，歸一化座標，比較函式。 在Metal程式裡初始化的取樣器必須使用constexpr修飾符宣告。 取樣器指標和引用是不支援的，將會導致編譯錯誤。

    constexpr sampler textureSampler (mag_filter::linear,
                                      min_filter::linear); // sampler是取樣器

運算子

• 矩陣相乘有一個運算元是標量，那麼這個標量和矩陣中的每一個元素相乘，得到一個和矩陣有相同行列的新矩陣。
• 右運算元是一個向量，那麼它被看做一個列向量，如果左運算元是一個向量，那麼他被看做一個行向量。這個也說明，為什麼我們要固定用mvp乘以position（左乘矩陣），而不能position乘以mvp！因為兩者的處理結果不一致。

三、Metal和OpenGL ES的差異

OpenGL的歷史已經超過25年。基於當時設計原則，OpenGL不支援多執行緒，非同步操作，還有著臃腫的特性。為了更好利用GPU，蘋果設計了Metal。 Metal的目標包括更高效的CPU&GPU互動，減少CPU負載，支援多執行緒執行，可預測的操作，資源控制和同非同步控制；介面與OpenGL類似，但更加切合蘋果設計的GPUs。

Metal的關係圖

Metal的關係圖如上，其中的Device是GPU裝置的抽象，負責管道相關物件的建立：

Device

Metal和OpenGL ES的程式碼對比

我們先看一段OpenGL ES的渲染程式碼，我們可以抽象為Render Targets的設定，Shaders繫結，設定Vertex Buffers、Uniforms和Textures，最後呼叫Draws指令。

glBindFramebuffer(GL_FRAMEBUFFER, myFramebuffer);
glUseProgram(myProgram);
glBindBuffer(GL_ARRAY_BUFFER, myVertexBuffer);
glBindBuffer(GL_UNIFORM_BUFFER, myUniforms);
glBindTexture(GL_TEXTURE_2D, myColorTexture);
glDrawArrays(GL_TRIANGLES, 0, numVertices);

再看Metal的渲染程式碼： Render Targets設定 是建立encoder； Shaders繫結 是設定pipelineState； 設定Vertex Buffers、Uniforms和Textures 是setVertexBuffer和setFragmentBuffer； 呼叫Draws指令 是drawPrimitives； 最後需要再呼叫一次endEncoding。

encoder = [commandBuffer renderCommandEncoderWithDescriptor:descriptor]; [encoder setPipelineState:myPipeline];
[encoder setVertexBuffer:myVertexData offset:0 atIndex:0];
[encoder setVertexBuffer:myUniforms offset:0 atIndex:1];
[encoder setFragmentBuffer:myUniforms offset:0 atIndex:1];
[encoder setFragmentTexture:myColorTexture atIndex:0];
[encoder drawPrimitives:MTLPrimitiveTypeTriangle vertexStart:0  vertexCount:numVertices];
[encoder endEncoding];

Metal和OpenGL ES的同異步處理

如下圖，是用OpenGL ES實現一段渲染的程式碼。CPU在Frame1的回撥中寫入資料到buffer，之後GPU會從buffer中讀取Frame1寫入的資料。

但在Frame2 CPU在往Buffer寫入資料時，Buffer仍儲存著Frame1的資料，且GPU還在使用該buffer，於是Frame2必須等待Frame1渲染完畢，造成阻塞。如下，會產生CPU的wait和GPU的idle。

Metal的處理方案會更加高效。如下圖，Metal會申請三個buffer對應三個Frame，然後根據GPU的渲染回撥，實時更新buffer的快取。 在Frame2的時候，CPU會操作Buffer2，而GPU會讀取Buffer1，並行操作以提高效率。

總結

Metal系列入門教程介紹了Metal的圖片繪製、三維變換、視訊渲染、天空盒、計算管道、Metal與OpenGL ES互動。結合本文的總結，能對Metal產生基本的認知，看懂大部分Metal渲染的程式碼。 接下來的學習方向是Metal進階，包括Metal濾鏡鏈的設計與實現、多重colorAttachments渲染、綠幕功能實現、更復雜的通用計算比如MPSImageHistogram，Shader的效能優化等。

相關閱讀
【每日課程推薦】機器學習實戰！快速入門線上廣告業務及CTR相應知識

此文已由作者授權騰訊雲+社群釋出，更多原文請點選

搜尋關注公眾號「雲加社群」，第一時間獲取技術乾貨，關注後回覆1024 送你一份技術課程大禮包！

海量技術實踐經驗，盡在雲加社群！