Metal语言规范

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

Metal语言的限制

• C++ 11.0 特性在Metal 语⾔中不⽀持之处
  • Lambda 表达式
  • 递归函数调⽤
  • 动态转换操作符
  • 类型识别
  • 对象创建new 和销毁delete 操作符
  • 操作符 noexcept
  • goto 跳转
  • 变量存储修饰符register 和 thread_local
  • 虚函数修饰符
  • 派⽣类
  • 异常处理
  • C++ 标准库在Metal 语⾔中也不可使⽤
• Metal 语⾔中对于指针使⽤的限制
  • Metal图形和并⾏计算函数⽤到的⼊参数; 如果是指针必须使⽤地址空间修饰符 (device,threadgroup,constant)
  • 不⽀持函数指针
  • 函数名不能出现main
• Metal 像素坐标系统: Metal 中纹理/帧缓存区attachment 的像素使⽤的坐标系统的原点是左上 ⻆;

Metal数据类型

• bool 布尔类型, true/false
• char 有符号8位整数
• unsigned char /uchar ⽆符号8-bit 整数
• short 有符号16-bit整数
• unsigned short / ushort ⽆符号32-bit 整数
• half 16位bit 浮点数（没有double类型）
• float 32bit 浮点数
• size_t 64 ⽆符号整数
• void 该类型表示⼀个空的值集合

Metal数据类型.png

Metal向量与矩阵

• 向量支持如下类型:
  • booln
  • charn
  • shortn
  • intn
  • ucharn
  • ushortn
  • uintn
  • halfn
  • floatn

向量中的n,指的是维度

例如,n=4的时候，我们可以用int4表示一个4维向量

// int型的4维向量
int4 test = int4(0,1,2,3);

• 向量的相关操作
  获取其中的元素

//向量可以通过x,y,z,w或者r,g,b,a来访问属性，但是x,y,z,w和r,g,b,a不能混合使用
int4 test = int4(0,1,2,3);
//x表示第一个元素
int a = test.x;
//y表示第二个元素
int b = test.y;
//z表示第三个元素
int c = test.z;
//w表示第四个元素
int d = test.w;

//r表示第一个元素
int e = test.r;
//g表示第二个元素
int f = test.g;
//b表示第三个元素
int g = test.b;
//a表示第四个元素
int h = test.a;

获取其中元素组成的向量

//2维向量
int2 test1 = test.xy;
//3维向量
int3 test1 = test.xyz;

向量赋值

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);
//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);
//pos = (3.0f,5.0f,9.0f,7.0f);
pos.xyz = float3(3.0f,5.0f,9.0f);
//非法,x出现2次
pos.xx = float2(3.0,4.0f);
//不合法-使用混合限定符
pos.xy = float4(1.0f,2.0,3.0,4.0);

• 矩阵支持如下类型:
  • halfnxm
  • floatnxm
    nxm分别指的是矩阵的⾏数和列数（最大为4*4）

float4x4 m;
//将第二排的值设置为0
m[1] = float4(2.0f);
//设置第一行/第一列为1.0f
m[0][0] = 1.0f;
//设置第三行第四列的元素为3.0f
m[2][3] = 3.0f;

纹理Textures 类型

纹理类型是⼀个句柄, 它指向⼀个⼀维/⼆维/三维纹理数据

枚举值: 定义了访问权利

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

 //当access为sample类型时，access a = access::sample可以省略
 void foo (texture2d<float> imgA [[ texture(0) ]] , texture2d<float, access::read> imgB [[ texture(1) ]], texture2d<float, access::write> imgC [[ texture(2) ]])
  {  ...  }

采样器Samplers类型

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

• 从纹理中采样时,纹理坐标是否需要归⼀化

enum class coord { normalized, pixel };

• 纹理采样过滤⽅式, 放⼤/缩⼩过滤模式

enum class filter { nearest, linear };

• 设置纹理采样的缩⼩过滤模式

enum class min_filter { nearest, linear };

• 设置纹理采样的放⼤过滤模式

enum class mag_filter { nearest, linear };

• 设置纹理s,t,r坐标的寻址模式

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

• 设置所有的纹理坐标的寻址模式

enum class address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat };

-设置纹理采样的mipMap过滤模式, 如果是none,那么只有⼀层纹理⽣效

