Metal Shader language (著色語言規(guī)范)

Metal簡述

  • Metal著色器語言是用來編寫 3D圖形渲染邏輯、并行Metal計算核心邏輯 的一門編程語言醉者，當你使用Metal框架來完成APP的實現(xiàn)時則需要使用Metal編程語言但狭。
  • Metal語言使用Clang 和LLVM進行編譯處理违寿，編譯器對于在GPU上的代碼執(zhí)行效率有更好的控制
  • Metal基于C++ 11.0語言設計的，在C++基礎(chǔ)上多了一些擴展和限制熟空，主要用來編寫在GPU上執(zhí)行的圖像渲染邏輯代碼以及通用并行計算邏輯代碼
  • Metal 像素坐標系統(tǒng)：Metal中紋理 或者 幀緩存區(qū)attachment的像素使用的坐標系統(tǒng)的原點是左上角

Metal語言的限制

  • Metal中不支持C++11.0的如下特性
    • Lambda表達式
    • 遞歸函數(shù)調(diào)用
    • 動態(tài)轉(zhuǎn)換操作符
    • 類型識別
    • 對象創(chuàng)建new和銷毀delete操作符
    • 操作符noexcept
    • go跳轉(zhuǎn)
    • 變量存儲修飾符 register 和thread_local
    • 虛函數(shù)修飾符
    • 派生類
    • 異常處理
  • C++標準庫在Metal語言中也不可使用
  • Metal語言對于指針使用的限制
    • Metal圖形和并行計算函數(shù)用到的入?yún)ⅲū热缰羔?/ 引用），如果是指針 / 引用必須使用地址空間修飾符（比如device搞莺、threadgroup息罗、constant）
    • 不支持函數(shù)指針
    • 函數(shù)名不能出現(xiàn)main

Metal 基本數(shù)據(jù)類型

基本數(shù)據(jù)類型主要有 ：

  • 標量
  • 向量
  • 矩陣

標量

Metal中的標量類型如下圖所示


標量類型.png
  • 常用的主要有 bool、int才沧、uint 迈喉、half
  • unsigned char可以簡寫為 uchar
  • unsigned short 可以簡寫為 ushort
  • unsigned int 可以簡寫為 uint
  • 其中half 相當于OC中的float，float 相當于OC中的double
    size_t用來表示內(nèi)存空間温圆， 相當于 OC中 sizeof
bool a = true;
char b = 5;
int  d = 15;
//用于表示內(nèi)存空間
size_t c = 1;
ptrdiff_t f = 2;

向量

向量支持如下類型

  • booln挨摸、charn、shortn岁歉、intn得运、ucharn、ushortn锅移、uintn熔掺、halfn、floatn非剃，其中 n 表示向量的維度置逻，最多不超過4維向量
//直接賦值初始化
bool2 A= {1,2};
//通過內(nèi)建函數(shù)float4初始化
float4 pos = float4(1.0,2.0,3.0,4.0);

//通過下標從向量中獲取某個值
float x = pos[0];
float y = pos[1];

//通過for循環(huán)對一個向量進行運算
float4 VB;
for(int i = 0; i < 4 ; i++)
{
    VB[i] = pos[i] * 2.0f;
}
  • 在OpenGL ES的GLSL語言中，例如2.0f备绽，在著色器中書寫時券坞，是不能加f，寫成2.0肺素，而在Metal中則可以寫成2.0f恨锚，其中f可以是大寫，也可以是小寫

向量的訪問規(guī)則

  • 通過向量字母獲取元素： 向量中的向量字母僅有2種倍靡，分別為xyzw眠冈、rgba 
int4 test = int4(0,1,2,3);
int a = test.x; //獲取的向量元素0
int b = test.y; //獲取的向量元素1
int c = test.z; //獲取的向量元素2
int d = test.w; //獲取的向量元素3

int e = test.r; //獲取的向量元素0
int f = test.g; //獲取的向量元素1
int g = test.b; //獲取的向量元素2
int h = test.a; //獲取的向量元素3
  • 多個分量同時訪問
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);
  • 多分量訪問可以亂序/重復
    • 賦值時分量不可重復，取值時分量可重復
    • 右邊是取值 和 左邊賦值都合法
    • xyzw與rgba不能混合使用
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);

