Metal 一、初識 Metal 及其語言規范


Metal 簡介 與 Metal 編程語言語法。

一、Metal 簡介

1、Metal 是 Apple 為了解決 3D 渲染而推出的框架。游戲開發者的新的技術平台,該技術能夠為3D圖像提高 10 倍的渲染性能。蘋果2018年推出 Metal,在此之前一直使用OpenGL ES。

Metal 的2個目的:1圖形渲染; 2高並發計算

蘋果文檔給出的 Metal 優化:CPU 低消耗、GPU 高利用、連續處理器並行、有效的資源利用

Metal 官方文檔

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 可以修飾結構體,其結構體成員可以有多個,類型可以為一個 整型/浮點型 的 標量/向量。

 


免責聲明!

本站轉載的文章為個人學習借鑒使用,本站對版權不負任何法律責任。如果侵犯了您的隱私權益,請聯系本站郵箱yoyou2525@163.com刪除。



 
粵ICP備18138465號   © 2018-2025 CODEPRJ.COM