Metal Shader Language语法规范

Doe豆
1,837 阅读9分钟

Metal 介绍

  • Metal基于C++ 11.0 语⾔设计,是⼀个⽤来编写3D图形渲染逻辑和并⾏计算核⼼逻辑的编程语⾔,编写Metal框架的APP需要使⽤Metal着⾊语⾔程序.
  • Metal着⾊器语⾔使⽤Clang和LLVM,编译器对于在GPU上的代码执⾏效率有更好的控制.

假如当需要利用GPU的高计算能力时,我们也可以使用Metal来帮助我们,因为它不仅仅可以做图形渲染,也拥有高并发计算的能力. 比如在AVFoundation做人脸识别、二维码识别、音视频的编解码等工作中，需要对每一帧的数据进行计算，这种情况下的计算量是非常大的，苹果就利用硬件加速器(指的就是GPU)来进行高并发计算。


Metal和C++11.0 异同

1. C++ 11.0 特性在Metal 语⾔中不⽀持之处

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

2.Metal 语⾔中对于指针使⽤的限制

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

3.Metal 像素坐标系统

  • Metal 中纹理/帧缓存区attachment的像素使⽤的坐标系统的原点是左上⻆

Metal数据类型

1.标量类型

Metal ⽀持后缀表示字⾯量类型, 例如 0.5F, 0.5f; 0.5h, 0.5H. 这在GLSL中是不被允许的.

bool a = true;
char b = 1;
int  c = 15;
size_t d = 1;

2.向量类型

向量支持如下类型:

n指的是维度

  • booln
  • charn
  • shortn
  • intn
  • ucharn
  • ushortn
  • uintn
  • halfn
  • floatn
语法规则:
bool2 a = [1, 2]; //布尔类型的二维向量,也可以理解为布尔数组使用
float4 pos = float4(1.0f,2.0f,3.0f,4.0f); 
// 获取pos中数据的方式
float x = pos[0];  //index下标

// 向量分量(字母)来获取元素有两种: 'xyzw' | 'rgba',这是索引对应的
int4 t = int4(0,1,2,3);
int a = t.x; // a = 0
int b = t.r; // b = 0

// 多个分量访问
float4 t = float4(0,1,2,3);
t.xyzw = float4(3,4,5,6); // 覆盖重新赋值
t.z = 7;
t.xy = float2(8,9);

// 多分量乱序/重复访问  -- GLSL是不允许乱序的
float4 t = float(0,1,2,3);
float4 tt = t.zywx;
float4 ttt = t.xxyy;
t.xw = float2(4,5);
t.wx = float2(6,7);
t.ww = float2(6,7); // !!! 重复赋值是不允许的,

float2 t = float2(1,2);
t.x = 3;
t.z = 4; // 这是不合法的,float2 => z 是数组越界了

// 不可以混用
float t = float4(0,1,2,3);
t.x = 4;
t.xg = float2(5,6); // 不合法,‘xyzw’和‘rgba’是不可以混用的

3.矩阵

nxm分别指的是矩阵的行数和列数,最大支持4行4列 例:4x4

  • halfnxm
  • floatnxm
语法规则
float4x4 m;
m[1] = float4(2.0f); // 指定第二行都为2.0f
m[1] = float4(1,2,3,4); // 指定第二行不同的数据
m[0][1] = 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);

纹理类型

纹理类型是⼀个句柄, 它指向⼀个⼀维/⼆维/三维纹理数据. 在⼀个函数中描述纹理对象的类型;

enum class access {sample ,read ,write};

  • sample:纹理对象可以被采样. 采样⼀维(sample2D/sample3D:采样二维/三维纹理数据),这是使⽤或不使⽤采样器从纹理中读取数据;
  • 不使⽤采样器, ⼀个图形渲染函数或者⼀个并⾏计算函数可以读取纹理对象;
  • ⼀个图形渲染函数或者⼀个并⾏计算函数可以向纹理对象写⼊数据;
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 函数名称 (texture2d<float> imgA [[ texture(0) ]] , 
texture2d<float, access::read> imgB [[ texture(1) ]], 
texture2d<float, access::write> imgC [[ texture(2) ]]) 
{
 ... 
}

修饰符以下详细说明.


采样器类型

在GLSL中可以对纹理的环绕方式、过滤方式等进行设置,同样在Metal中,可以使用采样器类型,来对一个纹理进行采样操作.在Metal框架中,有一个对应着色器语言的采样器的对象MTLSampleState,作为图形渲染着色器函数参数,或是并行计算函数的参数传递

enum class min_filter { nearest, linear }; 设置纹理采样的缩⼩过滤方式;
enum class mag_filter { nearest, linear }; 设置纹理采样的放⼤过滤方式;

设置纹理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 };

注意:在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)

函数修饰符