//可以僅對 xw / wx 修改
//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);

//可以僅對 xyz 進行修改
//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; //非法菌瘫，pos是二維向量蜗顽，沒有z這個索引

float3 pos2;
pos2.z = 1.0f; //合法
pos2.w = 1.0f; //非法

// 賦值 時 分量不可重復，取值 時 分量可重復
//非法,x出現(xiàn)2次
pos.xx = float2(3.0,4.0f);
pos.xy = swiz.xx;

//向量中xyzw與rgba兩組分量不能混合使用
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;
  • GLSL中向量不能亂序訪問雨让，只是和Metal中的向量相似雇盖，并不是等價。

矩陣

矩陣支持如下類型

  • halfnxm栖忠、floatnxm崔挖，其中 nxm表示矩陣的行數(shù)和列數(shù)贸街，最多4行4列，其中half狸相、float相當于OC中的float薛匪、double
  • 普通的矩陣其本質(zhì)就是一個數(shù)組
float4x4 m;
//將第二行的所有值都設置為2.0
m[1] = float4(2.0f);

//設置第一行/第一列為1.0f
m[0][0] = 1.0f;

//設置第三行第四列的元素為3.0f
m[2][3] = 3.0f;
  • float4 類型向量的構(gòu)造方式
    • 1個float構(gòu)成,表示一行都是這個值
    • 4個float構(gòu)成
    • 2個float2構(gòu)成
    • 1個float2+2個float構(gòu)成（順序可以任意組合）
    • 1個float2+1個float
    • 1個float4
//float4類型向量的所有可能構(gòu)造方式
//1個一維向量，表示一行都是x
float4(float x);/
//4個一維向量 --> 4維向量
float4(float x,float y,float z,float w);
//2個二維向量 --> 4維向量
float4(float2 a,float2 b);
//1個二維向量+2個一維向量 --> 4維向量
float4(float2 a,float b,float c);
float4(float a,float2 b,float c);
float4(float a,float b,float2 c);
//1個三維向量+1個一維向量 --> 4維向量
float4(float3 a,float b);
float4(float a,float3 b);
//1個四維向量 --> 4維向量
float4(float4 x);
  • float3 類型向量的構(gòu)造方式
    • 1個float構(gòu)成,表示一行都是這個值
    • 3個float
    • 1個float+1個float2（順序可以任意組合）
    • 1個float2
//float3類型向量的所有可能的構(gòu)造的方式
//1個一維向量
float3(float x);
//3個一維向量
float3(float x,float y,float z);
//1個一維向量 + 1個二維向量
float3(float a,float2 b);
//1個二維向量 + 1個一維向量
float3(float2 a,float b);
//1個三維向量
float3(float3 x);
  • float2 類型向量的構(gòu)造方式
    • 1個float構(gòu)成,表示一行都是這個值
    • 2個float
    • 1個float2
//float2類型向量的所有可能的構(gòu)造方式
//1個一維向量
float2(float x);
//2個一維向量
float2(float x,float y);
//1個二維向量
float2(float2 x);

Metal 其他類型

主要有以下兩種

  • 紋理
  • 采樣器

紋理類型

紋理類型是一個句柄脓鹃，指向一維/二維/三維紋理數(shù)據(jù)逸尖，而紋理數(shù)據(jù)對應一個紋理的某個level的mipmap的全部或者一部分

紋理的訪問權(quán)限

在一個函數(shù)中描述紋理對象的類型
access枚舉值由Metal定義，定義了紋理的訪問權(quán)利 enum class access {sample, read, write};瘸右，有以下3種訪問權(quán)利娇跟，當沒寫access時，默認的access 就是 sample

  • sample: 紋理對象可以被采樣（即使用采樣器去紋理中讀取數(shù)據(jù)太颤，相當于OpenGL ES的GLSL中sampler2D）苞俘，采樣一維這時使用 或者 不使用都可以從紋理中讀取數(shù)據(jù)（即可讀可寫可采樣）
  • read：不使用采樣器，一個圖形渲染函數(shù) 或者 一個并行計算函數(shù)可以讀取紋理對象（即僅可讀）
  • write：一個圖形渲染函數(shù) 或者 一個并行計算可以向紋理對象寫入數(shù)據(jù)（即 可讀可寫）

