Metal Shader Language

Metal Shading Language介绍

  Metal 着⾊语⾔ 是⽤来编写 3D 图形渲染逻辑 和 并⾏计算核⼼逻辑的⼀⻔编程语⾔. 当你使⽤Metal 框架来完成APP 的实现时则需要使⽤Metal 编程语⾔。

  Metal 语⾔使⽤Clang 和 LLVM 进⾏编译处理~Metal 基于C++ 11.0 语⾔设计.我们主要⽤来编写 在 GPU 上执⾏的图像渲染逻辑代码以及通用并⾏计算逻辑代码。

Metal 与 C++ 11.0

  Metal 这⻔语⾔是基于C++ 11.0标准设计的.它在C++基础是⾏多了⼀些拓展和限制.下⾯我们可以简单介绍介绍Metal着⾊语⾔与C++11.0 相⽐之下的修改和限制.

1. C++ 11.0 特性在Metal 语⾔中不⽀持的情况

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

2. Metal 语⾔中对于指针使⽤的限制

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

3. Metal 像素坐标系统

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

Metal数据类型

标量数据类型

Metal ⽀持后缀表示字⾯量类型, 例如 0.5F, 0.5f; 0.5h, 0.5H;

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

向量数据类型

向量中的n,指的是维度。

• booln
• charn
• shortn
• intn
• ucharn
• ushortn
• uintn
• halfn
• floatn纹理Textures 类型:纹理Textures 类型:

向量的访问

int4 test = int4(0,1,2,3);

1. 通过索引方式访问test[n]

2. 通过向量字母获取元素xyzw 和 rgba，只取部分、乱序取均可。

   例：

   test.x = 3;
   test.xyz = int3(1,2,3);
   test.rgb = int3(1,2,3);
   

矩阵数据类型

矩阵的行数n，矩阵的列数m

• halfnxm
• floatnxm

例：

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

1. float4类型向量的所有可能构造方式

   float4(float x);
   float4(float x,float y,float z,float w);
   float4(float2 a,float2 b);
   float4(float2 a,float b,float c);
   float4(float a,float2 b,float c);
   float4(float a,float b,float2 c);
   float4(float3 a,float b);
   float4(float a,float3 b);
   float4(float4 x);
   

2. float3类型向量的所有可能的构造的方式

   float3(float x);
   float3(float x,float y,float z);
   float3(float a,float2 b);
   float3(float2 a,float b);
   float3(float3 x);
   

3. float2类型向量的所有可能的构造的方式

   float2(float x);
   float2(float x,float y);
   float2(float2 x);
   

纹理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

  采取器类型决定了如何对⼀个纹理进⾏采样操作。在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); 
 constexpr sampler a(coord::normalized); 
 constexpr sampler b(address::repeat); 
 constexpr sampler s(address::clamp_to_zero,
                      filter::linear);

修饰符

函数修饰符

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

• kernel

  表示该函数是⼀个数据并⾏计算着⾊函数. 它可以被分配在⼀维/⼆维/三维线程组中去执⾏，其返回值类型必须为void。

  kernel void CCTestKernelFunctionA(int a,int b)
  {
  
      /*
       注意:
       1. 使用kernel 修饰的函数返回值必须是void 类型
       2. 一个被函数修饰符修饰过的函数,不允许在调用其他的被函数修饰过的函数. 非法
       3. 被函数修饰符修饰过的函数,只允许在客户端对齐进行操作. 不允许被普通的函数调用.
       */
  }
  

• vertex

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

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

• fragment

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

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

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

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

  Metal 着⾊器语⾔使⽤ 地址空间修饰符 来表示⼀个函数变量或者参数变量 被分配于那⼀⽚内存区域.

  所有的着⾊函数(vertex, fragment, kernel)的参数,如果是指针或是引⽤, 都必须带有地址空间修饰符号;

• device
• threadgrounp
• constant
• thread

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

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

Device Address Space(设备地址空间)

  Device是比较通用的访问模式，使用限制比较少，在设备地址空间(Device) 指向设备内存池分配出来的缓存对象，它是可读也是可写的; ⼀个缓存对象可以被声明成⼀个标量,向量或是⽤户⾃定义结构体的指针或是引⽤。

示例

device float4 *color; 
struct Foo { 
float a[3]; 
int b[2]; 
}; 
device Foo *my_info;

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

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

  线程组地址空间⽤于为 并⾏计算着⾊函数分配内存变量。 这些变量被⼀个线程组的所有线程共享。在线程组地址空间分配的变量不能被⽤于图形绘制着⾊函数(顶点着⾊函数, ⽚元着⾊函数)

  在并⾏计算着⾊函数中, 在线程组地址空间分配的变量为⼀个线程组使⽤, 声明周期和线程组相同;

示例

/*
 1. threadgroup 被并行计算计算分配内存变量, 这些变量被一个线程组的所有线程共享. 在线程组分配变量不能被用于图像绘制.
 2. thread 指向每个线程准备的地址空间. 在其他线程是不可见切不可用的
 */
kernel void CCTestFouncitionF(threadgroup float *a)
{
    //在线程组地址空间分配一个浮点类型变量x
    threadgroup float x;

    //在线程组地址空间分配一个10个浮点类型数的数组y;
    threadgroup float y[10];

}

constant Address Space( 常量地址空间)

  常量地址空间指向的缓存对象也是从设备内存池分配存储, 但是它是只读的，并且限定大小。在程序域的变量必须定义在常量地址空间并且声明的时候初始化，⽤来初始化的值必须是编译时的常量。

  在程序域的变量的⽣命周期和程序⼀样, 在程序中的并⾏计算着⾊函数或者图形绘制着⾊函数调⽤, 但是constant 的值会保持不变。

注意:

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

示例

constant float samples[] = { 1.0f, 2.0f, 3.0f, 4.0f };

thread Address Space (线程地址空间)

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

示例

kernel void CCTestFouncitionG(void)
{
    //在线程空间分配空间给x,p
    float x;
    thread float p = &x;

}

函数参数与变量

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

• 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_vectros ,实现2个设备地址空间中的缓存A与缓存B相加.然后将结果写入到缓存out.
//属性修饰符"(buffer(index))" 为着色函数参数设定了缓存的位置
//并行计算着色器函数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修饰符的⾃定义的结构体，其成员可以为⼀个整形或浮点标量，或是整形或浮点向量。