Metal Shading Language

作者: MonKey_Money | 来源:发表于2020-08-24 14:08 被阅读0次

Metal 语言介绍

Metal着色语⾔是⼀个⽤来编写3D图形渲染逻辑和并⾏计算核⼼逻辑的编程语⾔言,编写Metal框架的APP需要使用Metal 着⾊语⾔程序.
Metal 着⾊语⾔与 Metal 框架配合使用,Metal 框架管理Metal着⾊语言的运行和可选编译选项. Metal 着色器器语言使⽤Clang和LLVM,编译器对于在GPU上的代码执⾏效率有更好的控制.
Metal 这⻔语⾔是基于C++ 11.0标准设计的.它在C++基础增多了一些拓展和限制

Metal Shading Language语言规则

C++ 11.0 特性在Metal 语⾔中不⽀持之处

Lambda 表达式;
递归函数调⽤
动态转换操作符
类型识别
对象创建new 和销毁delete 操作符;
操作符 noexcept
goto 跳转
变量存储修饰符register 和 thread_local;
虚函数修饰符;
派⽣类
异常处理
C++ 标准库在Metal 语⾔中也不可使⽤;

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

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

Metal 像素坐标系统

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

Metal 数据类型

标量数据类型

类型	描述
bool	布尔类型, true/false,true可以拓展为整形常量1,false可以拓展为整形常量0
char	有符号8-bit整形
unsigned char /uchar	⽆符号8-bit 整数
short	有符号16-bit整数
unsigned short / ushort	⽆符号32-bit 整数
half	16 bit 浮点数
float	32bit 浮点数
size_t	64bit ⽆符号整数,表示sizeof操作符的结果
ptrdiff_t	64bit有符号整型,表示2个指针的差
void	该类型表示⼀个空的值集合

Metal ⽀持后缀表示字⾯量类型, 例如 0.5F, 0.5f; 0.5h, 0.5H;

向量

booln charn shortn intn ucharn ushortn uintn halfn floatn
表示n维度的对应数据类型，数据类型见上表。

pos = float4(1.0f, 2.0f, 3.0f, 4.0f);
float x = pos[0];// x = 1.0
float4 vA = float4(1.0f, 2.0f, 3.0f, 4.0f); float4 vB;
for (int i=0; i<4; i++)
vB[i] = vA[i] * 2.0f // vB = (2.0, 4.0, 6.0, 8.0);
//<vector_data_type>.xyzw 
//<vector_data_type>.rgba
int4 test = int4(0, 1, 2, 3); 
inta=test.x; // a=0 
intb=test.y; // b=1
float4 c;
c.xyzw = float4(1.0f, 2.0f, 3.0f, 4.0f);
float4 pos = float4(1.0f, 2.0f, 3.0f, 4.0f);
float4 swiz = pos.wzyx; // swiz = (4.0f, 3.0f, 2.0f, 1.0f)
float4 pos = float4(1.0f, 2.0f, 3.0f, 4.0f); // pos = (5.0, 2.0, 3.0, 6.0)
pos.xw = float2(5.0f, 6.0f);
// pos = (8.0, 2.0, 3.0, 7.0) 
pos.wx = float2(7.0f, 8.0f);

float2 pos;
pos.x = 1.0f; // x,y是合法的
pos.z = 1.0f; //z,w是不合法的
float3 pos;
pos.z = 1.0f; // x,y,z是合法的
pos.w = 1.0f; // w是不合法的
pos.xx = float2(3.0f, 4.0f); //不合法的，x被赋值两次
pos.xy = float4(1.0f, 2.0f, 3.0f, 4.0f); //不合法的，四维向量复制给二维向量
float4 pos = float4(1.0f, 2.0f, 3.0f, 4.0f);
pos.x = 1.0f; // 合法
pos.g = 2.0f; // 合法
pos.xg = float2(3.0f, 4.0f);//不合法，不能使用混合属性
float3 coord = pos.ryz;//不合法，不能使用混合属性
float4 pos = float4(1.0f, 2.0f, 3.0f, 4.0f); 
my_func(&pos.xy); // 不合法

初始化方法

//四维
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, float b, float2 c);
float4(float a, float2 b, float c);
float4(float3 a, float b);
float4(float a, float3 b);
float4(float4 x);
//三维
float3(float x);
float3(float x, float y, float z);
float3(float a, float2 b);
float3(float2 a, float b);
float3(float3 x);
//二维
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);

矩阵

halfnxm 表示n行m列的16位浮点数
floatnxm 表示n行m列的32位浮点数

float4x4(fval);
fval 0.0 0.0 0.0
0.0 fval 0.0 0.0
0.0 0.0 fval 0.0
0.0 0.0 0.0 fval

float2x2(float2, float2);
float3x2(float2, float2, float2);
float2x2(float, float, float, float);
float3x2(float, float, float, float, float, float);
float2x3(float2 a, float b, float2 c, float d); //不合法

纹理Textures

纹理类型是⼀个句柄, 它指向⼀个⼀维/⼆维/三维纹理数据. 在⼀个函数中描述纹理对象的类型;
enum class access {sample ,read ,write};
texture1d<T, access a = access::sample>
texture1d_array<T, access a = access::sample>
texture2d<T, access a = access::sample>
texture2d_array<T, access a = access::sample>
texture3d<T, access a = access::sample>
texturecube<T, access a = access::sample>
texture2d_ms<T, access a = access::read>
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 类型; **

只有图形着⾊函数才可以被 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) 指向设备内存池分配出来的缓存对象, 它是可读也是可写的; ⼀个缓存对象可以被声明成⼀个标量,向量或是⽤户⾃定义结构体的指针或是引⽤.

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