探索 M3 和 A17 Pro 中的 GPU 改进

欢迎大家 我叫 Jedd Haberstro 是 Apple GPU、图形 和显示软件团队的工程师

很高兴向大家介绍 全新 Apple 系列 9 GPU 架构中的 A17 Pro 和 M3 系列芯片 这些芯片是 iPhone 15 Pro 和新款 Mac 的核心

在 Apple 的整个产品系列中 GPU 为用户喜爱的丰富体验 提供了支持 无论是使用新款 iPhone 15 Pro 随时随地玩游戏

在新款 iMac 上为你的 App 提供顺畅的 UI 动画

还是利用机器学习 在新款 MacBook Pro 上 执行高级视频和图像处理

GPU 在支持这些 App 方面都 发挥着关键作用

Metal API 用于 利用 Apple GPU 的 计算能力 这些 App 共同运行一组不同的 Metal 着色语言程序 这些着色器程序的范围广泛 包括 从仅执行几行代码的小型简单着色器 到需要数十万行代码 框架和库的 大型复杂着色器 所有这些着色器的共同点 是海量数据并行性 这是通过并行运行 来极大提高 App 性能的 机会 此并行性通过在不同输入上 并行运行很多次 Metal 着色器程序

来实现 此类输入包括 3D 渲染 场景中的每个顶点 或屏幕的每个像素

最终 GPU 负责并行执行 这些着色器

每个 GPU 的心脏是其着色器核心 每个着色器核心都可以 并行运行数千个线程 为了进一步扩展性能 一个 GPU 上的多个着色器核心 也可以并行运行 为 App 提供数以万计的 并行执行线程

当今现有的 iPhone、iPad 和 Mac 中的 GPU 已经具有出色的性能

以及一套强大的开发者工具 让 App 开发者能够充分地 发挥 GPU 的潜力 但凭借全新 Apple 系列 9 GPU 我们将性能推到一个新高度 这要归功于几项令人兴奋的新改进

首先是全新的着色器核心架构 它可提高现有 App 的性能和能效 从而立即提升你交付的体验 同时帮助你应对构建新一代 App 所面临的严峻挑战

硬件加速光线追踪的提升 对已经使用 Metal 光线追踪 API 的 App 是透明的 同时带来更多使用光线追踪的机会 以出色的性能实现丰富的渲染效果

而且 通过硬件加速的网格着色 App 能够以前所未有的方式构建 高级几何处理管道

在我们更详细地讨论这些方面之前 我们来看看新的 Apple 9 系列 GPU 让 App 能够在没有任何变化的 情况下实现的出色性能 这是 Larian Studios 的 Baldur's Gate 3 上面是配备 M3 Mac 的新款 MacBook Pro 的运行效果 下面是配备 M2 Mac 的 MacBook Pro 的运行效果 每个渲染均采用 1800p 的 超视频质量设置 M3 Mac 能够提供 显著的性能提升 这得益于新一代着色器核心能够 以更高的线程占用率 运行游戏的 Metal 着色器 这是 Blender 渲染的 理发店场景图像 使用的是的 Cycles Path Tracer 后者利用了 M3 Mac 上的 Metal 光线追踪 两个渲染均在同一时间开始 但得益于硬件加速光线追踪 和新一代着色器核心 M3 Mac 上的渲染的 收敛速度明显更快

这是由 Pixar 的 Hydra Storm 渲染的 《玩具总动员 4》 Antiques Mall USD 的实时可视化效果 Hydra Storm 使用 M3 Mac 上的 Metal 网格着色 后者与硬件加速网格着色相结合时 运行速度比以往更快

现在 我们从新一代着色器核心开始 更详细地了解每项功能

Apple 9 系列 GPU 由 多个构建块组成 例如解析 Metal 命令缓冲区的 计算和顶点命令处理器 分派片元着色器执行的光栅化器

以及一系列缓存 包括为所有 GPU 内存流量 提供服务的 GPU 末级缓存

但任何 GPU 的核心都是 其着色器核心 这些是 执行 App Metal 着色器的 构建块

着色器核心还与一个纹理单元配对 后者可以采样和写入纹理资源 此外 着色器核心也与一个全新的 光线追踪单元配对 后者可加速光线相交请求

