Scale compute workloads across Apple GPUs

Highlight

Apple M1 GPU 从 8 核 iPad 到 64 核 Mac Studio 共享同一架构，但计算工作负载要获得线性扩展性能，需要解决工作分布不均、GPU 时间线空闲、原子操作竞争三大瓶颈。

核心内容

你的 Metal 计算应用在 M1 上跑得不错，但放到 M1 Ultra 上性能只提升了 30%，远没有达到核心数翻倍应有的效果。问题出在哪？

GPU 工作负载的可扩展性（scalability）指的是：随着 GPU 核心数增加，性能能否成比例提升。理想情况是线性扩展——核心数翻倍，性能也翻倍。现实中常见三种非理想情况：性能达到平台期、GPU 时间线存在大量空闲间隙、性能提升不均匀（在某些核心数区间遇到瓶颈）。

（00:59）Apple M1 GPU 家族从 8 核扩展到 64 核，同一套 GPU 架构支持所有 Metal 3 特性。Affinity Photo 和 DaVinci Resolve 等专业应用已经实现了优秀的跨设备扩展性能。

瓶颈分两类：计算受限（computation-bound）和带宽受限（bandwidth-bound）。优化过程中你可能在两者之间来回切换。如果计算是瓶颈，可以尝试把部分负载转移到内存访问；反之亦然。MPS 和 MPSGraph 框架已经针对所有硬件做了优化，但自定义计算内核仍需手动调优。

（04:37）工作分布（work distribution）是第一个需要检查的环节。工作负载以 3D 线程组网格（threadgroup grid）的形式分派。线程组被均匀分布到 GPU 核心，每个线程组可以访问容量有限但速度极快的线程组内存（threadgroup memory）。单个线程组进一步拆分为 SIMD-group（在 Apple GPU 上宽度固定为 32）。每个线程组最多 1024 个线程，最多 32KB 线程组内存。

如果线程组数量太少，GPU 无法饱和。经验法则是：每个 shader 核心需要 1K 到 2K 并发线程才能达到良好占用率（occupancy）。M1（8 核）需要至少 8K-16K 线程，M1 Max（32 核）需要 32K-64K，M1 Ultra（64 核）需要 64K-128K。

（06:43）另一个常见问题是线程组过大。过大的线程组会阻碍负载在 GPU 核心间的均匀分布。最佳实践是使用 SIMD 宽度（32）的最小倍数，且能很好映射到你的工作负载。

详细内容

消除 GPU 时间线间隙

（08:18）GPU 空闲是最大的性能杀手。假设一个工作负载因为 CPU/GPU 串行化只利用了 50% 的 GPU。CPU 做一部分工作，等 GPU 完成，再做下一部分。即使 GPU 核心翻倍，整体性能提升也有限：核心翻倍后 GPU 部分快了一倍，但 CPU 部分不变，总时间从 2 单位降到 1.5 单位，只提升了 33%。

根本原因通常是 waitUntilCompleted 导致的 CPU/GPU 同步。修复方法有几种：

使用 MTLSharedEvents 替代 waitUntilCompleted

// 不推荐：阻塞 CPU 等待 GPU
[commandBuffer waitUntilCompleted];

// 推荐：使用共享事件减少同步开销
id<MTLSharedEvent> sharedEvent = [device newSharedEvent];
[commandBuffer encodeSignalEvent:sharedEvent value:1];
// CPU 可以在需要时检查事件状态，不必阻塞

关键点：

waitUntilCompleted 会阻塞当前线程，造成 GPU 时间线间隙
MTLSharedEvents 开销更低，允许更灵活的同步策略
如果可能，完全移除 CPU/GPU 同步点

流水线化工作负载

（10:57）如果算法能提前知道下一批数据，可以在等待共享事件之前就编码好下一批工作。这样 GPU 不会干涸，始终有工作可处理。

使用并发分派（concurrent dispatches）

（12:54）对于处理多张图片的场景，原本一张接一张串行处理。使用 MTLDispatchTypeConcurrent 后，驱动可以把不同图片的独立工作交错执行，隐藏内核同步开销，同时填满各内核的启动和收尾阶段。

// 创建支持并发分派的命令缓冲区
id<MTLCommandBuffer> commandBuffer = [commandQueue commandBufferWithUnretainedReferences];

// 使用 concurrent dispatch type
// 需要手动插入 barrier 管理依赖关系

实验结果显示：两张图片交错处理提升 30%，三张并行提升 70%。

优化原子操作

（14:07）全局原子操作（global atomics）在整个 GPU 上保持一致性。当大量线程竞争同一个内存地址时，性能会急剧下降，而且核心数越多竞争越严重。

