WWDC Quick Look 💓 By SwiftGGTeam
中文 EN 日本語
Bring your Metal app to Apple silicon Macs

Bring your Metal app to Apple silicon Macs

MetalGraphicsApple Silicon
Watch original video

Highlight

Apple Silicon Mac brings the TBDR architecture of Apple GPUs to Mac, and Metal apps will need to instead use API-driven feature detection, explicit load/store actions, positional immutability, and threadgroup synchronization to avoid the performance cost of the Rosetta-compatible path.

Core Content

Many Mac graphics applications originally only worked with Intel, AMD, or Nvidia GPUs. They are used to the behavior of the Immediate Mode Renderer (immediate mode renderer): depth, color, and stencil buffers are frequently read and written in the video memory, and many incorrect usages are not immediately exposed on some drivers.

Apple Silicon Macs are replaced with Apple-developed GPUs. This GPU uses Tile-Based Deferred Rendering (TBDR). It first divides the geometry into screen tiles, and then completes depth testing, mixing and some color operations in the on-chip tile memory. A well-written renderer will use a lot less memory bandwidth, benefiting both power consumption and performance.

The migration path is divided into two steps. Old x86 applications can be run through Rosetta first, and the system will open some Metal compatibility workarounds to make the picture as correct as possible. Next, developers need to recompile the native Apple Silicon version and fix Metal code that relies on undefined behavior. Compatibility workarounds have a performance cost, and native versions cannot use them as a long-term solution.

The value of this session is that it does not stop at “recompile”. It directly points to four types of real issues: inability to judge by GPU name, incorrect load/store actions, lack of positional invariance, incorrect threadgroup memory synchronization assumptions, and sampling the current depth/stencil attachment in the same render pass. These problems will turn into picture errors or wasted bandwidth on the TBDR architecture.

Detailed Content

Use Metal API to detect GPU capabilities

(14:23) Apple explicitly requires apps to query the Metal GPU feature directly. Do not use operating system names or GPU names to infer capabilities. Apple Silicon Mac supports both Mac 2 family and Apple GPU family. The past judgment that “there are no Apple GPU features on macOS” has become invalid.

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

    self.simdgroupSize = computePipeline.threadExecutionWidth

    self.isLowPower = metalDevice.isLowPower
}

Key points:

  • supportsFamily(.apple5)Directly query Apple GPU family support to determine whether Apple GPU capabilities such as programmable blending, tile shaders, and local image blocks are available. -threadExecutionWidthGet the SIMD group width from the compute pipeline, which is needed for the subsequent threadgroup synchronization logic. -isLowPowerCode path used to distinguish integrated or discrete GPU style; session is specifically stated, on Apple GPU it returnsfalse, should be viewed in terms of the performance characteristics of a discrete GPU.

Correctly choose load/store action

(15:23) The Load/store action controls how color, depth, and stencil attachments enter or leave on-chip tile memory at the beginning and end of the render pass. The TBDR GPU is sensitive to this choice because it determines whether to load old content from system memory and whether to write tile memory back to system memory.

The judgment method is very straightforward: at the beginning of the render pass, select load action = load only if the previous content needs to be preserved; otherwise, avoid unnecessary load. At the end of the render pass, select store action = store only if the content will be consumed in a later pass; otherwise avoid store.

Key points:

  • An example of session is to render a texture in the previous pass, and then continue to draw on the old content in the last pass; if you useDontCareload action, Apple GPU will not upload the texture in the system memory to tile memory.
  • Uninitialized tile memory can leave graphics errors when subsequent draws do not cover the entire framebuffer; load should be selected when old content needs to be preserved.
  • The judgment of store action also comes from the resource flow: only when the attachment will be consumed by the subsequent render pass, the tile memory will be flushed back to the system memory.
  • Attachments that do not require subsequent reading avoid the store, which can reduce additional memory traffic.

When you need to accurately reuse positions across passes, explicitly turn on position invariance

(20:58) A common practice for multi-pass algorithms is to write depth in the first pass and useequaldepth test draws the same batch of geometry again. Even if two vertex shaders call the same position calculation function, compiler optimization may cause slight differences in position output. Apple GPUs do not guarantee position invariance by default.

// 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;
};

Key points:

  • MTLCompileOptions()Create Metal compilation options. -options.preserveInvariance = trueCompilers are required to preserve positional invariance. -makeLibrary(source:options:)Build a shader library using a compile configuration with this option. -[[position, invariant]]Marked on the vertex output position, tells the compiler that this position needs to be consistent across shaders.
  • This ability has a performance cost and should only be used for shaders that rely on the same depth value, e.g. usingequaldepth compare pass.