着色器核心可以进一步细分 其组成部分

每个着色器核心都有一组执行管道 用于执行不同类型的指令 例如 FP32、FP16 和整数计算 它们对应于对着色器中变量的操作 并使用 Metal 数据类型 (例如 float、half 和 int) 此外 着色器核心均具有 用于对纹理和缓冲区 进行读写操作的内存管道

保持所有这些执行管道忙碌 通常需要执行 来自多个 SIMD 组的指令 因此 有一个池可以追踪 在着色器核心上 运行的 SIMD 组 还有一个调度程序用于 选择下一个执行指令的 SIMD 组

通常还有一些片上内存 用于存储着色器程序可能使用的 不同类型的数据 例如用于存储变量值的寄存器 线程组和图块内存 用于存储在计算线程组中 共享的数据 或在分块中 共享的颜色附件数据

此外还有一个缓存用于提高 堆栈和缓冲区访问的性能

了解过着色器核心的作用 及其组成部分后 我想解释一下 Apple 系列 9 GPU 着色器核心 三项令人兴奋的新改进 这些改进将提高着色器的性能 而无需更改 App 但通过更好地了解 这个新的着色器核心的工作原理 你更能从中受益

第一个变化是动态着色器核心内存 它让 App 实现更高的线程占用率 因此通常可以实现更佳性能

第二个变化是灵活的片上内存 这将提高着色器访问 缓冲区、堆栈、线程组 和分块内存的效率

最后一个变化是着色器核心的 高性能 ALU 管道 其并行执行能力得到提升 这将提高执行浮点 或整数运算组合的 App 的性能

在进一步探索这些新功能之前 我们来更详细地了解 着色器核心如何保持其执行管道忙碌 以及在这方面线程占用率的重要性

假设你的 Metal 着色器 使用 ALU 管道 执行一些数学运算后读取一个缓冲区 其结果将在之后立即使用 访问缓冲区 可能需要一直访问到设备内存 这是一个较长的延迟操作 在此期间 SIMD 组无法执行其他操作 这会导致 ALU 管道闲置

为了缓解这种情况 着色器核心可以执行 来自不同 SIMD 组的指令 可能有一些自己的 ALU 指令 这减少了 ALU 运行和使用的时间 并允许 SIMD 组并行运行 从而提高了性能

如果有其他 SIMD 组 在着色器核心上运行 会多次采用这种方式 直到 ALU 和其他执行管道 不会缺乏要执行的指令

在一个着色器核心上 同时运行的 SIMD 组数量 称为其线程占用率

但你可能会问自己 是什么决定了 有多少个 SIMD 组 能在着色器核心上同时运行？

为了回答这个问题 我们来看一个例子 这是一个典型的光线追踪计算内核 可与具有加速结构的光线相交

检查相交结果

然后基于相交图元的材质 执行不同的着色函数 在这个例子中 它支持玻璃和皮革材质的着色

每行代码都会使用一定数量的寄存器 来存储程序的变量 在程序的不同阶段 将使用更多或更少的寄存器 具体取决于代码会做什么 在这个特定的例子中 shadeGlass 函数的实现 比程序的其余部分使用更多的寄存器

在 Apple 系列 9 GPU 之前 SIMD 组需要 其从片上寄存器文件分配寄存器后 才能在着色器核心上开始执行 分配的数量将等于 程序中任何时候的最大寄存器使用量 SIMD 组将在自身的 整个持续时间内 保留这些分配的寄存器 即使这些寄存器中的大多数 可能在程序的大部分中未得到使用 因此根据最大寄存器使用情况 我们可能只能在着色器核心上 一次运行四个 SIMD 组 因为如果更多的话将需要 比现有更多的片上寄存器字段内存 不过 得益于Apple 系列 9 GPU 的 全新动态着色器核心内存功能 最大寄存器使用量将不再限制 可以运行的 SIMD 组数量 现在 根据程序每个部分 实际使用的内容 片上寄存器内存 在着色器的生命周期内 得到动态分配和解除分配 这使得 SIMD 组能够更有效地 使用片上寄存器文件 从而释放原本无法可供使用的空间 这可能会对 App 的 线程占用率产生深远影响 并最终通过允许更多 SIMD 组 同时运行来有效影响其性能