enum class mip_filter { none, nearest, linear };

注意: 在Metal 程序中初始化的采样器必须使⽤ constexpr 修饰符声明

代码示例

constexpr sampler s(coord::pixel,address::clamp_to_zero, filter::linear);

函数修饰符

Metal 有以下3种函数修饰符:

• kernel , 表示该函数是⼀个数据并⾏计算着⾊函数. 它可以被分配在⼀维/⼆维/三维线程组中去执⾏;

 kernel void foo(...)  
{
  ...  
}

注意: 使⽤kernel 修饰的函数. 其返回值类型必须是void 类型;

• vertex , 表示该函数是⼀个顶点着⾊函数 , 它将为顶点数据流中的每个顶点数据执⾏⼀次然后为每个顶 点⽣成数据输出到绘制管线;

//顶点函数
vertex int CCTestVertexFunctionB(int a,int b)
{
    ...  
}

• fragment , 表示该函数是⼀个⽚元着⾊函数, 它将为⽚元数据流中的每个⽚元 和其关联执⾏⼀次然后 将每个⽚元⽣成的颜⾊数据输出到绘制管线中;

//片元函数
fragment int CCTestVertexFunctionB(int a,int b)
{
    ...
}

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

⼀个被函数修饰符修饰的函数不能在调⽤其他也被函数修饰符修饰的函数; 这样会导致编译失败;
即kernel、vertex、fragment之间无法相互调用

kernel void hello(...)
{ ... } 

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

⽤于变量或者参数的地址空间修饰符

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

• device
• threadgroup
• constant
• thread

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

Device Address Space(设备地址空间)

  在设备地址空间(Device)指向设备(此处设备指的是显存)内存池分配出来的缓存对象, 它是可读也是可写的;
  ⼀个缓存对象可 以被声明成⼀个标量,向量或是⽤户⾃定义结构体的指针或是引⽤。
  由于device修饰的变量存放在显存中，所以读取速度会比放在内存中要快。

 //device修饰4维向量的颜色
 device float4 *color;

struct Foo {  
float a[3]; 
int b[2];  
}; 
//device修饰结构体
device Foo *my_info;

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

threadgrounp Address Space (线程组地址空间)

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

 kernel void my_func(threadgroup float *a [[ threadgroup(0) ]], ...){
     threadgroup float x;
     threadgroup float b[10];
}

constant Address Space (常量地址空间)

  常量地址空间指向的缓存对象也是从设备内存池分配存储, 但是它是只读的;
  在程序域的变量必须定义在常量地址空间并且声明的时候初始化; ⽤来初始化的值必须是编译时的常量.
  在程序域的变量的⽣命周期和程序⼀样, 在程序中的并⾏计算着⾊函数或者图形绘制着⾊函数调⽤, 但 是constant 的值会保持不变;

注意: 常量地址空间的指针或是引⽤可以作为函数的参数. 向声明为常量的变量赋值会产⽣编译错误. 声明常量但是没有赋予初值也会产⽣编译错误;

  //定义一个常量samples
  constant float samples[] = { 1.0f, 2.0f, 3.0f, 4.0f };
  //对⼀个常量地址空间的变量进⾏修改也会失败,因为它只读的
  sampler[4] = {3,3,3,3}; //编译失败;
   //定义为常量地址空间声明时不赋初值也会编译失败
  constant float a; 

thread Address Space (线程地址空间)

  thread 地址空间指向每个线程准备的地址空间, 这个线程的地址空间定义的变量在其他线程不可⻅, 在 图形绘制着⾊函数或者并⾏计算着⾊函数中声明的变量thread 地址空间分配;

kernel void my_func(...)
{
      float x;
      thread float p = &x;
}

函数参数与变量

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

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

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

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

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

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 ]])
{
    //thread_position_in_grid : ⽤于表示当前节点在多线程⽹格中的位置;由Metal自行进行传递
    out[id] = inA[id] + inB[id];
}

[[ buffer(index) ]]中的index不能和其他的相同，否则会覆盖掉之前的数据

内建变量属性修饰符

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

  [[stage_in]]表示从顶点着色器中经过图元装配、光栅化等操作后传递给片元着色器的数据。类似于GLSL中的Varying。

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