Metal-02-Metal Shading Language

宇宙那么大
2,436 阅读11分钟

《目录-iOS & OpenGL & OpenGL ES & Metal》
我们今天来大致讲一下Metal Shading Language 着色语言规范。看一下平时写Metal 需要注意些什么~

一、Metal Shading Language

1、定义及作用

Metal着⾊语⾔ 是⽤来编写 3D 图形渲染逻辑 和 并⾏计算核⼼逻辑的⼀⻔编程语⾔. 当你使⽤ Metal 框架来完成APP 的实现时则需要使⽤Metal 编程语⾔;

Metal 语⾔使⽤Clang 和 LLVM 进⾏编译处理~

Metal 基于C++ 11.0 语⾔设计的，我们主要⽤来编写 在 GPU 上执⾏的图像渲染逻辑代码 以及 通⽤ 并⾏计算逻辑代码;

2、注意点

  1. 对于指针的限制：

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

  2. Metal 像素坐标系统: Metal中 纹理/帧缓存区attachment 的像素使⽤的坐标系统的原点是左上⻆（与OpenGL里的不同）。

3、Metal语⾔ 与 C++ 11.0 的不同之处

虽然 Metal 是基于 C++ 11.0 语⾔设计的，但是有一些 C++ 11.0 中的特性语法，在Metal中 不支持:

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

二、基础数据类型

与GLSL语言不同，float类型后面允许带f或者F。同样，half类型后面允许带h或者H。

1、标量

  • 注意 unsigned 修饰的可以简写成 u。例如：unsigned char 简写 uchar
bool a = true;
char b = 5;
int  c = 15;
//用于表示内存空间的大小
size_t d = 1;
ptrdiff_t e = 2;

2、向量

booln 、charn 、shortn 、intn 、ucharn 、ushortn 、uintn 、halfn 、floatn（n指的是维度 1~4）

//初始化类型1：直接赋值
bool2 A = {1,0};
float4 pos ={1.0,2.0,3.0,4.0};
//初始化类型2：通过内建变量赋值
bool2 A = bool2(1,0);
float4 pos = float4(1.0,2.0,3.0,4.0);

//根据下标取值
float x = pos[0];
float y = pos[1];

//通过for循环赋值
float4 VB;
for(int i = 0; i < 4 ; i++)
{
    VB[i] = pos[i] * 2.0f;
}

向量有且只有两套向量分量（xyzw 、rgba）用来获取元素，但是注意：
可以单套乱序取值，不能两套混合使用，赋值的时候分量不能重复，取值的时候分量可以重复

//通过向量字母来获取元素(向量分量)
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);


float4 pos = float4(1.0f,2.0f,3.0f,4.0f);
float4 swiz = pos.wxyz;  //swiz = (4.0,1.0,2.0,3.0);
float4 dup = pos.xxyy;  //dup = (1.0f,1.0f,2.0f,2.0f);

//pos = (5.0f,2.0,3.0,6.0)
pos.xw = float2(5.0f,6.0f);

//pos = (8.0f,2.0f,3.0f,7.0f)
pos.wx = float2(7.0f,8.0f);

//pos = (3.0f,5.0f,9.0f,7.0f);
pos.xyz = float3(3.0f,5.0f,9.0f);

pos.xr //非法的，不能混合使用

pos.xx = float2(2.0f,3.0f); //非法的，不能连续重复赋值

注意，不要越界

float2 pos;
pos.x = 1.0f; //合法
pos.z = 1.0f; //非法，只定义了2个元素，取第3个就越界了

float3 pos2;
pos2.z = 1.0f; //合法
pos2.w = 1.0f; //非法，只定义了3个元素，取第4个就越界了

列举一下，二维向量、三维向量、四维向量所有的构造方式

//float2类型向量的所有可能的构造方式
float2(float x);
float2(float x,float y);
float2(float2 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);

//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);

3、矩阵

halfnxm 、floatnxm （nxm分别指的是矩阵的⾏数和列数）

//定义一个4x4的矩阵m  注意，4行4列是 m[3][3]
float4x4 m;

//设置第一行/第一列为1.0f
m[0][0] = 1.0f;

//将第二排的4个值都设置为0
m[1] = float4(2.0f);

//设置第三行第四列的元素为3.0f
m[2][3] = 3.0f;

三、纹理类型和采样器类型

1、纹理类型 Textures

纹理类型是⼀个句柄（id）, 它指向⼀个⼀维/⼆维/三维纹理数据。是定义好的枚举值：

enum class access {
sample , //纹理对象可以被采样. 采样⼀维这是使⽤或不使⽤采样器从纹理中读取数据;
read ,   //不使⽤采样器, ⼀个图形渲染函数或者⼀个并⾏计算函数可以读取纹理对象
write    //⼀个图形渲染函数或者⼀个并⾏计算函数可以向纹理对象写⼊数据
};

//sample 用的最多，可读可写可采样