正如我刚才提到的 寄存器现在可以在 SIMD 组的生命周期内 得到动态分配和取消分配 这在一定程度上是可能的 因为寄存器文件现在是一个缓存 而不是寄存器的永久存储 这意味着可以使用的更多寄存器 能够存储在芯片上

灵活的片上内存功能将此处理 扩展到着色器核心的其余内存类型 例如线程组和分块内存 使其也成为缓存

现在寄存器、线程组、分块、堆栈 和缓冲区数据都缓存在芯片上 这使我们能够重新设计片上内存 将其分成数量更少更大的缓存 来服务所有这些内存类型 这种灵活性将有利于 并不大量使用每种内存类型的着色器 在过去 如果计算内核不使用 线程组内存 则其相应的片上存储 将完全未被使用 现在片上存储将动态分配给 着色器使用的内存类型 为它们提供比过去更多的片上存储 最终实现更出色的性能

例如 对于寄存器使用量较大的着色器 这可能意味着更高的占用率

对于重复访问大量工作缓冲区数据的 着色器来说 这意味着更高的缓存命中率 以及更低的缓冲区访问延迟 从而带来更出色的性能 对于大量使用非内联函数的 App (例如函数指针、可见函数表 和动态关联的着色器库) 这意味着有更多的片上堆栈空间 来传递函数参数 因此函数调用速度更快

但如果你的 App 使用的内存 仍然多于片上存储 会发生什么情况？ 如果不加以缓解 这些数据将溢出到下一个缓存级别 甚至溢出到主内存 幸运的是 着色器核心将动态监控 着色器的行为并调整占用水平 以防止这种情况发生 它会将数据保留在芯片上 最终执行管道将变得忙碌

不过这确实意味着 着色器的占用率会受到

着色器的访问线程组、分块 堆栈和缓冲区内存

以及动态寄存器使用情况的影响

这些新的硬件功能 提高了许多 App 的占用率 这意味着开发者 无需像过往那么频繁地 优化占用率的次数 但如果你需要进一步 优化 Apple 系列 9 GPU 的占用率 我们开发了一套分析工具来帮助你 要了解有关如何诊断和优化 占用率的更多信息 请参考这些讲座

关于 Apple 9 系列 GPU 着色器核心 我想讨论的最后一个特性 是其高性能 ALU 管道

Apple GPU 着色器核心具有 针对不同指令类型 (包括 FP16 指令) 的单独 ALU 管道 Apple GPU 高度优化 FP16 运算 我们建议你尽可能 使用 FP16 数据类型

FP16 数学指令以峰值吞吐量执行

与 FP32 同等类型相比 它们使用的寄存器更少 如果你的缓冲区以原生方式 通过 FP16 存储数据 它们会减少内存带宽 对于数学操作的源变量 或目标变量 还不是 FP16 的情况 无需开销即可进行相互转换

但如果你的 App 仍然 执行其他数学运算 例如 FP32 和整数 Apple 系列 9 GPU 着色器核心可以 比以往更高效地并行执行 所有三种数据类型的指令 与之前的 Apple GPU 相比 这可以提供 高达 2 倍的 ALU 性能 为了利用这种额外的并行性 指令必须从 多个 SIMD 组执行 这意味着增加占用率 可以提高 ALU 管道的利用率 让我们考虑一个例子 想象一下 有两个 SIMD 组同时运行 两者都执行 ALU 指令 过去这些 SIMD 组可能必须 相继运行

但如果它们在不同时间点 执行 FP32 和 FP16 指令 (如此处所示) 那么它们的执行就可以重叠 以提高并行性和性能

我们来回顾一下 新一代着色器核心的新增功能 它将在着色器的生命周期内 动态分配和回收寄存器 从而提高了线程占用率

它具有大型片上缓存 为寄存器、线程组、分块、堆栈 和缓冲内存提供服务 从而提高了 访问这些内存类型的性能

着色器核心将动态调整占用率 让芯片上的数据和执行管道保持忙碌

最后 FP16、FP32 和整数运算 可比以往更多地并行执行 从而提高了 ALU 性能

