Metal着色器语⾔规范

Metal语言介绍

Metal着⾊语⾔是⼀个⽤来编写3D图形渲染逻辑和并⾏计算核⼼逻辑的编程语⾔,编写Metal框架的APP需要使⽤Metal着⾊语⾔程序.
Metal着⾊语⾔与Metal 框架配合使⽤,Metal框架管理Metal着⾊语⾔的运
⾏和可选编译选项. Metal 着⾊器语⾔使⽤Clang和LLVM,编译器对于在GPU上的代码执⾏效率有更好的控制.

Metal Restrictions 限制

如下的C++11.0的特性在Metal 着⾊语⾔中是不⽀持的

Lambda表达式
递归函数调用
动态转换操作符
类型识别
对象创建(new)和释放(delloc)操作符
操作符 noexcept
goto跳转
虚函数修饰符
派生类
异常处理

c++的标准库也不可以在Metal着色器语言中使用
Metal着色器语言对于指针的使用限制

Metal图形和并行计算函数用到的入参如果是指针使用地址空间修饰符(device ,threadgroup ,constant)
不支持函数指针
Metal函数名不能命名为main函数

Metal数据类型

标量

Metal标量数据类型.png
向量

booln ( true/false)
charn (有符号8位整数)
shortn (有符号16位整数)
intn (有符号int型)
ucharn (⽆符号8位整数)
ushortn (⽆符号32位整数)
uintn (无符号int型)
halfn (16位bit浮点数)
floatn (32位bit浮点数)
向量中的n,指的是维度

矩阵

halfnxm
floatnxm
nxm分别指的是矩阵的⾏数和列数

缓存buffer

在Metal 中实现缓存靠的是⼀个指针.它指向⼀个在Device 或者constant 地址空间中的内建或是开发者⾃定义的数据块.缓存可以被定在程序域域中,或是当做函数的参数传递.

纹理Textures和采样器

纹理Textures
纹理数据类型是⼀个句柄,它是⼀个⼀维/⼆维/三维纹理数据.⽽纹理数据对
应这个⼀个纹理的某个level的mipmap的全部或者⼀部分。

enum class access {sample ,read ,write}; 
texture1d<T, access a = access::sample> 
texture1d_array<T, access a = access::sample> 
texture2d<T, access a = access::sample> 
texture2d_array<T, access a = access::sample> 
texture3d<T, access a = access::sample> 
texturecube<T, access a = access::sample> 
texture2d_ms<T, access a = access::read>

void foo (texture2d<float> imgA [[ texture(0) ]] , 
texture2d<float, access::read> imgB [[ texture(1) ]], 
texture2d<float, access::write> imgC [[ texture(2) ]]) 
{ 
    ... 
}

采样器

Metal采样器.png

enum class coord { normalized, pixel }; 
从纹理中采样时,纹理坐标是否需要归⼀化;

enum class filter { nearest, linear }; 
纹理采样过滤⽅式, 放⼤/缩⼩过滤模式; 

enum class min_filter { nearest, linear }; 
设置纹理采样的缩⼩过滤模式; 

enum class mag_filter { nearest, linear }; 
设置纹理采样的放⼤过滤模式; 

enum class s_address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat }; 
enum class t_address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat }; 
enum class r_address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat }; 
设置纹理s,t,r坐标的寻址模式; 

enum class address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat }; 
设置所有的纹理坐标的寻址模式; 

enum class mip_filter { none, nearest, linear }; 
设置纹理采样的mipMap过滤模式, 如果是none,那么只有⼀层纹理⽣效;

注意: 在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,

Metal函数修饰符

Metal着色器语言支持的函数修饰符：

kernel,表示该函数是一个数据并行计算着色函数，它将被分配在一个一维/二维/三维的线程组网格中执行.
vertex,表示该函数是⼀个顶点着⾊函数,它将为顶点数据流中的每个顶点数据执⾏⼀次然后为每个顶点⽣成数据输出到绘制管线中去.
fragment,表示该函数是⼀个⽚元着⾊函数,它将⽚元数据流中的每个⽚元和其关联的数据执⾏⼀次然后为每个⽚元⽣成数据输出到绘制管线中去.

注意: 使⽤kernel 修饰的函数. 其返回值类型必须是void 类型;
只有图形着⾊函数才可以被 vertex 和 fragment 修饰. 对于图形着⾊函数, 返回值类型可以辨认出它是为顶点做计算还是为每像素做计算. 图形着⾊函数的返回值可以为 void , 但是这也就意味着该函数不产⽣数据输出到绘制管线; 这是⼀个⽆意义的动作;

注意: ⼀个被函数修饰符修饰的函数不能在调⽤其他也被函数修饰符修饰的函数; 这样会导致编译失败;

kernel void hello2(...){  
} 
vertex float4 hello1(...){
     //⼀个被函数修饰符修饰的函数不能在调⽤其他也被函数修饰符修饰的函数; 这样会 导致编译失败; 
    hello2(...); //错误调⽤❌  
}

变量与参数的地址空间修饰符

Metal 着⾊语⾔使⽤"地址空间修饰符号",来表⽰⼀个函数变量或是参数被分于那块内存区域. 下⾯这些修饰符描述了不相交叠地址空间：

device— 设备地址空间
纹理对象总是在设备地址空间分配内存, device 地址空间修饰符不必出现在纹理类型定义中. ⼀个纹理对象的内容⽆法直接访问. Metal 提供读写纹理的内建函数;
threadgroup — 线程组地址空间
线程组地址空间⽤于为并⾏计算着⾊函数分配内存变量. 这些变量被⼀个线程组的所有线程共享. 在线程组地址空间分配的变量不能被⽤于图形绘制着⾊函数[顶点着⾊函数, ⽚元着⾊函数] . 在并⾏计算着⾊函数中, 在线程组地址空间分配的变量为⼀个线程组使⽤, 声明周期和线程组相同;
constant — 常量地址空间
常量地址空间的指针或是引⽤可以作为函数的参数. 向声明为常量的变量赋值会产⽣编译错误. 声明常量但是没有赋予初值也会产⽣编译错误;
thread — thread地址空间
thread 地址空间指向每个线程准备的地址空间, 这个线程的地址空间定义的变量在其他线程不可⻅, 在图形绘制着⾊函数或者并⾏计算着⾊函数中声明的变量thread 地址空间分配;

对于图形着⾊器函数, 其指针或是引⽤类型的参数必须定义为 device 或是 constant 地址空间; 对于并⾏计算着⾊函数, 其指针或是引⽤类型的参数必须定义为 device 或是 threadgrounp 或是 constant 地址空间;

函数参数和变量

图形绘制或是并⾏计算着⾊函数的输⼊/输出都需要通过参数传递(除了常量地址空间变量和程序域中定义的采样器以外).参数可以是如下之⼀:

device buffer — 设备缓存
constant buffer— 常量缓存
texture object — 纹理对象
sample object — 采样器对象
threadgroup — 线程共享的缓存

用于寻址缓存，纹理，采样器属性修饰符

为每个参数,指定⼀个属性修饰符. 指定明确的缓冲,纹理位置.

device 和 constant buffers — [[ buffer ( index ) ]]
texture — [[ texture ( index ) ]]
sampler — [[ sampler ( index ) ]]
threadgroup buffer — [[ threadgroup ( index ) ]]

注意: 被着⾊器函数的缓存(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))”为着⾊函数参数设定了缓存的位置。
 kernel void add_vectors(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];
}
//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”修饰符的⾃定义的结构体，其成员可以为⼀个整形或浮点标量，或是整形或浮点向量