Metal Shader Language-着色语言规范
1.简述
Metal 着⾊语⾔是⽤来编写 3D 图形渲染逻辑和并⾏计算核⼼逻辑的⼀⻔编程语⾔。当你使⽤Metal框架来完成APP的实现时则需要使⽤Metal 编程语⾔;
Metal 语⾔使⽤Clang 和 LLVM 进⾏编译处理~
Metal 基于C++ 11.0 语⾔设计.我们主要⽤来编写 在 GPU 上执⾏的图像渲染逻辑代码 以及 通⽤并⾏计算逻辑代码;
2.C++ 11.0 和 Metal 语⾔的不同
2.1 C++ 11.0 特性在Metal 语⾔中不⽀持之处
- Lambda 表达式;
- 递归函数调⽤
- 动态转换操作符
- 类型识别
- 对象创建new 和销毁delete 操作符;
- 操作符 noexcept
- goto 跳转
- 变量存储修饰符register 和 thread_local;
- 虚函数修饰符;
- 派⽣类
- 异常处理
- C++ 标准库在Metal 语⾔中也不可使⽤;
2.2 Metal 语⾔中对于指针使⽤的限制:
- Metal图形和并⾏计算函数⽤到的⼊参数; 如果是指针必须使⽤地址空间修饰符(device,threadgroup,constant)
- 不⽀持函数指针;
- 函数名不能出现main
2.3 像素坐标系统
- Metal 像素坐标系统: Metal 中纹理/帧缓存区attachment 的像素使⽤的坐标系统的原点是左上⻆;
3. Metal基本数据类型
3.1 标量
标量
3.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;
向量的值的访问
通过向量字母来获取元素：
int4 test = int4(0,1,2,3);
int a = test.x;
int b = test.y;
int c = test.z;
int d = test.w;
int e = test.r;
int f = test.g;
int g = test.b;
int h = test.a;
多个分量同时访问：
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);
非法访问：
float2 pos;
pos.x = 1.0f; //合法
pos.z = 1.0f; //非法
float3 pos2;
pos2.z = 1.0f; //合法
pos2.w = 1.0f; //非法
//非法,x出现2次
pos.xx = float2(3.0,4.0f);
//不合法-使用混合限定符
pos.xy = float4(1.0f,2.0,3.0,4.0);
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;
float4 pos5 = float4(1.0f,2.0f,3.0f,4.0f);
//非法,使用指针来指向向量/分量
my_func(&pos5.xy);
3.2 矩阵
矩阵支持如下类型
- halfnxm、floatnxm，其中 nxm表示矩阵的行数和列数，最多4行4列，其中half、float相当于OC中的float、double
- 普通的矩阵其本质就是一个数组
float4x4 m;
//将第二排的值设置为0
m[1] = float4(2.0f);
//设置第一行/第一列为1.0f
m[0][0] = 1.0f;
//设置第三行第四列的元素为3.0f
m[2][3] = 3.0f;
构造方式
//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);
4. 缓存 buffer
在Metal 中实现缓存靠的是一个指针。它指向一个在Device 或者 constant 地址空间中的内建或是开发者自定义的数据块。缓存可以被定在程序域中，或是当做函数的参数传递。
5. 纹理 Texture
纹理数据类型是一个句柄,它是一个一维/二维/三维纹理数据.而纹理数据对
应这个一个纹理的某个level的mipmap的全部或者一部分。
access枚举值定义了纹理的访问权利
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, short, int 等;
代码示例：
void foo (texture2d<float> imgA[[texture(0)]],
texture2d<float,access::read> imgB[[texture(1)]],
texture2d<float,access::write> imgC[[texture(2)]])
{
//...
}
类型
- 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
6.采样器类型 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)
7.函数修饰符
Metal 着色器语言支持下列的函数修饰符:
kernel：表示该函数是一个数据并行计算着色函数，它可以被分配在一维/二维/三维线程组中去执行，表示函数要并行计算，其返回值类型必须是void类型，是一个高并发函数
vertex：表示该函数是一个顶点着色函数，它将为顶点数据流中的每个顶点数据执行一次，然后为每个顶点生成数据输出到绘制管线
fragment：表示该函数是一个片元着色函数，它将为片元数据流中的每个片元 和其相关联的数据执行一次，然后将每个片元生成的颜色数据输出到绘制管线中
注意：
1、使用kernel修饰的函数，其返回值类型必须是void类型
2、一个被函数修饰符修饰的函数不能在调用其他也被函数修饰符修饰的函数，这样会导致编译失败，即Kernel、vertex、fragment修饰的函数不能相互调用，也不能同修饰符函数相互调用。但是可以调用普通函数
3、被函数修饰符修饰过的函数,只允许在客户端对其进行操作. 不允许被普通的函数调用
4、Metal中并不是所有函数都需要上述3个修饰符修饰，是可以在Metal中定义普通函数的，即不带任何修饰符的函数
5、只有图形着色函数才可以被vertex和fragment修饰，对于图形着色函数，通过返回值类型可以辨认出是为顶点计算还是像素计算，其返回值也可以是void，意味着不产生数据输出到绘制管线，是一个无意义的动作
代码示例：
//1.并行计算函数(kernel)
kernel void TestKernelFunctionA(int a,int b)
{
//不可以的!
//一个被函数修饰符修饰过的函数,不允许在调用其他的被函数修饰过的函数. 非法
TestKernelFunctionB(1,2);//非法
TestVertexFunctionB(1,2);//非法
//可以! 你可以调用普通函数.而且在Metal 不仅仅只有这3种被修饰过的函数.普通函数也可以存在
Test();
}
kernel void TestKernelFunctionB(int a,int b)
{
}
//顶点函数
vertex int TestVertexFunctionB(int a,int b){
}
//片元函数
fragment int TestVertexFunctionB(int a,int b){
}
//普通函数
void Test()
{
}
8. 变量与参数的地址空间修饰符
Metal 着⾊器语⾔使⽤ 地址空间修饰符 来表示⼀个函数变量或者参数变量 被分配于那⼀⽚内存区域。下面这些修饰符描述了不相交叠地址空间:
device — 设备地址空间
threadgrounp — 线程组地址空间
constant — 常量地址空间
thread — thread地址空间
所有的着⾊函数(vertex, fragment, kernel)的参数,如果是指针或是引⽤, 都必须带有地址空间修饰符号;
对于图形着⾊器函数, 其指针或是引⽤类型的参数必须定义为device或是 constant 地址空间;
对于并⾏计算着⾊函数, 其指针或是引⽤类型的参数必须定义为 device 或是 threadgrounp 或是constant 地址空间;
代码示例：
//变量/参数地址空间修饰符
void TestFouncitionE(device int *g_data,
threadgroup int *l_data,
constant float *c_data)
{
//...
}
8.1 Device Address Space(设备地址空间)
在设备地址空间(Device) 指向设备内存池分配出来的缓存对象, 它是可读也是可写的; ⼀个缓存对象可以被声明成⼀个标量、向量或是⽤户⾃定义结构体的指针或是引⽤。
// an array of a float vector with 4 components 3 device float4 *color;
struct Foo {
float a[3];
int b[2];
};
// an array of Foo elements
device Foo *my_info;
注意: 纹理对象总是在设备地址空间分配内存, device 地址空间修饰符不必出现在纹理类型定义中. ⼀个纹理对象的内容⽆法直接访问. Metal 提供读写纹理的内建函数;
8.2 threadgrounp Address Space 线程组地址空间
线程组地址空间⽤于为 并⾏计算着⾊函数分配内存变量。这些变量被⼀个线程组的所有线程共享。在线程组地址空间分配的变量不能被⽤于图形绘制着⾊函数[顶点着⾊函数, ⽚元着⾊函数]
在并⾏计算着⾊函数中, 在线程组地址空间分配的变量为⼀个线程组使⽤, 声明周期和线程组相同;
/*
1. threadgroup 被并行计算计算分配内存变量, 这些变量被一个线程组的所有线程共享. 在线程组分配变量不能被用于图像绘制.
2. thread 指向每个线程准备的地址空间. 在其他线程是不可见切不可用的
*/
kernel void TestFouncitionF(threadgroup float *a)
{
//在线程组地址空间分配一个浮点类型变量x
threadgroup float x;
//在线程组地址空间分配一个10个浮点类型数的数组y;
threadgroup float y[10];
}
8.3 constant Address Space 常量地址空间
常量地址空间指向的缓存对象也是从设备内存池分配存储, 但是它是只读的;
在程序域的变量必须定义在常量地址空间并且声明的时候初始化; ⽤来初始化的值必须是编译时的常量。
在程序域的变量的⽣命周期和程序⼀样, 在程序中的并⾏计算着⾊函数或者图形绘制着⾊函数调⽤, 但是constant 的值会保持不变;
注意: 常量地址空间的指针或是引⽤可以作为函数的参数。 向声明为常量的变量赋值会产⽣编译错误。声明常量但是没有赋予初值也会产⽣编译错误;
constant float samples[] = { 1.0f, 2.0f, 3.0f, 4.0f };
//对一个常量地址空间的变量进行修改也会失败,因为它只读的
sampler[4] = {3,3,3,3}; //编译失败;
//定义为常量地址空间声明时不赋初值也会编译失败
constant float a;
8.4 thread Address Space 线程地址空间
thread 地址空间指向每个线程准备的地址空间, 这个线程的地址空间定义的变量在其他线程不可⻅, 在图形绘制着⾊函数或者并⾏计算着⾊函数中声明的变量thread 地址空间分配;
kernel void TestFouncitionG(void)
{
//在线程空间分配空间给x,p
float x;
thread float p = &x;
}
9.函数的参数和修饰符
图形绘制或者并⾏计算着⾊器函数的输⼊输出都是通过参数传递。除了常量地址空间变量和程序域定义的采样器以外。参数可以是如下之一：
-device buffer- 设备缓存, ⼀个指向设备地址空间的任意数据类型的指针或者引⽤;
-constant buffer -常量缓存区, ⼀个指向常量地址空间的任意数据类型的指针或引⽤
-texture - 纹理对象;
-sampler - 采样器对象;
-threadGrounp - 在线程组中供各线程共享的缓存.
对于每个着色函数来说，一个修饰符是必须指定的，它用来设置一个缓存、纹理、采样器的位置(定位)，传递修饰符对应的写法如下：
device buffers/ constant buffer --> [[buffer (index)]]
texture -- [[texture (index)]]
sampler -- [[sampler (index)]]
threadgroup buffer -- [[threadgroup (index)]]
在代码中如何表现:
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 可以由开发者来指定.
index是⼀个unsigned integer类型的值，它表示了⼀个缓存、纹理、采样器参数的位置(在函数参数索引表中的位置)。 从语法上讲，属性修饰符的声明位置应该位于参数变量名之后。
//并行计算着色器函数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];
}
//着色函数的多个参数使用不同类型的属性修饰符的情况
kernel void my_kernel(device float4 *p [[buffer(0)]],
texture2d<float> img [[texture(0)]],
sampler sam [[sampler(0)]])
{
//.....
}
10.内建变量属性修饰符
- [[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]] ：片元着色函数使用的单个片元输入数据是由顶点着色函数输出然后经过光栅化生成的（即由顶点着色函数之后的颜色传递到片元着色函数），类似于GLSL中的varying传递纹理/颜色
顶点和片元着色器函数都只能有一个参数被声明为使用stage_in修饰符（即有且仅有一个）
对于一个使用了stage_in修饰符的自定义结构体，其成员可以为一个整型/浮点类型标量，或是整型/浮点类型向量