接下来 我们来看看硬件加速的光线追踪

借助 Metal 光线追踪 App 可以利用 Apple GPU 的 大规模并行性 让光线与其场景几何体相交 如果你不熟悉 Metal 光线追踪 希望了解更多信息 请观看《Metal 光线追踪指南》 以及《使用 Metal 光线追踪 增强你的 App》

Metal 光线追踪 API 的核心 是相交器对象 该对象负责确定 光线与加速结构中包含的图元的 交点 它经常被光线追踪 App 的 GPU 函数 (也称为着色器) 多次调用 因此它是 App 性能的核心

早些时候 我在讨论该 raytracingKernel 的 寄存器使用情况时 展示了该 GPU 功能 它创建了一个相交器对象 并通过调用该对象的 intersect 方法 找到交点

为了确定交点 相交器执行几个关键阶段 首先 它遍历加速结构 以找到候选图元 然后 它调用 相交函数 (可以由 App 提供) 来确定光线是否与图元相交

如果是 系统会将交点与之前的交点进行比较 并重复该过程 直到找到最接近的交点

然后最近的交点将 返回到调用 GPU 函数 以进行特定于 App 的进一步处理

Apple 系列 9 GPU 的新功能是 相交器对象的实现 是由硬件加速的 这大大提高了 这一关键操作的性能

硬件加速交点 不按照 GPU 函数执行 因此 为了促进光线 和在两者之间的光线有效负载的通信 数据将被读写入片上内存 你可以在新 Xcode 中使用 RT 暂存性能计数器来予以观察

我已经讨论了相交器的角色和职责 我们现在将通过一个示例 来剖析一次性通过 此相交循环的性能特征

想象一下 我们的 App 正在执行 两个 SIMD 组 每个均希望 通过加速结构使四条光线相交

在此示例中 我们的加速结构包含 经典的康奈尔盒子 具有一个长方体对象和一个球体对象

通过调用 intersect 方法 光线被投射到场景中 向其传递光线、加速结构 和相交函数表 每个 SIMD 组都有 两条与长方体相交的光线 和两条与球体相交的光线 在此示例中 通过使用 MTLAccelerationStructure TriangleGeometryDescriptor

并设置其不透明属性为 yes 将长方体定义为不透明三角形图元 因此相交可以使用 Metal 内置的 相交函数计算交集

但球体是使用相交必须调用的 自定义边界框相交函数 在程序上加以定义的

自定义 BoundingBoxIntersection 函数是使用 intersection 属性和 bounding_box 参数来声明的

正如我之前提到的 intersect 方法由每个正在 针对加速结构测试光线的线程调用 有鉴于此示例 我来看看在传统实现中 每个 intersect 如何调用遍历 以及相交测试如何执行

在典型用法中 并非所有遍历都会花费相同的时间 来定位要测试光线的图元 这造成了所谓的执行发散 导致 SIMD 组中的每个线程 需要等待该 SIMD 组的 最长遍历后才能进入下一阶段

事实证明 执行相交函数时 同样的开销也会增加 执行发散导致各类型 相交函数依次运行 进一步减少并行性 从两个阶段的总和来看 每个线程大部分运行时间 都处于空闲状态 等待 SIMD 组中的 其他线程完成 这是一个主要的性能瓶颈

有鉴于传统实现的情况 我们来讨论一下硬件加速光线追踪 如何优化这些低效情况

第一个主要改进 是硬件相交器 能够使用固定函数硬件 完全独立地运行每次遍历 这在一定程度上是可能的 因为数组被发送到硬件相交器 进行处理 而不是在 GPU 函数中执行 这大大减少了遍历所花费的时间 并且还消除了 传统遍历执行发散的开销

另一方面 相交函数 是 Metal 着色语言代码 因此它们仍须分组到 SIMD 组 才能在着色器核心上运行 但由于硬件相交器 独立执行每条光线 因此可以自由地将 来自不同 SIMD 组的光线 的相交函数调用组合在一起

这就是重新排序阶段的作用 当光线在非常接近的距离和时间内 到达此阶段时 相交函数调用将被分组到 相关的 SIMD 组 从而减少甚至完全消除 传统实现中存在的执行发散开销

