iOS Metal语言规范（二）
一.函数修饰符
Metal 有以下3种函数修饰符:
1）kernel , 表示该函数是⼀个数据并⾏计算着⾊函数. 它可以被分配在⼀维/⼆维/三维线程组中去执⾏
2）vertex , 表示该函数是⼀个顶点着⾊函数 , 它将为顶点数据流中的每个顶点数据执⾏⼀次然后为每个顶 点⽣成数据输出到绘制管线
3）fragment , 表示该函数是⼀个⽚元着⾊函数, 它将为⽚元数据流中的每个⽚元 和其关联执⾏⼀次然后 将每个⽚元⽣成的颜⾊数据输出到绘制管线中;
eg：
//1.并行计算函数(kernel)
kernelvoidCCTestKernelFunctionA(inta,intb)
{
/*
注意:
1. 使用kernel 修饰的函数返回值必须是void 类型
2. 一个被函数修饰符修饰过的函数,不允许在调用其他的被函数修饰过的函数. 非法
3. 被函数修饰符修饰过的函数,只允许在客户端对齐进行操作. 不允许被普通的函数调用.
*/
//不可以的!
//一个被函数修饰符修饰过的函数,不允许在调用其他的被函数修饰过的函数. 非法
CCTestKernelFunctionB(1,2);//非法
CCTestVertexFunctionB(1,2);//非法
//可以! 你可以调用普通函数.而且在Metal 不仅仅只有这3种被修饰过的函数.普通函数也可以存在
CCTest();
}
kernelvoidCCTestKernelFunctionB(inta,intb)
{
}
//顶点函数
vertexintCCTestVertexFunctionB(inta,intb){
}
//片元函数
fragmentintCCTestVertexFunctionB(inta,intb){
}
//普通函数
voidCCTest()
{
}
说明：使⽤kernel 修饰的函数. 其返回值类型必须是void 类型; 只有图形着⾊函数才可以被 vertex 和 fragment 修饰. 对于图形着⾊函数, 返回值类型可以辨认出它是为 顶点做计算还是为每像素做计算. 图形着⾊函数的返回值可以为 void , 但是这也就意味着该函数不产⽣数 据输出到绘制管线; 这是⼀个⽆意义的动作; ⼀个被函数修饰符修饰的函数不能在调⽤其他也被函数修饰符修饰的函数; 这样会导致编译失败;
二.⽤于变量或者参数的地址空间修饰符
Metal 着⾊器语⾔使⽤ 地址空间修饰符 来表示⼀个函数变量或者参数变量 被分配于那⼀⽚内存区域. 所有的着⾊函数(vertex, fragment, kernel)的参数,如果是指针或是引⽤, 都必须带有地址空间修饰符号;
1） device： 设备地址空间
Device Address Space(设备地址空间) ，在设备地址空间(Device) 指向设备内存池分配出来的缓存对象, 它是可读也是可写的; ⼀个缓存对象可 以被声明成⼀个标量,向量或是⽤户⾃定义结构体的指针或是引⽤.
eg：
// 设备地址空间: device 用来修饰指针.引用
//1.修饰指针变量
device float4 *color;
structCCStruct{
floata[3];
intb[2];
};
//2.修饰结构体类的指针变量
device CCStruct*my_CS;
注意: 纹理对象总是在设备地址空间分配内存, device 地址空间修饰符不必出现在纹理类型定义中. ⼀个纹 理对象的内容⽆法直接访问. Metal 提供读写纹理的内建函数;
2）threadgroup： 线程组地址空间
线程组地址空间⽤于为 并⾏计算着⾊函数分配内存变量. 这些变量被⼀个线程组的所有线程共享. 在线 程组地址空间分配的变量不能被⽤于图形绘制着⾊函数[顶点着⾊函数, ⽚元着⾊函数] 在并⾏计算着⾊函数中, 在线程组地址空间分配的变量为⼀个线程组使⽤, 声明周期和线程组相同;
eg：
/*
1. threadgroup 被并行计算计算分配内存变量, 这些变量被一个线程组的所有线程共享. 在线程组分配变量不能被用于图像绘制.
2. thread 指向每个线程准备的地址空间. 在其他线程是不可见切不可用的
*/
kernel void CCTestFouncitionF(threadgroup float *a)
{
//在线程组地址空间分配一个浮点类型变量x
threadgroup float x;
//在线程组地址空间分配一个10个浮点类型数的数组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 常量地址空间
常量地址空间指向的缓存对象也是从设备内存池分配存储, 但是它是只读的; 在程序域的变量必须定义在常量地址空间并且声明的时候初始化; ⽤来初始化的值必须是编译时的常 量. 在程序域的变量的⽣命周期和程序⼀样, 在程序中的并⾏计算着⾊函数或者图形绘制着⾊函数调⽤, 但 是constant 的值会保持不变;
注意: 常量地址空间的指针或是引⽤可以作为函数的参数. 向声明为常量的变量赋值会产⽣编译错误. 声明常量但是没有赋予初值也会产⽣编译错误;
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 地址空间指向每个线程准备的地址空间, 这个线程的地址空间定义的变量在其他线程不可⻅, 在 图形绘制着⾊函数或者并⾏计算着⾊函数中声明的变量thread 地址空间分配; 在图形绘制着色函数 或者 并行计算着色函数中声明的变量，在线程地址空间分配存储
eg：
kernel void CCTestFouncitionG(void){
//在线程空间分配空间给x,p
float x;
thread float p=&x;
}
对于图形着⾊器函数, 其指针或是引⽤类型的参数必须定义为 device 或是 constant 地址空间; 对于并⾏计算着⾊函数, 其指针或是引⽤类型的参数必须定义为 device 或是 threadgrounp 或是 constant 地址空间; 并不是所有的变量都需要修饰符，也可以定义普通变量（即无修饰符的变量）
三.函数参数与变量
图形绘制或者并⾏计算着⾊器函数的输⼊输出都是通过参数传递. 除了常量地址空间变量和程序域定义 的采样器以外.
device buffer- 设备缓存, ⼀个指向设备地址空间的任意数据类型的指针或者引⽤;
constant buffer -常量缓存区, ⼀个指向常量地址空间的任意数据类型的指针或引⽤
texture - 纹理对象;
sampler - 采样器对象;
threadGrounp - 在线程组中供各线程共享的缓存.
注意: 被着⾊器函数的缓存(device 和 constant) 不能重名;
Attribute Qualifiers to Locate Buffers, Textures, and Samplers ⽤于寻址缓存,纹理, 采样器的属性修饰符;对于每个着⾊器函数来说, ⼀个修饰符是必须指定的. 他⽤来设定⼀个缓存,纹理, 采样器的位置;
device buffers/ constant buffer --> [[buffer (index)]]
texture -- [[texture (index)]]
sampler -- [[sampler (index)]]
threadgroup buffer -- [[threadgroup (index)]]
index是⼀个unsigned integer类型的值，它表示了⼀个缓存、纹理、采样器参数的位置(在函数参数索引 表中的位置)。 从语法上讲，属性修饰符的声明位置应该位于参数变量名之后
例⼦中展示了⼀个简单的并⾏计算着⾊函数 add_vectors，它把两个设备地址空间中的缓存inA和inB相 加，然后把结果写⼊到缓存out。属性修饰符 “(buffer(index))”为着⾊函数参数设定了缓存的位置。
//属性修饰符
/*
1. device buffer(设备缓存)
2. constant buffer(常量缓存)
3. texture Object(纹理对象)
4. sampler Object(采样器对象)
5. 线程组 threadgroup
属性修饰符目的:
1. 参数表示资源如何定位? 可以理解为端口
2. 在固定管线和可编程管线进行内建变量的传递
3. 将数据沿着渲染管线从顶点函数传递片元函数.
在代码中如何表现:
1.已知条件:device buffer(设备缓存)/constant buffer(常量缓存)
代码表现:[[buffer(index)]]
解读:不变的buffer ,index 可以由开发者来指定.
2.已知条件:texture Object(纹理对象)
代码表现: [[texture(index)]]
解读:不变的texture ,index 可以由开发者来指定.
3.已知条件:sampler Object(采样器对象)
代码表示: [[sampler(index)]]
解读:不变的sampler ,index 可以由开发者来指定.
4.已知条件:threadgroup Object(线程组对象)
代码表示: [[threadgroup(index)]]
解读:不变的threadgroup ,index 可以由开发者来指定.
*/
//并行计算着色器函数add_vectros ,实现2个设备地址空间中的缓存A与缓存B相加.然后将结果写入到缓存out.
//属性修饰符"(buffer(index))" 为着色函数参数设定了缓存的位置
//并行计算着色器函数add_vectros ,实现2个设备地址空间中的缓存A与缓存B相加.然后将结果写入到缓存out.
//属性修饰符"(buffer(index))" 为着色函数参数设定了缓存的位置
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 : ⽤于表示当前节点在多线程⽹格中的位置;
四.内建变量属性修饰符
[[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]] : ⽚元着⾊函数使⽤的单个⽚元输⼊数据是由顶点着⾊函数输出然后经过光栅化⽣成的.顶 点和⽚元着⾊函数都是只能有⼀个参数被声明为使⽤“stage_in”修饰符，对于⼀个使⽤ 了“stage_in”修饰符的⾃ 定义的结构体，其成员可以为⼀个整形或浮点标量，或是整形或浮点向量
