利用 Metal 网格着色器改变几何形状

大家好 我是 Andrei 是 Metal Frameworks 团队的 GPU 软件工程师 今天 我很高兴能为大家介绍 我们新推出的 Metal 网格着色器 网格着色器是 Metal 中 全新灵活的管线 可用于 GPU 驱动的 几何元素生成和处理 它可以提高顶点/片元管线 增加灵活性 消除顶点处理限制 它有多重应用 包括但不限于 细粒度几何裁剪 GPU 上可扩展程序化几何体创建 允许自定义几何输入 如压缩顶点流 meshlet 和复杂程序化算法 这三点也是我今天想要分享的内容 首先 我会介绍 Metal 网格着色器是什么 然后 我会为大家展示两个 网格着色器的使用案例 网格着色器对生成 程序化几何体极为有效 如渲染程式化头发 网格着色器也可以帮助 改善场景处理和渲染 最常见的一个例子就是 使用网格着色器 来进行 GPU 驱动的 meshlet 裁剪 我们先介绍下网格着色器 这个是 Stanford Bunny 代表了您在 GPU 上 可以渲染的典型网格 为了渲染该网格 顶点和索引数据 首先要放入设备内存中 然后 您要使用渲染指令编码器 来执行绘制调用 传统渲染管线包含三个基础阶段 一个可编程的顶点着色阶段 一个固定函数光栅化阶段 以及一个可编程的片元着色阶段 顶点着色阶段从设备内存中 提取几何体 作为输入并进一步处理 光栅化程序生成屏幕空间片元 片元着色器将其着色 并生成最终图像 这种管线自始至终都能 很好地服务于这个目标 然而 它缺乏灵活性且有一定的限制 我们来看一个例子 假设您想要在 GPU 上 生成程序化几何体 比如您决定在这个兔子上 添加程序化毛发 我给大家演示下传统几何管线中 这个任务是如何完成的 传统方法中 要生成程序化几何体 您还要有计算指令编码器 来调度计算内核 计算内核将原始网格作为输入 生成程序化几何体 并将其输出存回设备内存 然后 您可以使用渲染指令编码器 来执行绘制调用 将程序化几何体作为输入 生成最终图像 这一方法需要两个指令编码器 且需要您分配额外内存 来存储程序化几何体 如果是间接绘制调用 或扩散系数高的情况下 该内存会非常高 且难以预测 两个编码器之间还有一个屏障 使 GPU 间的工作序列化 Metal 网格着色器 则可以解决以上所有问题 网格着色器是一个 全新的几何管线 用两个可编程的新阶段 替换了顶点着色阶段 对象着色阶段和网格着色阶段 在这个例子中 对象着色器 将几何体作为输入 对其进行处理 输出我们称为 “payload” 的数据 到网格着色器中 该数据可由您自行决定 网格着色器则会使用该数据 来生成程序化几何体 该程序化几何体仅在 绘制调用中存在 所以并不会要求您 分配任何设备内存 它是管线式的 直接到光栅化程序 然后到片元着色器 生成最终图像 网格绘制调用使用 与传统绘制调用一样类型的 渲染指令编码器来运行 网格绘制调用和传统绘制调用 可以混合和匹配使用 现在 我们来看下两个 新的可编程阶段

与顶点着色器相比 对象和网格着色器 都与计算内核相近 它们是在线程组栅格中启动的 每个线程组都是独立的线程栅格 如同计算线程 它们可以相互沟通 另外 每个对象线程组 可以生成一个网格栅格 以编程方式确定 它启动网格栅格的大小 提供足够的灵活性 每个对象线程组将 payload 数据 传送给它生成的网格栅格 如其名称所示 对象阶段处理的是对象 对象是一个抽象的概念 您可以根据需求来定义 可以是场景模型 场景模型的一部分 比如 您想要生成程序化几何体的 空间区域 网格阶段的设计是为了搭建网格 直接向光栅化程序 发送几何体数据 接下来两个例子能展示 对象和网格之间的关系 第一个使用了网格着色器 来进行毛发渲染 为了将任务简化一些 我用了简易的平面 而不是兔子模型 要生成一撮毛发 我将输入几何体划分为分块 每个分块会计算一层细节 以及需要生成的发丝数量 然后生成独立的发丝 我为大家演示下如何在这个平面上 用网格着色器 程序化地生成毛发 平面可以划分为分块 每个分块对应一个对象线程组 每个对象线程组可计算 发丝数量 为每根发丝生成曲线控制点 这个会成为 payload 然后 我们的对象线程组 启动网格栅格 每个网格线程组代表一根发丝 每个网格线程组输出网格 到光栅化程序 新的几何体管线让您可以 将几何处理更紧密地映射到硬件 并让您可以充分利用 GPU 提供的所有线程 在网格渲染管线中 输入几何体按对象着色栅格 划分分块 每个对象着色线程组 可以独立生成 payload 并启动网格栅格 每个来自栅格的 网格着色线程组生成 metal::mesh 它在剩余的渲染管线中 被进一步处理 我们仔细看看每个阶段 生成的数据 payload 在对象着色器中定义 每个对象线程组 向其生成的网格栅格 传送到自定义 payload 如果渲染头发 payload 包含 曲线控制点 同时 网格着色器 通过新的 metal::mesh 类型 输出顶点和图元数据 稍后我会详细解释

