Metal Shading Language

• Metal着⾊语⾔是⽤来编写 3D 图形渲染逻辑和并⾏计算核⼼逻辑的⼀⻔编程语⾔。当使⽤ Metal 框架来完成APP 的实现时则需要使⽤Metal 编程语⾔；
• Metal语言使用的是Clang 和LLVM进行编译处理的；
• Metal 基于C++ 11.0 语⾔设计。主要⽤来编写 在 GPU 上执⾏的图像渲染逻辑代码 以及 通⽤ 并⾏计算逻辑代码;

Metal 与C++ 11.0 相爱相杀

虽然Metal基于C++,但它很苹果，并不是支持C++11.0所以特性的，以下Metal暂时用不上，它就不支持：

• 递归函数调用；
• Lambda 表达式;
• 动态转换操作符
• 类型识别
• 对象创建new 和销毁delete 操作符;
• 操作符 noexcept；
• goto 跳转；
• 变量存储修饰符register 和 thread_local;
• 虚函数修饰符;
• 派⽣类；
• 异常处理；
• C++ 标准库在Metal 语⾔中也不可使⽤;

Metal使用指针的局限

• Metal图形和并⾏计算函数⽤到的⼊参数; 如果是指针必须使⽤地址空间修饰符 (device,threadgroup,constant)；
• 不支持函数指针；
• 函数名不能出现main；

Metal 像素坐标系统: Metal 中纹理/帧缓存区attachment 的像素使⽤的坐标系统的原点是左上 ⻆;

Metal 数据类型

标量数据类型

标量数据类型

纹理数据类型（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, short, int 等;

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

采样器类型（Samplers）

采取器类型决定了如何对⼀个纹理进⾏采样操作. 在Metal 框架中有⼀个对应着⾊器语⾔的采样器的对象 MTLSamplerState 这个对象作为图形渲染着⾊器函数参数或是并⾏计算函数的参数传递;

采样器状态及默认值

/*
constexpr：修饰符（必须写）
sampler：类型
s：采样器变量名称
coord: 是否需要归一化，不需要归一化，用的是像素pixel
address: 地址环绕方式
filter: 过滤方式
*/
 constexpr sampler s(coord::pixel,address::clamp_to_zero,filter::linear);
 constexpr sampler a(coord::normalized);
 constexpr sampler b(address::repeat);

函数修饰符

Metal的3种函数修饰符：

• Kernel：表示该函数是⼀个数据并⾏计算着⾊函数. 它可以被分配在⼀维/⼆维/三维线程组中去执⾏;
• vertex：表示该函数是⼀个顶点着⾊函数 , 它将为顶点数据流中的每个顶点数据执⾏⼀次然后为每个顶 点⽣成数据输出到绘制管线;
• fragment , 表示该函数是⼀个⽚元着⾊函数, 它将为⽚元数据流中的每个⽚元 和其关联执⾏⼀次然后 将每个⽚元⽣成的颜⾊数据输出到绘制管线中;

注意：

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

2. 一个被函数修饰符修饰过的函数,不允许在调用其他的被函数修饰过的函数.

3. 被函数修饰符修饰过的函数,只允许在客户端对齐进行操作. 不允许被普通的函数调用.

地址空间修饰符

它是⽤于变量或者参数，Metal 着⾊器语⾔使⽤它来表示⼀个函数变量或者参数变量 被分配于那⼀⽚内存区域. 所有的着⾊函数(vertex, fragment, kernel)的参数,如果是指针或是引⽤, 都必须带有地址空间修饰符号;

• 1.device
• 2.threadgroup
• 3.constant
• 4.thread

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

Device Address Space(设备地址空间)

在设备地址空间(Device) 指向设备内存池分配出来的缓存对象, 它是可读也是可写的; ⼀个缓存对象可 以被声明成⼀个标量,向量或是⽤户⾃定义结构体的指针或是引⽤.

注意: 纹理对象总是在设备地址空间分配内存, device 地址空间修饰符不必出现在纹理类型定义中. ⼀个纹 理对象的内容⽆法直接访问. Metal 提供读写纹理的内建函数;

线程组地址空间（threadgrounp Address Space）

线程组地址空间⽤于为 并⾏计算着⾊函数分配内存变量. 这些变量被⼀个线程组的所有线程共享. 在线 程组地址空间分配的变量不能被⽤于图形绘制着⾊函数[顶点着⾊函数, ⽚元着⾊函数] 在并⾏计算着⾊函数中, 在线程组地址空间分配的变量为⼀个线程组使⽤, 声明周期和线程组相同;

函数参数与变量

图形绘制或者并⾏计算着⾊器函数的输⼊输出都是通过参数传递. 除了常量地址空间变量和程序域定义 的采样器以外.

• device buffer- 设备缓存, ⼀个指向设备地址空间的任意数据类型的指针或者引⽤;
• constant buffer -常量缓存区, ⼀个指向常量地址空间的任意数据类型的指针或引⽤；
• texture - 纹理对象;
• sampler - 采样器对象;
• threadGrounp - 在线程组中供各线程共享的缓存.

注意: 被着⾊器函数的缓存(device 和 constant) 不能重名;

对于每个着⾊器函数来说, ⼀个修饰符是必须指定的. 他⽤来设定⼀个缓存,纹理, 采样器的位置：

• device buffers/ constant buffer --> [[buffer (index)]]
• texture -- [[texture (index)]]
• sampler -- [[sampler (index)]]
• threadgroup buffer -- [[threadgroup (index)]]

index是⼀个unsigned integer类型的值，它表示了⼀个缓存、纹理、采样器参数的位置(在函数参数索引 表中的位置)。 从语法上讲，属性修饰符的声明位置应该位于参数变量名之后

//并行计算着色器函数add_vectros ,实现2个设备地址空间中的缓存A与缓存B相加.然后将结果写入到缓存out.
//属性修饰符"(buffer(index))" 为着色函数参数设定了缓存的位置
// thread_position_in_grid : ⽤于表示当前节点在多线程⽹格中的位置;
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];
}

内建变量属性修饰符

• [[vertex_id]] 顶点id 标识符;
• [[position]] 顶点信息(float4) /� 述了⽚元的窗⼝相对坐标(x, y, z, 1/w) ；
• [[point_size]] 点的⼤⼩(float) ；
• [[color(m)]] 颜⾊, m编译前得确定;
• [[stage_in]] : ⽚元着⾊函数使⽤的单个⽚元输⼊数据是由顶点着⾊函数输出然后经过光栅化⽣成的.顶 点和⽚元着⾊函数都是只能有⼀个参数被声明为使⽤“stage_in”修饰符，对于⼀个使⽤ 了“stage_in”修饰符的⾃ 定义的结构体，其成员可以为⼀个整形或浮点标量，或是整形或浮点向量

//定义了片元输入的结构体，
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; 
}