Metal着色语言规范

Metal着色语言规范：Metal-Shading-Language-Specification.pdf，目前的版本是Version 3.2，这个pdf有298页。

官方文档英文的，硬看的话比较费时间，这里根据网上搜到的，整理一番。

1、Metal着色语言认识

Metal着色语言是用来编写3D图形渲染逻辑、并行Metal计算核心逻辑的一门编程语言，比你高切当使用Metal框架来完成APP的某些功能时也需要使用Metal编程语言。

Metal语言使用Clang和LLVM进行编译处理，编译器对于GPU上的代码执行效率有更好的控制。

Metal基于C++ 11.0语言设计，在C++基础上多了一些扩展和限制，目前版本更新到了C++ 14.0。

由于不懂C++语言，所以两者的差异就不做比较。

2、数据类型

2.1、基础数据类型

包括标量、向量、矩阵。

2.1.1、标量

标量数据类型：

2.1.2、向量

支持类型：booln、charn、shortn、intn、ucharn、ushortn、uintn、halfn、floatn

注意：这里的n表示向量的维度，最多不超过4维

使用规则

1. 直接通过下标来使用
2. 通过字母来使用，有 xyzw 和 rgba 两种表达形式， 分别代表顶点坐标和色值。

示例代码：

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.1.3、矩阵

有两种类型，halfnxm、floatnxm。nxm表示n行m列，最多就是4行4列。可以把矩阵看做一个二维数组来使用。

float2类型向量的构造方式

//float2类型向量的所有可能的构造方式
//1个一维向量
float2(float x);
//2个一维向量
float2(float x,float y);
//1个二维向量
float2(float2 x);

float3类型向量的构造方式

//float3类型向量的所有可能的构造的方式
//1个一维向量
float3(float x);
//3个一维向量
float3(float x,float y,float z);
//1个一维向量 + 1个二维向量
float3(float a,float2 b);
//1个二维向量 + 1个一维向量
float3(float2 a,float b);
//1个三维向量
float3(float3 x);

float4类型向量的构造方式

//float4类型向量的所有可能构造方式
//1个一维向量，表示一行都是x
float4(float x);/
//4个一维向量 --> 4维向量
float4(float x,float y,float z,float w);
//2个二维向量 --> 4维向量
float4(float2 a,float2 b);
//1个二维向量+2个一维向量 --> 4维向量
float4(float2 a,float b,float c);
float4(float a,float2 b,float c);
float4(float a,float b,float2 c);
//1个三维向量+1个一维向量 --> 4维向量
float4(float3 a,float b);
float4(float a,float3 b);
//1个四维向量 --> 4维向量
float4(float4 x);

2.2、其它类型

有两种，纹理类型和采样器类型。

2.2.1、纹理类型

纹理类型是一个句柄，指向一维 /  二维 / 三维纹理数据，而纹理数据对应一个纹理的某个level的mipmap的全部或者一部分。

纹理类型的定义

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

1、texture1d, texture2d, texture3d都表示这是一个纹理类型，分别定义的事一维 / 二维 / 三维。

2、T是一个泛型，表示从纹理中读取数据 或是 写入时的颜色类型，T可以是half、float、short、int等。

3、access表示纹理访问权限，默认是sample。

纹理访问权限

enum class access {
    sample, 
    read, 
    write
};

sample：纹理对象可以被采样，采样一维时使用 或者 不使用都可以从纹理中读取数据（即可读可写可采样）。

read：不适用采样器，一个图形渲染函数 或者 一个并行计算函数可以读取纹理对象（即仅可读）。

write：一个图形渲染函数 或者 一个并行计算可以向纹理写入数据（即 可读可写）。

2.2.2、采样器类型 Samples

对采样器设置采样类型，决定了对这个纹理进行采样时的操作方式。在Metal框架中通过采样器对象MTLSamplerState进行设置采样器类型，这个对象作为图形渲染着色器函数参数 或者 并行计算函数的参数传递。有以下几种状态：

1、coord

• 描述：从纹理中采样时，纹理坐标是否需要归一化
• 参数：enum class coord { nornalized, pixel };

2、filter

• 描述：纹理采样过滤方式，统一设置，包括访达 / 缩小两种过滤方式
• 参数：enum class filter { nearest, linear };

3、min_filter

• 描述：设置纹理采样的缩小过滤方式
• 参数：enum class min_filter { nearest, linear }; 邻近过滤、线性过滤

4、mag_filter

• 描述：设置纹理采样的放大过滤方式
• 参数：enum class mag_filter { nearest, linear }; 邻近过滤、线性过滤

5、s_address、t_address、r_address