对象和网格阶段输出接下来管线 所需的网格数据 与传统管线中的顶点输出相似 光栅化程序首先处理网格数据 然后执行片元着色器 我们详细看下如何配置 毛发渲染网格管线 首先 将要用毛发覆盖的平面 划分分块 每个分块对应一个对象线程组 对象线程组决定网格栅格大小 将向网格传送的 payload 数据初始化 在这个情况下 该分块中有 6 根发丝 与每根发丝的 曲线 payload 数据一起 生成 3x2 网格栅格 每个线程组生成独特的 网格栅格大小 下一个线程组只需要生成 4 根发丝 所以与那 4 根发丝初始化的 曲线 payload 数据一起 所以设置 2x2 网格 这就是这一实现方法的 对象着色器 object 属性已添加到 Metal 中 以指定对象着色器的代码 除了 payload 属性 和 object_data 地址空间 允许 payload 参数 使用于着色器之外

mesh_grid_properties 参数 用于编码网格栅格大小 下一步是管线初始化 首先 分配网格渲染管线描述符 然后初始化对象函数 并制定所需的 payload 长度 与每线程组的线程最大数量 对象着色器上有一些限制 Payload 格式和内容 全部可以自定义 然而 payload 大小不能超过 16KB 的限值 同时 每个对象线程组生成的 网格线程组 最大数量不能超过 1024 对象着色阶段准备好后 下一步就是 初始化网格着色阶段 网格着色器让用户定义 payload 以作为输入 在这个例子中 payload 是一套曲线控制点 每个网格线程组 生成一个 metal::mesh 也就是一根发丝 网格着色器的输出网格 必须有 metal::mesh 类型 metal::mesh 是 Metal 中的内置结构 可为您提供输出顶点和图元数据 到光栅化程序 和片元着色器的界面 每个 metal::mesh 定义一个 顶点数据类型 就像 顶点着色器的输出类型 一个图元数据类型 顶点的最大数值 图元的最大数值 最后 网格拓扑 不管是点 线还是三角 Metal 着色语言通过增加的 mesh 属性 指定什么代码是网格着色器 metal::mesh 用于 网格着色器中的输出结构

网格着色器对 GPU 驱动的 几何体处理很有效 因为它们能让您动态 生成这些 metal::mesh 用于光栅化程序 网格着色器用 metal::mesh 作为其优势 这样您可以在处理过程中 加入更多渲染指令 而无需额外计算通道 编码网格在同一个线程组的 线程之间完成 在这个例子中 线程组的 前 9 条线程可以编码发丝的顶点 索引和图元数据 线程 0 到 4 每个都在网格中 编码一个顶点 线程组中的剩余线程 不编码网格中的顶点 接下来 所有 9 条线程将 1 条索引 编码到网格目录中

接下来 前 3 条线程 为 3 个三角形编码图元数据 剩余的线程不编码任何 图元数据 最后 1 条线程会为网格 编码图元数量 我为大家演示下 这个网格着色器的源代码 网格着色器会组织为 尽可能避免 线程间的分化 遵循顶点 索引和图元数据的编码步骤 最后是图元数量的步骤

