一.函數(shù)修飾符

Metal 有以下3種函數(shù)修飾符:

? ? ? ? 1）kernel , 表示該函數(shù)是?個數(shù)據(jù)并?計算著?函數(shù). 它可以被分配在?維/?維/三維線程組中去執(zhí)?

? ? ? ? 2）vertex , 表示該函數(shù)是?個頂點著?函數(shù) , 它將為頂點數(shù)據(jù)流中的每個頂點數(shù)據(jù)執(zhí)??次然后為每個頂 點?成數(shù)據(jù)輸出到繪制管線

? ? ? ? 3）fragment , 表示該函數(shù)是?個?元著?函數(shù), 它將為?元數(shù)據(jù)流中的每個?元 和其關(guān)聯(lián)執(zhí)??次然后 將每個?元?成的顏?數(shù)據(jù)輸出到繪制管線中;?

eg：

//1.并行計算函數(shù)(kernel)

kernelvoidCCTestKernelFunctionA(inta,intb)

{

? ? /*

?? ? 注意:

?? ? 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();

}

kernelvoidCCTestKernelFunctionB(inta,intb)

{

}

//頂點函數(shù)

vertexintCCTestVertexFunctionB(inta,intb){

}

//片元函數(shù)

fragmentintCCTestVertexFunctionB(inta,intb){

}

//普通函數(shù)

voidCCTest()

{

}

????????說明：使?kernel 修飾的函數(shù). 其返回值類型必須是void 類型;?只有圖形著?函數(shù)才可以被 vertex 和 fragment 修飾. 對于圖形著?函數(shù), 返回值類型可以辨認出它是為 頂點做計算還是為每像素做計算. 圖形著?函數(shù)的返回值可以為 void , 但是這也就意味著該函數(shù)不產(chǎn)?數(shù) 據(jù)輸出到繪制管線; 這是?個?意義的動作;??個被函數(shù)修飾符修飾的函數(shù)不能在調(diào)?其他也被函數(shù)修飾符修飾的函數(shù); 這樣會導(dǎo)致編譯失敗;

二.?于變量或者參數(shù)的地址空間修飾符

? ??????Metal 著?器語?使? 地址空間修飾符 來表示?個函數(shù)變量或者參數(shù)變量 被分配于那??內(nèi)存區(qū)域. 所有的著?函數(shù)(vertex, fragment, kernel)的參數(shù),如果是指針或是引?, 都必須帶有地址空間修飾符號;?

? ? ? ? 1） device： 設(shè)備地址空間

????????Device Address Space(設(shè)備地址空間) 食侮，在設(shè)備地址空間(Device) 指向設(shè)備內(nèi)存池分配出來的緩存對象, 它是可讀也是可寫的; ?個緩存對象可 以被聲明成?個標量,向量或是?戶?定義結(jié)構(gòu)體的指針或是引?.?

eg：

????????// 設(shè)備地址空間: device 用來修飾指針.引用

????????//1.修飾指針變量

????????device float4 *color;

????????structCCStruct{

? ? ????????floata[3];

? ????????? intb[2];

????????};

????????//2.修飾結(jié)構(gòu)體類的指針變量

????????device ?CCStruct*my_CS;

????????注意: 紋理對象總是在設(shè)備地址空間分配內(nèi)存, device 地址空間修飾符不必出現(xiàn)在紋理類型定義中. ?個紋 理對象的內(nèi)容?法直接訪問. Metal 提供讀寫紋理的內(nèi)建函數(shù);?

? ? ? ? 2）threadgroup： 線程組地址空間

????????線程組地址空間?于為 并?計算著?函數(shù)分配內(nèi)存變量. 這些變量被?個線程組的所有線程共享. 在線 程組地址空間分配的變量不能被?于圖形繪制著?函數(shù)[頂點著?函數(shù), ?元著?函數(shù)] 在并?計算著?函數(shù)中, 在線程組地址空間分配的變量為?個線程組使?, 聲明周期和線程組相同;?

eg：

????????/*

?????????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;

}

? ? ? ?3） constant?常量地址空間

????????常量地址空間指向的緩存對象也是從設(shè)備內(nèi)存池分配存儲, 但是它是只讀的; 在程序域的變量必須定義在常量地址空間并且聲明的時候初始化; ?來初始化的值必須是編譯時的常 量. 在程序域的變量的?命周期和程序?樣, 在程序中的并?計算著?函數(shù)或者圖形繪制著?函數(shù)調(diào)?, 但 是constant 的值會保持不變;?

? ??????注意: 常量地址空間的指針或是引?可以作為函數(shù)的參數(shù). 向聲明為常量的變量賦值會產(chǎn)?編譯錯誤.?聲明常量但是沒有賦予初值也會產(chǎn)?編譯錯誤;?

eg：

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

????????2?//對?個常量地址空間的變量進?修改也會失敗,因為它只讀的?

????????3?sampler[4] = {3,3,3,3}; //編譯失敗;?

? ? ? ? 4 //定義為常量地址空間聲明時不賦初值也會編譯失敗?

? ? ? ? 5 constant float a;?

? ? ? ? 4）thread?線程地址空間

????????thread 地址空間指向每個線程準備的地址空間, 這個線程的地址空間定義的變量在其他線程不可?, 在 圖形繪制著?函數(shù)或者并?計算著?函數(shù)中聲明的變量thread 地址空間分配;?在圖形繪制著色函數(shù) 或者 并行計算著色函數(shù)中聲明的變量悉抵，在線程地址空間分配存儲

eg：

????????kernel void CCTestFouncitionG(void){

????????//在線程空間分配空間給x,p

????????float x;?

????????thread ? float ? p=&x;

????????}

? ??????對于圖形著?器函數(shù), 其指針或是引?類型的參數(shù)必須定義為 device 或是 constant 地址空間;?對于并?計算著?函數(shù), 其指針或是引?類型的參數(shù)必須定義為 device 或是 threadgrounp 或是?constant 地址空間;?并不是所有的變量都需要修飾符屁桑，也可以定義普通變量（即無修飾符的變量）

三.函數(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ù)變量名之后?

????????例?中展示了?個簡單的并?計算著?函數(shù) add_vectors，它把兩個設(shè)備地址空間中的緩存inA和inB相 加颜屠，然后把結(jié)果寫?到緩存out辰妙。屬性修飾符 “(buffer(index))”為著?函數(shù)參數(shù)設(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è)定了緩存的位置

kernelvoidadd_vectros(

? ? ? ? ? ? ? ? const device float4*inA [[buffer(0)]],

? ? ? ? ? ? ? ? const device float4*inB [[buffer(1)]],

? ? ? ? ? ? ? ? device float4*out [[buffer(2)]]

? ? ? ? ? ? ? ? uintid[[thread_position_in_grid]])

{

? ? out[id] = inA[id] + inB[id];

}

? ??? ??注意：thread_position_in_grid : ?于表示當(dāng)前節(jié)點在多線程?格中的位置;?

四.內(nèi)建變量屬性修飾符?

????????[[vertex_id]] 頂點id 標識符;?

????????[[position]] 頂點信息(float4) /??述了?元的窗?相對坐標(x, y, z, 1/w)?

????????[[point_size]] 點的??(float)?

????????[[color(m)]] 顏?, m編譯前得確定;?

????????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;

????????}

????????[[stage_in]] :??元著?函數(shù)使?的單個?元輸?數(shù)據(jù)是由頂點著?函數(shù)輸出然后經(jīng)過光柵化?成的.頂 點和?元著?函數(shù)都是只能有?個參數(shù)被聲明為使?“stage_in”修飾符甫窟，對于?個使? 了“stage_in”修飾符的? 定義的結(jié)構(gòu)體密浑，其成員可以為?個整形或浮點標量，或是整形或浮點向量

? ??????