简介

• Metal 着色语言是用来编写 '3D 图形渲染' 和 '并行计算核心逻辑' 的 ⼀⻔编程语⾔。 当你使⽤ Metal 框架来完成APP的实现时，则需要使⽤Metal 编程语⾔
• Metal 语言使用 Clang 和 LLVM 进行编译处理，编译器对于在 GPU 上的代码执行效率有更好的控制
• Metal 这⻔门语⾔言是基于 C++ 11.0 标准设计的。它在 C++ 基础上多了一些拓拓展和限制
• 主要⽤来编写在 GPU 上执⾏的图像渲染逻辑代码以及通⽤并⾏计算逻辑代码

C++ 11.0 和 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 ⽀持后缀表示字⾯量类型， 例如 0.5F, 0.5f; 0.5h, 0.5H。

//基本数据类型
bool a = true;
char b = 5;
int  d = 15;
size_t c = 1;
ptrdiff_t f = 2;

向量和矩阵

向量支持如下类型：booln、charn、shortn、intn、ucharn、ushortn、uintn、halfn、floatn，其中 n 表示向量的维度，最多不超过4维向量

//向量
bool2 A= {1,2};
float4 pos = float4(1.0,2.0,3.0,4.0);
float x = pos[0];
float y = pos[1];

float4 VB;
for(int i = 0; i < 4 ; I++)
    VB[i] = pos[i] * 2.0f;

向量访问规则

• 通过向量字母获取元素： 向量中的向量字母仅有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

• 多个分量同时访问

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

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

  • 赋值时分量不可重复，取值时分量可重复
  • 右边是取值 和 左边赋值都合法
  • 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;

矩阵支持以下类型：halfnxm 和 floatnxm，nxm分别指的是矩阵的⾏数和列数

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

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

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

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

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

缓冲buffer

• 在Metal 中实现缓存靠的是⼀个指针，它指向⼀个在Device 或者 constant 地址空间中的内建或是开发者⾃定义的数据块，缓存可以被定在程序域域中，或是当做函数的参数传递。

//缓存buffer
device float4 *device_buffer;
struct my_user_data{
    float4 a;
    float b;
    int2 c;
};
constant my_user_data *user_data;

纹理类型（Textures）

纹理类型是一个句柄，指向一维/二维/三维纹理数据

纹理的访问权限

在一个函数中描述纹理对象的类型，枚举值定义了纹理的访问权利 enum class access {sample, read, write};，有以下3种访问权利，当没写access时，默认的access 就是 sample

• sample：纹理对象可以被采样（即使用采样器去纹理中读取数据，相当于OpenGL ES的GLSL中sampler2D），采样一维这时使用 或者 不使用都可以从纹理中读取数据（即可读可写可采样）
• 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 表示纹理访问权限，当access没写时，默认是sample

//类型 变量 修饰符
/*
 类型
    - texture2d<float>，读取的数据类型是float，没写access，默认是sample
    - texture2d<float,access::read>，读取的数据类型是float，读取的方式是read
    - texture2d<float,access::write>，读取的数据类型是float，读取的方式是write
 变量名
    - imgA
    - imgB
    - imgC
 修饰符
    - [[texture(0)]] 对应纹理0
    - [[texture(1)]] 对应纹理1
    - [[texture(2)]] 对应纹理2
 */
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坐标（对应纹理坐标的x、y、z）的寻址方式
  • s坐标：enum class s_address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat };
  • t坐标：enum class t_address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat };
  • r坐标：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 }; 设置所有纹理坐标的寻址方式
• enum class mip_filter { none, nearest, linear }; 设置纹理采样的mipMap过滤模式, 如果是none,那么只有一层纹理生效

采样器所有状态如下所示：

⚠️注意：在 Metal 程序中初始化的采样器必须使用 constexpr 修饰符声明；openGL ES 中纹理坐标对应的是 stq，Metal 中纹理坐标对应是 str

