跟OpenGL ES类似，Metal也有自己的着色器语言，相比更接近C语言的GLSL，Metal Shading Language的语言更接近于C++。而且，Metal有着和OpenGL管道相同的渲染流程，即：顶点着色器->图元装配->光栅化->片元着色器->提交缓冲区，不同的是，Metal中将着色器描述为函数，但本质是相同的。

但是毕竟是苹果的亲儿子，Xcode不仅可以提供了创建Metal文件的入口，还能使用Clang+LLVM对Metal文件进行编译链接，这相比于要开发者手敲GLSL的OpenGL ES来说，无疑是比较优势的一点，毕竟仅仅是以字符串形式存在的着色器来说，查找错误是一个麻烦的过程。

但是Metal的语法仿佛继承了OC语言恶心的特质，学习起来还是需要一定成本的，本文将对Metal Shading Language做一个学习归纳，方便日后查看。

Metal Shading Language 基本语法

Metal Shading Language使用的是 Clang 和 LLVM 进行编译处理的，它基于C++ 11.0 语言设计，但它也不完全等同与C++的语言，Metal不支持以下特征：

• Lambda表达式
• 递归函数调用
• 动态转换操作符
• 类型识别
• 对象创建new和销毁delete操作符
• 操作符noexcept
• goto跳转
• 变量存储修饰符register和thread_local
• 虚函数修饰符
• 派生类
• 异常处理
• C++标准库

另外，Metal像素坐标系统使用的原点是左上角，这与OpenGL的左下角是不同的，所以我们使用Metal渲染纹理时不再需要翻转。

而且，Metal函数名不能命名为Main，这点和GLSL是不同的。

Metal支持结构体和枚举。

基本数据类型

Metal 有以下基本数据类型：

与OpenGL不同的是，Metal支持数字后缀表示字面量类型，如0.4f，0.5h。

指针类型

Metal Shading Language 也是支持指针类型的，同样是使用 * 表示指针,但是在Metal Shading Language中，对指针的使用以下限制：

• Metal图型和并行计算函数用的的入参如果是指针或者应用类型，必须使用地址空间修饰符（如device，threadgroup，constant）
• 不支持函数指针

向量和矩阵数据类型

向量

Metal Shading Language 使用基本数据类型+维度表示向量，如float3表示每个维度类型为float的3维向量，它的定义可以写为：

short4 vector1 = {1,2,3,4};
float3 vector2 = float3(1.f, 2.f, 3.f);

和GLSL类似，它也可以使用.xyzw或者.rgba读取或赋予对应的值：

short2 vector3 = vector1.xy; // vector3 = (1, 2)
float2 vector4 = vector2.rb; // vector4 = (1.f, 3.f)

向量读取的是后实际是使用索引值，所以向量的xyzw（或rgba）可以颠倒和重复：

vector3.xy = vector3.yx; // vector3 = (2, 1)
vector4.rg = vector4.rr; // vector4 = (1.f, 1,f)

另外，因为xyzw（或rgba）读取的索引值，所以不能读取越界的索引，也不能xyzw和rgba混用，如

vector3.zw // 错误，二维向量没有zw值
vector4.xg // 错误，xyzw和rgba不能混用

向量支持如下类型：

• booln
• charn
• shortn
• intn
• ucharn
• ushortn
• uintn
• halfn
• floatn

0<n≤4,表示维度

矩阵

矩阵支持如下类型：

• halfnxm
• floatnxm

n、m分别表示矩阵的行数和列数

它可以这么定义：

half2x2 matrix = half2x2(1.h,2.h,3.h,4.h);
matrix = {2.h,2.h,2.h,2.h};

纹理类型

纹理类型是一个句柄，它指向一个一维/二维/三维纹理数据。他是这么定义的:

texture1d<T, access a = access::sample>
texture2d<T, access a = access::sample>
texture3d<T, access a = access::sample>

T表示数据类型，设定了从纹理中读取的或是向纹理中写入时的颜色类型，可以是half，float，short，int等

access是一个枚举值，定义了访问权利：

enum class access {
    sample, // 默认值， 纹理对象可以被采样，采样一维时使用或不使用都从纹理中读取数据
    read, // 不使用采样器，一个图形渲染函数或一个并行计算函数可以读取纹理对象
    write // 一个图形渲染函数或者一个并行函数可以像纹理对象写入数据
}

示例如下：

void foo(texture2d<float> imgA[[ texture(0) ]],
         texture2d<float, access::read> imgB [[ texture(1) ]],
         texture2d<float, access::write> imgC [[ texture(2) ]])

采样器类型

采样器类型决定了如何对一个纹理进行采样操作，在Metal框架中有一个对应着着色器语言的采样器对象MTLSamplerState，这个对象作为图形渲染着色器函数参数或者并行计算函数的参数传递。
需要注意的是，Metal程序中，初始化采样器必须使用constexpr修饰符声明。

constexpr sampler s(coord::normalized)

例中括号内表示采样器的采样参数，是一个枚举值，Metal中采样器以下几种类型：

例：

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

修饰符