我们回到初始化网格管线 描述符 网格管线描述符设置了 网格函数 和每网格线程组的最大线程 metal::mesh 结构有一些 限制是需要严守的 Metal 的网格着色器 有以下限制 Metal::mesh 支持高达 256 个顶点 以及高达 512 个图元 metal::mesh 总大小 不可超过 16KB 现在网格栅格 已经生成 metal::mesh 它们随后会被带到光栅化程序中 最后运行片元着色器 所以 与传统渲染管线类似 片元函数在网格管线 描述符上设置 初始化描述符后 管线状态已通过 makeRenderPipelineState(descriptor:) 方法 在 Metal 设备创建 编码网格管线和编码 传统绘制调用很接近 管线在编码器上设置 管线的每个阶段都有资源绑定 在这个例子中 绑定资源有 对象阶段的对象缓冲区 网格阶段的纹理 片元阶段的片元缓冲区 接下来 我定义几个启动 网格管线所需的常量 对象栅格尺寸 每对象线程组的线程数量 每网格线程组的线程数量 以及通过新的 drawMeshThreadgroups 方法 使用这些常量来编码绘制 用于渲染平面毛发的相同方法 可以应用到整个兔子上 通过网格管线 程序化地生成毛发 接下来 我们看下使用 网格着色器的另一种方法 网格着色器可以用 Meshlet 裁剪高效处理 以及渲染大量几何体 这一技术的基础就在于 将场景网格划分为 称为 meshlet 的小块

将场景几何体划分为 meshlet 增加了场景的粒度 可以执行更高效 更深入的裁剪 这也可以极大地减少 几何体开销 利用 meshlet 粒度处理 可允许高效的遮挡和裁剪算法 如屏幕空间遮挡裁剪 和法线过滤程序 您可以使用网格着色器 来执行全 GPU 驱动的裁剪 和渲染管线 这是传统 GPU 驱动的管线 可用一种计算和一条渲染通道 运行场景处理和渲染 该场景数据可划分为 meshlet 并用于计算通道 从而负责视锥体剔除 LOD 选择 以及将绘制编码到设备内存中 然后渲染通道为场景 执行绘制指令 然后生成最终图像 通过使用网格着色器 可以移除同步点 通过将两条通道 合并到一个网格着色分发 避免中间绘制指令 我为大家演示下操作过程 这是单一渲染通道 可执行网格着色分发 对象着色器运行视锥体剔除 为每个可见的 meshlet 计算 LOD 网格着色器的 payload 是 应编码的 meshlet IDs 列表 然后 网格着色器编码用于光栅化 和着色的 metal::mesh 对象 然后最终图像在 片元着色器中着色 与传统管线一样 几何体处理完全在 网格线程组指令内 和单一编码器内完成 不再需要通过中间缓冲区 存储这些绘制指令 因为三角形数据 会通过网格着色器编码

我们现在来看下裁剪 尤其是 meshlet 裁剪的执行 该场景包含代表模型的形状 在这个操作中 场景的每个模型 都会变成对象栅格的一部分 网格栅格由对象着色线程组生成 包含 meshlet 模型表面的三角小块 新的几何体管线灵活性很高 您可以选择如何将场景 映射到对象栅格上 在这个例子中 我将每个模型 都映射到对象线程组上 但您可以使用更符合 您任务所需的映射 现在 对象着色器可以用视锥 来确定 meshlet 的可见性 只分派在最终图像上可显示的内容 我们来看场景中的两个模型 对象着色器基于确定的 可见度来启动网格栅格 然后 网格着色器处理 meshlet 并构建 metal::mesh 可编程的网格栅格大小 可以灵活调度 所以只有可见的 meshlet 能经由网格着色器处理 这可以减少在后续管线中 处理不可见几何体的时间 固定函数光栅化程序只接收 可见的表面 而且可以减少处理和裁剪 不可见几何体的时间 最后 可编程片元着色器被调用 生成最终图像 如您所见 新几何体管线 可供您处理 各种不同的问题 如创建程序化网格 或让您的绘制调用更高效 正如这个 meshlet 裁剪样本 展示的一样 Metal 现在包含了创新灵活的 全新几何体管线 创建程序化几何更为简单 正如头发渲染例子演示的一样 另外 单一渲染通道中 GPU 驱动的 工作可能性得以扩展 无需额外计算通道或中间缓冲区 正如在 meshlet 裁剪 demo 中 看到的一样 这个全新的几何体管线 在 Family7 和 Mac2 设备中均可用

为了帮助您开始 网格着色器的学习和操作 Apple 开发者网站 有示例代码 展示了如何使用新的 API 我很期待看到您 如何使用这一功能 利用 Apple GPU 强大的并行性 来适应您几何体处理的需求 感谢收看