/*
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 有三种函数修饰符，放在函数的返回值前面

• kernel：表示该函数是一个数据并行计算着色函数，它可以被分配在一维/二维/三维线程组中去执行
• vertex：表示该函数是一个顶点着色函数，它将为顶点数据流中的每个顶点数据执行一次，然后为每个顶点生成数据输出到绘制管线
• fragment：表示该函数是一个片元着色函数，它将为片元数据流中的每个片元 和其相关联的数据执行一次，然后将每个片元生成的颜色数据输出到绘制管线中

//并行计算函数(kernel)
kernel void CCTestKernelFunctionA(int a,int b)
{ 
    /*
     注意:
     1. 使用kernel 修饰的函数返回值必须是void 类型
     2. 一个被函数修饰符修饰过的函数,不允许在调用其他的被函数修饰过的函数. 非法
     3. 被函数修饰符修饰过的函数,只允许在客户端对其进行操作. 不允许被普通的函数调用.
     */
     
    //不可以的!
    //一个被函数修饰符修饰过的函数,不允许在调用其他的被函数修饰过的函数. 非法
    CCTestKernelFunctionB(1,2);//非法，错误调用！！！
    CCTestVertexFunctionB(1,2);//非法，错误调用！！！
    
    //可以! 你可以调用普通函数.而且在Metal 不仅仅只有这3种被修饰过的函数.普通函数也可以存在
    CCTest();
    
}

//并行计算函数
kernel void CCTestKernelFunctionB(int a,int b)
{
    .....
}

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

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

//普通函数
void CCTest()
{
    .....
}

⚠️注意

• 使⽤kernel 修饰的函数，其返回值类型必须是 void 类型
• 一个被函数修饰符修饰的函数不能在调用其他也被函数修饰符修饰的函数，这样会导致编译失败，即 Kernel、vertex、fragment 修饰的函数不能相互调用，也不能同修饰符函数相互调用。但是可以调用普通函数
• 被函数修饰符修饰过的函数,只允许在客户端对其进行操作. 不允许被普通的函数调用
• 只有图形着色函数才可以被 vertex 和 fragment 修饰，对于图形着色函数，通过返回值类型可以辨认出是为顶点计算还是像素计算，其返回值也可以是 void，意味着不产生数据输出到绘制管线，是一个无意义的动作

变量、参数的地址空间修饰符

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

• device 设备地址空间
• threadgroup 线程组地址空间
• constant 常量地址空间
• thread 线程地址空间

⚠️注意

• 所有的着⾊函数(vertex, fragment, kernel)的参数，如果是指针或是引⽤，都必须带有地址空间修饰符号
• 对于图形着色器函数（即vertex/fragment修饰的函数），其指针/引用类型的参数必须定义为 device、constant 地址空间
• 对于并行计算函数（即 kernel 修饰的函数），其指针/引用类型的参数必须定义为 device、threadgroup、constant
• 并不是所有的变量都需要修饰符（例如普通变量）

//变量/参数地址空间修饰符
void CCTestFouncitionE(device int *g_data,
                       threadgroup int *l_data,
                       constant float *c_data
                       )
{
}

设备地址空间修饰符（device）

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

// 设备地址空间: device 用来修饰指针.引用
//1.修饰指针变量
device float4 *color;

struct CCStruct{
    float a[3];
    int b[2];
};
//2.修饰结构体类的指针变量
device CCStruct *my_CS;

⚠️注意

• 纹理对象总是在设备地址空间分配内存，即纹理对象默认在GPU分配内存
• device 地址空间修饰符不必出现在纹理类型定义中
• 一个纹理对象的内容无法直接访问，Metal 提供读写纹理的内建函数

线程组地址空间修饰符（threadgroup）

• 线程组地址空间用于为并行计算着色器函数分配内存变量，这些变量被一个线程组的所有线程共享，在线程组地址空间分配的变量不能用于图形绘制着色函数（即顶点着色函数 / 片元着色函数），即在图形绘制着色函数中不能使用线程组
• 在并行计算着色函数中，在线程组地址空间分配的变量为一个线程组使用，生命周期和线程组相同

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

