Metal shading language简单介绍

一、Metal语⾔介绍

• Metal着⾊语⾔是⼀个⽤来编写3D图形渲染逻辑和并⾏计算核⼼逻辑的编程语⾔,编写Metal框架的APP需要使⽤Metal 着⾊语⾔程序.
  Metal 着⾊语⾔ 与 Metal 框架配合使⽤,Metal 框架管理Metal着⾊语⾔的运⾏和可选编译选项.
• Metal 着⾊器语⾔使⽤Clang和LLVM,编译器对于在GPU上的代码执⾏效率有更好的控制.
• Metal 与 C++ 11.0区别： Metal 这⻔语⾔是基于C++ 11.0标准设计的.它在C++基础是⾏多了⼀些拓展和限制.C++ 11.0 特性在

Metal 语⾔中不⽀持之处:

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

Metal Restrictions 限制

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

二、Metal 数据类型--标量数据类型 Metal数据类型

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

三、纹理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 等;

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

四、采样器类型 Samplers

采样去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 }; 设置纹理采样的放⼤过滤模式;

• 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 有以下3种函数修饰符:

• kernel , 表示该函数是⼀个数据并⾏计算着⾊函数. 它可以被分配在⼀维/⼆维/三维线程组中去执⾏; 注意：使⽤kernel 修饰的函数. 其返回值类型必须是void 类型;

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

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

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

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

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

• device

• threadgrounp

• constant

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

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

// an array of a float vector with  components 
 device float4 *color; 
 struct Foo { 
 float a[3];
  int b[2]; 
 }; 
 // an array of Foo elements 
 device Foo *my_info;

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

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

kernel void my_func(threadgroup float *a [[ threadgroup(0) ]], ...)  { 
 // A float allocated in threadgroup address space 
 threadgroup float x; 
 // An array of 10 floats allocated in 
 // threadgroup address space 
threadgroup float b[10]; 
 }

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

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 和 constant) 不能重名;

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

• 内建变量属性修饰符