以归约（reduction）算法为例：把所有 buffer 中的值求和。

不推荐的做法：每个线程直接对主内存做 atomic add。

// 性能差：所有线程竞争同一个全局地址
kernel void badReduction(device float* buffer [[buffer(0)]],
                         device atomic_float* result [[buffer(1)]],
                         uint gid [[thread_position_in_grid]])
{
    atomic_fetch_add_explicit(result, buffer[gid], memory_order_relaxed);
}

推荐的做法：分两级归约，先在线程组内用 SIMD 指令和 threadgroup 原子操作，最后每个线程组只写一个全局原子。

kernel void goodReduction(device float* buffer [[buffer(0)]],
                          device atomic_float* result [[buffer(1)]],
                          uint lid [[thread_index_in_threadgroup]],
                          uint tgid [[threadgroup_position_in_grid]])
{
    threadgroup float sharedSum;
    
    // 第一步：SIMD-group 内归约，使用寄存器操作，不访问内存
    float localSum = buffer[lid]; // 简化示例
    localSum = simd_sum(localSum);
    
    // 第二步：每个 SIMD-group 的最后一个线程写入 threadgroup 内存
    if (lid == 31) {
        atomic_fetch_add_explicit(&sharedSum, localSum, memory_order_relaxed);
    }
    
    threadgroup_barrier(mem_flags::mem_threadgroup);
    
    // 第三步：每个线程组只写一个全局原子
    if (lid == 0) {
        atomic_fetch_add_explicit(result, sharedSum, memory_order_relaxed);
    }
}

关键点：

simd_sum 等 SIMD-group 指令在寄存器间交换数据，无需内存往返
threadgroup 原子操作由每个 GPU 核心独立的 threadgroup 内存处理，随核心数线性扩展
最后一步全局原子操作的数量等于线程组数量，且各线程组完成时间不同，天然分散了竞争

优化内存访问模式

（17:15）即使 GPU 时间线没有间隙，扩展性能仍可能不理想。这时需要检查 GPU limiters。Xcode 和 Metal System Trace 提供了详细的计数器。

低效内存访问会导致高 Last Level Cache（LLC）或 MMU limiter，以及低利用率。两种优化方向：

重组数据布局：让数据访问更局部化
调整访问模式：让线程组的访问模式匹配数据布局

（18:12）假设内存中数据按行排列，但线程组以方形 2D 模式访问。第一个 SIMD-group 访问的数据被打包到 cache line 中，但大部分 cache line 空间被浪费了。解决方案之一是把数据重新组织成条带（stripes），让线程组能充分利用每个 cache line。

另一个方案是调整线程组形状。比如用更扁长的矩形线程组替代方形，让访问模式与行优先的数据布局对齐。

Blender Cycles 案例

（20:08）Blender Cycles 使用排序来减少线程发散：按材质类型对光线命中点排序。这减少了线程发散，但增加了空间内存发散，导致高 MMU limiter。

解决方案：在排序前对内存范围做分区（partitioning）。排序时确保同一分区的索引不会混合。这样 SIMD-group 访问的数据范围被限制在分区内，大幅降低 MMU 压力。实验结果：Top limiter 和 LLC limiter 降低约 20%，GPU 读带宽显著提升，整体性能提升 10% 到 30%（取决于场景）。

核心启发

优化你的图像处理管线

如果你有批量图像滤镜处理的 App，检查是否在用 waitUntilCompleted 串行化 CPU 和 GPU。改用 MTLSharedEvents 配合流水线化，可以让 GPU 利用率从 50% 提升到接近 100%，在 M1 Ultra 上获得接近线性的性能扩展。入口是 MTLSharedEvent 和 encodeSignalEvent:value:。

实现 GPU 上的并行粒子系统

粒子模拟通常需要每帧做位置更新、碰撞检测、渲染。用并发分派把不同粒子组的工作交错提交，隐藏内核同步开销。注意在依赖点手动插入 barrier。关键 API 是 MTLDispatchTypeConcurrent。

为机器学习推理优化内存布局

如果你的 App 用 Metal 做自定义神经网络层，检查内存访问模式。把权重数据从行优先重排成更适合线程组访问的格式，可以显著降低 LLC limiter。用 Xcode 14 的新增计数器（MMU Limiter、MMU Utilization、MMU TLB Miss Rate）来验证优化效果。

关联 Session

Discover Metal 3 — Metal 3 的新特性概览
Program Metal in C++ with metal-cpp — 用 C++ 编写 Metal 代码
Transform your geometry with Metal mesh shaders — Metal Mesh Shaders 几何处理管线