Bring your Metal app to Apple silicon Macs

Highlight

Apple Silicon Mac 把 Apple GPU 的 TBDR 架构带到 Mac，Metal 应用需要改用 API 驱动的特性检测、显式 load/store action、位置不变性和 threadgroup 同步来避免 Rosetta 兼容路径的性能成本。

核心内容

很多 Mac 图形应用原来只面对 Intel、AMD 或 Nvidia GPU。它们习惯了 Immediate Mode Renderer（立即模式渲染器）的行为：深度、颜色、模板缓冲区在显存中频繁读写，很多错误用法在某些驱动上没有立刻暴露。

Apple Silicon Mac 换成了 Apple 自研 GPU。这个 GPU 使用 Tile-Based Deferred Rendering（基于图块的延迟渲染，TBDR）。它先把几何划分到屏幕图块，再在片上 tile memory 里完成深度测试、混合和部分颜色操作。写得好的渲染器会少走很多内存带宽，功耗和性能都受益。

迁移路径分成两步。旧的 x86 应用可以先通过 Rosetta 运行，系统会打开一些 Metal 兼容性 workaround，让画面尽量正确。接下来，开发者需要重新编译原生 Apple Silicon 版本，并修复依赖未定义行为的 Metal 代码。兼容性 workaround 有性能成本，原生版本不能把它们当作长期方案。

这场 session 的价值在于，它没有停留在“重新编译”四个字。它直接指出四类真实问题：用 GPU 名称判断能力、错误的 load/store action、缺少位置不变性、threadgroup memory 同步假设错误，以及在同一个 render pass 里采样当前 depth/stencil attachment。这些问题都会在 TBDR 架构上变成画面错误或带宽浪费。

详细内容

用 Metal API 检测 GPU 能力

（14:23）Apple 明确要求应用直接查询 Metal GPU feature。不要用操作系统名称或 GPU 名称推断能力。Apple Silicon Mac 同时支持 Mac 2 family 和 Apple GPU family，过去“macOS 上没有 Apple GPU 特性”的判断已经失效。

{
    self.appleGPUFeatures = metalDevice.supportsFamily(.apple5)

    self.simdgroupSize = computePipeline.threadExecutionWidth

    self.isLowPower = metalDevice.isLowPower
}

关键点：

• supportsFamily(.apple5) 直接查询 Apple GPU family 支持情况，用来决定 programmable blending、tile shaders、local image blocks 这类 Apple GPU 能力是否可用。
• threadExecutionWidth 从 compute pipeline 获取 SIMD group 宽度，后面的 threadgroup 同步逻辑需要它。
• isLowPower 用来区分集成或独立 GPU 风格的代码路径；session 特别说明，Apple GPU 上它返回 false，应按独立 GPU 的性能特征看待。

正确选择 load/store action

（15:23）Load/store action 控制 color、depth、stencil attachment 在 render pass 开始和结束时如何进入或离开片上 tile memory。TBDR GPU 对这个选择很敏感，因为它决定是否从系统内存加载旧内容，是否把 tile memory 写回系统内存。

判断方式很直接：render pass 开始时，只有 previous content needs to be preserved，才选择 load action = load；否则避免不必要的 load。render pass 结束时，只有 content will be consumed in a later pass，才选择 store action = store；否则避免 store。

关键点：

• session 的例子是先在前一个 pass 渲染一张 texture，最后一个 pass 继续在这张旧内容上绘制；如果此时使用 DontCare load action，Apple GPU 不会把系统内存中的 texture 上传到 tile memory。
• 当后续绘制没有覆盖整个 framebuffer 时，未初始化的 tile memory 会留下图形错误；需要保留旧内容时应选择 load。
• store action 的判断同样来自资源流向：只有 attachment 会被后续 render pass 消费时，才把 tile memory flush 回系统内存。
• 不需要后续读取的 attachment 避免 store，可以减少额外内存流量。

需要跨 pass 精确复用位置时，显式打开 position invariance

（20:58）多 pass 算法常见做法是第一遍写 depth，第二遍用 equal depth test 再画同一批几何。两个 vertex shader 即使调用同一个位置计算函数，编译器优化也可能让 position 输出有微小差异。Apple GPU 默认不保证位置不变性。

// Renderer.swift
let options = MTLCompileOptions()
options.preserveInvariance = true

library = try device.makeLibrary(source: sourceString,
                                 options: options)


// vertex.metal
struct VertexOut
{
    float4 pos [[position, invariant]];
    float data;
};

关键点：

• MTLCompileOptions() 创建 Metal 编译选项。
• options.preserveInvariance = true 要求编译器保留位置不变性。
• makeLibrary(source:options:) 用带有该选项的编译配置生成 shader library。
• [[position, invariant]] 标在 vertex 输出位置上，告诉编译器这个 position 需要跨 shader 保持一致。
• 这项能力有性能成本，只应该用于依赖相同 depth 值的 shader，例如使用 equal depth compare 的 pass。

threadgroup memory 必须按 SIMD group 大小同步

（24:25）Compute shader 中的 threadgroup memory 由同一个 threadgroup 内的线程共享。不同 GPU 的 SIMD group 大小不同。Apple GPU 是 32；如果应用假设 64 个线程天然在一个 SIMD group 内执行，就可能漏掉跨 SIMD group 的同步。

// Correct synchronization

