Metal Shading Language着色语言规范

Potato_土豆
341 阅读12分钟

  • Metal Shading Language概述

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

  • Metal Shading Language语言的限制

    Metal 是基于C++ 11.0 语⾔设计，但是还是有部分不支持的地方和一些使用上的限制具体如下：
    • C++ 11.0 特性在Metal 语⾔中不⽀持之处
      • Lambda 表达式
      • 递归函数调⽤
      • 动态转换操作符
      • 类型识别
      • 对象创建new 和销毁delete 操作符
      • 操作符 noexcept
      • goto 跳转
      • 变量存储修饰符register 和 thread_local
      • 虚函数修饰符
      • 派⽣类
      • 异常处理
      • C++ 标准库在Metal 语⾔中也不可使⽤
    • Metal 语⾔中对于指针使⽤的限制
      • Metal图形和并⾏计算函数⽤到的⼊参数; 如果是指针必须使⽤地址空间修饰(device,threadgroup,constant)
      • 不⽀持函数指针
      • 函数名不能出现main
    • Metal 像素坐标系统: Metal 中纹理/帧缓存区attachment 的像素使⽤的坐标系统的原点是左上⻆（OpenGL中的纹理坐标实在左下角）

  • Metal Shading Language基本数据类型

    • 标量数据类型
    • 向量、矩阵数据类型
      • 向量数据类型
        1. booln
        2. charn
        3. shortn
        4. intn
        5. ucharn
        6. ushortn
        7. uintn
        8. halfn
        9. floatn
          向量中的n代表的是向量的维度最大不能超过4（也就是分别代表一维、二维、三维、四维向量）n前面代表的是向量中数据的数据类型
      • 矩阵数据类型
        1. halfnxm
        2. floatnxm
          nxm分别指的是矩阵的行数和列数，同样的前面值得是矩阵中数据的数据类型
    • 相关用法代码实现： 
      //基本数据类型
bool a = true;
char b = 5;
int  d = 15;
size_t c = 1;
ptrdiff_t f = 2;

//向量（所有的向量赋值方式都有两种一种是下面A的方式另外一种是pos的赋值方式）
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;

//通过向量字母来获取元素(两种方式xyzw、rgba分别代表四个元素)
int4 test = int4(0,1,2,3);
int a = test.x; //a=0
int b = test.y; //b=1
int c = test.z; //c=2
int d = test.w; //d=3

//下面同理
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);// 最后得到c = {3.0f,4.0f,5.0f,4.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; //非法(应为是二维向量也就是说是个申请了两个地址空间的数组，z代表访问第四个地址空间，所以会造成数组越界)

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;
//将第二排的值设置为0
m[1] = float4(2.0f);

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

//设置第三行第四列的元素为3.0f
m[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);

  • Metal Shading Language其他类型

    • 缓存buffer 在Metal 中实现缓存靠的是⼀个指针.它指向⼀个在Device 或者 constant 地址空间中的内建或是开发者⾃定义的数据块.缓存可以被定在 程序域域中,或是当做函数的参数传递.
      相关代码实现：
    //缓存buffer
device float4 *device_buffer;
struct my_user_data{
    float4 a;
    float b;
    int2 c;
};
constant my_user_data *user_data;

    • 纹理Textures 类型 纹理类型是⼀个句柄, 它指向⼀个⼀维/⼆维/三维纹理数据. 在⼀个函数中描述纹理对象的类型; 枚举值: 定义了访问权利;
      enum class access {sample ,read ,write};

      • sample : 纹理对象可以被采样. 采样⼀维这是使⽤或不使⽤采样器从纹理中读取数据;（不写access默认sample）
      • read : 不使⽤采样器,⼀个图形渲染函数或者⼀个并⾏计算函数可以读取纹理对象
      • write: ⼀个图形渲染函数或者⼀个并⾏计算函数可以向纹理对象写⼊数据;
        相关代码实现：
      //纹理texture
//T是指数据类型 设定了从纹理中读取或是向纹理中写⼊时的颜⾊类型. T可以是half, float, short, int 等;
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>

//带有深度格式的纹理必须被声明为下面纹理数据类型中的一个
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)]])
{
    //...
}

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

      1. enum class coord { normalized, pixel };
        从纹理中采样时,纹理坐标是否需要归⼀化;
      2. enum class filter { nearest, linear };
        纹理采样过滤⽅式, 放⼤/缩⼩过滤模式;
      3. enum class min_filter { nearest, linear };
        设置纹理采样的缩⼩过滤模式;
      4. enum class mag_filter { nearest, linear };
        设置纹理采样的放⼤过滤模式;
      5. enum class s_address { clamp_to_zero, clamp_to_edge,repeat, mirrored_repeat };
        设置纹理s坐标的寻址模式;
      6. enum class t_address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat };
        设置纹理t坐标的寻址模式;
      7. enum class r_address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat };
        设置纹理r坐标的寻址模式;
      8. enum class address { clamp_to_zero, clamp_to_edge, repeat, mirrored_repeat };
        设置所有的纹理坐标的寻址模式;
      9. enum class mip_filter { none, nearest, linear };
        设置纹理采样的mipMap过滤模式, 如果是none,那么只有⼀层纹理⽣效;

      注意: 在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种函数修饰符:

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

      注意项：

      1. 使⽤kernel 修饰的函数. 其返回值类型必须是void 类型; 只有图形着⾊函数才可以被 vertex 和 fragment 修饰. 对于着⾊函数, 返回值类型可以辨认出它是为顶点做计算还是为每像素做计算. 图形着⾊函数的返回值可以为 void , 但是这也就意味着该函数不产⽣数据输出到绘制管线; 这是⼀个⽆意义的动作;
      2. ⼀个被函数修饰符修饰的函数不能在调⽤其他也被函数修饰符修饰的函数; 这样会导致编译失败;
        相关代码实现：
      //函数修饰符.
