Metal 簡介 與 Metal 編程語言語法。
一、Metal 簡介
1、Metal 是 Apple 為了解決 3D 渲染而推出的框架。游戲開發者的新的技術平台,該技術能夠為3D圖像提高 10 倍的渲染性能。蘋果2018年推出 Metal,在此之前一直使用OpenGL ES。
Metal 的2個目的:1圖形渲染; 2高並發計算。
蘋果文檔給出的 Metal 優化:CPU 低消耗、GPU 高利用、連續處理器並行、有效的資源利用
2、圖形管道 graphic pipeline
二、Metal Shading Language - Metal 編程語言與規范
1、Metal 語言特點
1.1)Metal 着色器語言,用來編寫3D 圖形渲染邏輯和並行計算核心邏輯(高並發) 的一門編程語言。當我們需要使用 Metal 框架來完成 APP 的開發時(或 需要使用 Metal 的高並發計算能力時),就要使用 Metal 編程語言進行編程。
1.2)Metal 編程語言 使用 clang 和 LLVM 進行編譯處理
Metal的語言規則,是基於 C++ 11.0 標椎設計的,並在此基礎上進行了一定的擴展和限制。
而我們的實際業務開發場景中所需要做的工作就是: 編寫 在GPU 上執行的圖像渲染邏輯代碼 以及 通用的並行計算邏輯 代碼。
1.3)Metal 與 C++ 11.0
1.3.1)Metal 限制(不支持):
Lambda 表達式
遞歸函數調用
動態轉換操作符
類型識別
對象創建 new 和 銷毀 delete 操作符
操作符 noexcept
goto 跳轉
變量存儲修飾符 register 和 thread_local
虛函數修飾符
派生類
異常處理
C++ 標准函數庫 在 Metal 中不支持,不能使用
1.3.2)Metal 中指針的使用限制
Metal 圖形和並行計算函數用到的參數,如果是指針 必須使用地址空間修飾符(devide / threadgroup / constant)
不支持函數指針
函數名不能出現 main
1.3.3)Metal 像素坐標系
Metal 中紋理/幀緩沖區 attachment 的像素使用的坐標系的原點是在 左上角。--> 蘋果的坐標原點
2、Metal 的數據類型
2.1)標量類型
2.2)向量與矩陣 數據類型
2.2.1)向量
booln
charn
shortn
intn
ucharn / ushortn / uintn
halfn
floatn
n --> 向量中的 n 表示維度,最大為4
代碼示例:
bool2 b = [1, 2]; float4 f4 = float4(1.0, 2.0, 3.0, 4.0); float f = f4[0];// x = 1.0 --> 類似數組 // int4 --> 4個變量組成的4維向量 // xyzw rgba int4 test = int4(0, 1, 2, 3); int x = test.x;// x = 0 int y = text.y;// y = 1 int x1 = test.r;// x1 = 0 // 多個分量的訪問 float4 c = float4(0,0,0,0); c.xyzw = float4(1,2,3,4);// 重賦值 c = [1,2,3,4] c.xy = float2(6,0);// c = [6,0,3,4] c.yzw = float3(7,8,9);// c = [6,7,8,9] // 可亂序 --> // 注意!!!這里 Metal 不同於 GLSL --> GLSL 可以多分量,但是不可亂序 xyzw/rgba 順序是不可變的 float4 pos = c.wxyz;// pos = [9,6,7,8] float4 rep = c.xxwz;// rep = [6,6,9,8] rep.xw = float2(5,7);// rep = [5,6,9,7] // 不可混用 float4 m = float4(4,3,2,1); m.xg = float2(0,9);!error 非法, xyzw 和 rgba 不能混了,他倆只可選其一 m.rg = float(0,9);// m = [0,9,2,1] /// 構造方式 // float4 類型向量的所有可能構造方式 float4(float x); float4(float x,float y,float z,float w); float4(float2 a,float2 b); float4(float2 a,float b,float c); float4(float a,float2 b,float c); float4(float a,float b,float2 c); float4(float3 a,float b); float4(float a,float3 b); float4(float4 x); // float3 類型向量的所有可能的構造的方式 float3(float x); float3(float x,float y,float z); float3(float a,float2 b); float3(float2 a,float b); float3(float3 x); // float2 類型向量的所有可能的構造方式 float2(float x); float2(float x,float y); float2(float2 x); // 多個向量構造的使用 float x = 1.0f,y = 2.0f,z = 3.0f,w = 4.0f; float4 a = float4(0.0f); float4 b = float4(x,y,z,w); float2 c = float2(5.0f,6.0f); float2 a = float2(x,y); float2 b = float2(z,w); float4 x = float4(a.xy,b.xy);
2.2.2)矩陣
halfxnm / floatnxm
nxm 中 n m 分別指 矩陣的行數 和 列數 --> 最大 4 x 4 : 4行4列
float4x4 mix; // mix[1] = float4(2.0f);// 矩陣 第一行的值都是 2 mix[1] = float4(1,2,3,4);// 矩陣的第一行的值 mix[0][0] = 3;// 矩陣的第0行0列的值為 3 mix[3][2] = 7;// 矩陣 3行2列 位置的值
2.3)紋理 Texture 類型
紋理類型是一個句柄,指向一個 一維/二維/三維紋理數據。在函數中描述紋理對象的類型。
枚舉:
enum class access { sample, read, write } // 定義訪問權限
sample:紋理對象可以被采樣,采樣器可將紋理讀取出來,可讀可寫可采樣 --> 使用最多
read:不使用采樣器,一個圖形渲染函數或並行計算函數 可以讀取紋理對象
write:一個圖形渲染函數或並行計算函數 可以向紋理對象寫入數據。
texture1d<T, access a = access::sample> // 一維紋理 texture2d<T, access a = access::sample>// 二維紋理 texture3d<T, access a = access::sample>// 三維紋理
T:數據類型 ,指定從紋理中 讀取/寫入 時的顏色類型。T可以是 half、float、int 等;
access:讀寫方式(權限)
代碼示例:
void foo (texture2d<float> imgA [[ texture(0) ]] ,// texture2d<float>: 2 維紋理,類型 float,訪問權限 sample --> 默認權限就是 sample 可不寫 texture2d<float, access::read> imgB [[ texture(1) ]],// texture2d<float, access::read>: 類型 float,權限 read texture2d<float, access::write> imgC [[ texture(2) ]]) // 權限 write { ... }
2.4)采樣器類型 Samplers
采樣器類型 決定了如何對一個紋理進行采樣操作。
metal 框架中有一個對應 着色器語言的采樣器對象:MTLSamplerState, 此對象做為 圖像渲染着色器函數 or 並行計算函數 的參數進行傳遞。
枚舉 們:
// 從紋理中采樣時,紋理坐標是否歸一化 enum class coord { normalized, peixel } // 紋理采樣過濾方式 - 放大/縮小 enum class filter { nearest, linear } // 縮小過濾方式 enum class min_filter { nearest, linear } // 放大過濾方式 enum class mag_filter { nearest, linear } // 設置紋理 s t r 坐標的尋址模式 (str 即 xyz 環繞方式) enum class s_address { clapm_to_zero, clapm_to_edge, repeat, mirrored_repeat } enum class t_address { clapm_to_zero, clapm_to_edge, repeat, mirrored_repeat } enum class r_address { clapm_to_zero, clapm_to_edge, repeat, mirrored_repeat } // 設置所有紋理坐標的尋址模式 enum class address { clapm_to_zero, clapm_to_edge, repeat, mirrored_repeat } // 設置紋理采樣的 mipMap 過濾模式,如果是 none ,則只有一層紋理生效 enum class mag_filter { none, nearest, linear }
注意:在 Metal 中,初始化采樣器必須使用 constexpr 修飾符聲明
代碼示例:
// 初始化 創建 采樣器 屬性設置 constexpr sampler s(coord::pixel, address::clamp_to_zero, filter::linear); // constexpr sampler a(coord::normalized); constexpr sampler b(address::repeat); constexpr sampler s(address::clamp_to_zero, filter::linear);
3、修飾符
3.1)函數修飾符
kernel:表示該函數是一個 數據並行計算着色函數。我們要高效並發運算就用它。它可以被分配在 一維/二維/三維 線程組中去執行;--> 使用他修飾的函數 返回類型必須是 void
vertex:頂點着色函數。為頂點數據流中的 每個頂點數據 執行一次,然后為每個頂點生成數據輸出到繪制管線;
fragment:片元着色函數,為 片元數據流中的 每個片元 與其關聯 執行一次,然后將 每個片元生成的顏色數據 輸出到繪制管線中;
注意1:只有 圖形着色函數才能用 vertex/fragment 修飾。函數返回類型 可以用來辨認出它是為頂點 or 為每個像素 做計算的。返回 void 也可以但是無意義,因為頂點/片元函數本就是為了計算出相應數據 將數據傳到繪制管線的。
注意2:被函數修飾符修飾的函數不能再調用 '被修飾符修飾的函數',否則編譯失敗。即:被函數修飾符修飾的函數們不能相互調用。
例:kernel void func1 (...) {}; vertex float4 funcV1 (...) { func1(...) } --> 錯誤調用,無法編譯
注意2:特定函數修飾,普通函數隨意。
代碼示例:
kernel void foo(...) { ... }
3.2)變量 或 參數 的地址空間修飾符 Address Space
地址空間修飾符:用來指定 一個函數 參數/變量 被分配在內存中的哪塊區域。
device:設備地址空間
threadgroup:線程組地址空間
constant:常量地址空間
thread: thread 地址空間
a、對於 圖形着色器函數,是 指針 或 引用 類型的參數 必須定義為 device 或 const 地址空間
b、對於並行計算着色函數,對於是 指針 或 引用 的參數,必須使用 device 或 threadgroup 或 constant 修飾。
3.2.1)Device Address Space(設備地址空間)
device:設備地址空間 指向設備內存(顯存)池分配出來的緩存對象,它可以是可讀也可以是可寫的;一個緩存對象可以被聲明成一個 標量、向量、自定義結構體的指針或引用。
代碼示例:
// an array of a float vector with 4 components device float4 *color;
// 定義個結構體 struct Foo { float a[3]; int b[2]; } // an array of Foo elements device Foo *my_info;
注意1:紋理對象 總是在設備地址空間分配內存,device 地址空間修飾符不必出現在紋理類型定義中。一個紋理對象的內容無法直接訪問,Metal 提供了讀寫紋理的內建函數。
3.2.2)線程組地址空間 threadgroup
threadgroup:用於 為 並行計算着色函數 分配內存變量(在GPU里),這些變量被一個線程組的所有線程共享。在線程組地址空間 分配的變量 不能被用於圖形繪制着色函數。
在並行計算着色函數中,在線程組地址空間分配的變量 為一個線程組使用,生命周期和線程相同。
代碼示例:
// kernel 高速並行 kernel void my_func(threadgroup float *a [[ threadgroup(0) ]], ...) { // A float allocated in threadgroup address space threadgroup float x; // An array of 10 floats allocated in threadgroup address space threadgroup float b[10]; }
3.2.3)constant 常量地址空間
constant:指向的緩存對象也是從設備內存池 分配存儲,但是是只讀的。
在程序域的變量 必須定義在常量地址空間 並且在聲明的時候初始化;用來初始化的值 必須是編譯時的常量。此變量的生命周期和程序一樣,在程序中的 並行計算着色函數or圖形繪制着色函數 調用,但 constant 的值會保持不變。
注意:常量地址空間 的指針或引用 可以作為函數的參數(constant修飾的常量可作為函數的參數)。向聲明為常量的變量賦值會產生變異錯誤(代碼示例中sampler),聲明為常量但沒有賦予初始值也會產生變異錯誤(代碼示例中a)。
錯誤代碼示例:
constant float sampler[] = {1.0f, 2.0f, 3.0f,4.0f}; // 對一個常量地址空間的變量進行修改會失敗,因為它是只讀的 sampler[4] = {3,3,3,3};// 編譯失敗 // 定義常量地址空間但不初始化賦值 --> 也編譯失敗 const float a;// 編譯失敗
3.2.4)線程地址空間 thread
thread:指向每個線程准備的地址空間,這個線程的地址空間 定義的變量 在其他線程是不可見的,在圖形繪制着色函數or並行計算着色函數 中聲明的變量可以使用 thread 地址空間分配。
代碼示例:
kernel void func2 (...) { float x; thread float p = &x; ... }
3.3)函數參數與變量
圖形繪制/並行計算着色函數的 輸入/輸出 都需要通過參數傳遞 ( 除了常量地址空間變量和程序域中定義的采樣器 外)。參數如下:
device buffer:設備緩存 - 指向設備地址空間的任意數據類型的指針 or 引用
constant buffer:常量緩存 - 指向常量地址空間的任意數據類型的指針 or 引用
texture object:紋理對象
sample object:采樣器對象
threadgroup:線程共享的緩存
對於每個着色器函數來說,一個修飾符是必須指定的,它用來設定一個 緩存、紋理、采樣器的位置:
device buffer / constant buffer --> [[buffer(index)]]
texture --> [[texture(index)]]
sample --> [[sampler(index)]]
threadgroup buffer --> [[threadgroup(index)]]
index:一個 unsigned integer 類型的值,表示一個 緩存、紋理、采樣器的位置(在函數參數索引表中的位置)。語法上講,屬性修飾符的聲明位置 應該位於參數變量名之后。
通過示例理解:
// 一個簡單的並行計算着色函數 my_add ,它把兩個設備地址空間的魂村 inA、inB 相加,把結果寫入緩存 out。 // 屬性修飾符 “buffer(index)” 為着色函數參數 設定了緩存的位置 kernel void my_add (constant device float4 *inA [[ buffer(0) ]],// inA: 放在設備地址空間,緩存位置對應的是 buffer(0)這個ID , constant 修飾的不可變 constant device float4 *inB [[ buffer(1) ]], device float4 *out [[ buffer(2) ]], uit id [[ thread_position_in_grid ]],) { out[id] = inA[id] + inB[id]; }
thread_position_in_grid:用於表示當前節點,在多線程網格中的位置 --> 我們是無法知道當前在GPU的哪個運算單元里,thread_position_in_grid 知道,我們通過它獲取即可。
3.4)內建變量屬性修飾符
[[vertex_id]] -- 頂點ID標識符
[[position]] -- 1、當前頂點信息(float4 - xyzw) 2、也可描述 片元在窗口的相對坐標:當前這個像素點在屏幕上的哪個位置
[[point_size]] -- 點的大小
[[color(m)]] -- 顏色,m 編譯前要確定
如下代碼:
// 定義顏色結構體 struct myFragmentOutput { // 三組顏色,要知道使用時取哪一個 float4 clr_f [[color(0)]];// int4 clr_i [[color(1)]];// uint4 clr_ui [[color(2)]];// } fragment myFragmentOutput my_grag_shader (...) { myFragmentOutput f; ... f.clr_f = ...; ... return f; }
另補充一個:[[stage_in]] -- 其實就是:頂點着色器輸出 經過光柵化 生成的 傳給片元着色器的 每個片元數據。
頂點和片元着色函數都是 有且僅有 一個參數可以被聲明為 使用"stage_in"修飾符的。 stage_in 可以修飾結構體,其結構體成員可以有多個,類型可以為一個 整型/浮點型 的 標量/向量。