一、 Metal的语言编程规范
在 WWDC 2014 上，Apple为游戏开发者推出了新的平台技术 Metal，该技术能够为 3D 图像提高 10 倍的渲染性能，并支持大家熟悉的游戏引擎及公司。
二、Metal语言的作用
-
Metal着色语言是用来编写3D图形渲染逻辑和并行计算核心逻辑编程语言。当使用Metal框架来完成APP的实现时则需要使用Metal编程语言。
-
Metal语言使用Clang和LLVM进行编译处理。
-
Metal基于C++ 11.0 语言设计。我们主要用来编写在GPU上执行的图像渲染逻辑代码并进行计算逻辑代码。
那么C++11.0和Metal语言有什么异同之处呢？
- Lambda表达式。
- 递归函数调用
- 动态转换操作符
- 类型识别
- 对象创建new和销毁delete操作符
- 操作符noexcept
- goto跳转
- 变量存储修饰符register和thread_local;
- 虚函数修饰符
- 派生类
- 异常处理
- C++标准库
以上在Metal语言中也不可使用。
Metal语言对于指针使用的限制 - Metal图形和并行计算函数用到的入参数，如果是指针必须使用地址空间修饰符(device,threadgroup,constant)
- 不支持函数指针
- 函数名不能出现main
Metal像素坐标系统中：Metal中纹理/桢缓存区attachment的像素使用的坐标系统的原点是左上角
三、Metal的数据类型
3.1 标量数据类型
| 类型 | 描述 |
| bool | 布尔类型true/false|
| char | 有符号的8-bit整数 |
| unsign char/uchar | 无符号的8-bit整数 |
| short | 有符号的16-bit整数 |
| unsign short/ushort | 无符号的16-bit整数 |
| int | 有符号的32-bit整数 |
| unsign int/uint | 无符号的32-bit整数 |
| half | 一个16-bit浮点数 |
| float | 一个32-bit浮点数 |
| size-t | 64-bit无符号整数，表示sizeof操作符的结果 ｜
| ptrdiff_t | 64-bit有符号整数，表示两个指针的差|
| void | 表示一个空的值的集合 |
3.2 向量和矩阵数据类型
3.2.1向量支持如下类型：
- booln
- charn
- shortn
- intn
- ucharn
- ushortn
- uintn
- halfn
- floatn
向量中的n，指的是维度，最大是4
3.2.2矩阵支持的类型：
- halfnxm
- floatnxm
nxm分别指的是矩阵的行数和列数
3.3纹理Textures类型
纹理类型是一个句柄，它指向一个一维/二维/三维的纹理数据，在一个函数中描述纹理对象的类型。
枚举值：定义了访问权利
enum class access {sample ,read ,write};
texture1d<T, access a = access::sample> texture1d_array<T, access a = access::sample>
texture2d<T, access a = access::sample> texture2d_array<T, access a = access::sample>
texture3d<T, access a = access::sample>
texturecube<T, access a = access::sample>
texture2d_ms<T, access a = access::read>
sample：纹理对象可以被采样，采样一维这是使用或不使用采样器从纹理中读取数据
read：不使用采样器，一个图形渲染函数或者一个并行计算函数可以读取纹理对象；
write：个图形渲染函数或者一个并行计算函数可以读取纹理对象写入数据
T：数据类型 设定了从纹理中读取或者是向纹理中写入时的颜色类型。T可以是half，float，short，int等
代码示例：
void foo (texture2d<float> imgA [[ texture(0) ]] ,
texture2d<float, access::read> imgB [[ texture(1) ]],
texture2d<float, access::write> imgC [[ texture(2) ]])
{
...
}
3.4.采样器类型 Samplers
采样器类型决定了如何对一个纹理进行采样操作，在Metal框架中有一个对应着色器语言的采样器等对象MTSamplerState，这个对象作为图形渲染着色器函数参数或者是并行计算函数的参数传递。
image.png
注意：在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);
四、修饰符
4.1函数修饰符
- kernel,表示该函数是⼀个数据并行计算着⾊函数,它将被分配在一个⼀一维/二维/三维的线 程组⽹网格中执⾏行行.
- vertex,表示该函数是一个顶点着⾊函数,它将为顶点数据流中的每个顶点数据执行一次 然后为每个顶点生成数据输出到绘制管线中去.
- fragment,表示该函数是一个片元着⾊色函数,它将片元数据流中的每个片元和其关联的数 据执⾏一次然后为每个片元⽣成数据输出到绘制管线中去.
注意：使用kernel修饰的函数，其返回值类型必须是void类型。
kernel void foo(...)
{
...
}
能被vertex和fragment修饰的只有图形着色函数，对于图形着色函数，返回值可以辨认为它是为定点做计算还是为像素做计算，图形着色函数的返回值可以为void，但是这意味着该函数不产生数据输出到绘制管线，没有任何意义。
一个被函数修饰符修饰的函数不能在调用其他也被函数修饰符修饰的函数。
kernel void hello2(...)
{
}
vertex float4 hello1(...)
{
hello2(...); //这样是错误的
}
4.2用于变量/参数的地址空间修饰符
Metal着色器语言使用地址空间修饰符来表示一个函数变量或者参数变量被分配于哪一片内存区域。如果是指针或者引用，都必须带有地址空间修饰符
- device
- threadgrounp
- constant
- thread
图形着色器函数，其指针/引用类型的参数必须定义为device/constant地址空间。
并行着色器函数，其指针/引用类型的参数必须定义为device/constant/threadgrounp地址空间
4.2.1Device Address Space(设备地址空间)
device指向内存池分配出来的缓存对象，它是可读可写的；一个缓存对象可以被声明为一个标量，向量或者用户自定义结构体的指针或者引用
// an array of a float vector with 4 components device float4 *color;
struct Foo {
float a[3];
int b[2];
};
// an array of Foo elements
device Foo *my_info;
// an array of a float vector with 4 components device float4 *color;
struct Foo *my_info;
注意：纹理对象总是在设备地址空间分配内存，device地址空间修饰符不必出现在纹理类型定义中，纹理对象的内容无法直接访问，Metal提供读写纹理的内建函数。
4.2.2threadgrounp Address Space(线程组地址空间)
线程组地址空间用于并行计算着色函数分配内存变量，这些变量被一个线程组共享，在线程组地址空间分配的变量不能用于图形绘制着色函数(顶点着色函数，片元着色函数)
在并行着色函数中，在一个线程地址空间分配的变量为一个线程组使用，声明周期和线程组相同。
kernel void my_func(threadgroup float *a [[ 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];
}
4.2.3constant Address Space(常量地址空间)
常量地址空间指向的缓存对象也是从设备内存池分配存储的，只能用来读。
在声明的时候都初始化，其值不可变。向声明为常量的变量赋值或者声明常量没有赋予初始值都会产生编译错误。
constant float samples[] = { 1.0f, 2.0f, 3.0f, 4.0f };
//不可对常量地址空间的变量进行修改，只读，
sampler[4] = {3,3,3,3}; //编译失败
//声明常量没有赋予初始值都会产生编译错误
constant float a;
4.2.4thread Address Space(线程地址空间)
thread地址指向每个线程准备的地址空间，其变量在其他线程不可见，在图形绘制着色函数或者并行计算着色函数中声明的变量thread地址空间分配；
kernel void my_func(...)
{
float x;
thread float p = &x;
...
}
4.3函数参数与变量的修饰符
- device buffer — 设备缓存
- constant buffer— 常量量缓存
- texture object — 纹理理对象
- sample object — 采样器器对象
- threadgroup — 线程共享的缓存
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];
}
该例子为一个简单的并行计算着色函数add_vectors，它把两个设备地址空间中缓存的inA和inB相加，然后把结果写入到缓存out中，属性修饰符[ buffer(x) ]为着色函数参数设定了缓存的位置
thread_position_in_grid 表示当前节点在多线程网格中的位置。
对于每个着色器函数，一个修饰符是必须指定的。
- [[buffer(x)]]
- [[texture(x)]]
- [[sampler(x)]]
- [[threadgroup(x)]]
x是一个unit的类型。表示缓存、纹理、采样器和线程组参数的位置，
4.4内建变量属性修饰符
- [[vertex_id]] 顶点id标志符
- [[position]] 顶点信息(float4)描述了片元窗口的相对坐标(x,y,z,1/w)
- [[point-size]] 点的大小
- [[color(m)]]颜色，m在编译前得确定。
struct MyFragmentOutput {
// color attachment 0
float4 clr_f [[color(0)]]; // color attachment 1
int4 clr_i [[color(1)]]; // color attachment 2
uint4 clr_ui [[color(2)]]; };
fragment MyFragmentOutput my_frag_shader( ... )
{
MyFragmentOutput f;
....
f.clr_f = ...;
...
return f;
}
- [[stage_in]] 片元着色函数使用的单个片元输入数据是有顶点着色函数输出，然后经过光栅化生成的，顶点和片元着色函数都是只能有一个参数被声明为使用stage_in修饰符，对于一个使用了stage_in修饰符的自定义结构体，其成员可以为一个整型或者符点标量/向量。
网友评论