Metal(二) Metal语法规范

1. Metal简介

Metal 着⾊语⾔ 是⽤来编写3D 图形渲染逻辑和并⾏计算核⼼逻辑的⼀⻔编程语⾔. 当你使⽤ Metal 框架来完成APP 的实现时则需要使⽤Metal 编程语⾔;Metal 语⾔使⽤Clang 和 LLVM 进⾏编译处理~Metal 基于C++ 11.0 语⾔设计.我们主要⽤来编写 在 GPU 上执⾏的图像渲染逻辑代码 以及 通⽤ 并⾏计算逻辑代码;

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

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

Metal 像素坐标系统:

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

2. 标量数据类型

示例代码：

bool a = true;char b = 5;int d = 15;size_t c = 1;ptrdiff_t f = 2;

3. 向量数据类型

示例代码：

//向量bool2 A= {1,2};float4 pos = float4(1.0,2.0,3.0,4.0);float x = pos[0];float y = pos[1];float4 VB;for(int i = 0; i < 4 ; i++)VB[i] = pos[i] * 2.0f;//通过向量字母来获取元素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);float2 pos;pos.x = 1.0f; //合法pos.z = 1.0f; //非法float3 pos2;pos2.z = 1.0f; //合法pos2.w = 1.0f; //非法//非法,x出现2次pos.xx = float2(3.0,4.0f);//不合法-使用混合限定符pos.xy = float4(1.0f,2.0,3.0,4.0);float4 pos4 = float4(1.0f,2.0f,3.0f,4.0f);pos4.x = 1.0f;pos4.y = 2.0f;//非法,.rgba与.xyzw 混合使用pos4.xg = float2(2.0f,3.0f);非法,.rgba与.xyzw 混合使用float3 coord = pos4.ryz;float4 pos5 = float4(1.0f,2.0f,3.0f,4.0f);//非法,使用指针来指向向量/分量my_func(&pos5.xy);float4x4 m;//将第二排的值设置为0m[1] = float4(2.0f);//设置第一行/第一列为1.0fm[0][0] = 1.0f;//设置第三行第四列的元素为3.0fm[2][3] = 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);//多个向量构造器的使用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);

3. 纹理 Textures 类型

纹理类型是⼀个句柄，它指向⼀个⼀维/⼆维/三维纹理数据。⽽纹理数据对应这样⼀个纹理的某个level的mipmap的全部或者⼀部分。

枚举值: 定义了访问权利;

enum class access {sample ,read ,write};

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

示例代码：

//纹理textureenum 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>//带有深度格式的纹理必须被声明为下面纹理数据类型中的一个enum class depth_forma {depth_float};depth2d<T,access a = depth_format::depth_float>depth2d_array<T,access a = access::sample,depth_format d = depth_format::depth_float>depthcube<T,access a = access::sample,depth_format d = depth_format::depth_float>depth2d_ms<T,access a = access::read,depth_format d = depth_format::depth_float>void foo (texture2d<float> imgA[[texture(0)]],texture2d<float,access::read> imgB[[texture(1)]],texture2d<float,access::write> imgC[[texture(2)]]){//...}

4. 采样器类型 Samplers

采取器类型决定了如何对⼀个纹理进⾏采样操作. 在Metal 框架中有⼀个对应着⾊器语⾔的采样器的对象 MTLSamplerState 这个对象作为图形渲染着⾊器函数参数或是并⾏计算函数的参数传递;

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

constexpr sampler s(coord::pixel）

示例代码：

enum class coord {normalized, pixel }; 从纹理中采样时,纹理坐标是否需要归⼀化;enum class filter {nearest, linear };纹理采样过滤⽅式, 放⼤/缩⼩过滤模式;enum class min_filter {nearest, linear };设置纹理采样的缩⼩过滤模式; enum class mag_filter {nearest, linear }; 设置纹理采样的放⼤过滤模式; 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 }; 设置纹理s,t,r坐标的寻址模式; enum class address {clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat }; 设置所有的纹理坐标的寻址模式; enum class mip_filter {none, nearest, linear }; 设置纹理采样的mipMap过滤模式, 如果是none,那么只有⼀层纹理⽣效;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);

5. 函数修饰符

Metal 有以下3种函数修饰符:

kernel：表示该函数是⼀个数据并⾏计算着⾊函数，它可以被分配在⼀维/⼆维/三维线程组中去执⾏;vertex：表示该函数是⼀个顶点着⾊函数，它将为顶点数据流中的每个顶点数据执⾏⼀次然后为每个顶 点⽣成数据输出到绘制管线;fragment：表示该函数是⼀个⽚元着⾊函数，它将为⽚元数据流中的每个⽚元 和其关联执⾏⼀次然后 将每个⽚元⽣成的颜⾊数据输出到绘制管线中;

注意：

1. 使用kernel修饰的函数，其返回值类型必须是void类型2.一个被函数修饰符修饰的函数，不能再调用其他也被函数修饰符修饰的函数，这样会导致编译失败。就是不能相互调用，递归调用也不行。

示例代码：

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

6. ⽤于变量或者参数的地址空间修饰符

Metal 着⾊器语⾔使⽤地址空间修饰符来表示⼀个函数变量或者参数变量被分配于那⼀⽚内存区域. 所有的着⾊函数(vertex, fragment, kernel)的参数,如果是指针或是引⽤, 都必须带有地址空间修饰符号;

device : 设备地址空间threadgrounp : 线程组地址空间constant : 常量地址空间thread : thread地址空间

