Metal Shading Language 语法使用

2020-08-22 本文已影响0人泽泽伐木类

Metal Shading Language 是什么

Apple生态下的着色语言，其目的类似于OpenGL 的GL Shading Language, 用来编写3D图形渲染逻辑 和 并行计算核心逻辑 的一门编程语言。在iOS开发中，应用场景涉及到Metal框架时，需要通过Metal Shading Language来编写Metal文件。Metal语言通过Clang和LLVM进行编译处理(这里不同于OpenGL的手动编译和链接)。Metal语法基于C++11.0设计。

Metal 数据类型

Type	Description
`bool`	A conditional data type that has the value of either `true` or `false`. The value true expands to the integer constant `1`, and the value false expands to the integer constant `0`.
`char`	A signed two’s complement `8-bit integer`.
`unsigned char uchar`	An `unsigned 8-bit integer`.
`short`	A signed two’s complement `16-bit integer`.
`unsigned short ushort`	An `unsigned 16-bit integer`.
`int`	A signed two’s complement `32-bit integer`.
`unsigned int uint`	An `unsigned 32-bit integer`.
`half`	A `16-bit floating-point`. The half data type must conform to the IEEE 754 binary16 storage format.
`float`	A `32-bit floating-point`. The float data type must conform to the IEEE 754 single precision storage format.
`size_t`	An unsigned integer type of the result of the sizeof operator. This is a `64-bit unsigned integer`.
`ptrdiff_t`	A signed integer type that is the result of subtracting two pointers. This is a `64-bit signed integer`.
`void`	The `void` type comprises an empty set of values; it is an incomplete type that cannot be completed.

Metal 向量和矩阵数据类型

向量类型: booln charn shortn intn ucharn ushortn uintn halfn floatn(‘n’表示向量维度，eg. int3,float4)
矩阵类型: halfNxM floatNxM ('NxM'表示矩阵的行和列,eg. float4x4 half2x3)

纹理Textures

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

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>

void foo (texture2d<float> imgA[[texture(0)]],
          texture2d<float,access::read> imgB[[texture(1)]],
          texture2d<float,access::write> imgC[[texture(2)]])
{}

采样器Samplers

这个类似于GLSL的uniform sampler2D,只是Metal的Samplers融合了一些纹理的一些属性的配置(比如环绕方式,过滤方式)

枚举名称	有效值	描述
`coord`	`normalized`,`pixel`	从纹理中采样,纹理坐标是否需要归一化
`address`	`clamp_to_edge`,`clamp_to_zero`, `mirrored_repeat`,`repeat`	设置所有纹理的寻址模式
`s_address` `t_address` `r_address`	`clamp_to_edge`,`clamp_to_zero`, `mirrored_repeat`,`repeat`	设置某一个纹理坐标的寻址模式
`filter`	`nearest`,`linear`	设置纹理的采样缩放过滤方式
`mag_filter`	`nearest`,`linear`	设置纹理的采样放大过滤方式
`min_filter`	`nearest`,`linear`	设置纹理的采样缩小过滤方式
`mip_filter`	`none`,`nearest`,`linear`	设置纹理的采样minmap过滤方式
`compare_func`	`none`,`less`,`less_qual`,`greater`,`greate_qual`,`equal`,`not_equal`	设置比较测试逻辑,这个状态值的设置只能在Metal着色语言程序中完成

constexpr sampler s(coord::pixel, filter::linear, mip_filter::nearest)

#warning: 在Metal程序中初始化采样器必须使用constexpr关键字修饰

函数修饰符

kernel,表示该函数是一个数据并行计算着色函数,它将被分配在一个一维/二维/三维的线程组网格中执行.
vertex,表示该函数是一个顶点着色函数,它将为顶点数据流中的每个顶点数据执⾏一次, 然后为每个顶点生成数据输出到绘制管线中去.
fragment,表示该函数是一个片元着色函数,它将片元数据流中的每个片元和其关联的数据执行一次然后为每个⽚元生成数据输出到绘制管线中去.

//并行计算函数
kernel void ZZTestKernelFunctionA(int a,int b) { ... }
//顶点函数
vertex int ZZTestVertexFunctionB(int a,int b){ ... }
//片元函数
fragment int ZZTestVertexFunctionB(int a,int b){ ... }

注意

使用kernel 修饰的函数返回值必须是void类型
一个被函数修饰符修饰过的函数,不允许在调用其他的被函数修饰过的函数
被函数修饰符修饰过的函数,只允许在客户端对齐进行操作. 不允许被普通的函数调用.

地址空间修饰符

Metal 着色器语言使用，地址空间修饰符 来表示一个函数变量或者参数被分配到那块内存区域，所有的着色函数(vertex,fragment,kernel)的参数,如果是指针或者是引用,都必须带有地址空间修饰符号。

device ：设备地址空间
threadgroup ：线程组地址空间
constant ：常量地址空间
thread：thread 地址空间
对于图形着色器函数，其指针或者引用类型的参数必须定义为device 或者 constant地址空间；对于并行计算着色函数,其指针或者引用类型的参数必须为device 或者 threadgroup 或者 constant 地址空间；

device float4  *color;
struct Foo {
  float a[3];
  int b[2];
}
device Foo *my_info;
void ZZTestFouncitionE(device int *g_data,
                       threadgroup int *l_data,
                       constant float *c_data
                       )
{....}
kernel void ZZTestFouncitionF(threadgroup float *a)
{
    //在线程组地址空间分配一个浮点类型变量x
    threadgroup float x;
    //在线程组地址空间分配一个10个浮点类型数的数组y;
    threadgroup float y[10];
}
constant float sampler[] = {1.0f,2.0f,3.0f,4.0f};
kernel void ZZTestFouncitionG(void)
{
    //在线程空间分配空间给x,p
    float x;
    thread float p = &x;   
}

#warning：纹理对象总是在设备地址空间(device)分配内存，地址空间修饰符不必出现在纹理类型定义中，一个纹理对象的内容无法直接访问,Metal 提供了读写纹理的内建函数；

函数参数与变量

图形绘制或是并行计算着色函数的输入/输出都需要通过参数传递(除了常量地址空间变量和程序域中定义的采样器以外).参数可以是如下之一:

device buffer — 设备缓存
constant buffer — 常量量缓存
texture object — 纹理理对象
sample object — 采样器器对象
threadgroup — 线程共享的缓存

用于寻址缓存,纹理,采样器属性修饰符

为每个参数,指定一个属性修饰符. 指定明确的缓冲,纹理位置.

device 和 constant buffers — [[ buffer ( index ) ]]
texture — [[ texture ( index ) ]]
sampler — [[ sampler ( index ) ]]
threadgroup buffer — [[ threadgroup ( 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];
}
//着色函数的多个参数使用不同类型的属性修饰符的情况
kernel void my_kernel(device float4 *p [[buffer(0)]],
                      texture2d<float> img [[texture(0)]],
                      sampler sam [[sampler(0)]])
{
    //.....
}

通过唯一标识的索引值来实现参数传递。这里类似于GLSL中的(
GLuint positionSlot = glGetAttribLocation(program, "Position");
glEnableVertexAttribArray(positionSlot);
glVertexAttribPointer(positionSlot,...);
....
)

参考文献

Metal Shading Language 开发指南
 MetalKit