定義紋理類型

描述一個紋理對象/類型龄章，有以下三種方式吃谣，分別對應一維/二維/三維，

  • 其中T代表泛型做裙，設定了從紋理中讀取數(shù)據(jù)或是寫入時的顏色類型基协，T可以是half、float菇用、short澜驮、int等
    • access表示紋理訪問權(quán)限，當access沒寫時惋鸥，默認是sample
    • texture1d<T, access a = access::sample>
    • texture2d<T, access a = access::sample>
    • texture3d<T, access a = access::sample>
//類型 變量 修飾符
/*
 類型
    - texture2d<float>杂穷，讀取的數(shù)據(jù)類型是float，沒寫access卦绣，默認是sample
    - texture2d<float,access::read>耐量，讀取的數(shù)據(jù)類型是float，讀取的方式是read
    - texture2d<float,access::write>滤港，讀取的數(shù)據(jù)類型是float廊蜒，讀取的方式是write
 變量名
    - imgA
    - imgB
    - imgC
 修飾符
    - [[texture(0)]] 對應紋理0
    - [[texture(1)]] 對應紋理1
    - [[texture(2)]] 對應紋理2
 */
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ù)傳遞山叮，有以下幾種狀態(tài)：

-coord：從紋理中采樣時，紋理坐標是否需要歸一化

  • enum class coord { normalized, pixel };
  • filter：紋理采樣過濾方式添履，放大/縮小過濾方式
    • enum class filter { nearest, linear };
  • min_filter：設置紋理采樣的縮小過濾方式
    • enum class min_filter { nearest, linear };
  • mag_filter：設置紋理采樣的放大過濾方式
    • enum class mag_filter { nearest, linear };
  • s_address屁倔、t_address、r_address：設置紋理s暮胧、t锐借、r坐標（對應紋理坐標的x问麸、y、z）的尋址方式
    • s坐標：enum class s_address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat };
    • t坐標：enum class t_address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat };
    • r坐標：enum class r_address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat };
  • address：設置所有紋理坐標的尋址方式
    • enum class address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat };
  • mip_filter：設置紋理采樣的mipMap過濾模式, 如果是none,那么只有一層紋理生效;
    • enum class mip_filter { none, nearest, linear };

采樣器所有狀態(tài)如圖所示

采樣器狀態(tài).png

1钞翔、openGL ES中紋理坐標對應的是stq严卖，Metal中紋理坐標對應是str
2、在Metal程序中初始化的采樣器必須使用constexpr修飾符聲明

/*
constexpr：修飾符（必須寫）
sampler：類型
s：采樣器變量名稱
參數(shù)
    - coord: 是否需要歸一化布轿，不需要歸一化哮笆，用的是像素pixel
    - address: 地址環(huán)繞方式
    - filter: 過濾方式
*/
constexpr sampler s(coord::pixel, address::clamp_to_zero, filter::linear);

constexpr sampler a(coord::normalized);

constexpr sampler b(address::repeat);

函數(shù)修飾符

Metal有以下3中函數(shù)修飾符，放在函數(shù)的最前面驮捍，即位于函數(shù)返回值的前面

  • kernel：表示該函數(shù)是一個數(shù)據(jù)并行計算著色函數(shù)，它可以被分配在一維/二維/三維線程組中去執(zhí)行脚曾，表示函數(shù)要并行計算东且，其返回值類型必須是void類型，是一個高并發(fā)函數(shù)
  • vertex：表示該函數(shù)是一個頂點著色函數(shù)本讥，它將為頂點數(shù)據(jù)流中的每個頂點數(shù)據(jù)執(zhí)行一次珊泳，然后為每個頂點生成數(shù)據(jù)輸出到繪制管線
  • fragment：表示該函數(shù)是一個片元著色函數(shù)，它將為片元數(shù)據(jù)流中的每個片元 和其相關(guān)聯(lián)的數(shù)據(jù)執(zhí)行一次拷沸，然后將每個片元生成的顏色數(shù)據(jù)輸出到繪制管線中

