[toc]

前言

Metal Shadeing Language是做什么用的?

1. Metal着色语言是用来编写3D图形渲染逻辑和并行计算核心逻辑的一门语言,当时使用Metal框架来完成APP的实现时则需要使用Metal编程语言;
2. Metal语言使用Clang和LLVM进行编译处理~
3. Metal基于C++11.0语言设计,我们主要用来编写在GPU上

提问
C++ 11.0 和 Metal 语⾔的异同之处?

1. C++ 11.0 特性在Metal 语⾔中不⽀持之处
2. Lambda 表达式;
3. 递归函数调⽤
4. 动态转换操作符
5. 类型识别
6. 对象创建new和销毁delete 操作符;
7. 操作符 noexcept
8. goto 跳转
9. 变量存储修饰符register 和 thread_local;
10. 虚函数修饰符;
11. 派⽣类
12. 异常处理
13. C++ 标准库在Metal 语⾔中也不可使⽤;

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

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

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

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

image.png

1. bool 布尔类型, true/false
2. char 有符号8位整数;
3. unsigned char /(简写)uchar ⽆符号8-bit 整数;
4. short 有符号16-bit整数;
5. unsigned short / (简写)ushort ⽆符号32-bit 整数;
6. half 16位bit 浮点数; //OC中的float
7. float 32bit 浮点数;
8. size_t 64 ⽆符号整数;
9. void 该类型表示⼀个空的值集合

Metal 数据类型--向量数据类型

向量支持如下类型

• booln、charn、shortn、intn、ucharn、ushortn、uintn、halfn、floatn，其中 n 表示向量的维度，最多不超过4维向量

//直接赋值初始化
bool2 A= {1,2};
//通过内建函数float4初始化
float4 pos = float4(1.0,2.0,3.0,4.0);

//通过下标从向量中获取某个值
float x = pos[0];
float y = pos[1];

//通过for循环对一个向量进行运算
float4 VB;
for(int i = 0; i < 4 ; i++)
{
    VB[i] = pos[i] * 2.0f;
}

在OpenGL ES的GLSL语言中，例如2.0f，在着色器中书写时，是不能加f，写成2.0，而在Metal中则可以写成2.0f，其中f可以是大写，也可以是小写

向量的访问规则

1. 通过向量字母获取元素： 向量中的向量字母仅有2种，分别为xyzw、rgba

int4 test = int4(0,1,2,3);
int a = test.x; //获取的向量元素0
int b = test.y; //获取的向量元素1
int c = test.z; //获取的向量元素2
int d = test.w; //获取的向量元素3

int e = test.r; //获取的向量元素0
int f = test.g; //获取的向量元素1
int g = test.b; //获取的向量元素2
int h = test.a; //获取的向量元素3

2. 多个分量同时访问

float4 c;
c.xyzw = float4(1.0f,2.0f,3.0f,4.0f);
c.z = 1.0f;
c.xy = float2(3.0f,4.0f);
c.xyz = float3(3.0f,4.0f,5.0f);

3. 多分量访问可以乱序/重复

   • 赋值时分量不可重复，取值时分量可重复
   • 右边是取值和左边赋值都合法
   • xyzw与rgba不能混合使用

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

//可以仅对 xw / wx 修改
//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);

//可以仅对 xyz 进行修改
//pos = (3.0f,5.0f,9.0f,7.0f);
pos.xyz = float3(3.0f,5.0f,9.0f);

float2 pos;
pos.x = 1.0f; //合法
pos.z = 1.0f; //非法，pos是二维向量，没有z这个索引

float3 pos2;
pos2.z = 1.0f; //合法
pos2.w = 1.0f; //非法

// 赋值 时 分量不可重复，取值 时 分量可重复
//非法,x出现2次
pos.xx = float2(3.0,4.0f);
pos.xy = swiz.xx;

//向量中xyzw与rgba两组分量不能混合使用
float4 pos4 = float4(1.0f,2.0f,3.0f,4.0f);
pos4.x = 1.0f;
pos4.y = 2.0f;
//非法,.rgba与.xyzw 混合使用
pos4.xg = float2(2.0f,3.0f);
////非法,.rgba与.xyzw 混合使用
float3 coord = pos4.ryz;

GLSL中向量不能乱序访问，只是和Metal中的向量相似，并不是等价。

Metal 数据类型--矩阵数据类型矩阵

矩阵支持如下类型

1. halfnxm、floatnxm，其中 nxm表示矩阵的行数和列数，最多4行4列，其中half、float相当于OC中的float、double

2. 普通的矩阵其本质就是一个数组

float4x4 m;
//将第二行的所有值都设置为2.0
m[1] = float4(2.0f);

//设置第一行/第一列为1.0f
m[0][0] = 1.0f;

//设置第三行第四列的元素为3.0f
m[2][3] = 3.0f;

3. float4 类型向量的构造方式

• 1个float构成,表示一行都是这个值
• 4个float构成
• 2个float2构成
• 1个float2+2个float构成（顺序可以任意组合）
• 1个float2+1个float
• 1个float4

//float4类型向量的所有可能构造方式
//1个一维向量，表示一行都是x
float4(float x);/
//4个一维向量 --> 4维向量
float4(float x,float y,float z,float w);
//2个二维向量 --> 4维向量
float4(float2 a,float2 b);
//1个二维向量+2个一维向量 --> 4维向量
float4(float2 a,float b,float c);
float4(float a,float2 b,float c);
float4(float a,float b,float2 c);
//1个三维向量+1个一维向量 --> 4维向量
float4(float3 a,float b);
float4(float a,float3 b);
//1个四维向量 --> 4维向量
float4(float4 x);

4. float3 类型向量的构造方式

• 1个float构成,表示一行都是这个值
• 3个float
• 1个float+1个float2（顺序可以任意组合）
• 1个float2

//float3类型向量的所有可能的构造的方式
//1个一维向量
float3(float x);
//3个一维向量
float3(float x,float y,float z);
//1个一维向量 + 1个二维向量
float3(float a,float2 b);
//1个二维向量 + 1个一维向量
float3(float2 a,float b);
//1个三维向量
float3(float3 x);

5. float2 类型向量的构造方式

• 1个float构成,表示一行都是这个值
• 2个float
• 1个float2

//float2类型向量的所有可能的构造方式
//1个一维向量
float2(float x);
//2个一维向量
float2(float x,float y);
//1个二维向量
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 }; 设置纹理采样的放⼤过滤模式;

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

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

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

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

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

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

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

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

• device
• threadgrounp
• constant
• thread

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

Device Address Space(设备地址空间)

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

// an array of a float vector with 4 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) 不能重名;

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”修饰符的⾃定义的结构体，其成员可以为⼀个整形或浮点标量,或是整形或浮点向量.

参考:

Metal Shader language (着色语言规范)总结