Metal Shading Language - 语法小结Metal Shading Language - 语法小结

Metal Shading Language简述

• Metal 着色语言是用来编写3D图形渲染逻辑 和 并行计算核心逻辑的一门编程语言，如果需要使用Metal框架来实现某些逻辑则需要使用该语言
• Metal是通过Xcode的Clang 和 LLVM进行编译、链接，无需在手动编译
• Metal基于C++ 11.0语言设计的，在C++基础上多了一些扩展和限制
• Metal 像素坐标系统：Metal中纹理 或者 帧缓存区attachment的像素原点是在左上角

Metal语言的部分限制

• 递归函数
• C++标准库在Metal语言中也不可使用
• Metal图形和并行计算函数用到的入参，如果是指针 / 引用必须使用地址空间修饰符（比如device、threadgroup、constant）,否则编译报错
• 无法使用函数指针，也就是方法入参、返回值。
• Metal文件中不得出现main()函数
• 无法进行异常的捕捉和处理

Metal语言-数据类型

包含：

1. 基本数据类型
   1. 标量类型
   2. 向量类型
   3. 矩阵类型
2. 纹理类型
3. 采样器类型

基本数据类型

标量类型

• 常用的：bool,char,int,uint,float

向量类型

向量支持如下类型:

• 最多支持4维向量

bool2 A= {1,2};
float4 pos = float4(1.0,2.0,3.0,4.0);
pos = float4(float2(1.0,2.0),float2(3.0,4.0));
float x = pos[0];
x = pos.x;
float y = pos[1];
y = pos.y;

矩阵类型

矩阵类型有且只有两种：halfnxm,floatnxm n:行数 m:列数

half4x4 n;
float4x4 m;
m[0] = float4(1,1,1,1);
m[1][1] = 1;
m[3][3] = 3;

纹理类型

纹理类型是一个句柄，指向一个一维/二维/三维的纹理数据;相当于OpenGL中的textureBufferID.

纹理数据类型

限制从纹理中读取或是向纹理中写入是的颜色类型. 可以有：half,float,short,int等类型，一般使用：float

访问权限

• sample：可读可写，纹理可以被采样。权限默认值
• read：只读，不使用采样器。渲染函数、并行计算函数可以读取纹理数据
• write：可读可写，不使用采样器。渲染函数、并行计算函数可以修改纹理数据

//这两种写法等价
texture1d<float, access::sample>
texture1d<float>

texture2d<float, access::read>
texture3d<float, access::write>

kernel void my_kernel(texture2d<float> img [[texture(0)]])
{
    //.....
}

采样器类型

采样器类型决定了如何对一个纹理进行采样。

在Metal中有一个对应着色器语言色采样器对象MTLSamplerState这个对象作为：图形渲染着色器函数参数 、 并行计算函数的参数。

在Metal程序中初始化的采样器必须使用constexpr修饰符声明

constexpr sampler s(coord::pixel,
                  address::clamp_to_zero,
                  filter::linear); 
constexpr sampler a(coord::normalized);
constexpr sampler b(address::repeat);

Metal语言-函数修饰符

• kernel：表示该函数是一个数据并行计算着色函数，它可以被分配在一维/二维/三维线程组中去执行
• vertex：表示该函数是一个顶点着色函数，它将为顶点数据流中的每个顶点数据执行一次，然后为每个顶点生成数据输出到绘制管线
• fragment：表示该函数是一个片元着色函数，它将为片元数据流中的每个片元 和其相关联的数据执行一次，然后将每个片元生成的颜色数据输出到绘制管线中

注意点：

1. 被函数修饰符修饰的函数体内不能调用任何被函数修饰符修饰的函数
2. 被函数符修饰的函数系统会自动调用，开发者不允许调用
3. kernel修饰符修饰的函数返回值必须为void，其他两种（vertex、fragment）返回值也可以为void,但是这样会导致函数执行无效
4. 被函数符修饰的函数可以调用普通函数

//并行计算函数(kernel)
kernel void KernelFunction(int a,int b)
{ 
    KernelFunctionB(1,2);//非法调用！！！
}
kernel void KernelFunctionB(int a)
{ ... }


//顶点函数
vertex int VertexFunction(int a,int b)
{ 
    Test();//合法调用！！！
}

//片元函数
fragment int VertexFunction(int a,int b)
{ ..... }

//普通函数
void Test()
{ ..... }

Metal语言-地址空间修饰符

Metal着色器语言使用地址空间修饰符来表示一个 函数变量或者 参数变量 被分配于哪一片内存区域.所有被函数符修饰的函数其参数如果是指针、引用，就必须使用地址空间修饰符 包含以下4种：