3种纹理：
texture1d<T, access a = access::sample>
texture2d<T, access a = access::sample>
texture3d<T, access a = access::sample>
T : 数据类型 设定了从纹理中读取或是向纹理中写⼊时的颜⾊类型. T可以是half, float, short, int 等

代码示例：


/*
类型 变量 修饰符

类型：二维纹理
	 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
 
*/
void foo (  texture2d<float> imgA [[ texture(0) ]] , 
			texture2d<float, access::read> imgB [[ texture(1) ]], 
            texture2d<float, access::write> imgC [[ texture(2) ]]) 
{ 
	... 
}

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 address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat };
  • 设置纹理s,t,r坐标的寻址模式;
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 };
  • 设置纹理采样的mipMap过滤模式（如果是none,那么只有⼀层纹理⽣效）
enum class mip_filter { none, nearest, linear };

注意: 在Metal 程序中初始化的采样器必须使⽤ constexpr 修饰符声明。 代码示例：

//设置归一化、设置寻址模式、过滤方式
constexpr sampler s ( coord::pixel,
 					  address::clamp_to_zero, 
                      filter::linear );
                      
//如果都是用默认也可以不写，或者只写需要设定的
constexpr sampler a ( coord::normalized ); 
constexpr sampler b ( address::repeat );

四、函数修饰符

Metal 有以下3种函数修饰符：kernel、vertex、fragment。

1、kernel

表示该函数是⼀个数据并⾏计算着⾊函数. 它可以被分配在⼀维/⼆维/三维线程组中去执⾏

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

2、vertex

表示该函数是⼀个顶点着⾊函数 , 它将为顶点数据流中的每个顶点数据执⾏⼀次然后为每个顶 点⽣成数据输出到绘制管线;

3、fragment

表示该函数是⼀个⽚元着⾊函数, 它将为⽚元数据流中的每个⽚元 和其关联执⾏⼀次然后 将每个⽚元⽣成的颜⾊数据输出到绘制管线中;

注意：

  • 被上面三种修饰符修饰的函数中，不能调用 同样被这三个修饰符修饰的函数，否则会直接编译失败。可以调用普通函数。
  • 使⽤kernel 修饰的函数. 其返回值类型必须是 void 类型
  • 只有图形着⾊函数才可以被 vertex 和 fragment 修饰，返回值类型可以辨认出它是为 顶点做计算还是为每像素做计算
  • 图形着⾊函数的返回值可以为 void , 但是这也就意味着该函数不产⽣数 据输出到绘制管线; 这是⼀个⽆意义的动作

五、地址空间修饰符（⽤于变量或者参数）

地址空间修饰符有4种： device、threadgrounp、constant、thread

使⽤ 地址空间修饰符 来表示⼀个函数变量或者参数变量 被分配于那⼀⽚内存区域

  • 所有的着⾊函数(vertex, fragment, kernel)的参数,如果是指针或是引⽤, 都必须带有地址空间修饰符号
  • 对于图形着⾊器函数, 其指针或是引⽤类型的参数必须定义为 device 或是 constant 地址空间
  • 对于并⾏计算着⾊函数, 其指针或是引⽤类型的参数必须定义为 device 或是 threadgrounp 或是 constant 地址空间

1、device Address Space(设备地址空间)

device（设备地址空间）指向设备内存池分配出来的缓存对象（这里设备指的的显存，GPU）。它是可读也是可写的。⼀个缓存对象可 以被声明成⼀个标量、向量、⽤户⾃定义结构体或者结构体的指针

  • 放在显存中，是为了读取更快
  • 纹理对象，不用device修饰，也会默认放在 显存中
//1.修饰指针变量。表示：这个 四维向量颜色 的指针 放在显存中
device float4 *color;
 
//2、修饰结构体指针。表示：把这个结构体的指针 放在显存中
struct Foo { 
	float a[3]; 
	int b[2]; 
};

device Foo *my_info;

2、threadgrounp Address Space 线程组地址空间

threadgrounp（线程组地址空间）⽤于为 并⾏计算着⾊函数分配内存变量。 这些变量被⼀个线程组的所有线程共享. 在线 程组地址空间分配的变量不能被⽤于图形绘制着⾊函数（顶点着⾊函数, ⽚元着⾊函数），也就是在顶点、 ⽚元着⾊函数中不能使用threadgrounp修饰变量

在并⾏计算着⾊函数中, 在线程组地址空间分配的变量为⼀个线程组使⽤, 声明周期和线程组相同

/*
传入参数：
在线程组地址空间分配一个浮点类型变量a的指针
*/
kernel void my_func(threadgroup float *a [[ threadgroup(0) ]], ...)
{
	//在线程组地址空间分配一个浮点类型变量x
    threadgroup float x;
    
    //在线程组地址空间分配一个10个浮点类型数的数组y;
    threadgroup float y[10];
    
}

