• Metal 一、初识 Metal 及其语言规范


    Metal 简介 与 Metal 编程语言语法。

    一、Metal 简介

    1、Metal 是 Apple 为了解决 3D 渲染而推出的框架。游戏开发者的新的技术平台,该技术能够为3D图像提高 10 倍的渲染性能。苹果2018年推出 Metal,在此之前一直使用OpenGL ES。

    Metal 的2个目的:1图形渲染; 2高并发计算

    苹果文档给出的 Metal 优化:CPU 低消耗、GPU 高利用、连续处理器并行、有效的资源利用

    Metal 官方文档

    2、图形管道 graphic pipeline

    二、Metal Shading Language - Metal 编程语言与规范

    1、Metal 语言特点

    1.1)Metal 着色器语言,用来编写3D 图形渲染逻辑和并行计算核心逻辑(高并发) 的一门编程语言。当我们需要使用 Metal 框架来完成 APP 的开发时(或 需要使用 Metal 的高并发计算能力时),就要使用 Metal 编程语言进行编程。

    1.2)Metal 编程语言 使用 clang 和 LLVM 进行编译处理

    Metal的语言规则,是基于 C++ 11.0 标椎设计的,并在此基础上进行了一定的扩展和限制。

    而我们的实际业务开发场景中所需要做的工作就是: 编写 在GPU 上执行的图像渲染逻辑代码 以及 通用的并行计算逻辑 代码。

    1.3)Metal 与 C++ 11.0 

    1.3.1)Metal 限制(不支持):

      Lambda 表达式

      递归函数调用

      动态转换操作符

      类型识别

      对象创建 new 和 销毁 delete 操作符

      操作符 noexcept

      goto 跳转

      变量存储修饰符 register 和 thread_local

      虚函数修饰符

      派生类

      异常处理

      C++ 标准函数库 在 Metal 中不支持,不能使用

    1.3.2)Metal 中指针的使用限制

      Metal 图形和并行计算函数用到的参数,如果是指针 必须使用地址空间修饰符(devide / threadgroup / constant) 

      不支持函数指针

      函数名不能出现 main

    1.3.3)Metal 像素坐标系

      Metal 中纹理/帧缓冲区 attachment 的像素使用的坐标系的原点是在 左上角。--> 苹果的坐标原点

    2、Metal 的数据类型

    2.1)标量类型

     

    2.2)向量与矩阵 数据类型

    2.2.1)向量

    booln

    charn

    shortn

    intn

    ucharn / ushortn / uintn

    halfn

    floatn

    n --> 向量中的 n 表示维度,最大为4

    代码示例:

    bool2 b = [1, 2]; 
    
    float4 f4 = float4(1.0, 2.0, 3.0, 4.0);
    float f = f4[0];// x = 1.0 --> 类似数组
    
    // int4 --> 4个变量组成的4维向量
    // xyzw  rgba
    int4 test = int4(0, 1, 2, 3);
    int x = test.x;// x = 0
    int y = text.y;// y = 1
    int x1 = test.r;// x1 = 0
    
    // 多个分量的访问
    float4 c = float4(0,0,0,0);
    c.xyzw = float4(1,2,3,4);// 重赋值 c = [1,2,3,4]
    c.xy = float2(6,0);// c = [6,0,3,4]
    c.yzw = float3(7,8,9);// c = [6,7,8,9]
    
    // 可乱序 --> 
    // 注意!!!这里 Metal 不同于 GLSL --> GLSL 可以多分量,但是不可乱序 xyzw/rgba 顺序是不可变的
    float4 pos = c.wxyz;// pos = [9,6,7,8]
    
    float4 rep = c.xxwz;// rep = [6,6,9,8]
    rep.xw = float2(5,7);// rep = [5,6,9,7]
    
    // 不可混用
    float4 m = float4(4,3,2,1);
    m.xg = float2(0,9);!error 非法, xyzw 和 rgba 不能混了,他俩只可选其一
    m.rg = float(0,9);// m = [0,9,2,1]
    
    
    /// 构造方式
    // 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);

    2.2.2)矩阵

    halfxnm / floatnxm

    nxm 中 n m 分别指 矩阵的行数 和 列数 --> 最大 4 x 4 : 4行4列

    float4x4 mix;
    // mix[1] = float4(2.0f);// 矩阵 第一行的值都是 2
    mix[1] = float4(1,2,3,4);// 矩阵的第一行的值
    mix[0][0] = 3;// 矩阵的第0行0列的值为 3
    mix[3][2] = 7;// 矩阵 3行2列 位置的值

    2.3)纹理 Texture 类型

    纹理类型是一个句柄,指向一个 一维/二维/三维纹理数据。在函数中描述纹理对象的类型。 

    枚举: 

    enum class access { sample, read, write } // 定义访问权限

    sample:纹理对象可以被采样,采样器可将纹理读取出来,可读可写可采样 --> 使用最多

    read:不使用采样器,一个图形渲染函数或并行计算函数 可以读取纹理对象

    write:一个图形渲染函数或并行计算函数 可以向纹理对象写入数据。

    texture1d<T, access a = access::sample> // 一维纹理
    texture2d<T, access a = access::sample>// 二维纹理
    texture3d<T, access a = access::sample>// 三维纹理

    T:数据类型 ,指定从纹理中 读取/写入 时的颜色类型。T可以是 half、float、int 等;

    access:读写方式(权限)

    代码示例:

    void foo (texture2d<float> imgA [[ texture(0) ]] ,// texture2d<float>: 2 维纹理,类型 float,访问权限 sample --> 默认权限就是 sample 可不写
    texture2d<float, access::read> imgB [[ texture(1) ]],// texture2d<float, access::read>: 类型 float,权限 read
    texture2d<float, access::write> imgC [[ texture(2) ]]) // 权限 write
    { 
     ... 
    }

    2.4)采样器类型 Samplers

    采样器类型 决定了如何对一个纹理进行采样操作。

    metal 框架中有一个对应 着色器语言的采样器对象:MTLSamplerState, 此对象做为 图像渲染着色器函数 or 并行计算函数 的参数进行传递。

    枚举 们:

    // 从纹理中采样时,纹理坐标是否归一化
    enum class coord { normalized, peixel }
    
    // 纹理采样过滤方式 - 放大/缩小
    enum class filter { nearest, linear }
    
    // 缩小过滤方式
    enum class min_filter { nearest, linear }
    
    // 放大过滤方式
    enum class mag_filter { nearest, linear }
    
    // 设置纹理 s t r 坐标的寻址模式 (str 即 xyz 环绕方式)
    enum class s_address { clapm_to_zero, clapm_to_edge, repeat, mirrored_repeat }
    enum class t_address { clapm_to_zero, clapm_to_edge, repeat, mirrored_repeat }
    enum class r_address { clapm_to_zero, clapm_to_edge, repeat, mirrored_repeat }
    
    // 设置所有纹理坐标的寻址模式
    enum class address { clapm_to_zero, clapm_to_edge, repeat, mirrored_repeat }
    
    // 设置纹理采样的 mipMap 过滤模式,如果是 none ,则只有一层纹理生效
    enum class mag_filter { none, nearest, linear }

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

    3、修饰符

    3.1)函数修饰符

    kernel:表示该函数是一个 数据并行计算着色函数。我们要高效并发运算就用它。它可以被分配在 一维/二维/三维 线程组中去执行;--> 使用他修饰的函数 返回类型必须是 void

    vertex:顶点着色函数。为顶点数据流中的 每个顶点数据 执行一次,然后为每个顶点生成数据输出到绘制管线;

    fragment:片元着色函数,为 片元数据流中的 每个片元 与其关联 执行一次,然后将 每个片元生成的颜色数据 输出到绘制管线中;

    注意1:只有 图形着色函数才能用 vertex/fragment 修饰。函数返回类型 可以用来辨认出它是为顶点 or 为每个像素 做计算的。返回 void 也可以但是无意义,因为顶点/片元函数本就是为了计算出相应数据 将数据传到绘制管线的。

    注意2:被函数修饰符修饰的函数不能再调用 '被修饰符修饰的函数',否则编译失败。即:被函数修饰符修饰的函数们不能相互调用

    例:kernel void func1 (...) {}; vertex float4 funcV1 (...) { func1(...) } --> 错误调用,无法编译

    注意2:特定函数修饰,普通函数随意。

    代码示例:

    kernel void foo(...) {
            ...
    }

    3.2)变量 或 参数 的地址空间修饰符 Address Space

    地址空间修饰符:用来指定 一个函数 参数/变量 被分配在内存中的哪块区域。

    device:设备地址空间

    threadgroup:线程组地址空间

    constant:常量地址空间

    thread: thread 地址空间

    a、对于 图形着色器函数,是 指针 或 引用 类型的参数 必须定义为 device 或 const 地址空间

    b、对于并行计算着色函数,对于是 指针 或 引用 的参数,必须使用 device 或 threadgroup 或 constant 修饰。

    3.2.1)Device 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;

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

    3.2.2)线程组地址空间 threadgroup

    threadgroup:用于 为 并行计算着色函数 分配内存变量(在GPU里),这些变量被一个线程组的所有线程共享。在线程组地址空间 分配的变量 不能被用于图形绘制着色函数。

      在并行计算着色函数中,在线程组地址空间分配的变量 为一个线程组使用,生命周期和线程相同。

    代码示例:

    // kernel 高速并行
    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];
    }

    3.2.3)constant 常量地址空间

    constant:指向的缓存对象也是从设备内存池 分配存储,但是是只读的。

    在程序域的变量 必须定义在常量地址空间 并且在声明的时候初始化;用来初始化的值 必须是编译时的常量。此变量的生命周期和程序一样,在程序中的 并行计算着色函数or图形绘制着色函数 调用,但 constant 的值会保持不变。

    注意:常量地址空间 的指针或引用 可以作为函数的参数(constant修饰的常量可作为函数的参数)。向声明为常量的变量赋值会产生变异错误(代码示例中sampler),声明为常量但没有赋予初始值也会产生变异错误(代码示例中a)。

    错误代码示例:

    constant float sampler[] = {1.0f, 2.0f, 3.0f,4.0f};
    
    // 对一个常量地址空间的变量进行修改会失败,因为它是只读的
    sampler[4] = {3,3,3,3};// 编译失败
    
    // 定义常量地址空间但不初始化赋值 --> 也编译失败
    const float a;// 编译失败

    3.2.4)线程地址空间 thread

    thread:指向每个线程准备的地址空间,这个线程的地址空间 定义的变量 在其他线程是不可见的,在图形绘制着色函数or并行计算着色函数 中声明的变量可以使用 thread 地址空间分配。

    代码示例:

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

    3.3)函数参数与变量

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

    device buffer:设备缓存 - 指向设备地址空间的任意数据类型的指针 or 引用

    constant buffer:常量缓存 指向常量地址空间的任意数据类型的指针 or 引用

    texture object:纹理对象

    sample object:采样器对象

    threadgroup:线程共享的缓存

    对于每个着色器函数来说,一个修饰符是必须指定的,它用来设定一个 缓存、纹理、采样器的位置:

    device buffer / constant buffer --> [[buffer(index)]]

    texture --> [[texture(index)]]

    sample --> [[sampler(index)]]

    threadgroup buffer --> [[threadgroup(index)]]

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

    通过示例理解:

    // 一个简单的并行计算着色函数 my_add ,它把两个设备地址空间的魂村 inA、inB 相加,把结果写入缓存 out。
    // 属性修饰符 “buffer(index)” 为着色函数参数 设定了缓存的位置
    kernel void my_add (constant device float4 *inA [[ buffer(0) ]],// inA: 放在设备地址空间,缓存位置对应的是 buffer(0)这个ID , constant 修饰的不可变
                                    constant device float4 *inB [[ buffer(1) ]],
                                    device float4 *out [[ buffer(2) ]],
                                    uit id [[ thread_position_in_grid ]],) {
    
        out[id] = inA[id] + inB[id];
    }            

    thread_position_in_grid:用于表示当前节点,在多线程网格中的位置 --> 我们是无法知道当前在GPU的哪个运算单元里,thread_position_in_grid 知道,我们通过它获取即可。

    3.4)内建变量属性修饰符

    [[vertex_id]] -- 顶点ID标识符

    [[position]] -- 1、当前顶点信息(float4 - xyzw)  2、也可描述 片元在窗口的相对坐标:当前这个像素点在屏幕上的哪个位置

    [[point_size]] -- 点的大小

    [[color(m)]] -- 颜色,m 编译前要确定

    如下代码:

    // 定义颜色结构体
    struct myFragmentOutput {
    
        // 三组颜色,要知道使用时取哪一个
        float4 clr_f [[color(0)]];// 
        int4 clr_i [[color(1)]];//
        uint4 clr_ui [[color(2)]];//
    }
    
    fragment myFragmentOutput my_grag_shader (...) {
    
        myFragmentOutput f;
        ...
        f.clr_f = ...;
        ...
        return f;
    }

    另补充一个:[[stage_in]] -- 其实就是:顶点着色器输出 经过光栅化 生成的 传给片元着色器的 每个片元数据。

    顶点和片元着色函数都是 有且仅有 一个参数可以被声明为 使用"stage_in"修饰符的。 stage_in 可以修饰结构体,其结构体成员可以有多个,类型可以为一个 整型/浮点型 的 标量/向量。

  • 相关阅读:
    Pytorch 随机数种子设置
    python 利用 dictionary 的 .get() 操作,避免写 if-else
    PEP-8 or google 风格 python 代码风格和注释规范
    Vim 多文件切换使用
    Shell 变量及脚本使用
    python numpy 大矩阵运算容易内存爆炸
    Ubuntu 配置 Pytorch on Graph (PoG) 环境
    Markdown 学习笔记
    Linux-saltstack
    Python字符串详解
  • 原文地址:https://www.cnblogs.com/zhangzhang-y/p/13549930.html
Copyright © 2020-2023  润新知