資深程序員的Metal入門教程總結

歡迎你們前往騰訊雲+社區,獲取更多騰訊海量技術實踐乾貨哦~javascript

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

正文

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

1、Metal

Metal 是一個和 OpenGL ES 相似的面向底層的圖形編程接口,能夠直接操做GPU;支持iOS和OS X,提供圖形渲染和通用計算能力。(不支持模擬器)github

img

圖片來源 www.invasivecode.com/weblog/meta…web

MTLDevice 對象表明GPU,一般使用MTLCreateSystemDefaultDevice獲取默認的GPU; MTLCommandQueue由device建立,用於建立和組織MTLCommandBuffer,保證指令(MTLCommandBuffer)有序地發送到GPU;MTLCommandBuffer會提供一些encoder,包括編碼繪製指令的MTLRenderCommandEncoder、編碼計算指令的MTLComputeCommandEncoder、編碼緩存紋理拷貝指令的MTLBlitCommandEncoder。對於一個commandBuffer,只有調用encoder的結束操做,才能進行下一個encoder的建立,同時能夠設置執行完指令的回調。 每一幀都會產生一個MTLCommandBuffer對象,用於填放指令; GPUs的類型不少,每一種都有各自的接收和執行指令方式,在MTLCommandEncoder把指令進行封裝後,MTLCommandBuffer再作聚合到一次提交裏。 MTLRenderPassDescriptor 是一個輕量級的臨時對象,裏面存放較多屬性配置,供MTLCommandBuffer建立MTLRenderCommandEncoder對象用。編程

img

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 表示渲染管道,最主要的三個過程:頂點處理、光柵化、片元處理:數據結構

img

轉換幾何形狀數據爲幀緩存中的顏色像素,叫作點陣化(rasterizing),也叫光柵化。其實就是根據頂點的數據,檢測像素中心是否在三角形內,肯定具體哪些像素須要渲染。 對開發者而言,頂點處理和片元處理是可編程的,光柵化是固定的(不可見)。 頂點函數在每一個頂點被繪製時都會調用,好比說繪製一個三角形,會調用三次頂點函數。頂點處理函數返回的對象裏,必須有帶[[position]]描述符的屬性,表面這個屬性是用來計算下一步的光柵化;返回值沒有描述符的部分,則會進行插值處理。

img

插值處理

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

在編譯的時候,Xcode會單獨編譯.metal的文件,但不會進行連接;須要在app運行時,手動進行連接。 在包裏,能夠看到default.metallib,這是對metal shader的編譯結果。

img

MTLFunction能夠用來建立MTLRenderPipelineState對象,MTLRenderPipelineState表明的是圖形渲染的管道; 在調用device的newRenderPipelineStateWithDescriptor:error接口時,會進行頂點、像素函數的連接,造成一個圖像處理管道; MTLRenderPipelineDescriptor包括名稱、頂點處理函數、片元處理函數、輸出顏色格式。

setVertexBytes:length:atIndex:這接口的長度限制是4k(4096bytes),對於超過的場景應該使用MTLBuffer。MTLBuffer是GPU可以直接讀取的內存,用來存儲大量的數據;(經常使用於頂點數據) newBufferWithLength:options:方法用來建立MTLBuffer,參數是大小和訪問方式;MTLResourceStorageModeShared是默認的訪問方式。

紋理

Metal要求全部的紋理都要符合MTLPixelFormat上面的某一種格式,每一個格式都表明對圖像數據的不一樣描述方式。 例如MTLPixelFormatBGRA8Unorm格式,內存佈局以下:

img

每一個像素有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),以下:

img

有時候,紋理的座標會超過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

img

常見的圖形渲染管道

2、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是索引
複製代碼

img

地址空間的修飾符共有四個,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的有符號整數 一般用於指針的差值 half二、half三、half四、float二、float三、float4等,是向量類型,表達方式爲基礎類型+向量維數。矩陣相似half4x四、half3x三、float4x四、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

數據對齊 char三、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!由於二者的處理結果不一致。

3、Metal和OpenGL ES的差別

OpenGL的歷史已經超過25年。基於當時設計原則,OpenGL不支持多線程,異步操做,還有着臃腫的特性。爲了更好利用GPU,蘋果設計了Metal。 Metal的目標包括更高效的CPU&GPU交互,減小CPU負載,支持多線程執行,可預測的操做,資源控制和同異步控制;接口與OpenGL相似,但更加切合蘋果設計的GPUs。

img

Metal的關係圖

Metal的關係圖如上,其中的Device是GPU設備的抽象,負責管道相關對象的建立:

img

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);
複製代碼

img

再看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];
複製代碼

img

Metal和OpenGL ES的同異步處理

以下圖,是用OpenGL ES實現一段渲染的代碼。CPU在Frame1的回調中寫入數據到buffer,以後GPU會從buffer中讀取Frame1寫入的數據。

img

但在Frame2 CPU在往Buffer寫入數據時,Buffer仍存儲着Frame1的數據,且GPU還在使用該buffer,因而Frame2必須等待Frame1渲染完畢,形成阻塞。以下,會產生CPU的wait和GPU的idle。

img

Metal的處理方案會更加高效。以下圖,Metal會申請三個buffer對應三個Frame,而後根據GPU的渲染回調,實時更新buffer的緩存。 在Frame2的時候,CPU會操做Buffer2,而GPU會讀取Buffer1,並行操做以提升效率。

img

總結

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

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

此文已由做者受權騰訊雲+社區發佈,更多原文請點擊

搜索關注公衆號「雲加社區」,第一時間獲取技術乾貨,關注後回覆1024 送你一份技術課程大禮包!

海量技術實踐經驗,盡在雲加社區

相關文章
相關標籤/搜索