1色查、使用kernel修飾的函數(shù)，其返回值類型必須是void類型
2撞芍、一個被函數(shù)修飾符修飾的函數(shù)不能在調(diào)用其他也被函數(shù)修飾符修飾的函數(shù)秧了，這樣會導致編譯失敗，即Kernel序无、vertex验毡、fragment修飾的函數(shù)不能相互調(diào)用，也不能同修飾符函數(shù)相互調(diào)用帝嗡。但是可以調(diào)用普通函數(shù)
3晶通、被函數(shù)修飾符修飾過的函數(shù),只允許在客戶端對其進行操作. 不允許被普通的函數(shù)調(diào)用
4、Metal中并不是所有函數(shù)都需要上述3個修飾符修飾哟玷，是可以在Metal中定義普通函數(shù)的狮辽，即不帶任何修飾符的函數(shù)
5、只有圖形著色函數(shù)才可以被vertex和fragment修飾巢寡，對于圖形著色函數(shù)喉脖，通過返回值類型可以辨認出是為頂點計算還是像素計算，其返回值也可以是void抑月，意味著不產(chǎn)生數(shù)據(jù)輸出到繪制管線动看，是一個無意義的動作

//并行計算函數(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);//非法，錯誤調(diào)用Ｗ谩Ａ饨浴须误！
    CCTestVertexFunctionB(1,2);//非法，錯誤調(diào)用３鹎帷＞┝　！
    
    //可以! 你可以調(diào)用普通函數(shù).而且在Metal 不僅僅只有這3種被修飾過的函數(shù).普通函數(shù)也可以存在
    CCTest();
    
}

//并行計算函數(shù)
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()
{
    .....
}

變量篷店、參數(shù)的地址空間修飾符

Metal著色器語言使用地址空間修飾符來表示一個 函數(shù)變量或者參數(shù)變量 被分配于哪一片內(nèi)存區(qū)域祭椰，有以下4中地址空間修飾符

  • device： 設備地址空間
  • threadgroup： 線程組地址空間
  • constant: 常量地址空間
  • thread: 線程地址空間

1、所有的著色函數(shù)（vertex疲陕、fragment方淤、kernel）的參數(shù)，如果是指針/引用蹄殃，都必須帶有地址空間修飾符號
2携茂、對于圖形著色器函數(shù)（即vertex/fragment修飾的函數(shù)），其指針/引用類型的參數(shù)必須定義為 device诅岩、constant地址空間
3讳苦、對于并行計算函數(shù)（即kernel修飾的函數(shù)），其指針/引用類型的參數(shù)必須定義為device吩谦、threadgroup鸳谜、constant
4、并不是所有的變量都需要修飾符式廷，也可以定義普通變量（即無修飾符的變量）

/*
 注意:
 1. 所有被(kernel,vertex,fragment)所修飾的參數(shù)變量,如果其類型是指針/引用 都必須帶有地址空間修飾符.
 2. 被fragment修飾的片元函數(shù), 指針/引用必須被device/constant/threadgroup
 */

//變量/參數(shù)地址空間修飾符
void CCTestFouncitionE(device int *g_data,
                       threadgroup int *l_data,
                       constant float *c_data
                       )
{
    //...
    
}

1咐扭、紋理對象總是在設備地址空間分配內(nèi)存，即紋理對象默認在GPU分配內(nèi)存
2滑废、device地址空間修飾符不必出現(xiàn)在紋理類型定義中
3草描、一個紋理對象的內(nèi)容無法直接訪問，Metal提供讀寫紋理的內(nèi)建函數(shù)策严，通過內(nèi)建函數(shù)訪問紋理對象