修饰符是Metal Shading Language中比较难理解的点，Metal中的修饰符有函数修饰符、地址空间修饰符和传递修饰符3种

函数修饰符

函数修饰符位于函数之前，例如：

kernel void foo(...) {
 ...
}

Metal有以下3种函数修饰符

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

只有图形着色器函数才可以被vertex和fragment修饰，对于图形着色函数，返回值类型可以辨别出他是为顶点做计算还是为每个像素做计算，图型函数返回值可以为void，但这意味着该函数不产生数据输出到绘制管线，这是一个无意义的动作。

函数修饰符有以下注意点：

1. 使用kernel修饰的函数，其返回值类型必须是void类型；
2. 一个被函数修饰符修饰的函数不能调用其他被函数修饰符修饰的函数；

地址空间修饰符

地址空间修饰符用于变量或者参数，位于变量或参数类型之前，例如：

device float4 *color;

空间修饰符表示一个变量被分配在哪一片内存区域，所有着色器函数（vertex，fragment，kernel）的参数，如果是指针或者应用，都必须带有地址空间修饰符。

空间修饰符有以下4种。

• device：设备地址空间
• threadgroup：线程组地址空间
• constant：常量地址空间
• thread：线程地址空间

对于图形着色器函数，其指针或者应用类型参数必须定位为device或constant地址空间；
对于并行计算函数，其指针或者应用类型参数必须定义为device或threadgroup或constant地址空间；

1.设备地址空间

指的是设备内存池分配出来的缓存对象，它是可读可写的，一个缓存对象可以被声明为一个标量、向量或者是用户自定义结构体的指针或应用。
另外注意的是，纹理对象总是在设备地址空间分配内存，device地址空间修饰符不必出现在纹理类型定义中，一个纹理对象的内容无法直接访问，Metal提供读写纹理的内建函数。

例：

device float4 *color;
struct Foo {
    float a[3];
    int b[2];
}
device Foo *my_info;

2.线程组地址空间

指的是用于为并行计算着色函数分配内存变量的空间，这些变量被一个线程组的所有线程共享。
在线程组地址空间分配的变量不能用与图形绘制着色函数。
在并行计算着色函数中，在线程组地址空间分配的变量为一个线程组使用，声明周期和线程组相同。
例：

kernel void my_func(threadgroup float *a [[threadgroup(0)]]) {
    threadgroup float x;
    
    threadgroup float b[10];
}

3.常量地址空间

指向的缓存对象也是设备内存池分配，但他是只读的，它修饰的变量必须在定义的时候初始化，用来初始化的值必须是编译时的常量，修饰的变量在程序域的生命周期和程序一样，在程序中的并行计算着色函数或者图形绘制着色函数调用时，它的的值都会保持不变。
值得注意的是，常量地址空间的指针或引用也可以作为函数的参数，为声明的常量赋值会产生编译错误，声明常量但是没有赋予初值也会编译错误；
例如：

constant float samplers[] = {1.f, 2.f, 3.f, 3.f};

sampler[4] = {3,3,3,3}; // 编译失败

constant float a; // 编译失败

4.线程地址空间

指向每个线程准备的地址空间，这个线程地址空间定义的变量在其他线程不可见。可以在图形绘制函数或者并行计算函数中使用。
例：

kernel void my_func(...) {
    float x;
    thread float = &x;
}

传递修饰符

对于图形绘制或并行着色器函数来说，输入/输出都需要通过参数传递（除了常量地址空间变量和程序域中定义的采样器以外），参数可以以下之一：

• device buffer - 设备缓存，一个指向设备地址空间的任意数据类型的指针或者引用
• constant buffer - 常量缓存区，一个指向常量地址空间的任意数据类型的引用
• texture - 纹理对象
• sampler - 采样器对象
• threadgroup - 在线程组中供各线程共享的缓存

注意的是，着色器的缓存（device 和 constant）不能重名。

对于每个着色器来说，一个变量传递修饰符是必须指定的，它用来设置一个缓冲、纹理、采样器的位置。
传递修饰符位于参数变量之后，用[[]]表示，并在()内指定它的位置。
通常我们有以下几种类型可以从外部传入着色器：

• [[buffer(index)]] // device buffer / constant buffer
• [[texture(index)]] // texture
• [[sampler(index)]] // sampler
• [[threadgroup(index)]] // threadgroup

index是一个unsigned integer类型的值，它表示一个缓存、纹理、采样器参数的位置。

参考代码如下：

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: 用于表示当前节点在多线程网络中的位置,不需要用户传入

内建变量属性修饰符

除了以上的传递修饰符，Metal还提供了内建的传递修饰符。常见的有以下几种：

• [[vertex_id]] 顶点id表示符，即当前是第几个顶点
• [[position]] 顶点信息，表述片元的窗口相对坐标（x, y, z, 1/w）
• [[point_size]] 点的大小
• [[color(m)]] 颜色,m编译前得确定，表示传入颜色附着点m
• [[stage_in]] 表示由顶点着色函数输出经过光栅化生成传入片元函数的数据，一个顶点和片元函数都只能有一个参数被声明为stage_in