Metal 語言介紹
- Metal 著?語? 是?來編寫
3D 圖形渲染邏輯
和并?計算核?邏輯
的 ??編程語?漫贞。 當你使? Metal 框架來完成APP的實現(xiàn)時,則需要使?Metal 編程語?; - Metal 語?使?
Clang 和 LLVM
進?編譯處理伐弹; - Metal 基于C++ 11.0 語?設(shè)計披摄,我們主要?來編寫
在 GPU 上執(zhí)?的圖像渲染邏輯代碼 以及 通? 并?計算邏輯代碼
;
C++ 11.0 和 Metal 語?的異同之處
- C++ 11.0 特性在Metal 語?中不?持之處
- Lambda 表達式;
- 遞歸函數(shù)調(diào)? 動態(tài)轉(zhuǎn)換操作符 類型識別 對象創(chuàng)建
new
和銷毀delete
操作符; - 操作符
noexcept
跑揉、goto
跳轉(zhuǎn) 變量存儲修飾符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)的
原點是左上?
;
數(shù)據(jù)類型
標量類型
Metal ?持后綴表示字?量類型锅睛, 例如 0.5F, 0.5f; 0.5h, 0.5H;
//基本數(shù)據(jù)類型
bool a = true;
char b = 5;
int d = 15;
size_t c = 1;
ptrdiff_t f = 2;
向量和矩陣
向量?持如下類型:
- booln
- charn
- shortn
- intn
- ucharn
- ushortn
- uintn
- halfn
- floatn
向量中的n埠巨,指的是維度
矩陣?持如下類型:
- halfnxm
- floatnxm
nxm分別指的是矩陣的?數(shù)和列數(shù)
//向量
bool2 A= {1,2};
float4 pos = float4(1.0,2.0,3.0,4.0);
float x = pos[0];
float y = pos[1];
float4 VB;
for(int i = 0; i < 4 ; I++)
VB[i] = pos[i] * 2.0f;
//通過向量字母來獲取元素
int4 test = int4(0,1,2,3);
int a = test.x;
int b = test.y;
int c = test.z;
int d = test.w;
int e = test.r;
int f = test.g;
int g = test.b;
int h = test.a;
float4 c;
c.xyzw = float4(1.0f,2.0f,3.0f,4.0f);
c.z = 1.0f;
c.xy = float2(3.0f,4.0f);
c.xyz = float3(3.0f,4.0f,5.0f);
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);
float2 pos;
pos.x = 1.0f; //合法
pos.z = 1.0f; //非法
float3 pos2;
pos2.z = 1.0f; //合法
pos2.w = 1.0f; //非法
//非法,x出現(xiàn)2次
pos.xx = float2(3.0,4.0f);
//不合法-使用混合限定符
pos.xy = float4(1.0f,2.0,3.0,4.0);
float4 pos4 = float4(1.0f,2.0f,3.0f,4.0f);
pos4.x = 1.0f;
pos4.y = 2.0f;
//非法,.rgba與.xyzw 混合使用
pos4.xg = float2(2.0f,3.0f);
////非法,.rgba與.xyzw 混合使用
float3 coord = pos4.ryz;
float4 pos5 = float4(1.0f,2.0f,3.0f,4.0f);
//非法,使用指針來指向向量/分量
my_func(&pos5.xy);
//矩陣
float4x4 m;
//將第二排的值設(shè)置為0
m[1] = float4(2.0f);
//設(shè)置第一行/第一列為1.0f
m[0][0] = 1.0f;
//設(shè)置第三行第四列的元素為3.0f
m[2][3] = 3.0f;
//float4類型向量的所有可能構(gòu)造方式
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類型向量的所有可能的構(gòu)造的方式
float3(float x);
float3(float x,float y,float z);
float3(float a,float2 b);
float3(float2 a,float b);
float3(float3 x);
//float2類型向量的所有可能的構(gòu)造方式
float2(float x);
float2(float x,float y);
float2(float2 x);
//多個向量構(gòu)造器的使用
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);
緩存buffer
- 在Metal 中實現(xiàn)緩存靠的是?個指針,它指向?個在
Device
或者constant
地址空間中的內(nèi)建或是開發(fā)者?定義的數(shù)據(jù)塊现拒,緩存可以被定在程序域域中辣垒,或是當做函數(shù)的參數(shù)傳遞。
//緩存buffer
device float4 *device_buffer;
struct my_user_data{
float4 a;
float b;
int2 c;
};
constant my_user_data *user_data;
紋理Textures
紋理類型是?個句柄印蔬,它指向?個?維/?維/三維紋理數(shù)據(jù)勋桶。?紋理數(shù)據(jù)對應(yīng)這樣?個紋理的某個level的mipmap的全部或者?部分。
枚舉值: 定義了訪問權(quán)利;
- sample:紋理對象可以被采樣. 采樣?維這是使?或不使?采樣器從紋理中讀取數(shù)據(jù);
- read:不使?采樣器, ?個圖形渲染函數(shù)或者?個并?計算函數(shù)可以讀取紋理對象;
- write:?個圖形渲染函數(shù)或者?個并?計算函數(shù)可以向紋理對象寫?數(shù)據(jù);
//紋理texture
enum class access {sample,read,write};
texture1d<T,access a = access::sample>
texture1d_array<T,access a = access::sample>
texture2d<T,access a = access::sample>
texture2d_array<T,access a = access::sample>
texture3d<T,access a = access::sample>
texturecube<T,access a = access::sample>
texture2d_ms<T,access a = access::read>
//帶有深度格式的紋理必須被聲明為下面紋理數(shù)據(jù)類型中的一個
enum class depth_forma {depth_float};
depth2d<T,access a = depth_format::depth_float>
depth2d_array<T,access a = access::sample,depth_format d = depth_format::depth_float>
depthcube<T,access a = access::sample,depth_format d = depth_format::depth_float>
depth2d_ms<T,access a = access::read,depth_format d = depth_format::depth_float>
void foo (texture2d<float> imgA[[texture(0)]],
texture2d<float,access::read> imgB[[texture(1)]],
texture2d<float,access::write> imgC[[texture(2)]])
{
//...
}
采樣器類型 Samplers
采取器類型決定了如何對?個紋理進?采樣操作扛点。在Metal 框架中有?個對應(yīng)著?器語?的采樣器的對象 MTLSamplerState
這個對象作為圖形渲染著?器函數(shù)參數(shù)或是并?計算函數(shù)的參數(shù)傳遞;
enum class coord { normalized, pixel };
從紋理中采樣時,紋理坐標是否需要歸?化;
enum class filter { nearest, linear };
紋理采樣過濾?式, 放?/縮?過濾模式;
enum class min_filter { nearest, linear };
設(shè)置紋理采樣的縮?過濾模式;
enum class mag_filter { nearest, linear };
設(shè)置紋理采樣的放?過濾模式;
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 };
設(shè)置紋理s,t,r坐標的尋址模式;
enum class address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat };
設(shè)置所有的紋理坐標的尋址模式;
enum class mip_filter { none, nearest, linear };
設(shè)置紋理采樣的mipMap過濾模式, 如果是none,那么只有?層紋理?效;
注意: 在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)
函數(shù)修飾符
Metal 有以下3種函數(shù)修飾符:
- kernel:表示該函數(shù)是?個
數(shù)據(jù)并?計算著?函數(shù)
哥遮,它可以被分配在?維/?維/三維線程組中去執(zhí)?; - vertex:表示該函數(shù)是?個
頂點著?函數(shù)
,它將為頂點數(shù)據(jù)流中的每個頂點數(shù)據(jù)執(zhí)??次然后為每個頂 點?成數(shù)據(jù)輸出到繪制管線; - fragment:表示該函數(shù)是?個
?元著?函數(shù)
陵究,它將為?元數(shù)據(jù)流中的每個?元 和其關(guān)聯(lián)執(zhí)??次然后 將每個?元?成的顏?數(shù)據(jù)輸出到繪制管線中;
//函數(shù)修飾符.
/*
3個函數(shù)修飾符:
1. kernel : 并行計算函數(shù)
2. vertex : 頂點函數(shù)
3. fragment : 片元函數(shù)
*/
//1.并行計算函數(shù)(kernel)
kernel void CCTestKernelFunctionA(int a,int b)
{
/*
注意:
1. 使用kernel 修飾的函數(shù)返回值必須是void 類型
2. 一個被函數(shù)修飾符修飾過的函數(shù),不允許在調(diào)用其他的被函數(shù)修飾過的函數(shù). 非法
3. 被函數(shù)修飾符修飾過的函數(shù),只允許在客戶端對齊進行操作. 不允許被普通的函數(shù)調(diào)用.
*/
//不可以的!
//一個被函數(shù)修飾符修飾過的函數(shù),不允許在調(diào)用其他的被函數(shù)修飾過的函數(shù). 非法
CCTestKernelFunctionB(1,2);//非法
CCTestVertexFunctionB(1,2);//非法
//可以! 你可以調(diào)用普通函數(shù).而且在Metal 不僅僅只有這3種被修飾過的函數(shù).普通函數(shù)也可以存在
CCTest();
}
kernel void CCTestKernelFunctionB(int a,int b)
{
}
//頂點函數(shù)
vertex int CCTestVertexFunctionB(int a,int b){
}
//片元函數(shù)
fragment int CCTestVertexFunctionB(int a,int b){
}
//普通函數(shù)
void CCTest()
{
}
注意
使?kernel 修飾的函數(shù)眠饮,其返回值類型必須是void 類型;
只有圖形著?函數(shù)才可以被 vertex 和 fragment 修飾. 對于圖形著?函數(shù), 返回值類型可以辨認出它是為 頂點做計算還是為每像素做計算. 圖形著?函數(shù)的返回值可以為 void , 但是這也就意味著該函數(shù)不產(chǎn)?數(shù) 據(jù)輸出到繪制管線; 這是?個?意義的動作;?個被函數(shù)修飾符修飾的函數(shù)不能在調(diào)?其他也被函數(shù)修飾符修飾的函數(shù); 這樣會導(dǎo)致編譯失敗;
kernel void hello2(...){
}
vertex float4 hello1(...){
//?個被函數(shù)修飾符修飾的函數(shù)不能在調(diào)?其他也被函數(shù)修飾符修飾的函數(shù); 這樣會 導(dǎo)致編譯失敗;
hello2(...); //錯誤調(diào)??
}
?于變量或者參數(shù)的地址空間修飾符
Metal 著?器語?使? 地址空間修飾符
來表示?個函數(shù)變量或者參數(shù)變量 被分配于哪??內(nèi)存區(qū)域,所有的著?函數(shù)(vertex, fragment, kernel)
的參數(shù)铜邮,如果是指針或是引?仪召,都必須帶有地址空間修飾符號
;
- device— 設(shè)備地址空間
- threadgroup — 線程組地址空間
- constant — 常量地址空間
- thread — thread地址空間
注意
- 對于
圖形著?器函數(shù)
, 其指針或是引?類型的參數(shù)必須定義為 device 或是 constant 地址空間; - 對于
并?計算著?函數(shù)
松蒜,其指針或是引?類型的參數(shù)必須定義為 device 或是 threadgrounp 或是 constant 地址空間;
//變量/參數(shù)地址空間修飾符
void CCTestFouncitionE(device int *g_data,
threadgroup int *l_data,
constant float *c_data
)
{
//...
}
Device Address Space (設(shè)備地址空間)
在設(shè)備地址空間(Device) 指向設(shè)備內(nèi)存池分配出來的緩存對象扔茅,它是可讀也是可寫的;?個緩存對象可以被聲明成?個標量秸苗、向量或是?戶?定義結(jié)構(gòu)體的指針或是引?召娜。
注意: 紋理對象總是在設(shè)備地址空間分配內(nèi)存(即顯存), device 地址空間修飾符不必出現(xiàn)在紋理類型定義中. ?個紋理對象的內(nèi)容?法直接訪問. Metal 提供讀寫紋理的內(nèi)建函數(shù);
// 設(shè)備地址空間: device 用來修飾指針.引用
//1.修飾指針變量
device float4 *color;
struct CCStruct{
float a[3];
int b[2];
};
//2.修飾結(jié)構(gòu)體類的指針變量
device CCStruct *my_CS;
threadgrounp Address Space (線程組地址空間)
- 線程組地址空間?于
為并?計算著?函數(shù)分配內(nèi)存變量
,這些變量被?個線程組的所有線程共享惊楼。在線程組地址空間分配的變量不能被?于圖形繪制著?函數(shù)
[頂點著?函數(shù), ?元著?函數(shù)] - 在并?計算著?函數(shù)中玖瘸,在線程組地址空間分配的變量為?個線程組使?,聲明周期和線程組相同;
/*
1. threadgroup 被并行計算計算分配內(nèi)存變量, 這些變量被一個線程組的所有線程共享. 在線程組分配變量不能被用于圖像繪制.
2. thread 指向每個線程準備的地址空間. 在其他線程是不可見切不可用的
*/
kernel void CCTestFouncitionF(threadgroup float *a)
{
//在線程組地址空間分配一個浮點類型變量x
threadgroup float x;
//在線程組地址空間分配一個10個浮點類型數(shù)的數(shù)組y;
threadgroup float y[10];
}
constant float sampler[] = {1.0f,2.0f,3.0f,4.0f};
kernel void CCTestFouncitionG(void)
{
//在線程空間分配空間給x,p
float x;
thread float p = &x;
}
constant Address Space (常量地址空間)
- 常量地址空間指向的緩存對象也是從設(shè)備內(nèi)存池分配存儲檀咙,但是它是只讀的雅倒;
- 在程序域的變量必須定義在常量地址空間并且聲明的時候初始化;
- ?來初始化的值必須是編譯時的常量弧可;
- 在程序域的變量的?命周期和程序?樣蔑匣,在程序中的并?計算著?函數(shù)或者圖形繪制著?函數(shù)調(diào)?,但是constant 的值會保持不變棕诵;
注意: 常量地址空間的指針或是引?可以作為函數(shù)的參數(shù). 向聲明為常量的變量賦值會產(chǎn)?編譯錯誤. 聲明常量但是沒有賦予初值也會產(chǎn)?編譯錯誤;
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è)備緩存, ?個指向設(shè)備地址空間的任意數(shù)據(jù)類型的指針或者引?;
- constant buffer -常量緩存區(qū), ?個指向常量地址空間的任意數(shù)據(jù)類型的指針或引?
- texture - 紋理對象;
- sampler - 采樣器對象;
- threadGrounp - 在線程組中供各線程共享的緩存。
注意: 被著?器函數(shù)的緩存(device 和 constant) 不能重名;
Attribute Qualifiers to Locate Buffers, Textures, and Samplers ?于尋址緩存,紋理, 采樣器的屬性修飾符;
對于每個著?器函數(shù)來說搔确,?個修飾符是必須指定的彼棍。 他?來設(shè)定?個緩存灭忠、紋理、 采樣器的位置;
- device buffers/ constant buffer --> [[buffer (index)]]
- texture -- [[texture (index)]]
- sampler -- [[sampler (index)]]
- threadgroup buffer -- [[threadgroup (index)]]
index是?個unsigned integer類型的值座硕,它表示了?個緩存弛作、紋理、采樣器參數(shù)的位置(在函數(shù)參數(shù)索引 表中的位置)华匾。 從語法上講映琳,屬性修飾符的聲明位置應(yīng)該位于參數(shù)變量名之后
//屬性修飾符
/*
1. device buffer(設(shè)備緩存)
2. constant buffer(常量緩存)
3. texture Object(紋理對象)
4. sampler Object(采樣器對象)
5. 線程組 threadgroup
屬性修飾符目的:
1. 參數(shù)表示資源如何定位? 可以理解為端口
2. 在固定管線和可編程管線進行內(nèi)建變量的傳遞
3. 將數(shù)據(jù)沿著渲染管線從頂點函數(shù)傳遞片元函數(shù).
在代碼中如何表現(xiàn):
1.已知條件:device buffer(設(shè)備緩存)/constant buffer(常量緩存)
代碼表現(xiàn):[[buffer(index)]]
解讀:不變的buffer ,index 可以由開發(fā)者來指定.
2.已知條件:texture Object(紋理對象)
代碼表現(xiàn): [[texture(index)]]
解讀:不變的texture ,index 可以由開發(fā)者來指定.
3.已知條件:sampler Object(采樣器對象)
代碼表示: [[sampler(index)]]
解讀:不變的sampler ,index 可以由開發(fā)者來指定.
4.已知條件:threadgroup Object(線程組對象)
代碼表示: [[threadgroup(index)]]
解讀:不變的threadgroup ,index 可以由開發(fā)者來指定.
*/
//并行計算著色器函數(shù)add_vectros ,實現(xiàn)2個設(shè)備地址空間中的緩存A與緩存B相加.然后將結(jié)果寫入到緩存out.
//屬性修飾符"(buffer(index))" 為著色函數(shù)參數(shù)設(shè)定了緩存的位置
//并行計算著色器函數(shù)add_vectros ,實現(xiàn)2個設(shè)備地址空間中的緩存A與緩存B相加.然后將結(jié)果寫入到緩存out.
//屬性修飾符"(buffer(index))" 為著色函數(shù)參數(shù)設(shè)定了緩存的位置
kernel void add_vectros(
const device float4 *inA [[buffer(0)]],
const device float4 *inB [[buffer(1)]],
device float4 *out [[buffer(2)]]
uint id[[thread_position_in_grid]])
{
out[id] = inA[id] + inB[id];
}
// thread_position_in_grid : ?于表示當前節(jié)點在多線程?格中的位置;
//著色函數(shù)的多個參數(shù)使用不同類型的屬性修飾符的情況
kernel void my_kernel(device float4 *p [[buffer(0)]],
texture2d<float> img [[texture(0)]],
sampler sam [[sampler(0)]])
{
//.....
}
內(nèi)建變量屬性修飾符
- [[vertex_id]] 頂點id 標識符;
- [[position]] 頂點信息(float4) /? 述了?元的窗?相對坐標(x, y, z, 1/w)
- [[point_size]] 點的??(float)
- [[color(m)]] 顏?, m編譯前得確定;
struct MyFragmentOutput {
float4 clr_f [[color(0)]]; // color attachment 0
int4 clr_i [[color(1)]];// color attachment 1
uint4 clr_ui [[color(2)]]; }; // color attachment 2
fragment MyFragmentOutput my_frag_shader( ... ) {
MyFragmentOutput f;
....
f.clr_f = ...;
...
return f;
}
}
[[stage_in]] : ?元著?函數(shù)使?的單個?元輸?數(shù)據(jù)是由頂點著?函數(shù)輸出然后經(jīng)過光柵化?成的。頂點和?元著?函數(shù)都是只能有?個參數(shù)被聲明為使?“stage_in”修飾符蜘拉,對于?個使? 了“stage_in”修飾符的?定義的結(jié)構(gòu)體萨西,其成員可以為?個整形或浮點標量,或是整形或浮點向量旭旭。