Metal 一、初识 Metal 及其语言规范
Metal 简介 与 Metal 编程语言语法。
一、Metal 简介
1、Metal 是 Apple 为了解决 3D 渲染而推出的框架。游戏开发者的新的技术平台，该技术能够为3D图像提高 10 倍的渲染性能。苹果2018年推出 Metal，在此之前一直使用OpenGL ES。
Metal 的2个目的：1图形渲染； 2高并发计算。
苹果文档给出的 Metal 优化：CPU 低消耗、GPU 高利用、连续处理器并行、有效的资源利用
2、图形管道 graphic pipeline
二、Metal Shading Language - Metal 编程语言与规范
1、Metal 语言特点
1.1）Metal 着色器语言，用来编写3D 图形渲染逻辑和并行计算核心逻辑(高并发) 的一门编程语言。当我们需要使用 Metal 框架来完成 APP 的开发时(或 需要使用 Metal 的高并发计算能力时)，就要使用 Metal 编程语言进行编程。
1.2）Metal 编程语言 使用 clang 和 LLVM 进行编译处理
Metal的语言规则，是基于 C++ 11.0 标椎设计的，并在此基础上进行了一定的扩展和限制。
而我们的实际业务开发场景中所需要做的工作就是： 编写 在GPU 上执行的图像渲染逻辑代码 以及 通用的并行计算逻辑 代码。
1.3）Metal 与 C++ 11.0
1.3.1）Metal 限制(不支持)：
Lambda 表达式
递归函数调用
动态转换操作符
类型识别
对象创建 new 和 销毁 delete 操作符
操作符 noexcept
goto 跳转
变量存储修饰符 register 和 thread_local
虚函数修饰符
派生类
异常处理
C++ 标准函数库 在 Metal 中不支持，不能使用
1.3.2）Metal 中指针的使用限制
Metal 图形和并行计算函数用到的参数，如果是指针 必须使用地址空间修饰符(devide / threadgroup / constant)
不支持函数指针
函数名不能出现 main
1.3.3）Metal 像素坐标系
Metal 中纹理/帧缓冲区 attachment 的像素使用的坐标系的原点是在 左上角。--> 苹果的坐标原点
2、Metal 的数据类型
2.1）标量类型
2.2）向量与矩阵 数据类型
2.2.1）向量
booln
charn
shortn
intn
ucharn / ushortn / uintn
halfn
floatn
n --> 向量中的 n 表示维度，最大为4
代码示例：
bool2 b = [1, 2]; float4 f4 = float4(1.0, 2.0, 3.0, 4.0); float f = f4[0];// x = 1.0 --> 类似数组 // int4 --> 4个变量组成的4维向量 // xyzw rgba int4 test = int4(0, 1, 2, 3); int x = test.x;// x = 0 int y = text.y;// y = 1 int x1 = test.r;// x1 = 0 // 多个分量的访问 float4 c = float4(0,0,0,0); c.xyzw = float4(1,2,3,4);// 重赋值 c = [1,2,3,4] c.xy = float2(6,0);// c = [6,0,3,4] c.yzw = float3(7,8,9);// c = [6,7,8,9] // 可乱序 --> // 注意！！！这里 Metal 不同于 GLSL --> GLSL 可以多分量,但是不可乱序 xyzw/rgba 顺序是不可变的 float4 pos = c.wxyz;// pos = [9,6,7,8] float4 rep = c.xxwz;// rep = [6,6,9,8] rep.xw = float2(5,7);// rep = [5,6,9,7] // 不可混用 float4 m = float4(4,3,2,1); m.xg = float2(0,9);!error 非法, xyzw 和 rgba 不能混了，他俩只可选其一 m.rg = float(0,9);// m = [0,9,2,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); // float3 类型向量的所有可能的构造的方式 float3(float x); float3(float x,float y,float z); float3(float a,float2 b); float3(float2 a,float b); float3(float3 x); // float2 类型向量的所有可能的构造方式 float2(float x); float2(float x,float y); float2(float2 x); // 多个向量构造的使用 float x = 1.0f,y = 2.0f,z = 3.0f,w = 4.0f; float4 a = float4(0.0f); float4 b = float4(x,y,z,w); float2 c = float2(5.0f,6.0f); float2 a = float2(x,y); float2 b = float2(z,w); float4 x = float4(a.xy,b.xy);
2.2.2）矩阵
halfxnm / floatnxm
nxm 中 n m 分别指 矩阵的行数 和 列数 --> 最大 4 x 4 : 4行4列
float4x4 mix; // mix[1] = float4(2.0f);// 矩阵 第一行的值都是 2 mix[1] = float4(1,2,3,4);// 矩阵的第一行的值 mix[0][0] = 3;// 矩阵的第0行0列的值为 3 mix[3][2] = 7;// 矩阵 3行2列 位置的值
2.3）纹理 Texture 类型
纹理类型是一个句柄，指向一个 一维/二维/三维纹理数据。在函数中描述纹理对象的类型。
枚举:
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、int 等；
access：读写方式(权限)
代码示例：
void foo (texture2d<float> imgA [[ texture(0) ]] ,// texture2d<float>: 2 维纹理，类型 float，访问权限 sample --> 默认权限就是 sample 可不写 texture2d<float, access::read> imgB [[ texture(1) ]],// texture2d<float, access::read>: 类型 float，权限 read texture2d<float, access::write> imgC [[ texture(2) ]]) // 权限 write { ... }
2.4）采样器类型 Samplers
采样器类型 决定了如何对一个纹理进行采样操作。
metal 框架中有一个对应 着色器语言的采样器对象：MTLSamplerState, 此对象做为 图像渲染着色器函数 or 并行计算函数 的参数进行传递。
枚举 们：
// 从纹理中采样时，纹理坐标是否归一化 enum class coord { normalized, peixel } // 纹理采样过滤方式 - 放大/缩小 enum class filter { nearest, linear } // 缩小过滤方式 enum class min_filter { nearest, linear } // 放大过滤方式 enum class mag_filter { nearest, linear } // 设置纹理 s t r 坐标的寻址模式 （str 即 xyz 环绕方式） enum class s_address { clapm_to_zero, clapm_to_edge, repeat, mirrored_repeat } enum class t_address { clapm_to_zero, clapm_to_edge, repeat, mirrored_repeat } enum class r_address { clapm_to_zero, clapm_to_edge, repeat, mirrored_repeat } // 设置所有纹理坐标的寻址模式 enum class address { clapm_to_zero, clapm_to_edge, repeat, mirrored_repeat } // 设置纹理采样的 mipMap 过滤模式，如果是 none ，则只有一层纹理生效 enum class mag_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);
3、修饰符
3.1）函数修饰符
kernel：表示该函数是一个 数据并行计算着色函数。我们要高效并发运算就用它。它可以被分配在 一维/二维/三维 线程组中去执行；--> 使用他修饰的函数 返回类型必须是 void
vertex：顶点着色函数。为顶点数据流中的 每个顶点数据 执行一次，然后为每个顶点生成数据输出到绘制管线；
fragment：片元着色函数，为 片元数据流中的 每个片元 与其关联 执行一次，然后将 每个片元生成的颜色数据 输出到绘制管线中；
注意1：只有 图形着色函数才能用 vertex/fragment 修饰。函数返回类型 可以用来辨认出它是为顶点 or 为每个像素 做计算的。返回 void 也可以但是无意义，因为顶点/片元函数本就是为了计算出相应数据 将数据传到绘制管线的。
注意2：被函数修饰符修饰的函数不能再调用 '被修饰符修饰的函数'，否则编译失败。即：被函数修饰符修饰的函数们不能相互调用。
例：kernel void func1 (...) {}; vertex float4 funcV1 (...) { func1(...) } --> 错误调用，无法编译
注意2：特定函数修饰，普通函数随意。
代码示例：
kernel void foo(...) { ... }
3.2）变量 或 参数 的地址空间修饰符 Address Space
地址空间修饰符：用来指定 一个函数 参数/变量 被分配在内存中的哪块区域。
device：设备地址空间
threadgroup：线程组地址空间
constant：常量地址空间
thread： thread 地址空间
a、对于 图形着色器函数，是 指针 或 引用 类型的参数 必须定义为 device 或 const 地址空间
b、对于并行计算着色函数，对于是 指针 或 引用 的参数，必须使用 device 或 threadgroup 或 constant 修饰。
3.2.1）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;
注意1：纹理对象 总是在设备地址空间分配内存，device 地址空间修饰符不必出现在纹理类型定义中。一个纹理对象的内容无法直接访问，Metal 提供了读写纹理的内建函数。
3.2.2）线程组地址空间 threadgroup
threadgroup：用于 为 并行计算着色函数 分配内存变量(在GPU里)，这些变量被一个线程组的所有线程共享。在线程组地址空间 分配的变量 不能被用于图形绘制着色函数。
在并行计算着色函数中，在线程组地址空间分配的变量 为一个线程组使用，生命周期和线程相同。
代码示例：
// kernel 高速并行 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]; }
3.2.3）constant 常量地址空间
constant：指向的缓存对象也是从设备内存池 分配存储，但是是只读的。
在程序域的变量 必须定义在常量地址空间 并且在声明的时候初始化；用来初始化的值 必须是编译时的常量。此变量的生命周期和程序一样，在程序中的 并行计算着色函数or图形绘制着色函数 调用，但 constant 的值会保持不变。
注意：常量地址空间 的指针或引用 可以作为函数的参数(constant修饰的常量可作为函数的参数)。向声明为常量的变量赋值会产生变异错误(代码示例中sampler)，声明为常量但没有赋予初始值也会产生变异错误(代码示例中a)。
错误代码示例：
constant float sampler[] = {1.0f, 2.0f, 3.0f,4.0f}; // 对一个常量地址空间的变量进行修改会失败，因为它是只读的 sampler[4] = {3,3,3,3};// 编译失败 // 定义常量地址空间但不初始化赋值 --> 也编译失败 const float a;// 编译失败
3.2.4）线程地址空间 thread
thread：指向每个线程准备的地址空间，这个线程的地址空间 定义的变量 在其他线程是不可见的，在图形绘制着色函数or并行计算着色函数 中声明的变量可以使用 thread 地址空间分配。
代码示例：
kernel void func2 (...) { float x; thread float p = &x; ... }
3.3）函数参数与变量
图形绘制/并行计算着色函数的 输入/输出 都需要通过参数传递 ( 除了常量地址空间变量和程序域中定义的采样器 外)。参数如下：
device buffer：设备缓存 - 指向设备地址空间的任意数据类型的指针 or 引用
constant buffer：常量缓存 - 指向常量地址空间的任意数据类型的指针 or 引用
texture object：纹理对象
sample object：采样器对象
threadgroup：线程共享的缓存
对于每个着色器函数来说，一个修饰符是必须指定的，它用来设定一个 缓存、纹理、采样器的位置：
device buffer / constant buffer --> [[buffer(index)]]
texture --> [[texture(index)]]
sample --> [[sampler(index)]]
threadgroup buffer --> [[threadgroup(index)]]
index：一个 unsigned integer 类型的值，表示一个 缓存、纹理、采样器的位置(在函数参数索引表中的位置)。语法上讲，属性修饰符的声明位置 应该位于参数变量名之后。
通过示例理解：
// 一个简单的并行计算着色函数 my_add ，它把两个设备地址空间的魂村 inA、inB 相加，把结果写入缓存 out。 // 属性修饰符 “buffer(index)” 为着色函数参数 设定了缓存的位置 kernel void my_add (constant device float4 *inA [[ buffer(0) ]],// inA： 放在设备地址空间，缓存位置对应的是 buffer(0)这个ID ， constant 修饰的不可变 constant device float4 *inB [[ buffer(1) ]], device float4 *out [[ buffer(2) ]], uit id [[ thread_position_in_grid ]],) { out[id] = inA[id] + inB[id]; }
thread_position_in_grid：用于表示当前节点，在多线程网格中的位置 --> 我们是无法知道当前在GPU的哪个运算单元里，thread_position_in_grid 知道，我们通过它获取即可。
3.4）内建变量属性修饰符
[[vertex_id]] -- 顶点ID标识符
[[position]] -- 1、当前顶点信息(float4 - xyzw) 2、也可描述 片元在窗口的相对坐标：当前这个像素点在屏幕上的哪个位置
[[point_size]] -- 点的大小
[[color(m)]] -- 颜色，m 编译前要确定
如下代码：
// 定义颜色结构体 struct myFragmentOutput { // 三组颜色，要知道使用时取哪一个 float4 clr_f [[color(0)]];// int4 clr_i [[color(1)]];// uint4 clr_ui [[color(2)]];// } fragment myFragmentOutput my_grag_shader (...) { myFragmentOutput f; ... f.clr_f = ...; ... return f; }
另补充一个：[[stage_in]] -- 其实就是：顶点着色器输出 经过光栅化 生成的 传给片元着色器的 每个片元数据。
顶点和片元着色函数都是 有且仅有 一个参数可以被声明为 使用"stage_in"修饰符的。 stage_in 可以修饰结构体，其结构体成员可以有多个，类型可以为一个 整型/浮点型 的 标量/向量。
