Metal語言規(guī)范
??Metal著?語?是?個?來編寫3D圖形渲染邏輯和并?計算核?邏輯的編程語?,編寫Metal框架的APP需要使?Metal 著?語?程序。
??Metal 著?語? 與 Metal 框架配合使?,Metal 框架管理Metal著?語?的運?和可選編譯選項. Metal 著?器語?使?Clang和LLVM,編譯器對于在GPU上的代碼執(zhí)?效率有更好的控制。
Metal語言的限制
- C++ 11.0 特性在Metal 語?中不?持之處
- Lambda 表達式
- 遞歸函數(shù)調?
- 動態(tài)轉換操作符
- 類型識別
- 對象創(chuàng)建new 和銷毀delete 操作符
- 操作符 noexcept
- goto 跳轉
- 變量存儲修飾符register 和 thread_local
- 虛函數(shù)修飾符
- 派?類
- 異常處理
- C++ 標準庫在Metal 語?中也不可使?
- Metal 語?中對于指針使?的限制
- Metal圖形和并?計算函數(shù)?到的?參數(shù); 如果是指針必須使?地址空間修飾符 (device,threadgroup,constant)
- 不?持函數(shù)指針
- 函數(shù)名不能出現(xiàn)main
- Metal 像素坐標系統(tǒng): Metal 中紋理/幀緩存區(qū)attachment 的像素使?的坐標系統(tǒng)的原點是左上 ?;
Metal數(shù)據(jù)類型
- bool 布爾類型, true/false
- char 有符號8位整數(shù)
- unsigned char /uchar ?符號8-bit 整數(shù)
- short 有符號16-bit整數(shù)
- unsigned short / ushort ?符號32-bit 整數(shù)
- half 16位bit 浮點數(shù)(沒有double類型)
- float 32bit 浮點數(shù)
- size_t 64 ?符號整數(shù)
- void 該類型表示?個空的值集合

