Metal Shading Language簡(jiǎn)述
- Metal 著色語言是用來編寫
3D圖形渲染邏輯 和 并行計(jì)算核心邏輯的一門編程語言饵隙，如果需要使用Metal框架來實(shí)現(xiàn)某些邏輯則需要使用該語言- Metal是通過Xcode的Clang 和 LLVM進(jìn)行編譯、鏈接沮脖，無需在手動(dòng)編譯
- Metal基于C++ 11.0語言設(shè)計(jì)的金矛，在C++基礎(chǔ)上多了一些擴(kuò)展和限制
- Metal 像素坐標(biāo)系統(tǒng)：Metal中
紋理 或者 幀緩存區(qū)attachment的像素原點(diǎn)是在左上角
Metal語言的部分限制
- 遞歸函數(shù)
- C++標(biāo)準(zhǔn)庫在Metal語言中也不可使用
- Metal圖形和并行計(jì)算函數(shù)用到的入?yún)⑿炯保绻?code>指針 / 引用必須使用地址空間修飾符（比如device、threadgroup驶俊、constant）,否則編譯報(bào)錯(cuò)
- 無法使用函數(shù)指針娶耍，也就是方法入?yún)ⅰ⒎祷刂怠?/li>
- Metal文件中不得出現(xiàn)
main()函數(shù) - 無法進(jìn)行異常的捕捉和處理
Metal語言-數(shù)據(jù)類型
包含：
- 基本數(shù)據(jù)類型
- 標(biāo)量類型
- 向量類型
- 矩陣類型
- 紋理類型
- 采樣器類型
基本數(shù)據(jù)類型
標(biāo)量類型
- 常用的：bool,char,int,uint,float
向量類型
向量支持如下類型:
| 類型 | 舉例 |
|---|---|
| booln | bool,bool2,bool3,bool4 |
| charn | char,char2,char3,char4 |
| shortn | ... |
| intn | ... |
| ucharn | ... |
| ushortn | ... |
| uintn | ... |
| halfn | ... |
| floatn | ... |
- 最多支持4維向量
bool2 A= {1,2};
float4 pos = float4(1.0,2.0,3.0,4.0);
pos = float4(float2(1.0,2.0),float2(3.0,4.0));
float x = pos[0];
x = pos.x;
float y = pos[1];
y = pos.y;
矩陣類型
矩陣類型有且只有兩種：halfnxm,floatnxm
n:行數(shù)
m:列數(shù)
half4x4 n;
float4x4 m;
m[0] = float4(1,1,1,1);
m[1][1] = 1;
m[3][3] = 3;
紋理類型
紋理類型是一個(gè)句柄废睦，指向一個(gè)一維/二維/三維的紋理數(shù)據(jù);相當(dāng)于OpenGL中的textureBufferID.
紋理數(shù)據(jù)類型
限制從紋理中讀取或是向紋理中寫入是的顏色類型.
可以有：half,float,short,int等類型伺绽，一般使用：float
訪問權(quán)限
- sample：可讀可寫，紋理可以被采樣嗜湃。權(quán)限默認(rèn)值
- read：只讀奈应，不使用采樣器。渲染函數(shù)购披、并行計(jì)算函數(shù)可以讀取紋理數(shù)據(jù)
- write：可讀可寫杖挣，不使用采樣器。渲染函數(shù)刚陡、并行計(jì)算函數(shù)可以修改紋理數(shù)據(jù)
//這兩種寫法等價(jià)
texture1d<float, access::sample>
texture1d<float>
texture2d<float, access::read>
texture3d<float, access::write>
kernel void my_kernel(texture2d<float> img [[texture(0)]])
{
//.....
}
采樣器類型
采樣器類型決定了如何對(duì)一個(gè)紋理進(jìn)行采樣惩妇。
在Metal中有一個(gè)對(duì)應(yīng)著色器語言色采樣器對(duì)象
MTLSamplerState這個(gè)對(duì)象作為：圖形渲染著色器函數(shù)參數(shù) 、 并行計(jì)算函數(shù)的參數(shù)筐乳。
在Metal程序中初始化的采樣器必須使用constexpr修飾符聲明
constexpr sampler s(coord::pixel,
address::clamp_to_zero,
filter::linear);
constexpr sampler a(coord::normalized);
constexpr sampler b(address::repeat);
Metal語言-函數(shù)修飾符
- kernel：表示該函數(shù)是一個(gè)
數(shù)據(jù)并行計(jì)算著色函數(shù)歌殃，它可以被分配在一維/二維/三維線程組中去執(zhí)行 - vertex：表示該函數(shù)是一個(gè)
頂點(diǎn)著色函數(shù)，它將為頂點(diǎn)數(shù)據(jù)流中的每個(gè)頂點(diǎn)數(shù)據(jù)執(zhí)行一次蝙云，然后為每個(gè)頂點(diǎn)生成數(shù)據(jù)輸出到繪制管線 - fragment：表示該函數(shù)是一個(gè)
片元著色函數(shù)氓皱，它將為片元數(shù)據(jù)流中的每個(gè)片元 和其相關(guān)聯(lián)的數(shù)據(jù)執(zhí)行一次，然后將每個(gè)片元生成的顏色數(shù)據(jù)輸出到繪制管線中
注意點(diǎn)：
- 被函數(shù)修飾符修飾的函數(shù)體內(nèi)不能調(diào)用任何被函數(shù)修飾符修飾的函數(shù)
- 被函數(shù)符修飾的函數(shù)系統(tǒng)會(huì)自動(dòng)調(diào)用勃刨，開發(fā)者不允許調(diào)用
- kernel修飾符修飾的函數(shù)返回值必須為
void波材，其他兩種（vertex、fragment）返回值也可以為void,但是這樣會(huì)導(dǎo)致函數(shù)執(zhí)行無效 - 被函數(shù)符修飾的函數(shù)可以調(diào)用普通函數(shù)
//并行計(jì)算函數(shù)(kernel)
kernel void KernelFunction(int a,int b)
{
KernelFunctionB(1,2);//非法調(diào)用Ｉ硪Ｍ⑶！
}
kernel void KernelFunctionB(int a)
{ ... }
//頂點(diǎn)函數(shù)
vertex int VertexFunction(int a,int b)
{
Test();//合法調(diào)用＜致痢Ｏ肚帷！
}
//片元函數(shù)
fragment int VertexFunction(int a,int b)
{ ..... }
//普通函數(shù)
void Test()
{ ..... }
Metal語言-地址空間修飾符
Metal著色器語言使用地址空間修飾符來表示一個(gè) 函數(shù)變量或者 參數(shù)變量 被分配于哪一片內(nèi)存區(qū)域.所有被函數(shù)符修飾的函數(shù)其參數(shù)如果是指針垢揩、引用大脉，就必須使用地址空間修飾符
包含以下4種：
-
device：設(shè)備（GPU緩存）地址空間 -
constant：常量地址空間 -
threadgrounp：線程組地址空間 -
thread：線程地址空間
Device Address Space
- 設(shè)備地址空間指向GPU緩存分配出來的緩存對(duì)象，該值
可讀可寫水孩，一個(gè)緩存對(duì)象可以被聲明成一個(gè)標(biāo)量、向量或是用戶自定義結(jié)構(gòu)體的指針/引用 - device放在變量類型之前
- 紋理對(duì)象總是在設(shè)備地址空間分配內(nèi)存琐驴，此處的
device可以省略俘种。紋理對(duì)象無法直接訪問秤标，需要通過紋理的內(nèi)建變量（ [[texture(0)]] ）來獲取
//1. 修飾指針變量
device float4 *color;
struct Struct{
float a[3];
int b[2];
};
//2.修飾結(jié)構(gòu)體類的指針變量
device CCStruct *my_CS;
constant Address Space
- 常量地址空間指向的緩存對(duì)象也是從設(shè)備內(nèi)存池分配存儲(chǔ)，僅
可讀 - 在程序域的變量必須定義在常量地址空間并且聲明時(shí)初始化.不初始化會(huì)導(dǎo)致編譯錯(cuò)誤
- 在程序域的變量的生命周期和程序一樣宙刘，在程序中的并行計(jì)算著色函數(shù) 或者 圖形繪制著色函數(shù)調(diào)用苍姜，但是constant的值會(huì)保持不變
constant float samples[] = { 1.0f, 2.0f, 3.0f, 4.0f };
threadgroup Address Space
- 線程組地址空間用于為并行計(jì)算著色器函數(shù)分配內(nèi)存變量，這些變量被
一個(gè)線程組的所有線程共享 - 在線程組地址空間分配的變量不能用于圖形繪制著色函數(shù)（即頂點(diǎn)著色函數(shù) / 片元著色函數(shù)）悬包，即在圖形繪制著色函數(shù)中不能使用線程組
- 在并行計(jì)算著色函數(shù)中衙猪，在線程組地址空間分配的變量為一個(gè)線程組使用，生命周期和線程組相同
kernel void KernelFouncition(threadgroup float *a)
{
//在線程組地址空間分配一個(gè)浮點(diǎn)類型變量x
threadgroup float x;
}
thread Address Space
- 線程地址空間指向每個(gè)線程準(zhǔn)備的地址空間布近，也是在GPU中垫释，該線程的地址空間定義的變量在
其他線程不可見（即變量不共享） - 在圖形繪制著色函數(shù) 或者 并行計(jì)算著色函數(shù)中聲明的變量，
在線程地址空間分配存儲(chǔ)
kernel void CCTestFouncitionG(void)
{
//在線程空間分配空間給x,p
float x;
thread float p = &x;
}
注意：
- 在
圖形著色器函數(shù)(頂點(diǎn)函數(shù)片元函數(shù)),其指針/引用類型的參數(shù)必須定義為device撑瞧、constant地址空間 - 在
并行計(jì)算函數(shù)(kernel函數(shù))其指針/引用類型的參數(shù)必須定義為device棵譬、threadgroup、constant - 被
thread修飾的變量無法共享预伺，所以只能在三類函數(shù)體內(nèi)進(jìn)行使用
函數(shù)參數(shù)與變量的傳遞修飾符订咸，即屬性修飾符
圖形繪制 或者 并行計(jì)算著色器函數(shù)的輸入輸出都是通過參數(shù)傳遞，除了常量地址空間變量和程序域定義的采樣器之外, 其他參數(shù)修飾的可以是如下之一酬诀，常用的有以下5種屬性修飾符：
-
device buffer設(shè)備緩存：一個(gè)指向設(shè)備地址空間的任意數(shù)據(jù)類型的指針/引用 -
constant buffer常量緩存：一個(gè)指向常量地址空間的任意數(shù)據(jù)類型的指針/引用 -
texture紋理對(duì)象 -
sampler采樣器對(duì)象 -
threadGroup在線程組中供線程共享的緩存
參數(shù)表示資源的位置句柄脏嚷，可以理解為端口，相當(dāng)于OpenGl ES中的location
vertex int foo (const device float4 *inA [[ buffer(0) ]],
device float4 *out [[ buffer(1) ]]
texture2d<float> imgA[[texture(0)]],
texture2d<float> imgB[[texture(1)]],
threadgroup td [[threadgroup(0)]])
{
//...
}
常見的內(nèi)建變量修飾符
[[vertex_id]]:頂點(diǎn)函數(shù)中的index瞒御，并不由開發(fā)者傳遞[[position]]:在頂點(diǎn)著色函數(shù)中父叙，表示當(dāng)前的頂點(diǎn)信息，類型是float4葵腹、
在片元函數(shù)中還可以描述該像素點(diǎn)在窗口的相對(duì)坐標(biāo)（x高每，y，z践宴，1/w）鲸匿，即該像素點(diǎn)在屏幕上的位置信息[[point_size]]：點(diǎn)的大小，類型是float[[color(m)]]：顏色阻肩，m在編譯前就必須確定[[ thread_position_in_grid ]]: 用于表示當(dāng)前節(jié)點(diǎn)在多線程網(wǎng)格中的位置带欢，只能用在kernel函數(shù)中[[stage_in]]：片元著色函數(shù)使用的單個(gè)片元輸入數(shù)據(jù)是由頂點(diǎn)著色函數(shù)輸出然后經(jīng)過光柵化生成的，也就是片元函數(shù)的入?yún)⒂糜趯?duì)應(yīng)頂點(diǎn)函數(shù)的返回值.只允許在片元函數(shù)的參數(shù)中出現(xiàn)1次烤惊；可以使用各種標(biāo)量乔煞、自定義類型.
//頂點(diǎn)函數(shù)
vertex float4 VertexFunction(uint vertexId [[vertex_id]],
constant int *i [[buffer(0)]])
{
return float4(1,2,3,4);
}
//片元函數(shù)
fragment int VertexFunction(float4 param [[stage_in]])
{ ..... }