// launched with threadgroup size = 64
kernel void kernelMain(uint tid [[ thread_index_in_threadgroup ]],
                       uint simd_size [[ threads_per_simdgroup ]],
                       device uint * res [[ buffer(0) ]])
{
    threadgroup uint buf[64];

    buf[tid] = initBuffer(tid);

    if (simd_size == 64u)
        simdgroup_barrier(mem_flags::mem_threadgroup);
    else
        threadgroup_barrier(mem_flags::mem_threadgroup);

    uint index = (tid < 32) ? tid + 32 : tid - 32;
    res[tid] = buf[tid] + buf[index];
}

关键点：

• thread_index_in_threadgroup 给出当前线程在 threadgroup 内的索引。
• threads_per_simdgroup 在 shader 内取得当前 GPU 的 SIMD group 大小。
• threadgroup uint buf[64] 是 64 个线程共享的 threadgroup memory。
• buf[tid] = initBuffer(tid) 写入共享内存，后续其他线程会读取。
• simdgroup_barrier 只同步同一个 SIMD group 内的线程。
• threadgroup_barrier 同步整个 threadgroup，适合一个 threadgroup 含多个 SIMD group 的情况。
• res[tid] = buf[tid] + buf[index] 读取另一个线程写入的值，因此前面的 barrier 是正确性的必要条件。

不要在同一个 render pass 里采样当前 depth/stencil attachment

（26:11）session 最后指出一个 depth/stencil 绑定错误：同一个 texture 作为当前 render pass 的 attachment 时，不能同时作为采样 texture 使用。这会制造对同一底层 texture 的并发读写。

修正方式是把 current depth/stencil attachment 从 sampled texture 绑定中移走。确实需要读取当前 attachment 内容时，先在 render pass 之前创建第二份 copy，再让 shader 读取这份 copy。

关键点：

• session 描述的错误是 current depth attachment 同时被 fragment shader 采样，造成同一底层 texture 的并发读写。
• 对 opaque geometry，深度会在 fragment shader 执行前累积；到 render pass 结束时，Apple GPU 会把片上 depth/stencil memory flush 回系统内存，此时并发读写会触发正确性问题。
• 这个 race condition 可以发生在 render pass 中任意 draw，不只发生在最后一个 draw。
• 不要用 texture barrier 或 memory barrier 绕过这个问题；session 明确说这类 workaround 在 TBDR 架构上很昂贵。
• 如果算法确实需要采样当前 attachment，先创建第二份 copy，再让 shader 读取这份 copy。

核心启发

1. 给渲染器加一个 Apple GPU capability 层

• 做什么：启动时集中记录 supportsFamily、threadExecutionWidth 和 isLowPower。
• 为什么值得做：同一套 Mac 代码需要同时面对 Intel/AMD/Nvidia 和 Apple GPU，名称判断会很快失效。
• 怎么开始：封装 RendererCapabilities，把 programmable blending、tile shader、simdgroup size、GPU 性能类别都放进去，渲染路径只读这个对象。

2. 做一次 render pass load/store 审计

• 做什么：逐个检查 color、depth、stencil attachment 的 load/store action。
• 为什么值得做：TBDR 架构把 attachment 放进 tile memory，错误的 .dontCare 会导致画面错误，多余的 .load/.store 会浪费带宽。
• 怎么开始：按 pass 画一张资源流向表，标出哪些 attachment 被后续 pass 消费；只有这些 attachment 使用 .store。

3. 为多 pass depth equal 流程标记 position invariance

• 做什么：找出第一遍写 depth、第二遍用 equal 比较的渲染流程。
• 为什么值得做：不同 vertex shader 的优化差异会让 depth 值不完全相同，结果是部分像素被错误丢弃。
• 怎么开始：只对相关 library 打开 preserveInvariance，并在需要精确匹配的 vertex output 上加 [[position, invariant]]。

4. 给 compute shader 建立 SIMD group 版本策略

• 做什么：为依赖 threadgroup memory 的 kernel 准备按 SIMD group 大小选择的实现。
• 为什么值得做：Apple GPU 的 SIMD group 是 32，旧代码如果假设 64，会漏掉 threadgroup_barrier。
• 怎么开始：先用 threads_per_simdgroup 修正同步，再把高频 kernel 拆成针对 32 和其他宽度的两个版本，运行时按 threadExecutionWidth 选择。

5. 把 depth/stencil 采样改成双纹理流程

• 做什么：当前 pass 写 depth/stencil，shader 只读取上一份 depth/stencil copy。
• 为什么值得做：同一 render pass 内采样当前 attachment 是未定义行为，Apple 的兼容快照机制只适用于旧 SDK 路径且有性能成本。
• 怎么开始：为需要读取深度的效果在 pass 之前创建可采样 copy，禁止把当前 depth/stencil attachment 同时绑定为 sampled texture。

关联 Session

• Optimize Metal Performance for Apple silicon Macs — 这场 session 的直接后续，继续讲如何让 Metal 渲染器利用 Apple GPU 的带宽和 tile memory。
• Harness Apple GPUs with Metal — 从 Apple GPU 能力出发，补充 programmable blending、tile shaders 等特性适合的渲染场景。
• Optimize Metal apps and games with GPU counters — 用 GPU counters 验证 load/store、barrier 和带宽优化是否真的减少瓶颈。
• Discover ray tracing with Metal — 适合继续了解 2020 年 Metal 图形管线中更高级的渲染能力。