threadgroup：線程組地址空間修飾符

  • 線程組地址空間用于為并行計算著色器函數(shù)分配內(nèi)存變量穗慕，這些變量被一個線程組的所有線程共享，在線程組地址空間分配的變量不能用于圖形繪制著色函數(shù)（即頂點著色函數(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：常量地址空間修飾符

  • 常量地址空間指向的緩存對象也是從設備內(nèi)存池分配存儲倔韭，僅可讀
  • 在程序域的變量必須定義在常量地址空間并且聲明時初始化术浪，用來初始化的值必須是編譯時的常量
  • 在程序域的變量的生命周期和程序一樣，在程序中的并行計算著色函數(shù) 或者 圖形繪制著色函數(shù)調(diào)用寿酌，但是constant的值會保持不變

1胰苏、常量地址空間的指針/引用可以作為函數(shù)的參數(shù)，向聲明為常量的變量賦值會產(chǎn)生編譯錯誤
2醇疼、聲明常量但是沒有賦予初值也會產(chǎn)生編譯錯誤

constant float samples[] = { 1.0f, 2.0f, 3.0f, 4.0f };

//對一個常量地址空間的變量進行修改也會失敗,因為它只讀的
sampler[4] = {3,3,3,3}; //編譯失敗; 

//定義為常量地址空間聲明時不賦初值也會編譯失敗
constant float a;

thread：線程地址空間修飾符

  • 線程地址空間指向每個線程準備的地址空間硕并，也是在GPU中法焰，該線程的地址空間定義的變量在其他線程不可見（即變量不共享）
  • 在圖形繪制著色函數(shù) 或者 并行計算著色函數(shù)中聲明的變量，在線程地址空間分配存儲
kernel void CCTestFouncitionG(void)
{
    //在線程空間分配空間給x,p
    float x;
    thread float p = &x;
}

函數(shù)參數(shù)與變量的傳遞修飾符倔毙，即屬性修飾符

圖形繪制 或者 并行計算著色器函數(shù)的輸入輸出都是通過參數(shù)傳遞埃仪，除了常量地址空間變量和程序域定義的采樣器之外, 其他參數(shù)修飾的可以是如下之一，有以下5種屬性修飾符：

  • device buffer 設備緩存：一個指向設備地址空間的任意數(shù)據(jù)類型的指針/引用
  • constant buffer 常量緩存：一個指向常量地址空間的任意數(shù)據(jù)類型的指針/引用
  • texture 紋理對象
  • sampler 采樣器對象
  • threadGroup 在線程組中供線程共享的緩存

為什么需要屬性修飾符陕赃？

  • 參數(shù)表示資源的定位卵蛉，可以理解為端口，相當于OpenGl ES中的location
  • 在固定管線和可編程管線進行內(nèi)建變量的傳遞
  • 將數(shù)據(jù)沿著渲染管線從頂點函數(shù)傳遞到片元函數(shù)

傳遞修飾符在代碼中的體現(xiàn)

對于每個著色函數(shù)來說么库，一個修飾符是必須指定的傻丝，它用來設置一個緩存、紋理诉儒、采樣器的位置葡缰，傳遞修飾符對應的寫法如下：

  • device buffer ---> [[buffer(index)]]
  • constant buffer ---> [[buffer(index)]]
  • texture---> [[texture(index)]]
  • sampler --->[[sampler(index)]]
  • threadGroup ---> [[threadGroup(index)]]
    在代碼中的表現(xiàn)如下：
在代碼中如何表現(xiàn):
 1.已知條件:device buffer(設備緩存)/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ā)者來指定.

1、index是一個unsigned interger類型的值允睹，表示了一個緩存运准、紋理幌氮、采樣器參數(shù)的位置（即在函數(shù)參數(shù)索引表中的位置缭受，相當于OpenGl ES中的location）
2、從語法上來說该互，屬性修飾符的聲明位置應該位于參數(shù)變量名之后

//并行計算著色器函數(shù)add_vectros ,實現(xiàn)2個設備地址空間中的緩存A與緩存B相加.然后將結(jié)果寫入到緩存out.
//屬性修飾符"(buffer(index))" 為著色函數(shù)參數(shù)設定了緩存的位置
//thread_position_in_grid：用于表示當前節(jié)點在多線程網(wǎng)格中的位置,并不需要開發(fā)者傳遞米者，是Metal自帶的。
/*
 kernel：并行計算函數(shù)修飾符
 void：函數(shù)返回值類型
 add_vectros：函數(shù)名
 const device float4 *inA [[buffer(0)]]：定義了一個float4類型的指針宇智，指向一個4維向量空間蔓搞，放在設備內(nèi)存空間（即顯存GPU中）
    - const device：只決定放在哪里
    - inA：變量名
    - [[buffer(0)]] 對應 buffer中0這個id
 */
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];
}