注：对于图形着⾊器函数, 其指针或是引⽤类型的参数必须定义为 device 或是 constant 地址空间;

对于并⾏计算着⾊函数, 其指针或是引⽤类型的参数必须定义为 device 或是 threadgrounp 或是 constant 地址空间;

6.1 Device Address Space(设备地址空间)

在设备地址空间(Device) 指向设备内存池（显存）中分配出来的缓存对象, 它是可读也是可写的; ⼀个缓存对象可 以被声明成⼀个标量,向量或是⽤户⾃定义结构体的指针或是引⽤。

注: 纹理对象总是在设备地址空间分配内存, device 地址空间修饰符不必出现在纹理类型定义中. ⼀个纹 理对象的内容⽆法直接访问. Metal 提供读写纹理的内建函数;

示例代码：

// 设备地址空间: device 用来修饰指针.引用//1.修饰指针变量device float4 *color;struct CCStruct{float a[3];int b[2];};//2.修饰结构体类的指针变量device CCStruct *my_CS;

6.2 constant Address Space 常量地址空间

常量地址空间指向的缓存对象也是从设备内存池分配存储，但是它是只读的；在程序域的变量必须定义在常量地址空间并且声明的时候初始化；⽤来初始化的值必须是编译时的常量；在程序域的变量的⽣命周期和程序⼀样，在程序中的并⾏计算着⾊函数或者图形绘制着⾊函数调⽤，但是constant 的值会保持不变；

注: 常量地址空间的指针或是引⽤可以作为函数的参数. 向声明为常量的变量赋值会产⽣编译错误. 声明常量但是没有赋予初值也会产⽣编译错误;

示例代码：

constant float samples[] = {1.0f, 2.0f, 3.0f, 4.0f }; //对⼀个常量地址空间的变量进⾏修改也会失败,因为它只读的 sampler[4] = {3,3,3,3}; //编译失败; //定义为常量地址空间声明时不赋初值也会编译失败 constant float a;

6.3 threadgrounp Address Space (线程组地址空间)

线程组地址空间⽤于为并⾏计算着⾊函数分配内存变量. 这些变量被⼀个线程组的所有线程共享. 在线 程组地址空间分配的变量不能被⽤于图形绘制着⾊函数[顶点着⾊函数, ⽚元着⾊函数]。在并⾏计算着⾊函数中, 在线程组地址空间分配的变量为⼀个线程组使⽤, 声明周期和线程组相同;

实例代码：

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

6.4 thread Address Space (线程地址空间)

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

示例代码：

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

7. 函数参数与变量

图形绘制或者并⾏计算着⾊器函数的输⼊输出都是通过参数传递， 除了常量地址空间变量和程序域定义的采样器以外。

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

注意: 被着⾊器函数的缓存(device 和 constant) 不能重名;

Attribute Qualifiers to Locate Buffers, Textures, and Samplers ⽤于寻址缓存,纹理, 采样器的属性修饰符;

对于每个着⾊器函数来说, ⼀个修饰符是必须指定的. 他⽤来设定⼀个缓存,纹理, 采样器的位置;

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

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

示例代码：

//属性修饰符/*1. device buffer(设备缓存)2. constant buffer(常量缓存)3. texture Object(纹理对象)4. sampler Object(采样器对象)5. 线程组 threadgroup属性修饰符目的:1. 参数表示资源如何定位? 可以理解为端口2. 在固定管线和可编程管线进行内建变量的传递3. 将数据沿着渲染管线从顶点函数传递片元函数.在代码中如何表现: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 可以由开发者来指定.*///并行计算着色器函数add_vectros ,实现2个设备地址空间中的缓存A与缓存B相加.然后将结果写入到缓存out.//属性修饰符"(buffer(index))" 为着色函数参数设定了缓存的位置//并行计算着色器函数add_vectros ,实现2个设备地址空间中的缓存A与缓存B相加.然后将结果写入到缓存out.//属性修饰符"(buffer(index))" 为着色函数参数设定了缓存的位置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];}// thread_position_in_grid : ⽤于表示当前节点在多线程⽹格中的位置;//着色函数的多个参数使用不同类型的属性修饰符的情况kernel void my_kernel(device float4 *p [[buffer(0)]],texture2d<float> img [[texture(0)]],sampler sam [[sampler(0)]]){//.....}

8. 内建变量属性修饰符

[[vertex_id]] 顶点id 标识符;[[position]] 顶点信息(float4) /描 述了⽚元的窗⼝相对坐标(x, y, z, 1/w)[[point_size]] 点的⼤⼩(float)[[color(m)]] 颜⾊, m编译前得确定;[[stage_in]] ：⽚元着⾊函数使⽤的单个⽚元输⼊数据是由顶点着⾊函数输出然后经过光栅化⽣成的。顶点和⽚元着⾊函数都是只能有⼀个参数被声明为使⽤“stage_in”修饰符，对于⼀个使⽤ 了“stage_in”修饰符的⾃定义的结构体，其成员可以为⼀个整形或浮点标量，或是整形或浮点向量。

示例代码：

struct MyFragmentOutput {float4 clr_f [[color(0)]]; // color attachment 0 int4 clr_i [[color(1)]];// color attachment 1 uint4 clr_ui [[color(2)]]; }; // color attachment 2 fragment MyFragmentOutput my_frag_shader( ... ) {MyFragmentOutput f; ....f.clr_f = ...; ... return f; }}