Metal函数修饰符有3种

  • kernel:表示该函数是⼀个数据并⾏计算着⾊函数. 它可以被分配在⼀维/⼆维/三维线程组中去执⾏.(当某段代码需要高效计算时,可以使用kernel进行修饰)
  • vertex:表示该函数是⼀个顶点着⾊函数, 它将为顶点数据流中的每个顶点数据执⾏⼀次,然后为每个顶点⽣成数据输出到绘制管线,(可以理解为在GLSL种,顶点着色源码main函数中,返回的gl_Position);
  • fragment: 表示该函数是⼀个⽚元着⾊函数, 它将为⽚元数据流中的每个⽚元和其关联执⾏⼀次,然后将每个⽚元⽣成的颜⾊数据输出到绘制管线中(可以理解为在GLSL中,片元着色源码main函数中,返回的gl_FragColor)

注意:

  • 使⽤kernel 修饰的函数. 其返回值类型必须是void类型;
  • 在函数修饰符修饰过的函数中,不可以调用被函数修饰符修饰的函数,否则编译失败
kernel void test1(...){...}

vertex float4 test2(...)
{
	test1(...); //编译失败
}

void test3(...)
{
	test1(...); //编译通过
}

变量或参数的地址空间修饰符

Metal着色器语言使用地址空间修饰符来表示一个函数变量或参数变量,被分配在哪一片内存区域.

  • device
  • threadgrounp线程组
  • constant
  • thread

注意:

  • 对于图形着色器函数(vertex、fragment修饰的函数),当参数涉及指针变量或者引用类型时,必须使用device或constant的地址空间修饰符
  • 对于并行计算着色函数(kernel修饰的函数),当参数涉及指针变量或引用类型时,必须使用device或treadgrounp或这constant的地址空间修饰符

1.Device Address Space设备地址空间

被device修饰的参数,会指向设备内存池分配出来的缓存对象,即显存.它是可读写的,而且读取速度快.


device float4 *color; 

struct foo{
	float2[3];
    int b[2];
};
device foo *info;

纹理对象是默认分配在设备内存空间; ⼀个纹理对象的内容⽆法直接访问. Metal 提供读写纹理的内建函数;

2.threadgroup Address Space线程组地址空间

线程组地址空间⽤于为并⾏计算着⾊函数分配内存变量. 这些变量被⼀个线程组的所有线程共享. 在线程组地址空间分配的变量不能被⽤于图形绘制着⾊函数顶点着⾊函数, ⽚元着⾊函数

kernel void funcName(threadgroup float *name [[ threadgroup(0) ]])
{
	// A float allocated in threadgroup address space 
	threadgroup float x; 
    // An array of 10 floats allocated in threadgroup address space 
    threadgroup float b[10];
}

3.constant Address Space常量地址空间

  • 常量地址空间指向的缓存对象也是从设备内存池分配存储, 但是它是只读的
  • 在程序域的变量必须定义在常量地址空间并且声明的时候初始化; ⽤来初始化的值必须是编译时的常量.

注意:

声明为常量的变量赋值会产⽣编译错误,声明常量但是没有赋初值也会编译错误.

constant float samples[] = {1,2,3};
samples[] = {3,3,3}; // 编译错误

constant float a; 	 // 编译错误

可以理解Swift中let,OC中const

4.thread Address Space线程地址空间

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

希望变量可以在单独的线程中进行高效的计算,但又不想被共享时,多用于并发运算中,而不是图形处理

kernel void func(...)
{
	float x; // 普通变量
    thread float b; //分配在线程地址空间
}

函数参数与变量的传递修饰

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

被着色器函数的缓存(device 和 constant) 不能重名

对于每个着⾊器函数来说, ⼀个修饰符是必须指定的. 他⽤来设定⼀个缓存,纹理, 采样器的位置(location:如GLSL中的glGetAttribLocation)

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

index是⼀个unsigned integer类型的值，它表示了⼀个缓存、纹理、采样器参数的位置(在函数参数索引表中的位置)。 从语法上讲，属性修饰符的声明位置应该位于参数变量名之后

//代码案例:并⾏计算着⾊函数add_vectors,把两个设备地址空间中的缓存inA和inB相 加，然后把结果写⼊到缓存out
 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];
 }


内建变量属性修饰符

  • [[vertex_id]]顶点id标识符
  • [[position]] 在顶点函数中:表示顶点信息(float4) / 在片元函数中:表示像素点(片元)在屏幕窗⼝的相对位置(x, y, z, 1/w)
  • [[point_size]] 点的⼤⼩(float)
  • [[color(m)]] 颜⾊, m编译前得确定;
当结构体中定义多个颜色,需要做区分时
struct myColors{
	
   	float4 clr_f [[ color(0) ]];	// color attachment 1 附着点1
    int4 clr_i [[ color(1) ]];		// color attachment 2 附着点2
	uint clr_ui [[ color(2) ]];
}
  • [[stage_in]] ⽚元着⾊函数使⽤的单个⽚元输⼊数据,是由顶点着⾊函数输出然后经过光栅化⽣成的.在顶点/片元函数中只能有一个参数被[[stage_in]]修饰,它可以时整型\浮点标量、整型\浮点向量、结构体.
avatar
文章
阅读
粉丝

友情链接：