1. device：设备（GPU缓存）地址空间
2. constant：常量地址空间
3. threadgrounp：线程组地址空间
4. thread：线程地址空间

Device Address Space

• 设备地址空间指向GPU缓存分配出来的缓存对象，该值可读可写，一个缓存对象可以被声明成一个标量、向量或是用户自定义结构体的指针/引用
• device放在变量类型之前
• 纹理对象总是在设备地址空间分配内存，此处的device可以省略。纹理对象无法直接访问，需要通过纹理的内建变量（ [[texture(0)]] ）来获取

//1. 修饰指针变量
device float4 *color;

struct Struct{
    float a[3];
    int b[2];
};
//2.修饰结构体类的指针变量
device CCStruct *my_CS;

constant Address Space

• 常量地址空间指向的缓存对象也是从设备内存池分配存储，仅可读
• 在程序域的变量必须定义在常量地址空间并且声明时初始化.不初始化会导致编译错误
• 在程序域的变量的生命周期和程序一样，在程序中的并行计算着色函数 或者 图形绘制着色函数调用，但是constant的值会保持不变

constant float samples[] = { 1.0f, 2.0f, 3.0f, 4.0f };

threadgroup Address Space

• 线程组地址空间用于为并行计算着色器函数分配内存变量，这些变量被一个线程组的所有线程共享
• 在线程组地址空间分配的变量不能用于图形绘制着色函数（即顶点着色函数 / 片元着色函数），即在图形绘制着色函数中不能使用线程组
• 在并行计算着色函数中，在线程组地址空间分配的变量为一个线程组使用，生命周期和线程组相同

kernel void KernelFouncition(threadgroup float *a)
{
    //在线程组地址空间分配一个浮点类型变量x
    threadgroup float x;
}

thread Address Space

• 线程地址空间指向每个线程准备的地址空间，也是在GPU中，该线程的地址空间定义的变量在其他线程不可见（即变量不共享）
• 在图形绘制着色函数 或者 并行计算着色函数中声明的变量，在线程地址空间分配存储

kernel void CCTestFouncitionG(void)
{
    //在线程空间分配空间给x,p
    float x;
    thread float p = &x;
}

注意：

1. 在图形着色器函数(顶点函数 片元函数),其指针/引用类型的参数必须定义为device、constant地址空间
2. 在并行计算函数(kernel函数)其指针/引用类型的参数必须定义为 device、threadgroup、constant
3. 被thread修饰的变量无法共享，所以只能在三类函数体内进行使用

函数参数与变量的传递修饰符，即属性修饰符

图形绘制 或者 并行计算着色器函数的输入输出都是通过参数传递，除了常量地址空间变量和程序域定义的采样器之外, 其他参数修饰的可以是如下之一，常用的有以下5种属性修饰符：

• device buffer 设备缓存：一个指向设备地址空间的任意数据类型的指针/引用
• constant buffer 常量缓存：一个指向常量地址空间的任意数据类型的指针/引用
• texture 纹理对象
• sampler 采样器对象
• threadGroup 在线程组中供线程共享的缓存 参数表示资源的位置句柄，可以理解为端口，相当于OpenGl ES中的location

vertex int foo (const device float4 *inA [[ buffer(0) ]],
               device float4 *out [[ buffer(1) ]]
               texture2d<float> imgA[[texture(0)]],
               texture2d<float> imgB[[texture(1)]],
               threadgroup td [[threadgroup(0)]])
{
    //...
}

常见的内建变量修饰符

• [[vertex_id]]:顶点函数中的index，并不由开发者传递
• [[position]]:在顶点着色函数中，表示当前的顶点信息，类型是float4、 在片元函数中还可以描述该像素点在窗口的相对坐标（x，y，z，1/w），即该像素点在屏幕上的位置信息
• [[point_size]] ：点的大小，类型是float
• [[color(m)]] ：颜色，m在编译前就必须确定
• [[ thread_position_in_grid ]]: 用于表示当前节点在多线程网格中的位置，只能用在kernel函数中
• [[stage_in]]：片元着色函数使用的单个片元输入数据是由顶点着色函数输出然后经过光栅化生成的，也就是片元函数的入参用于对应顶点函数的返回值.只允许在片元函数的参数中出现1次；可以使用各种标量、自定义类型.

//顶点函数
vertex float4 VertexFunction(uint vertexId [[vertex_id]],
                           constant int *i [[buffer(0)]])
{ 
    return float4(1,2,3,4);
}

//片元函数
fragment int VertexFunction(float4 param [[stage_in]])
{ ..... }