我已经向大家展示了 硬件加速光线追踪如何提高 App 光线相交器调用的性能 现在我们来了解一下 你的 App 可以实施的一些最佳实践 以充分利用其优势

我们的第一个建议是 尽可能使用相交器对象 API Metal 还允许使用相交查询 API 执行光线追踪 但该 API 会增加 必须读取和写入的 光线追踪暂存内存量 并会禁用重新排序阶段

我们还建议 在编写自定义相交函数时 避免创建一个能够执行 许多不同逻辑相交例程的 超级函数 相反 为每个逻辑相交例程 创建一个 Metal 相交函数 这增加了重新排序阶段的好处

同样重要的是 尝试最小化 已传递到相交器对象 并从该对象返回的 光线负载结构的大小 这将减少着色器的延迟 并可能增加其线程占用率

有关如何优化光线追踪 App 的 更多详细信息和指南 请观看这些讲座

我们来总结一下 Apple 系列 9 GPU 通过具有固定函数遍历块 以及相交函数重新排序阶段的 新硬件加速 极大地提高了 光线追踪的性能

尽管这款新硬件将提高 所有 Metal 光线追踪 App 的性能 但为了充分地让 App 从中受益 最好尽可能使用相交 API 而不是相交查询 API

接下来我将介绍 Apple 系列 9 GPU 的最后一项改进 这就是硬件加速网格着色

网格着色是渲染管道中 GPU 驱动的灵活几何处理阶段 它以两个类似计算的着色器 取代了传统 顶点着色器阶段

对象着色器在第一阶段执行 可用于对特定于 App 的输入 (例如整个网格对象) 执行粗粒度处理 每个对象线程组都可以选择生成一个 网格组来执行后续更细粒度的处理 网格着色器构成第二阶段 通常 网格线程组将处理 父对象的组成部分 后者通常称为 meshlet

网格线程组的输出 是一个 Metal 网格对象 该对象包含了要由 传统图形管道的其余部分处理的 顶点和图元列表

网格着色有许多应用 例如细粒度几何调用

程序几何生成 自定义 App 特定的几何表示 例如压缩格式 这还包括从其他图形 API 移植几何和曲面细分着色器

如果你不熟悉 Metal 中的网格着色 建议查看下面的两个讲座

借助 Apple 系列 9 GPU 上的 硬件加速网格着色 你将观察到的最显著改进 是现有的网格着色代码 在性能上得到大幅提升

Apple 系列 9 GPU 能够更有效地安排对象 和网格线组 以将中间网格数据保存在芯片上 因而减少了内存流量

新的硬件 还附带了一些 Metal API 增强功能 首先是支持将绘制网格命令 编码到间接命令缓冲区中 这让 GPU 驱动的渲染管道 可利用网格着色 以及传统的顶点着色器

第二个 API 增强功能 扩展了每个网格的最大线程组数 从 1024 增加到超过 100 万

现在我们来了解一些最佳实践 确保可实现更佳的网格着色性能

网格线程组输出的 metal::mesh 对象 有几个模板参数 其尺寸尽可能小很重要

对于网格的顶点和图元数据类型 可以删除未使用属性来完成 此类属性因与其他不相关顶点或网格 函数共享这些数据类型而可能出现 网格类型还必须指定 可以输出的图元和顶点 的最大数量 这些不应设置为大于 App 的几何结构、管道和资源 实际需要的值 请注意 其大小将减少内存流量 并可能增加占用率

若在网格着色器中执行每个图元调用剔除 我们不建议将顶点位置 写入网格对象 只为硬件调用随后的调用阶段 更好的方法是完全省略 写入此类图元 因为这可以在硬件 几何处理阶段的其余部分中 节省大量处理时间

我们来回顾一下本次讲座介绍的 有关 Apple 系列 9 GPU 的内容

新一代着色器核心 通过动态分配寄存器存储 并在多种内存类型之间共享片上内存 来提高片上内存利用率 从而提高线程占用率和性能

硬件加速光线追踪大幅提升了 使用 Metal 光线追踪 API 的 App 的性能 实现了新的高保真视觉效果 最后 由于硬件加速 网格着色性能得到极大提高 使更多 App 能够自定义 其几何处理管道

感谢你的观看