/*
 3个函数修饰符:
 1. kernel : 并行计算函数
 2. vertex : 顶点函数
 3. fragment : 片元函数
 */
//1.并行计算函数(kernel)
kernel void TDTestKernelFunctionA(int a,int b)
{

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

    //不可以的!
    //一个被函数修饰符修饰过的函数,不允许在调用其他的被函数修饰过的函数. 非法
    TDTestKernelFunctionB(1,2);//非法
    TDTestVertexFunctionB(1,2);//非法

    //可以! 你可以调用普通函数.而且在Metal 不仅仅只有这3种被修饰过的函数.普通函数也可以存在
    TDTest();

}
kernel void TDTestKernelFunctionB(int a,int b)
{

}

//顶点函数
vertex int TDTestVertexFunctionB(int a,int b){

}

//片元函数
fragment int TDTestVertexFunctionB(int a,int b){

}

//普通函数
void TDTest()
{

}

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

      1. device — 设备地址空间
      2. threadgrounp — 线程组地址空间
      3. constant（constant=const device） — 常量地址空间
      4. thread — thread地址空间
        **注意：对于图形着⾊器函数, 其指针或是引⽤类型的参数必须定义为 device 或是 constant 地址空间;对于并⾏计算着⾊函数, 其指针或是引⽤类型的参数必须定义为 device 或是 threadgrounp 或是constant 地址空间; **
        详细说明及代码使用示例：
      • Device Address Space(设备地址空间) 在设备地址空间(Device) 指向设备内存池分配出来的缓存对象, 它是可读也是可写的; ⼀个缓存对象可以被声明成⼀个标量,向量或是⽤户⾃定义结构体的指针或是引⽤.
        注意：纹理对象总是在设备地址空间分配内存, device 地址空间修饰符不必出现在纹理类型定义中.⼀个纹理对象的内容⽆法直接访问. Metal 提供读写纹理的内建函数;
        代码示例
      // 设备地址空间: device 用来修饰指针.引用
//1.修饰指针变量
device float4 *color;

struct TDStruct{
    float a[3];
    int b[2];
};
//2.修饰结构体类的指针变量
device TDStruct *my_CS;
      • threadgrounp Address Space 线程组地址空间 线程组地址空间⽤于为并⾏计算着⾊函数分配内存变量.这些变量被⼀个线程组的所有线程共享. 在线程组地址空间分配的变量不能被⽤于图形绘制着⾊函数[顶点着⾊函数, ⽚元着⾊函数]
        在并⾏计算着⾊函数中, 在线程组地址空间分配的变量为⼀个线程组使⽤, 声明周期和线程组相同;
        代码示例 
        /*
 1. threadgroup 被并行计算计算分配内存变量, 这些变量被一个线程组的所有线程共享. 在线程组分配变量不能被用于图像绘制.
 2. thread 指向每个线程准备的地址空间. 在其他线程是不可见切不可用的
 */
kernel void TDTestFouncitionF(threadgroup float *a)
{
    //在线程组地址空间分配一个浮点类型变量x
    threadgroup float x;

    //在线程组地址空间分配一个10个浮点类型数的数组y;
    threadgroup float y[10];

}
      • constant Address Space 常量地址空间
        常量地址空间指向的缓存对象也是从设备内存池分配存储, 但是它是只读的;
        在程序域的变量必须定义在常量地址空间并且声明的时候初始化; ⽤来初始化的值必须是编译时的常量.
        在程序域的变量的⽣命周期和程序⼀样, 在程序中的并⾏计算着⾊函数或者图形绘制着⾊函数调⽤, 但是constant 的值会保持不变;
        注意: 常量地址空间的指针或是引⽤可以作为函数的参数. 向声明为常量的变量赋值会产⽣编译错误.声明常量但是没有赋予初值也会产⽣编译错误;
        代码示例：
      constant float sampler[] = {1.0f,2.0f,3.0f,4.0f};

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

        kernel void TDTestFouncitionG(void)
{
    //在线程空间分配空间给x,p
    float x;
    thread float p = &x;

}

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

      1. device buffer- 设备缓存,⼀个指向设备地址空间的任意数据类型的指针或者引⽤
      2. constant buffer -常量缓存区,⼀个指向常量地址空间的任意数据类型的指针或引用
      3. texture - 纹理对象;
      4. sampler - 采样器对象;
      5. threadGrounp - 在线程组中供各线程共享的缓存.
        注意：被着⾊器函数的缓存(device 和 constant) 不能重名;

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

      • 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))" 为着色函数参数设定了缓存的位置
//thread_position_in_grid⽤于表示当前节点在多线程⽹格中的位置
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)]])
{
    //.....

}

    • 内建变量属性修饰符

      1. [[vertex_id]] 顶点id 标识符

      2. [[position]] 顶点信息(float4) /述了⽚元的窗⼝相对坐标(x, y, z, 1/w)

      3. [[point_size]] 点的⼤⼩(float）

      4. [[color(m)]] 颜⾊, m编译前得确

        //定义了片元输入的结构体，
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; 
}

      5. [[stage_in]] : ⽚元着⾊函数使⽤的单个⽚元输⼊数据是由顶点着⾊函数输出然后经过光栅化⽣成的.顶点和⽚元着⾊函数都是只能有⼀个参数被声明为使⽤“stage_in”修饰符，对于⼀个使⽤了“stage_in”修饰符的⾃ 定义的结构体，其成员可以为⼀个整形或浮点标量，或是整形或浮点向量

avatar
文章
阅读
粉丝