Metal向量與矩陣
- 向量支持如下類型:
- booln
- charn
- shortn
- intn
- ucharn
- ushortn
- uintn
- halfn
- floatn
向量中的n,指的是維度
例如,n=4的時候,我們可以用int4表示一個4維向量
// int型的4維向量
int4 test = int4(0,1,2,3);
- 向量的相關操作
獲取其中的元素
//向量可以通過x,y,z,w或者r,g,b,a來訪問屬性,但是x,y,z,w和r,g,b,a不能混合使用
int4 test = int4(0,1,2,3);
//x表示第一個元素
int a = test.x;
//y表示第二個元素
int b = test.y;
//z表示第三個元素
int c = test.z;
//w表示第四個元素
int d = test.w;
//r表示第一個元素
int e = test.r;
//g表示第二個元素
int f = test.g;
//b表示第三個元素
int g = test.b;
//a表示第四個元素
int h = test.a;
獲取其中元素組成的向量
//2維向量
int2 test1 = test.xy;
//3維向量
int3 test1 = test.xyz;
向量賦值
float4 pos = float4(1.0f,2.0f,3.0f,4.0f);
//向量讀取可以亂序
float4 swiz = pos.wxyz; //swiz = (4.0,1.0,2.0,3.0);
//向量讀取可以重復
float4 dup = pos.xxyy; //dup = (1.0f,1.0f,2.0f,2.0f);
//pos = (5.0f,2.0,3.0,6.0)
pos.xw = float2(5.0f,6.0f);
//pos = (8.0f,2.0f,3.0f,7.0f)
pos.wx = float2(7.0f,8.0f);
//pos = (3.0f,5.0f,9.0f,7.0f);
pos.xyz = float3(3.0f,5.0f,9.0f);
//非法,x出現(xiàn)2次
pos.xx = float2(3.0,4.0f);
//不合法-使用混合限定符
pos.xy = float4(1.0f,2.0,3.0,4.0);
- 矩陣支持如下類型:
- halfnxm
- floatnxm
nxm分別指的是矩陣的?數(shù)和列數(shù)(最大為4*4)
float4x4 m;
//將第二排的值設置為0
m[1] = float4(2.0f);
//設置第一行/第一列為1.0f
m[0][0] = 1.0f;
//設置第三行第四列的元素為3.0f
m[2][3] = 3.0f;
紋理Textures 類型
紋理類型是?個句柄, 它指向?個?維/?維/三維紋理數(shù)據(jù)
枚舉值: 定義了訪問權利
enum class access {sample ,read ,write};
- sample : 紋理對象可以被采樣. 采樣?維這是使?或不使?采樣器從紋理中讀取數(shù)據(jù);
- read : 不使?采樣器, ?個圖形渲染函數(shù)或者?個并?計算函數(shù)可以讀取紋理對象;
- write: ?個圖形渲染函數(shù)或者?個并?計算函數(shù)可以向紋理對象寫?數(shù)據(jù);
紋理表達式
//一維紋理
texture1d<T, access a = access::sample>
//二維紋理
texture2d<T, access a = access::sample>
//三維紋理
texture3d<T, access a = access::sample>
T : 數(shù)據(jù)類型 設定了從紋理中讀取或是向紋理中寫?時的顏?類型. T可以是half, float, short, int 等;
//當access為sample類型時,access a = access::sample可以省略
void foo (texture2d<float> imgA [[ texture(0) ]] , texture2d<float, access::read> imgB [[ texture(1) ]], texture2d<float, access::write> imgC [[ texture(2) ]])
{ ... }
采樣器Samplers類型
采取器類型決定了如何對?個紋理進?采樣操作. 在Metal 框架中有?個對應著?器語?的采樣器的對象 MTLSamplerState 這個對象作為圖形渲染著?器函數(shù)參數(shù)或是并?計算函數(shù)的參數(shù)傳遞;
- 從紋理中采樣時,紋理坐標是否需要歸?化
enum class coord { normalized, pixel };
- 紋理采樣過濾?式, 放?/縮?過濾模式
enum class filter { nearest, linear };
- 設置紋理采樣的縮?過濾模式
enum class min_filter { nearest, linear };
- 設置紋理采樣的放?過濾模式
enum class mag_filter { nearest, linear };
- 設置紋理s,t,r坐標的尋址模式
enum class s_address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat }; enum class t_address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat }; enum class r_address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat };
- 設置所有的紋理坐標的尋址模式
enum class address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat };
-設置紋理采樣的mipMap過濾模式, 如果是none,那么只有?層紋理?效
enum class mip_filter { none, nearest, linear };
注意: 在Metal 程序中初始化的采樣器必須使? constexpr 修飾符聲明
代碼示例
constexpr sampler s(coord::pixel,address::clamp_to_zero, filter::linear);
函數(shù)修飾符
Metal 有以下3種函數(shù)修飾符:
- kernel , 表示該函數(shù)是?個數(shù)據(jù)并?計算著?函數(shù). 它可以被分配在?維/?維/三維線程組中去執(zhí)?;
kernel void foo(...)
{
...
}
注意: 使?kernel 修飾的函數(shù). 其返回值類型必須是void 類型;
- vertex , 表示該函數(shù)是?個頂點著?函數(shù) , 它將為頂點數(shù)據(jù)流中的每個頂點數(shù)據(jù)執(zhí)??次然后為每個頂 點?成數(shù)據(jù)輸出到繪制管線;
//頂點函數(shù)
vertex int CCTestVertexFunctionB(int a,int b)
{
...
}
- fragment , 表示該函數(shù)是?個?元著?函數(shù), 它將為?元數(shù)據(jù)流中的每個?元 和其關聯(lián)執(zhí)??次然后 將每個?元?成的顏?數(shù)據(jù)輸出到繪制管線中;
//片元函數(shù)
fragment int CCTestVertexFunctionB(int a,int b)
{
...
}
只有圖形著?函數(shù)才可以被 vertex 和 fragment 修飾.
對于圖形著?函數(shù), 返回值類型可以辨認出它是為 頂點做計算還是為每像素做計算.
圖形著?函數(shù)的返回值可以為 void , 但是這也就意味著該函數(shù)不產?數(shù) 據(jù)輸出到繪制管線; 這是?個?意義的動作;
?個被函數(shù)修飾符修飾的函數(shù)不能在調?其他也被函數(shù)修飾符修飾的函數(shù); 這樣會導致編譯失敗;
即kernel、vertex、fragment之間無法相互調用
kernel void hello(...)
{ ... }
vertex float4 hello1(...)
{
//?個被函數(shù)修飾符修飾的函數(shù)不能在調?其他也被函數(shù)修飾符修飾的函數(shù); 這樣會 導致編譯失敗;
hello(...);
//錯誤調??
}
?于變量或者參數(shù)的地址空間修飾符
??Metal 著?器語?使?地址空間修飾符來表示?個函數(shù)變量或者參數(shù)變量 被分配于那??內存區(qū)域.
??所有的著?函數(shù)(vertex, fragment, kernel)的參數(shù),如果是指針或是引?, 都必須帶有地址空間修飾符號;
- device
- threadgroup
- constant
- thread
對于圖形著?器函數(shù), 其指針或是引?類型的參數(shù)必須定義為 device 或是 constant 地址空間;
對于并?計算著?函數(shù), 其指針或是引?類型的參數(shù)必須定義為 device 或是 threadgroup 或是 constant 地址空間;
Device Address Space(設備地址空間)
??在設備地址空間(Device)指向設備(此處設備指的是顯存)內存池分配出來的緩存對象, 它是可讀也是可寫的;
???個緩存對象可 以被聲明成?個標量,向量或是?戶?定義結構體的指針或是引?。
??由于device修飾的變量存放在顯存中,所以讀取速度會比放在內存中要快。
//device修飾4維向量的顏色
device float4 *color;
struct Foo {
float a[3];
int b[2];
};
//device修飾結構體
device Foo *my_info;
注意: 紋理對象總是在設備地址空間分配內存, device 地址空間修飾符不必出現(xiàn)在紋理類型定義中. ?個紋 理對象的內容?法直接訪問. Metal 提供讀寫紋理的內建函數(shù);
threadgrounp Address Space (線程組地址空間)
??線程組地址空間?于為 并?計算著?函數(shù)分配內存變量. 這些變量被?個線程組的所有線程共享. 在線 程組地址空間分配的變量不能被?于圖形繪制著?函數(shù)[頂點著?函數(shù), ?元著?函數(shù)]
??在并?計算著?函數(shù)中, 在線程組地址空間分配的變量為?個線程組使?, 聲明周期和線程組相同;
kernel void my_func(threadgroup float *a [[ threadgroup(0) ]], ...){
threadgroup float x;
threadgroup float b[10];
}
constant Address Space (常量地址空間)
??常量地址空間指向的緩存對象也是從設備內存池分配存儲, 但是它是只讀的;
??在程序域的變量必須定義在常量地址空間并且聲明的時候初始化; ?來初始化的值必須是編譯時的常量.
??在程序域的變量的?命周期和程序?樣, 在程序中的并?計算著?函數(shù)或者圖形繪制著?函數(shù)調?, 但 是constant 的值會保持不變;
注意: 常量地址空間的指針或是引?可以作為函數(shù)的參數(shù). 向聲明為常量的變量賦值會產?編譯錯誤. 聲明常量但是沒有賦予初值也會產?編譯錯誤;
//定義一個常量samples
constant float samples[] = { 1.0f, 2.0f, 3.0f, 4.0f };
//對?個常量地址空間的變量進?修改也會失敗,因為它只讀的
sampler[4] = {3,3,3,3}; //編譯失敗;
//定義為常量地址空間聲明時不賦初值也會編譯失敗
constant float a;
thread Address Space (線程地址空間)
??thread 地址空間指向每個線程準備的地址空間, 這個線程的地址空間定義的變量在其他線程不可?, 在 圖形繪制著?函數(shù)或者并?計算著?函數(shù)中聲明的變量thread 地址空間分配;
kernel void my_func(...)
{
float x;
thread float p = &x;
}
函數(shù)參數(shù)與變量
??圖形繪制或者并?計算著?器函數(shù)的輸?輸出都是通過參數(shù)傳遞. 除了常量地址空間變量和程序域定義 的采樣器以外.
- device buffer- 設備緩存, ?個指向設備地址空間的任意數(shù)據(jù)類型的指針或者引?;
- constant buffer -常量緩存區(qū), ?個指向常量地址空間的任意數(shù)據(jù)類型的指針或引?
- texture - 紋理對象;
- sampler - 采樣器對象;
- threadGrounp - 在線程組中供各線程共享的緩存.
對于每個著?器函數(shù)來說, ?個修飾符是必須指定的. 他?來設定?個緩存,紋理, 采樣器的位置;
- device buffers/ constant buffer --> [[buffer (index)]]
- texture -- [[texture (index)]]
- sampler -- [[sampler (index)]]
- threadgroup buffer -- [[threadgroup (index)]]
index是?個unsigned integer類型的值,它表示了?個緩存、紋理、采樣器參數(shù)的位置(在函數(shù)參數(shù)索引 表中的位置,相當于GLSL中的Location)。 從語法上講,屬性修飾符的聲明位置應該位于參數(shù)變量名之后
kernel void add_vectors(const device float4 *inA [[ buffer(0) ]],
const device float4 *inB [[ buffer(1) ]],
device float4 *out [[ buffer(2) ]],
uint id [[ thread_position_in_grid ]])
{
//thread_position_in_grid : ?于表示當前節(jié)點在多線程?格中的位置;由Metal自行進行傳遞
out[id] = inA[id] + inB[id];
}
[[ buffer(index) ]]中的index不能和其他的相同,否則會覆蓋掉之前的數(shù)據(jù)
內建變量屬性修飾符
- [[vertex_id]] 頂點id 標識符;
- [[position]] 頂點信息(float4) / 在片元著色器中描述了?元的窗?相對坐標(x, y, z, 1/w)
- [[point_size]] 點的??(float)
- [[color(m)]] 顏?, m編譯前得確定;(有可能在一個函數(shù)中,會出現(xiàn)多種顏色,用m進行區(qū)分)
-
[[stage_in]] : ?元著?函數(shù)使?的單個?元輸?數(shù)據(jù)是由頂點著?函數(shù)輸出然后經過光柵化?成的.頂點和?元著?函數(shù)都是只能有?個參數(shù)被聲明為使?“stage_in”修飾符,對于?個使? 了“stage_in”修飾符的? 定義的結構體,其成員可以為?個整形或浮點標量,或是整形或浮點向量
[[stage_in]]表示從頂點著色器中經過圖元裝配、光柵化等操作后傳遞給片元著色器的數(shù)據(jù)。類似于GLSL中的Varying。
struct MyFragmentOutput {
// color attachment 0
float4 clr_f [[color(0)]];
// color attachment 1
int4 clr_i [[color(1)]];
// color attachment 2
uint4 clr_ui [[color(2)]];
};
fragment MyFragmentOutput my_frag_shader( ... )
{
MyFragmentOutput f;
....
f.clr_f = ...;
....
return f;
}