//著色函數(shù)的多個參數(shù)使用不同類型的屬性修飾符的情況
//紋理讀取的方式的sampler，即采樣器随橘，[[sampler(0)]]表示采樣器的緩存id
kernel void my_kernel(device float4 *p [[buffer(0)]],
                      texture2d<float> img [[texture(0)]],
                      sampler sam [[sampler(0)]])
{
    //.....
    
}

常見的內(nèi)建變量修飾符

  • [[vertex_id]] ：頂點id標識符喂分，并不由開發(fā)者傳遞
  • [[position]]
    • 在頂點著色函數(shù)中，表示當前的頂點信息机蔗，類型是float4蒲祈、
    • 還可以表示描述了片元的窗口的相對坐標（x，y萝嘁，z梆掸，1/w），即該像素點在屏幕上的位置信息
  • [[point_size]] ：點的大小牙言，類型是float
  • [[color(m)]]：顏色酸钦，m在編譯前就必須確定
//定義了片元輸入的結(jié)構(gòu)體，
struct MyFragmentOutput {
      // color attachment 0 顏色附著點0
     float4 clr_f [[color(0)]]; 
     // color attachment 1 顏色附著點1
     int4 clr_i [[color(1)]]; 
     // color attachment 2 顏色附著點2
     uint4 clr_ui [[color(2)]]; 
};

fragment MyFragmentOutput my_frag_shader( ... ) 
{
    MyFragmentOutput f;
    ....
    f.clr_f = ...;
    ....
    return f; 
}
  • [[stage_in]]：片元著色函數(shù)使用的單個片元輸入數(shù)據(jù)是由頂點著色函數(shù)輸出然后經(jīng)過光柵化生成的（即由頂點著色函數(shù)之后的顏色傳遞到片元著色函數(shù)）咱枉，類似于GLSL中的varying傳遞紋理/顏色
    • 頂點和片元著色器函數(shù)都只能有一個參數(shù)被聲明為使用stage_in修飾符（即有且僅有一個）
    • 對于一個使用了stage_in修飾符的自定義結(jié)構(gòu)體卑硫，其成員可以為一個整型/浮點類型標量徒恋，或是整型/浮點類型向量
