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、对于 图形着色器函数,是 指针 或 引用
b、对于并行计算着色函数,对于是 指针 或 引用
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]]
[[position]]
[[point_size]]
[[color(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 可以修饰结构体,其结构体成员可以有多个,类型可以为一个 整型/浮点型 的 标量/向量。