Threadgroup memory must be synchronized according to SIMD group size

(24:25) Threadgroup memory in Compute shader is shared by threads in the same threadgroup. Different GPUs have different SIMD group sizes. The Apple GPU is 32; if the application assumes that 64 threads naturally execute within a SIMD group, synchronization across SIMD groups may be missed.

// 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];
}

Key points:

  • thread_index_in_threadgroupGives the index of the current thread within the threadgroup. -threads_per_simdgroupGet the SIMD group size of the current GPU in the shader. -threadgroup uint buf[64]Is threadgroup memory shared by 64 threads. -buf[tid] = initBuffer(tid)Write to shared memory and other threads will read it later. -simdgroup_barrierOnly synchronize threads within the same SIMD group. -threadgroup_barrierSynchronize the entire threadgroup, suitable for situations where a threadgroup contains multiple SIMD groups. -res[tid] = buf[tid] + buf[index]Reads a value written by another thread, so the preceding barrier is necessary for correctness.

Do not sample the current depth/stencil attachment in the same render pass

(26:11) The session finally pointed out a depth/stencil binding error: when the same texture is used as the attachment of the current render pass, it cannot be used as a sampling texture at the same time. This creates concurrent reads and writes to the same underlying texture.

The fix is ​​to remove the current depth/stencil attachment from the sampled texture binding. When you really need to read the contents of the current attachment, create a second copy before the render pass, and then let the shader read this copy.

Key points:

  • The error described by session is that the current depth attachment is sampled by the fragment shader at the same time, causing concurrent reading and writing of the same underlying texture.
  • For opaque geometry, the depth will be accumulated before the fragment shader is executed; at the end of the render pass, the Apple GPU will flush the on-chip depth/stencil memory back to the system memory. At this time, concurrent reading and writing will trigger correctness issues.
  • This race condition can occur in any draw in the render pass, not just in the last draw.
  • Don’t use texture barriers or memory barriers to get around this problem; the session explicitly states that such workarounds are expensive on TBDR architectures.
  • If the algorithm really needs to sample the current attachment, create a second copy first, and then let the shader read this copy.

Core Takeaways

1. Add an Apple GPU capability layer to the renderer

  • What: Centralized logging on startupsupportsFamily、threadExecutionWidthandisLowPower.
  • Why it’s worth doing: The same set of Mac code needs to face Intel/AMD/Nvidia and Apple GPUs at the same time, and name judgment will quickly become invalid.
  • How ​​to start: EncapsulationRendererCapabilities, put programmable blending, tile shader, simdgroup size, GPU performance category into it, and the rendering path only reads this object.

2. Do a render pass load/store audit

  • What to do: Check the load/store actions of color, depth, and stencil attachment one by one.
  • Why is it worth doing: TBDR architecture puts attachment into tile memory, which is wrong.dontCarewill cause screen errors, redundant.load/.storeBandwidth will be wasted.
  • How ​​to start: Draw a resource flow table according to pass, marking which attachments are consumed by subsequent passes; only these attachments are used.store。

3. Mark position invariance for multiple pass depth equal processes

  • What to do: Find out how to write depth in the first pass and use it in the second passequalComparative rendering processes.
  • Why it’s worth doing: The optimization differences of different vertex shaders will make the depth values ​​​​not exactly the same, and the result is that some pixels are incorrectly discarded.
  • How ​​to start: Only open for related librariespreserveInvariance, and add it to the vertex output that requires exact matching[[position, invariant]]。

4. Establish SIMD group version strategy for compute shader

  • What: Prepare an implementation of selection by SIMD group size for kernels that rely on threadgroup memory.
  • Why it’s worth doing: The SIMD group of Apple GPU is 32. If the old code assumes 64, it will be missed.threadgroup_barrier.
  • How ​​to start: Use firstthreads_per_simdgroupCorrect the synchronization, and then split the high-frequency kernel into two versions for 32 and other widths. Press when runningthreadExecutionWidthchoose.

5. Change depth/stencil sampling to dual texture process

  • What to do: The current pass writes depth/stencil, and the shader only reads the previous depth/stencil copy.
  • Why it’s worth doing: Sampling the current attachment within the same render pass is undefined behavior, Apple’s compatible snapshot mechanism only works with old SDK paths and has a performance cost.
  • How ​​to start: Create a sampleable copy before pass for effects that need to read depth. It is forbidden to bind the current depth/stencil attachment as a sampled texture at the same time.

Comments

GitHub Issues · utterances