• 描述：设置纹理 s、t、r 坐标（对应纹理坐标的  x 、y、z ）寻址方式
• 参数s_address：enum class s_address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat };
• 参数t_address：enum class t_address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat };
• 参数r_address：enum class r_address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat };

6、address

• 描述：设置所有纹理坐标的寻址方式
• 参数：enum class address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat };

7、mip_filter

• 描述：设置纹理采样的mipMap过滤方式，如果是none，那么只有一层纹理生效；
• 参数：enum class mip_filter { none, nearest, linear };

定义：

/*
constexpr：修饰符（必须写）
sampler：类型
s：采样器变量名称
参数
    - coord: 是否需要归一化，不需要归一化，用的是像素pixel
    - address: 地址环绕方式
    - filter: 过滤方式
*/
constexpr sampler s(coord::pixel, address::clamp_to_zero, filter::linear);

constexpr sampler a(coord::normalized);

constexpr sampler b(address::repeat);

注意：constexpr作为修饰符必须写

3、函数修饰符

函数修饰符用来修饰函数，放在函数的最前面，即位于函数返回值的前面。有三种，kernel、vertex、fragment。

kernel：表示该函数是一个数据并行计算着色函数，它可以被分配在一维 / 二维 / 三维线程组中去执行，表示函数要并行计算，其返回值类型必须是void类型，是一个高并发函数。

vertex：表示该函数是一个顶点着色函数，它将为顶点数据流中的每个顶点数据执行一次，然后为每个顶点生成数据输出到绘制管线。

fragment：表示该函数是一个片元着色函数，它将为片元数据流中的每个片元 和 其相关联的数据执行一次，然后将每个片元生成的颜色数据输出到绘制管线中。

示例：

//并行计算函数(kernel)
kernel void CCTestKernelFunctionA(int a,int b)
{ 
    /*
     注意:
     1. 使用kernel 修饰的函数返回值必须是void 类型
     2. 一个被函数修饰符修饰过的函数,不允许在调用其他的被函数修饰过的函数. 非法
     3. 被函数修饰符修饰过的函数,只允许在客户端对其进行操作. 不允许被普通的函数调用.
     */
     
    //不可以的!
    //一个被函数修饰符修饰过的函数,不允许在调用其他的被函数修饰过的函数. 非法
    CCTestKernelFunctionB(1,2);//非法，错误调用！！！
    CCTestVertexFunctionB(1,2);//非法，错误调用！！！
    
    //可以! 你可以调用普通函数.而且在Metal 不仅仅只有这3种被修饰过的函数.普通函数也可以存在
    CCTest();
    
}

//并行计算函数
kernel void CCTestKernelFunctionB(int a,int b)
{
    .....
}

//顶点函数
vertex int CCTestVertexFunctionB(int a,int b)
{
    .....
}

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

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

注意：

1、使用kernel修饰的函数，其返回值类型必须是void类型

2、Metal中并不是所有函数都需要上述3个修饰符修饰，是可以在Metal定义普通函数的，即不带任何修饰符的函数。

3、被函数修饰符修饰的函数不能相互调用，只能调用普通函数，普通函数也不能调用被修饰符修饰的函数。这样容易理解，它们各自有其特殊的含义，是要被系统调用的。

4、只有图形着色函数才可以被 vertex 和 fragment 修饰，对于图形着色函数，通过返回值类型可以辨认出是为顶点计算还是像素计算，其返回值也可以是void，意味着不产生数据输出到绘制管线，是一个无意义的动作。

4、变量的地址空间修饰符

地址空间修饰符用来表示一个变量 或 参数要分配在哪一片区域，有device、threadgroup、constant、thread四种。

注意事项：

1、所有的着色函数（vertex、fragment、kernel）的参数，如果是指针 / 引用，都必须带有地址空间修饰符号。

2、对于图形着色器函数（即vertex / fragment修饰的函数），其指针 / 引用类型的参数必须定义为device、constant地址空间。

3、对于并行计算函数（即kernel修饰的函数），其指针 / 引用类型的参数必须定义为device、threadgroup、constant。

4、并不是所有的变量都需要修饰符，也可以定义普通变量（即无修饰符的变量）。

示例代码：

//变量/参数地址空间修饰符
void CCTestFouncitionE(device int *g_data,
                       threadgroup int *l_data,
                       constant float *c_data
                       )
{
    //...
    
}

4.1、device：设备地址空间修饰符

设备地址空间指向设备内存池分配出来的缓存对象（设备指显存，即GPU），即GPU空间分配的缓存对象，它是可读可写的。这个缓存对象可以存储变量和用户自定义结构体的指针 / 引用。

示例代码：