constant float sampler[] = {1.0f,2.0f,3.0f,4.0f};
kernel void CCTestFouncitionG(void)
{
    //在线程空间分配空间给x,p
    float x;
    thread float p = &x;
    
}

常量地址空间修饰符（constant）

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

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

 //定义为常量地址空间声明时不赋初值也会编译失败 
 constant float a;

⚠️注意

• 常量地址空间的指针/引用可以作为函数的参数，向声明为常量的变量赋值会产生编译错误
• 声明常量但是没有赋予初值也会产生编译错误

线程地址空间修饰符（thread）

• 线程地址空间指向每个线程准备的地址空间，也是在 GPU 中，该线程的地址空间定义的变量在其他线程不可见（即变量不共享）
• 在图形绘制着色函数或者并行计算着色函数中声明的变量，在线程地址空间分配存储

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

函数参数与变量

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

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

为什么需要属性修饰符？

• 参数表示资源的定位，可以理解为端口，相当于OpenGl ES中的location
• 在固定管线和可编程管线进行内建变量的传递
• 将数据沿着渲染管线从顶点函数传递到片元函数

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

对于每个着色函数来说，一个修饰符是必须指定的，它用来设置一个缓存、纹理、采样器的位置

• device buffer ---> [[buffer(index)]]
• constant buffer ---> [[buffer(index)]]
• texture ---> [[texture(index)]]
• sampler ---> [[sampler(index)]]
• threadGroup ---> [[threadGroup(index)]]

⚠️注意

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

//属性修饰符
/*
 1. device buffer(设备缓存)
 2. constant buffer(常量缓存)
 3. texture Object(纹理对象)
 4. sampler Object(采样器对象)
 5. 线程组 threadgroup
 
 属性修饰符目的:
 1. 参数表示资源如何定位? 可以理解为端口
 2. 在固定管线和可编程管线进行内建变量的传递
 3. 将数据沿着渲染管线从顶点函数传递片元函数.
 
 在代码中如何表现:
 1.已知条件:device buffer(设备缓存)/constant buffer(常量缓存)
 代码表现:[[buffer(index)]]
 解读:不变的buffer ,index 可以由开发者来指定.
 
 2.已知条件:texture Object(纹理对象)
 代码表现: [[texture(index)]]
 解读:不变的texture ,index 可以由开发者来指定.
 
 3.已知条件:sampler Object(采样器对象)
 代码表示: [[sampler(index)]]
 解读:不变的sampler ,index 可以由开发者来指定.
 
 4.已知条件:threadgroup Object(线程组对象)
 代码表示: [[threadgroup(index)]]
 解读:不变的threadgroup ,index 可以由开发者来指定.
 */

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

//着色函数的多个参数使用不同类型的属性修饰符的情况
kernel void my_kernel(device float4 *p [[buffer(0)]],
                      texture2d<float> img [[texture(0)]],
                      sampler sam [[sampler(0)]])
{
    //.....
    
}

内建变量修饰符

• [[vertex_id]] 顶点id 标识符
• [[position]]
  • 在顶点着色函数中，表示当前的顶点信息，类型是 float4
  • 还可以表示描述了片元的窗口的相对坐标（x，y，z，1/w），即该像素点在屏幕上的位置信息
• [[point_size]] 点的大小，类型是 float
• [[color(m)]] 颜色，m 在编译前就必须确定

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

• [[stage_in]] ⽚元着⾊函数使⽤的单个⽚元输⼊数据是由顶点着⾊函数输出然后经过光栅化⽣成的。
  • 顶点和⽚元着⾊函数都是只能有⼀个参数被声明为使⽤“stage_in”修饰符，
  • 对于⼀个使⽤ 了“stage_in”修饰符的⾃定义的结构体，其成员可以为⼀个整形或浮点标量，或是整形或浮点向量