Metal 一、初识 Metal 及其语言规范

Metal 简介 与 Metal 编程语言语法。

一、Metal 简介

1、Metal 是 Apple 为了解决 3D 渲染而推出的框架。游戏开发者的新的技术平台,该技术能够为3D图像提高 10 倍的渲染性能。苹果2018年推出 Metal,在此之前一直使用OpenGL ES。

Metal 的2个目的:1图形渲染; 2高并发计算

苹果文档给出的 Metal 优化:CPU 低消耗、GPU 高利用、连续处理器并行、有效的资源利用

Metal 官方文档

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 可以修饰结构体,其结构体成员可以有多个,类型可以为一个 整型/浮点型 的 标量/向量。

原文地址:https://www.cnblogs.com/zhangzhang-y/p/13549930.html