// 设备地址空间: device 用来修饰指针.引用
//1.修饰指针变量
device float4 *color;

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

注意事项：

1. 纹理对象总是在设备地址空间分配内存，即纹理对象默认分配在显存中。

2. device地址空间修饰符不必出现在纹理类型定义中。

3. 一个纹理对象的内容无法直接访问，Metal提供读写纹理的内建函数，通过 内建函数访问纹理对象。

4.2、constant：常量地址空间修饰符

constant指向的缓存对象也是存储在显存中，但是仅可读。

示例代码：

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

//对一个常量地址空间的变量进行修改也会失败,因为它只读的
sampler[4] = {3,3,3,3}; //编译失败; 

//定义为常量地址空间声明时不赋初值也会编译失败
constant float a;

4.3、threadgroup：线程组地址空间修饰符

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

/*
 1. threadgroup 被并行计算计算分配内存变量, 这些变量被一个线程组的所有线程共享. 在线程组分配变量不能被用于图像绘制.
 2. thread 指向每个线程准备的地址空间. 在其他线程是不可见切不可用的
 */
kernel void CCTestFouncitionF(threadgroup float *a)
{
    //在线程组地址空间分配一个浮点类型变量x
    threadgroup float x;
    
    //在线程组地址空间分配一个10个浮点类型数的数组y;
    threadgroup float y[10];
    
}

4.4、thread：线程地址空间修饰符

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

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

5、变量的属性修饰符

在函数的传递参数中，除了常量地址空间变量和程序预定义的采样器以外，也即从外界传入的参数需要使用使用属性修饰符。

作用：标志从客户端传递资源到服务器端的定位。也就是OpenGL ES中的通道location。

属性修饰符类型有5种：

• device buffer 设备缓存：一个指向设备地址空间的任意数据类型的指针 / 引用
• 常量缓存：一个指向常量地址空间的任意数据类型的指针 / 引用
• 纹理对象
• 采样器对象
• 在线程组中仅供线程共享的缓存

示例代码：

在代码中如何表现:
 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 可以由开发者来指定.

注意事项：

1、index是一个unsigned interger类型的值，表示了一个缓存、纹理、采样器参数的位置（即在函数参数索引表中的位置，相当于OpenGL ES中的location）

2、从语法上来说，属性修饰符的声明位置应该位于参数变量名之后

示例代码：

//并行计算着色器函数add_vectros ,实现2个设备地址空间中的缓存A与缓存B相加.然后将结果写入到缓存out.
//属性修饰符"(buffer(index))" 为着色函数参数设定了缓存的位置
//thread_position_in_grid：用于表示当前节点在多线程网格中的位置,并不需要开发者传递，是Metal自带的。
/*
 kernel：并行计算函数修饰符
 void：函数返回值类型
 add_vectros：函数名
 const device float4 *inA [[buffer(0)]]：定义了一个float4类型的指针，指向一个4维向量空间，放在设备内存空间（即显存GPU中）
    - const device：只决定放在哪里
    - inA：变量名
    - [[buffer(0)]] 对应 buffer中0这个id
 */
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];
}

//着色函数的多个参数使用不同类型的属性修饰符的情况
//纹理读取的方式的sampler，即采样器，[[sampler(0)]]表示采样器的缓存id
kernel void my_kernel(device float4 *p [[buffer(0)]],
                      texture2d<float> img [[texture(0)]],
                      sampler sam [[sampler(0)]])
{
    //.....
    
}

6、内建变量修饰符

对于特殊的变量提供了内建的修饰符直接使用，有4种。

[[vertex_id]]：顶点id标识符，并不由开发者传递

[[position]]：在顶点着色函数中，表示当前的顶点信息，类型是float4，还可以表示描述了片元的窗口的相对坐标（x， y， z， 1/w），即该像素点在屏幕上的位置信息

[[point_size]]：点的大小，类型是float

[[color(m)]]：颜色，m在编译前就必须确定

[[stage_in]]：片元着色函数使用的单个片元输入数据是由顶点着色函数输出，然后经过光栅化生成的（即由顶点着色函数之后的颜色传递到片元着色函数），类似GLSL中的varying传递纹理/颜色

注意事项：

1. 顶点和片元着色器函数都只能有一个参数被声明为使用stage_in修饰符（即有且仅有一个）

2. 对于一个使用了stage_in修饰符的自定义结构体，其成员可以为一个整型 / 浮点数类型标量，或是整型 / 浮点类型向量 

//定义了片元输入的结构体，
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; 
}

本文主要参考一下文章：

十四、Metal - Metal Shader language (着色语言规范)总结

Metal的语言编程规范

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