?著作權(quán)歸作者所有,轉(zhuǎn)載或內(nèi)容合作請聯(lián)系作者
  • 人面猴
    序言：七十年代末，一起剝皮案震驚了整個濱河市拔恰，隨后出現(xiàn)的幾起案子因谎，更是在濱河造成了極大的恐慌，老刑警劉巖颜懊，帶你破解...
    沈念sama閱讀 219,490評論 6贊 508
  • 死咒
    序言：濱河連續(xù)發(fā)生了三起死亡事件财岔，死亡現(xiàn)場離奇詭異，居然都是意外死亡河爹，警方通過查閱死者的電腦和手機匠璧，發(fā)現(xiàn)死者居然都...
    沈念sama閱讀 93,581評論 3贊 395
  • 救了他兩次的神仙讓他今天三更去死
    文/潘曉璐 我一進店門，熙熙樓的掌柜王于貴愁眉苦臉地迎上來咸这，“玉大人夷恍，你說我怎么就攤上這事∠蔽” “怎么了酿雪？”我有些...
    開封第一講書人閱讀 165,830評論 0贊 356
  • 道士緝兇錄：失蹤的賣姜人
    文/不壞的土叔 我叫張陵，是天一觀的道長侄刽。 經(jīng)常有香客問我指黎，道長，這世上最難降的妖魔是什么州丹？ 我笑而不...
    開封第一講書人閱讀 58,957評論 1贊 295
  • ?港島之戀（遺憾婚禮）
    正文 為了忘掉前任醋安，我火速辦了婚禮，結(jié)果婚禮上墓毒，老公的妹妹穿的比我還像新娘吓揪。我一直安慰自己，他們只是感情好所计，可當我...
    茶點故事閱讀 67,974評論 6贊 393
  • 惡毒庶女頂嫁案：這布局不是一般人想出來的
    文/花漫 我一把揭開白布柠辞。 她就那樣靜靜地躺著，像睡著了一般主胧。 火紅的嫁衣襯著肌膚如雪叭首。 梳的紋絲不亂的頭發(fā)上，一...
    開封第一講書人閱讀 51,754評論 1贊 307
  • 城市分裂傳說
    那天讥裤，我揣著相機與錄音放棒，去河邊找鬼。 笑死己英，一個胖子當著我的面吹牛间螟，可吹牛的內(nèi)容都是我干的。 我是一名探鬼主播，決...
    沈念sama閱讀 40,464評論 3贊 420
  • 雙鴛鴦連環(huán)套：你想象不到人心有多黑
    文/蒼蘭香墨 我猛地睜開眼厢破，長吁一口氣：“原來是場噩夢啊……” “哼荣瑟！你這毒婦竟也來了？” 一聲冷哼從身側(cè)響起摩泪，我...
    開封第一講書人閱讀 39,357評論 0贊 276
  • 萬榮殺人案實錄
    序言：老撾萬榮一對情侶失蹤笆焰，失蹤者是張志新（化名）和其女友劉穎，沒想到半個月后见坑，有當?shù)厝嗽跇淞掷锇l(fā)現(xiàn)了一具尸體嚷掠，經(jīng)...
    沈念sama閱讀 45,847評論 1贊 317
  • ?護林員之死
    正文 獨居荒郊野嶺守林人離奇死亡，尸身上長有42處帶血的膿包…… 初始之章·張勛 以下內(nèi)容為張勛視角 年9月15日...
    茶點故事閱讀 37,995評論 3贊 338
  • ?白月光啟示錄
    正文 我和宋清朗相戀三年荞驴，在試婚紗的時候發(fā)現(xiàn)自己被綠了不皆。 大學時的朋友給我發(fā)了我未婚夫和他白月光在一起吃飯的照片。...
    茶點故事閱讀 40,137評論 1贊 351
  • 活死人
    序言：一個原本活蹦亂跳的男人離奇死亡熊楼，死狀恐怖霹娄，靈堂內(nèi)的尸體忽然破棺而出，到底是詐尸還是另有隱情鲫骗，我是刑警寧澤犬耻，帶...
    沈念sama閱讀 35,819評論 5贊 346
  • ?日本核電站爆炸內(nèi)幕
    正文 年R本政府宣布，位于F島的核電站执泰，受9級特大地震影響枕磁，放射性物質(zhì)發(fā)生泄漏。R本人自食惡果不足惜坦胶，卻給世界環(huán)境...
    茶點故事閱讀 41,482評論 3贊 331
  • 男人毒藥：我在死后第九天來索命
    文/蒙蒙 一透典、第九天 我趴在偏房一處隱蔽的房頂上張望晴楔。 院中可真熱鬧顿苇，春花似錦、人聲如沸税弃。這莊子的主人今日做“春日...
    開封第一講書人閱讀 32,023評論 0贊 22
  • 一樁弒父案，背后竟有這般陰謀
    文/蒼蘭香墨 我抬頭看了看天上的太陽则果。三九已至幔翰，卻和暖如春，著一層夾襖步出監(jiān)牢的瞬間西壮，已是汗流浹背遗增。 一陣腳步聲響...
    開封第一講書人閱讀 33,149評論 1贊 272
  • 情欲美人皮
    我被黑心中介騙來泰國打工， 沒想到剛下飛機就差點兒被人妖公主榨干…… 1. 我叫王不留款青，地道東北人做修。 一個月前我還...
    沈念sama閱讀 48,409評論 3贊 373
  • 代替公主和親
    正文 我出身青樓，卻偏偏與公主長得像，于是被迫代替她去往敵國和親饰及。 傳聞我的和親對象是個殘疾皇子蔗坯，可洞房花燭夜當晚...
    茶點故事閱讀 45,086評論 2贊 355