3、constant Address Space 常量地址空间

  • constant（常量地址空间）指向的缓存对象也是从设备内存池（显存中）分配存储, 但是它是只读的
  • 在程序域中，constant 修饰的变量，必须在声明的时候就初始化并赋值
  • 在程序域中，constant 修饰的变量的生命周期和程序一样，在程序中的并⾏计算着⾊函数或者图形绘制着⾊函数调⽤, 但是 constant 的值会保持不变（也就是一旦初始化就是常量，不允许修改）

注意：constant修饰的变量，如果不赋值、或者后面进行修改，都会产生编译错误

//这样声明是正确的 ✅
constant float samples[] = { 1.0f, 2.0f, 3.0f, 4.0f };

//但是，一旦修改，就编译错误了 ❌
sampler[4] = {3,3,3,3};

//这样声明是错误的 ❌  没有赋值，直接编译失败
constant float a;

4、thread Address Space 线程地址空间

thread（线程地址空间）指向每个线程准备的地址空间，这个线程的地址空间定义的变量在其他线程不可⻅，在 图形绘制着⾊函数或者并⾏计算着⾊函数中声明的变量thread 地址空间分配

  • threadgroup修饰的变量，线程之间可以共享。thread修饰的就不行，
  • 在着色函数中，不能用threadgroup， 可以用thread
kernel void my_func(...)
{
	//在线程空间分配空间给x,p
    float x;
	thread float p = &x;
}

六、传递修饰符（函数参数与变量的属性修饰符）

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

  • device buffer ：设备缓存，一个指向设备地址空间的任意数据类型的指针/引用

  • constant buffer ：常量缓存，一个指向常量地址空间的任意数据类型的指针/引用

  • texture ：纹理对象

  • sampler ：采样器对象

  • threadGroup ：在线程组中供线程共享的缓存

那么，为什么需要属性修饰符？

  • 参数表示资源的定位，可以理解为端口，相当于OpenGl ES中的location
  • 在固定管线和可编程管线进行内建变量的传递
  • 将数据沿着渲染管线从顶点函数传递到片元函数

对于每个着色函数来说，一个修饰符是必须指定的，它用来设置一个缓存、纹理、采样器的位置，传递修饰符对应的写法如下：

  • device buffer ---> [[buffer(index)]]
  • constant buffer ---> [[buffer(index)]]
  • texture ---> [[texture(index)]]
  • sampler ---> [[sampler(index)]]
  • threadGroup ---> [[threadGroup(index)]]
index 可以由开发者来指定，它可以是一个unsigned interger类型的值，也可以是一个自定义的枚举值（枚举放在桥接.h文件中）
表示了一个缓存、纹理、采样器参数的位置，即在函数参数索引表中的位置。
属性修饰符一般放在变量名后面。

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

代码示例

//并行计算着色器函数，属性修饰符"[[buffer(index)]]" 为着色函数参数设定了缓存的位置
/*
kernel :并行计算函数修饰符
void :函数的返回值 （kernel修饰，返回值必须是void）
add_vectors：函数名
const device float4 *inA [[buffer(0)]]：
	const ：修饰的这个变量不能修改
    device：修饰的这个变量，放在显存中
    float4：这个变量的类型：四维向量
    inA：变量名
    [[buffer(0)]]：表示这个变量在 buffer中对应的0这个id
    
thread_position_in_grid：用于表示当前节点在多线程网格中的位置,并不需要开发者传递，是Metal自带的。
*/
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];
}

七、常见的内建变量属性修饰符

  1. [[vertex_id]] ：顶点id标识符，并不由开发者传递
  2. [[position]] ：
    • 在顶点函数中，表示顶点信息，类型是float4
    • 在片元函数中，表示片元的窗口相对坐标（x,y,z,w）,即像素在屏幕上的位置信息
  3. [[point_size]] ：点的大小，float类型
  4. [[color(m)]] ：颜色，m在编译前必须确定
  5. [[stage_in]] ：由顶点函数输出 经过光栅化生成 传入片元函数的数据。 需要注意：无论在顶点还是片元函数中，只能声明一个被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; 
}
avatar
文章
阅读
粉丝
相关推荐
Metal Shader Language 语言规范
2.0k阅读
 · 
8点赞
Metal-01-初识Metal
2.3k阅读
 · 
7点赞
Metal 入门知识
3.5k阅读
 · 
13点赞
探索Metal Petal框架
4.4k阅读
 · 
6点赞
[MetalKit]47-Introducing Metal 2.2 Metal 2.2简介
3.5k阅读
 · 
1点赞
Flutter+Metal实现图像处理
10k阅读
 · 
8点赞
Metal入门
1.2k阅读
 · 
4点赞
Metal 框架之使用 Metal Debugger 查看 GPU 工作负载
1.6k阅读
 · 
5点赞
iOS 利用 Metal 实现 LUT 滤镜
4.3k阅读
 · 
40点赞

友情链接：