Ocarina of Time: state on main (2026-11-08)

[seedds]

The Legend of Zelda: Ocarina of Time (0400486024DEC000, the Switch 2 release, update v131072): where it stands on main as of 04419bc (2026-11-08).

Status

• Playable on the Mac. No iPhone build yet: its memory use makes an iPhone unlikely as it stands.
• The latest push changes the shader chain (89f44f9b6bb0): rebuild the game's shader libraries (run the NSP script again) to get the fixes below.

What's new in 04419bc

• Full resolution by default. The game renders its 3D view at 1536x864. A new launcher setting, Dynamic resolution (DM_DYNAMIC_RESOLUTION=1), hands the choice back to the game's own controller. That runs at 60 fps on more Macs, at a lower resolution that follows the load.

• The resolution no longer stays at its lowest after a loading screen. Two host bugs made the game's frame timer read more GPU time than the frame took:

  • a frame-start timestamp that read a stale value
  • event waits that left the GPU idle about 2 ms a frame

  With them fixed, the game's own dynamic resolution climbs back to full size.

• Less GPU work per frame. The mesh shaders' 128-bit memory loads are now one load instead of four, and occlusion-culled draws are skipped once their query result has landed.

Measured on an M4 Pro (20-core GPU), Kokiri Forest:

	Resolution	fps	GPU per frame
Before	stuck at 683x384	30–52	16–17 ms
Now, default	1536x864	~48	15.0 ms
Now, Dynamic resolution on	follows load	60	12.6 ms

At full size the GPU needs almost a whole 60 Hz frame on that Mac, so it often holds 30 for stretches. On slower GPUs, turn Dynamic resolution on.

Fixed earlier

• The title movie's network is about 25% cheaper (14.8 → 11.4 ms on an M6).
• Dotted horizontal rows in fog and light (Ocarina of Time (0400486024DEC000 v131072, Mac): dotted horizontal lines across the picture in Kokiri Forest fog and light shafts #1537).
• The pale outline around Link and lighting flashes: event waits between GPU queues are honoured.
• Blue holes in the ground and wrong terrain materials.

Known open reports

• GPU hangs / "GPU access revoked": Ocarina of Time (v131072, Mac, M2 16 GB): repeatable GPU hang at the Saria/ocarina cutscene #1690 (Saria's ocarina cutscene), Title: Ocarina of Time Switch 2 — Repeated GPU hang on M1 Pro 16GB (macOS 27.2) #1736, [0400486024DEC000] Ocarina of Time (macOS): GPU Hang -> 'GPU access revoked', game quits #1788
• Crash entering a new area (Deku Tree, Hyrule Field): Game crashes when entering a new area, such as the Deku Tree, or Hyrule Field. And when you load the save after the crash it a black screen. (Ocarina of Time) #1692
• Music: none in gameplay (Ocarina of Time (0400486024DEC000 v131072, Mac): no background music in gameplay or pause menu, sound effects play #1725); stops when a Pro Controller connects (Ocarina of Time [0400486024DEC000]: BGM stops when Pro Controller connects #1719)
• Playing the ocarina: Ocarina of Time (0400486024DEC000 v131072, Mac): host crash sig=11 (sdk+0x1f7050, addr 0x340) when playing the ocarina at the closed Door of Time #1525 (Door of Time), Ocarina of Time (0400486024DEC000): playing a note on the Ocarina aborts with 0xe15 — nn::hid::IsOunceVibrationDeviceMounted (SDK 22) is not answered by the host; tested fix #1433
• Picture: rainbow speckles on M1 (Ocarina of Time (0400486024DEC000): scattered rainbow speckle artifacts on M1, survive DM_UPSCALER=metalfx #1568), dotted lines on an M5 MacBook Air (Ocarina of Time: dotted line artifacts across the scene on M5 MacBook Air, with either upscaler #1626), less glow and haze than the console (Ocarina of Time: fire has little glow and the air no haze compared with console video #1499)

Ideas wanted: performance

At full resolution the GPU needs about 15 ms a frame on an M4 Pro, just short of 60 fps, and slower Macs fall well below. Suggestions are welcome: Metal techniques, profiling approaches, or things you've seen work in other translation layers.

Where the time goes (M4 Pro, Kokiri Forest, Metal System Trace; measured at 683x384 before the latest fixes, so the shares shift at full size):

• Geometry, about 9 ms: vertex and mesh stages. The foliage task/mesh pass alone took 4.6 ms of vertex-stage time against 0.2 ms of fragment. The game's task/mesh shaders are translated to Metal mesh shaders; the 128-bit load fix took about 3 ms off this.
• Compute, about 5 ms: mostly the game's NDSR upscaler, which writes a 2048x1152 target whatever the internal resolution, and a culling group.
• Fragment, about 4 ms.

Already tried, no gain:

• making the mesh shaders' threadgroup memory 50–77% smaller, and dropping its zero fill
• the pool-table reads and the mesh shaders' primitive output (about 0.3 ms)
• DM_NVN_TBDR=1
• a separate Metal queue for presentation (it removed 6 ms of idle per frame but made pacing worse)
• the GPU clock: it holds about 1,560 MHz at 94% busy during play

Open questions:

1. Mesh shaders on Apple GPUs: what makes a translated task/mesh pipeline expensive in the vertex stage when it does little fragment work? Threadgroup size, payload size, how many primitives each meshlet emits? Would converting those passes to compute plus indirect draws be faster?
2. The upscaler is a fixed cost at any internal resolution. Any experience with large full-screen compute dispatches on Apple GPUs, or with running such an upscaler at a lower output size?
3. Anything in a Metal System Trace or Xcode GPU capture of this game that stands out to you. A profile from another Mac (M1, M2, M3, M5) would help too: the bottleneck may not be the same everywhere.

If you want to dig in: scripts/bench.py runs a fixed route (kokiri-deck in the profile's bench.tsv). DM_NVN_GPUPROF=<first>-<last> logs every encoder's GPU start and duration over those frames.

How to help

If you play on the latest build, reply with:

• your Mac model and memory
• the frame rate and whether Dynamic resolution was on
• for a crash or hang: the log.txt and, if you can, the save from just before it

[aksmith-git]

Measured answers from our Ocarina of Time work. Two machines: Mac = M4 Pro (20-core GPU), iPad = iPad Pro M5. All runs Docked 1536x864, Kokiri Forest.

Results

Change	Device	Before → after
Grass task/mesh draws → compute pre-pass + indexed draws	Mac	grass 3.5 → 1.6 ms; frame GPU p50 19.2 → 16.7 ms
Float atomic add CAS loops → simd_sum + one native float atomic per SIMD group (6 compute shaders)	Mac	1bba… 361 → 9 µs, 6e28… 289 → 70 µs; frame-start compute block −0.69 ms
NDSR upscaler: B operand kept in registers + 2x2 gather merge	Mac	864 → 678 µs a dispatch, output bit-identical
#1752 applied to every shader (full rebuild)	iPad	GPU 19.3 → 16.6 ms
Completion writes moved off the woken SyncWait (SEG_POST + SEG_RB)	iPad	frame −0.7 ms
Translate at submit, GPU waits on a shared event	Mac	frame p50 19.37 → 18.38 ms

Mesh shaders: why the vertex stage is expensive

• Stage split: the grass pass is 89% vertex phase (3.4 of 3.9 ms on the Mac); fragment is 0.4–0.6 ms.
• Cause is occupancy, not ALU. Each translated mesh threadgroup needs ~13 KB of threadgroup memory for its stage array, plus up to 20 KB of mesh output. Apple's mesh stage then runs about 1 threadgroup per core. ~1,300 mesh threadgroups × ~75 µs ÷ 20 cores ≈ 4.9 ms, which matches what we measured.
• Answer to "compute + indirect draws?": yes. The stage array becomes a per-threadgroup slice of a device buffer, so the kernel uses no threadgroup memory and the GPU fills its cores. A pass-through vertex function reads the vertices straight from that slice. Task shaders append (task, mesh) entries to a list under an atomic counter, and a small kernel turns the count into the indirect dispatch and draw arguments. All draws in the pass share one compute encoder, so their threadgroups overlap.
  • Mac: vertex phase 3.1 → 0.23 ms; the grass costs 1.6 ms in total instead of 3.5.
  • Same kernel with the stage array still in threadgroup memory: 3.4 ms of compute. So the device-memory staging is the win, not the move to compute.
  • One dispatch per draw (11 encoders, no overlap): 2.2 ms. Worse.
• Host-side switches that changed nothing (Mac, A/B):
  • direct draws instead of indirect (MESH_DIRECT): no pass split, same frame time
  • pipeline hints (MESH_PSO: execution-width multiple, grid cap): 4.0 vs 4.1 µs per task threadgroup
  • no stage-array zero fill; setting only the referenced vertices
  • turf density 150 → 100 / 50: threadgroup counts do not move

Pitfalls if you build the pre-pass

• Depth EQUAL. The colour pass draws the grass with depth EQUAL and no depth write, so the pre-pass positions must be bit-identical to the depth pass's. Invariant mesh outputs compile without contract/reassoc; the kernel copy compiled with fused multiply-adds, so some positions came out 1 ulp off and those blades rendered white. Fix: #pragma clang fp contract(off) reassociate(off) on the kernel sources. Result: 0 differing words in 36 captures, and the picture matches.
• iOS relaxed math. On iOS the same kernel still mismatches the invariant depth on nearly every blade, even with the pragma. Workaround under test: LEQUAL with a small constant depth bias for those draws. On the Mac it matches the mesh path. Not yet confirmed on the device.
• iOS link failure. A vertex pipeline with a pass-through vertex function will not link if the vertex outputs are narrower (float, float3) than the fragment inputs (float4). Mesh pipelines accept the mismatch. Fix: widen every user output to float4.

Upscaler (fixed cost)

• NDSR "vr light", grid 32x36, 128 threads: 177 HMMA.16816 per thread, emulated with simdgroup_matrix and operands staged through threadgroup memory. ~550 µs of the 842 µs is the HMMA emulation, of which ~135 µs each is the A and B round trips through threadgroup memory.
• B operands repeated by later calls stay in registers (fa30: 118 of 176 calls reuse one). Four offset gathers that form a 2x2 footprint become one gather plus a swizzle. Result: 864 → 678 µs (−22%), bit-identical. Across the trimmed library 56 of 62 HMMA shaders change; the movie networks gain 7–27% too.

Compute: float atomics

• NVN float atomic add is translated as a per-thread compare-and-swap loop. 256 threads on one word serialise at the L2. Where the old value is unused, simd_sum followed by one atomic_fetch_add_explicit on atomic_float per SIMD group gives 32x fewer atomics and no retry loop. Pixel shaders are excluded (helper lanes).

iPad specifics

• Metal System Trace, per frame: Vertex 11.0 / Compute 5.8 / Fragment 4.4 ms. On the M5 geometry is 2.5x the pixel work.
• GPU performance state was "Medium" ~70% of the time (warm, charging). Absolute iPad numbers are inflated; relative gains hold.
• CPU side: translating a frame takes 9.5 ms on the iPad vs 2.8 ms on the Mac. The translating thread averages ~1.36 GHz and 58% of the process's CPU runs on efficiency cores. A time-constraint policy did not raise that.
• The woken SyncWait was running the completion writes (readbacks + ~1,800 reports): 1.5 ms per frame on the critical path. Moving them to the segments' feedback handlers: frame −0.7 ms.
• Translating at submit with the GPU waiting on a shared event (safe here: the CPU writes between submit and signal touch no bytes the translation reads): frame −1.0 ms on the Mac. iPad run pending.

Profiling that worked

• Metal System Trace channel totals (Vertex / Compute / Fragment) per frame, with every encoder labelled so each interval maps to a pass or shader key.
• Per-dispatch GPU timing with the shader key attached, ranked across two runs.
• A standalone harness that runs one compute shader alone, plus stubbed variants (HMMA off, gathers merged, stores removed) to price each part.
• In-frame A/B (switch flipped every N frames) instead of separate runs: the scene changes more than most of these effects.
• An empty-fragment mesh pipeline as a diagnostic: about 40% of the mesh stage was vertex output and 60% the body.

Branches

Each change is a commit on our tree (zelda-grass pre-pass, zelda-v20 atomics rule, zelda-v21 HMMA reuse, zelda-v19p pacing). Happy to push them to my fork or split any into a PR. Ask here.

[cover93-95]

@seedds Thanks. I'll post iPhone numbers from my build, measured with DM_NVN_GPUPROF=every:150 as you asked. The device is an iPhone Duo (iPhone19,4, iOS 27.2).

The build has d00defb in it, together with caba365 and 519f082 from the same branch. The shader chain is 2f63616dff08, and the shader library is still compiling tonight.

Before, on the same device: docked, 30 fps, the game's own NDSR, main pass pinned at 2048x1152, inside the Deku Tree. The scene colour passes on the main target take about 12.4 ms of about 25 ms of GPU time at full clock, spread over about 18 passes a frame. That 12.4 ms is the total for all of them, not one pass.

Because the other two commits are in the same build, the after numbers won't isolate the plain load on its own. I'll give the per-pass times at the same spot and say clearly that they come from the combined build.

[seedds]

NDSR gather merge is on main (0c8a3ac, chain eb68aba755a5; credited to @aksmith-git)

When the four offsets of a textureGatherOffsets call form a 2x2 square, the translator now emits one gather and a swizzle instead of four gathers. The result is exact: a GPU test compares it with the four gathers bit for bit, clamped and repeating. This changes the shader chain, so rebuild the game's shaders (run the NSP script again) to get it.

In Ocarina of Time it changes 27 compute shaders, the upscaler's networks. On an M6, standing on the Kokiri Forest deck, the compute encoder that holds the network went from 959–961 µs to 935.5 µs (two runs each in one session). That's too small to show in the frame's total GPU time. Your M4 Pro figure was 845 → 779 µs.

The HMMA B-matrix reuse is not in it. Reusing B only when a later call matches removes part of the scratch round trip. I'd rather look at translating HMMA so B never goes through scratch at all.

[cover93-95]

iPhone result for the build with d00defb, caba365 and 519f082: a GPU hang that ended in a kernel panic

These are the results I said I'd post above. The timing comparison has to wait, because this build hangs the GPU on a specific in-game action.

Device: iPhone Duo (iPhone19,4, T8160), iOS 27.2.1.

Build: my v5.2d tree (chain b70f92017083), plus:

• upstream 09dcb62 (the mesh object function) and e4e5ff6 (-fwrapv);
• 519f082 (shader_redsum on), caba365 and d00defb from pub/oot-ipad-m5.

The new chain is 2f63616dff08, with all 42,848 shaders built. Settings: Metal 4 pipelines on, DM_GUEST_FILE=1, encode thread off, docked, 30 fps.

What happened

• It's 52 s into a session in Kokiri Forest. Up to then the session was clean: 30 fps, 0 dropped frames between frames 600 and 1200. The thermal state went from nominal to fair at 34 s.
• Link kills a Deku Baba and picks up the Deku Stick it drops. Frame 1351 creates one new pipeline (93.5 ms). No frame is presented after it.
• The process keeps running for at least 20 s after that: the guest-file and audio lines keep coming every 10 s, and a pipeline archive is written. There is no [RED] line and no command-buffer error.
• When the app went to the background, graphics suspended: in-flight GPU work drained never appeared.
• Relaunching gave a black screen. 65 s later the phone rebooted with this panic:
  wait for termination failed (ret = 0xe00002d6) timeout 65s @AppleCentauriManager.cpp:1190
• In the stackshot, the old DmShell (3.1 GB resident) is terminating, with every thread in TH_UNINT. The only thread with a kernel stack is NVN queue replay, in TH_WAIT|TH_UNINT with TH_SFLAG_ABORT. So the submission thread is stuck in the GPU driver, waiting on work that never completes, and that blocks the process from exiting.

The v5.2d build has never done this on the same device with the same settings, over sessions of 15,000 to 20,000 frames. I'm repeating the same Deku Stick pickup on v5.2d now and will post the result.

Suspects

All of them are baked into the shaders, so none can be switched off at run time:

• redsum: 6 compute shaders, SIMD-group sums in the reduction loops;
• the mesh object-function launch (09dcb62);
• caba365;
• d00defb.

My next test is a rebuild with only DMSH_REDSUM=0, unless one of you has a better guess. I'm happy to run a specific build or environment variable on the device.

[cover93-95]

Follow-up: v5.2d does not hang on the same action

On the same iPhone with the same settings, I went back to v5.2d (chain b70f92017083) and did the same thing: kill a Deku Baba and pick up the Deku Stick in Kokiri Forest. There was no hang. The session ran to frame 12,000, and when the app went to the background, graphics suspended: in-flight GPU work drained appeared as usual.

So the hang comes with one of the changes in 2f63616dff08. I'm now building the same tree with only DMSH_REDSUM=0. The chain stays the same, and the Metal compile cache keeps it to about a quarter of a full build. I'll post whether the pickup still hangs without redsum.

[seedds]

HMMA B-matrix reuse is on main (013c875, chain caf9df51cf82; credited to @aksmith-git)

An HMMA whose B repeats now takes it from a slot of matrices kept in the SIMD-group, with no store, load or barrier for B. Two calls share a B when their B arguments come, through copies, from the same variables, each written once at the function's top level. The output is the same bits: a GPU test compares a network-shaped kernel with itself with every slot off. It changes the shader chain, so rebuild the game's shaders to get it. In Ocarina of Time, 10,293 of 16,620 HMMAs in 56 of 71 shaders take B from a slot; fa3033d05c0a41d6 takes it in 118 of 176, the same counts you found.

Why reuse rather than a different lowering: I tried building the matrices straight from the lanes' registers, through thread_elements() and SIMD shuffles. The M6's 8x8 lane layout is fixed and matches the one published for the M1. It gives the same bits, but it's slower. On a latency-bound chain:

HMMA body	Time
B through the scratch block	1.448 ms
B built from registers	1.716 ms
Every operand built from registers	8.9 ms

So B's trip can't get cheaper per call; it can only be skipped where B repeats.

Measured on an M6:

• Register pressure: chains with 10 to 160 weight pairs ran 16–28 % faster, even at Ocarina's largest slot counts (up to 168 slots).
• In the game, standing on the Kokiri Forest deck: the compute encoder holding the upscaler's network went from 984 / 979 µs to 951 / 946 µs (two runs each).

Your M4 Pro figure was 779 → 655 µs. An M4 or iPhone measurement on a build from current main would be very welcome.

One more finding: the exact Ampere rounding (356a7132) is now the larger cost. An HMMA chain takes about twice as long as with f16 accumulation: 1.448 against 0.734 ms, which gives different bits.

[sleepyeyedjoe]

Two things from an M4 Pro MacBook Pro (Mac16,8, 24 GB, macOS 27.0.1) that may affect how the fps figures read:

1. ProMotion. In fullscreen on the built-in screen, my sessions had long stretches at about 56 fps: every 11th frame took 25 ms while the CPU was near idle and GPU p50 was 9–12 ms. That's the display-link shift in Present on the Mac: MetalFX scales at the snapshot, and pacing leaves the display link when the screen shifts under it #1045. DM_PRESENT_LINK=0, or the display fixed at 60 Hz, removes it. If your M4 Pro numbers come from a MacBook Pro in fullscreen, part of the gap to 60 may be this rather than the GPU. Data in Present on the Mac: MetalFX scales at the snapshot, and pacing leaves the display link when the screen shifts under it #1045.
2. Holds at 30. On my build (main 1794a35), the 30 fps stretches in real play came with "CPU 98–99 % on performance cores" from the frame step. ea0c98d's submit+present was 6.7–9.0 ms a frame against 0.4–0.6 ms with Honor NVN event dependencies across guest queues #1245 in the same place. Numbers in Honor NVN event dependencies across guest queues #1245. I haven't measured 04419bc yet.

I'll post fps with Dynamic resolution on and off once I've rebuilt on 04419bc.

[kartur]

your Mac model and memory:
Base Mac Mini M4 with 16GB

the frame rate and whether Dynamic resolution was on:
With a locked 30fps is is mostly constant with occasional slowdown and shader compilation stutters
With dynamic resolution (Upscaling standard / games own) it is mostly in the 30s. See video

c78d7375bea9c862679ec294a6a69ff3e22e18d0.mp4

//EDIT: Actually have to correct myself. With these settings after playing more than a minute the game starts eventually stuttering badly. If you pause the game and then contiue the FPS recover and the same scene runs smoothly again. Perhaps some memory issue.

However, if I use

• 30FPS + DISABLE dynamic resolution + Upscaling: standaring + Game's upscaler: MetalFX

the game runs pretty great. Very good frame pacing and no longer stutters or slowdown even after playing for a few minutes. Example:

c78d7375bea9c862679ec294a6a69ff3e22e18d0_30fps_nodyn.mp4

[LeonardoBai12]

Mac: MacBook Pro 14", M1 Pro, 16 GB, macOS 26.6.1. Fullscreen on the built-in ProMotion screen (the host switched it to 120 Hz).

Build: main 04419bc plus #1331, #1453, #1377, #1379 and an SDK guard; chain 89f44f9b6bb0, all 42,848 shaders. v131072, Handheld, DM_UPSCALER=game.

Settings: 60 fps target, hold frames on, Dynamic resolution off (full 1536x864), Metal 4 pipelines off.

Frame rate: about 30 fps most of the time, smooth to play.

• The host's per-window fps, over two sessions: p10 21–23, p50 31, p90 31–60.
• Probes up to 60 fail at full resolution on this GPU, so it mostly holds 2 vsyncs.
• Game time ran at 0.99–1.00 of real time.

Two sessions, 111 min in all, both ended by quitting:

Session	Length	Where	GPU hangs
1	73 min	Kakariko, the graveyard, the mask quest, Castle Town, Hyrule Field, Death Mountain	22 failed submissions, all recovered
2	38 min	Death Mountain, Goron City, Lost Woods	1, recovered

• Hangs: the earlier builds lost the GPU (GPU access revoked) after 7–44 min, and these two sessions never did. In session 1 the ground in Hyrule Field and Kakariko went missing for moments after about 30 min, which I think were the failed buffers. Details are in Ocarina of Time on an M1 (16 GB): GPU hang (kIOGPUCommandBufferCallbackErrorHang) under memory pressure, about 24 fps at 720p / 30 fps #1605.
• Pipelines: with Metal 4 pipelines off, pipeline prepare takes 5.6 s; it took 146 s on the first start after the chain change. With them on, it took about 6 minutes on every start.
• Other issues:
  • In session 1, a broken-looking piece of geometry floated in the air on Death Mountain (passable). It wasn't there in session 2, on the same mountain.
  • Apple's compiler crashes on one task pipeline (task a7f22b4a795ff4d1) in every session, with no visible effect.
  • The grain dots of Ocarina of Time (0400486024DEC000 v131072, Mac M4 Pro): docked, the 3D picture outdoors turns into flat patches; handheld is fine, MetalFX and HDR change nothing #1653 are hidden here with DM_NVN_SKIP_FS=5ad65c0eefd09c7f.

Dynamic resolution on (a third session, same build): no gain here.

• Per-window fps p10 21, p50 31, p90 32, the same as at full size, and the picture was visibly softer. Probes to 60 failed as before (146 probes, 147 failed trials).
• That fits the geometry share in the breakdown above: on this M1 Pro the frame seems bound by the task/mesh and vertex work, not by the pixel count. This is only a guess; I haven't profiled it.
• The session ended in GPU access revoked (exit 70) at 1,444 s. Three earlier episodes (339, 1,066 and 1,270 s) recovered. One session isn't enough to tie that to the setting.

I can send a DM_NVN_GPUPROF window from the kokiri-deck route if an M1 profile helps.

[clong119-debug]

plus 1 for m1 pro 16gb !

[seedds]

Update: the flickering dark specks in the mist are fixed (877c7ec)

The ground mist tests each pixel against a depth its own earlier pass stored for the same surface. On the Switch's NVIDIA GPU those two values are always equal. Apple's GPU sometimes stores the depth one rounding step lower, at 0.2–3.8% of pixels, so the mist rejected its own surface there. The fix is general, in the shader translator: a shader comparing a pixel's depth with a depth it read treats them as equal when they're within that gap. It touches 39 of Ocarina's shaders (the mist and some foliage).

The shaders rebuild on your next build (new chain). If you still see specks or anything new in the mist or grass, please post a screenshot and your Mac's chip.

[seedds]

Follow-up: the gap depends on the GPU generation. On an M4 Pro, 9,511 of 718,400 test pixels have their stored depth below the pixel's own depth. On an M6 none do: it stores exactly the depth the pixel shader reads, as the Switch does. So the specks showed on M4-generation Macs, never on an M6, where the fix changes nothing.

M1–M3 aren't measured yet. If you have one and a clone of the repository, you can run this one test (a few seconds, no game files needed):

make -C dm/gfx/tests GAME=0400486024DEC000 check_t_zsnap

Please post its last two lines (the Apple M…: gl_FragCoord.z above the stored depth … at N pixels line and t_zsnap: ok) with your chip.

[quartz-forge]

M1 Max (Apple M1 generation), macOS 27.2. Last two lines of
make -C dm/gfx/tests GAME=0400486024DEC000 check_t_zsnap:

Apple M1 Max: gl_FragCoord.z above the stored depth of its own surface at 9511 pixels in all (the rasterizer's own; 0 would mean this GPU stores the z its fragments read)
t_zsnap: ok

[Prince-Mendiratta]

M3 Max, macOS 27.0.

~/Documents/startup/nxnative on main
% make -C dm/gfx/tests GAME=0400486024DEC000 check_t_zsnap
clang -O1 -g -Wall -fobjc-arc -o t_zsnap t_zsnap.m -framework Metal -framework Foundation
2026-11-08 23:14:07.934 t_zsnap[83448:8398292] Metal API Validation Enabled
perspective: 143680 pixels, gl_FragCoord.z above the stored depth at 321 (0.22%), through dm_zsnap at 0
perspective, 4e-6 farther: z through dm_zsnap above the stored depth at 143680 of 143680 pixels
w = 1: 143680 pixels, gl_FragCoord.z above the stored depth at 522 (0.36%), through dm_zsnap at 0
w = 1, 4e-6 farther: z through dm_zsnap above the stored depth at 143680 of 143680 pixels
constant depth: 143680 pixels, gl_FragCoord.z above the stored depth at 2839 (1.98%), through dm_zsnap at 0
constant depth, 4e-6 farther: z through dm_zsnap above the stored depth at 143680 of 143680 pixels
depth < 0.001: 143680 pixels, gl_FragCoord.z above the stored depth at 404 (0.28%), through dm_zsnap at 0
depth < 0.001, 4e-6 farther: z through dm_zsnap above the stored depth at 143680 of 143680 pixels
depth > 0.999: 143680 pixels, gl_FragCoord.z above the stored depth at 5425 (3.78%), through dm_zsnap at 0
depth > 0.999, 4e-6 farther: z through dm_zsnap above the stored depth at 143680 of 143680 pixels
Apple M3 Max: gl_FragCoord.z above the stored depth of its own surface at 9511 pixels in all (the rasterizer's own; 0 would mean this GPU stores the z its fragments read)
t_zsnap: ok

[Amost-]

MacBook Pro M2 Pro (19-core GPU, 32 GB), macOS 27.2. Last two lines of make -C dm/gfx/tests GAME=0400486024DEC000 check_t_zsnap on main eddae74f:

Apple M2 Pro: gl_FragCoord.z above the stored depth of its own surface at 9511 pixels in all (the rasterizer's own; 0 would mean this GPU stores the z its fragments read)
t_zsnap: ok

So the M2 matches the M1 Max and the M4 Pro (9,511), and in the same run dm_zsnap brings it to 0.

[seedds]

Why dynamic resolution is off by default now

Since 04419bc the launcher's Dynamic resolution setting (DM_DYNAMIC_RESOLUTION) is off by default, and off means the game renders at full resolution (1536x864 main pass). Turning it on gives the game its own controller back.

The game's controller got stuck at its lowest resolution (683x384) on the Mac and never climbed, even with GPU time to spare. Two host bugs made its timer read high: a stale frame-start timestamp, and event waits that idled the GPU. Those are fixed, and it still stayed at the bottom. The rest is in the game's own cost model:

• Each frame it takes the measured GPU time, subtracts a small fixed cost, and assumes all the rest grows with the pixel count. It then extrapolates to each resolution level by the square of the scale.
• That fits the Switch 2 it was tuned for. On the Mac the frame hardly depends on resolution: on an M4 Pro, 683x384 costs 14.3 ms of GPU and 1536x864 costs 14.7 ms. Most of the time is mesh shading, compute and the upscaler, which writes a fixed 2048x1152 at any internal resolution.
• So at the lowest level the measured load was under its threshold, but the extrapolation put full resolution at 4-5 times over, and it never stepped up.

Since the low resolution saved only about 0.4 ms and looked much worse, full resolution is the better default. A controller that steps one level at a time on the measured load, instead of the extrapolation, is being worked on but isn't on main yet.

[Amost-]

MacBook Pro M2 Pro (19-core GPU, 32 GB), macOS 27.2: Kokiri Forest on main 028a2631 + my picks

Frame rate, on my own build. It is not plain main: main 028a2631 (with 04419bc2) plus local picks (#1331, #1368, #1176 HDR, #1377, #1379, DM_TEX_RELEASE (off for this game), a dispatch-overlap pick and a pool-table copy fix). Dynamic resolution off, docked. Measured headless on scripted 6-minute runs: a Kokiri Forest save loads about 99 s in, then a fixed walk. This is not bench.py's kokiri-deck.

• At the game's own 1536x864 (two runs, one of which hung from 119 s and lost GPU access): median frame time 33.3 ms (30 fps) in both, GPU p50 21–23 ms. The frame step held 2 vsyncs for 62–83 % of frames, 3 (20 fps) for 3–10 %, and 1 for the rest (14–28 %, mostly its failed probes up to 60). So at full resolution 60 fps is out of reach on this GPU (GPU p50 above 16.7 ms). I haven't measured with Dynamic resolution on.
• With a RomFS render-scale overlay of mine at 0.93 (1906x1072, above the native size): GPU p50 24–30 ms, the same 30 fps median, and 20 fps for 12–24 % of frames.
• Before 04419bc2 (main 1794a35a with the same picks and the overlay), GPU p50 was 26–40 ms (30–36 ms in most runs; that build also counted about 2 ms a frame of event-wait idle as GPU time), and the frame step mostly held 20 fps.

Hangs on this build: 1 of 6 full runs (no overlay, from 119 s, GPU access revoked), plus a 200 s run that hung at 193 s. Details on #1326: #1326 (comment)

[PaulEins]

MacBook Pro 14" (M4 Pro, 20-core GPU), 48 GB, macOS 26.7.2, Xcode 27.2

Build: main at 028a2631, Mac app via build_ipa_from_nsp.py --platform mac-app, full shader map (42,848 shaders, 39:37), pipeline archive 12:20, about one hour in total. Launcher: Docked, 60 fps, Upscaling Standard, game's own upscaler, Metal 4 pipelines and pipeline prebuild on, encode thread off, hold frames when late on. Display set to a fixed 60 Hz. macOS Game Mode off (the default).

Window (Composited), Dynamic resolution on: 60 fps, GPU 11.8–13.5 ms. The resolution drops on every short hitch, so the picture is low resolution most of the time, although the GPU has room to spare.

Fullscreen (Direct), Dynamic resolution off:

• Kokiri Forest, standing: 60.00 fps, GPU 13.72 ms
• Hyrule Field, walking with camera turns: 57.6 fps, GPU 16.06 ms, frame interval 17.36 ms, with short single spikes

Full resolution in fullscreen feels smooth and is the setup preferred here. No visual errors and no crashes in this session; saves load fine.

Screenshots below.

_{Written with AI: Claude Opus 5.5 · Responsibility: PaulEins}

Problem:

[ZhaowTao]

M5 MacBook Air (10-core GPU, 24 GB), a 45,300-frame run, zero hangs

Adding the first M5 laptop data point to the thread. It is not a full-resolution run: this build predates 04419bc2, so it still has the two host timer bugs and the game's own dynamic resolution is in charge.

Machine: MacBook Air Mac17,3, Apple M5, 10-core GPU (4 performance + 6 efficiency cores), 24 GB unified, macOS 26.6.2, built-in Liquid Retina 2560x1664, windowed.

Build: main 1794a35a (2026-11-08 17:01), --platform mac-app, shader chain 579d06bf3932. The shader library was built from a seen-shader list, 11,815 of 42,848 keys. I mention the consequences of that at the end.

Session: one continuous run, 45,300 frames, about 24 minutes of play after start-up. Route: Kokiri Forest → Lost Woods → Sacred Forest Meadow (beat the Wolfos there), then quit normally (the host exited with 0).

Frame timing

151 reports of 300 frames each:

	min	p25	median	p75	max
frame time p50 (ms)	16.66	16.67	16.67	16.74	33.33
GPU p50 (ms)	9.33	13.70	14.07	14.67	27.73

• 25 of the 151 reports (17%) sat at 33.3 ms, all of them with GPU p50 at 20–27 ms. Averaged over the session that is about 50 fps.
• Dropped frames: 496 of 45,300 (1.1%).
• The 30 fps stretches clustered in Hyrule Field and in the Sacred Forest Meadow transition; the Lost Woods and the Forest Meadow proper held 16.67 ms.

Memory

	value
footprint, median	12,087 MiB
footprint, max	12,507 MiB
at frame 45,000	12,412 MiB

Flat. It settles by about frame 18,000 and does not creep after that, which I hope is useful next to #1831 (a 16 GB M4 Air reaching a 15 GB GPU allocation on a full-resolution build). 12.2 GB resident on 24 GB was never under pressure here, so this build does not reproduce that climb.

Stability

• kIOGPUCommandBufferCallbackErrorHang: 0
• GPU access revoked: 0
• nn::diag abort: 0

That is one run, so it is not evidence that the hangs in #1326 / #1690 / #1830 are gone, only that a 45,300-frame session through the Lost Woods did not hit one.

Audio

audio #0 device callbacks reached 368,000, 240–241 frames a callback, 4.0 blocks queued on average. 183 of the 184 reports say 0 underruns in their 2,000; the one exception is at 196,000, with 17 underruns and 78.8 ms zero-padded. One audible gap in 24 minutes, and I have not looked at what the frame was doing at that moment.

Two notes

Resolution. I cannot state the main-view scale from this log, so please read the GPU figure with that in mind. The log shows a 1536x864 attachment sampled in a feedback pass and a scissored clear: first case 683x384 @0,0 (target 960x540 level 0), and WindowCrop 0,0 2048x1152 followed by a 2048x1152 → 1280x720 converting copy. Since this build is pre-04419bc2, the expectation is that it sits where the game's controller left it, as you describe for the M4 Pro before the fix. A median GPU p50 of 14.07 ms on a 10-core M5 is then close to the 15.0 ms you quote for an M4 Pro (20 cores) at 1536x864, which would put the two GPUs roughly level at the same setting — but that comparison needs a pass dump on this build, and I would rather flag it than assert it.

Start-up. Time to frame 58 was 586 s; frame 169, 772 s; frame 1031, 851 s. So about 14 minutes before the route starts, which is well past the ~5 minutes in #1747. On this app the shader library came from a partial build and the host says so twice at start:

[RED] shader library .../msl/b-579d06bf3932 has no stamp (inputs.tsv): no full dmsh_build.py run made it as it is, so c1_full.tsv and cb_extent.tsv are not taken
[RED] c1_full.tsv is stale or missing (...); every c1 gets its own 64 KiB

288 distinct keys then failed to load at runtime (failed to load shader <key>_b: cannot open .../<key>_b.metallib), 599 lines in all. Nothing on this route looked wrong with them missing, and the game kept running, but if the start-up figure is of interest I would want it measured on a full library before drawing a conclusion from it.

[zephleggett]

Hey, I’m edge inference expert and this project nerd sniped me.

In this PR:
#1837

I have re-implemented NVDE movie network in pure metal kernels. This takes the title screen to 36ms a frame to 10ms a frame on an M1 Pro.

The output of the reimplemented network is exact to the original network run on CPU.

[seedds]

Dynamic resolution update (7de11da) and a request for tests on other Macs

Thanks for the reports. @PaulEins, what you saw with Dynamic resolution on (low resolution most of the time with GPU to spare) was the game's own controller. It predicts each resolution's cost from the pixel count, which on a Mac overstates full resolution four to five times.

The launcher's Dynamic resolution setting is still off by default (full resolution). Turned on, the game's controller is now replaced by one that climbs on what it measures. It tries each higher resolution, keeps it while the measured GPU load fits the game's own budget, and drops a level only when the load stays over for 8 frames. On an M4 Pro it reaches 1536x864 within seconds and holds 60 fps.

One weakness is still open. On an M4 Pro the GPU time in heavier scenes hardly depends on resolution: at 683x384 the load was about the same as at full size. Even so, the controller still steps all the way down, and the picture goes soft for no gain. The next change will make it step down only to a resolution that measured cheaper. Until then, full resolution stays the default.

Could you test it? Slower Macs are what this is for. @Amost-, your M2 Pro (30 fps at full resolution, GPU 21–23 ms) is exactly the case. On a build from current main:

1. Turn on Performance HUD in the launcher's settings.
2. Play the same stretch of Kokiri Forest (walking around the village, turning the camera) for a minute or two, once with Dynamic resolution off and once on.
3. Post your Mac model and chip, and for each run roughly what frame rate the HUD showed and whether the picture went visibly soft.

A base M4 or an M1–M3 would help as well. If it brings a slower Mac from about 30–45 fps up to 60, it can become the default there.

[Amost-]

MacBook Pro M2 Pro (19-core GPU, 32 GB), macOS 27.2: Dynamic resolution off and on, main 7f502220

Built from main 7f502220, docked, without my RomFS render-scale overlay. The numbers come from the host's log rather than the HUD: headless runs of a scripted walk around Kokiri Forest's cave pond (the save loads about 99 s in, then a fixed walk with a camera turn), 6 minutes each, with warm pipeline caches, using a plain-main host on the app's files.

Dynamic resolution	fps	GPU p50 after the load
off (2 runs)	30 (frame time p50 33.3 ms)	22.2 / 22.2 ms
on (2 runs)	30 (33.3 ms)	18.6 / 20.0 ms

• It steps down when on. With it off the log shows only the full 1536x864. With it on it also shows 960x540, 763x429 and 683x384 render targets (sampled log lines, not a full trace).
• The picture goes visibly soft. At the same spot of the walk, the rock texture, the ferns and Link's tunic lose their fine detail.
• It doesn't reach 60 here. A frame still takes more than one vsync (about 1.5–1.9 when the frame step tries 60), so it stays at 30 either way. Frames held at 3 vsyncs (20 fps) came and went in both modes, 0–13 % per run, with no clear difference.
• It shows the weakness you described. It goes down to 683x384, about a fifth of the pixels of 1536x864, yet GPU time drops only 10–16 %.
• No GPU hang in any of the 7 runs: 5 on the plain-main host (3 off, 2 on) and 2 on my own app with local picks (1 off, 1 on). My app also stayed at 30 fps with it on.

[nawangwe]

Bright band on Deku Tree walls / pale patch at the Kokiri exit: cause found, fix in #1844

M4 Pro (20-core GPU), 24 GB, macOS 27.2; v131072, Docked, 1536x864 (tested on penis6's oot2 fork at 18cdc00, which compiles its shaders with the same flags as main).

Before (stock shaders):

In Gohma's room one section of the dark wall is lit brightly, and a pale patch sits on the ground at the Kokiri Forest exit; both shimmer while the camera moves and settle when it stops. Both are task/mesh-shader models; the vertex-shader models next to them are fine.

The mesh shaders decode packed 16-bit vertex data by hand, and mesh 35785cddc971290c (task a5288d0a6c4971fd, pixel f4f7e6955d37f40c) extracts the low half of a word with as_type<int>(w) * 65536 then >> 16. Signed overflow is undefined in MSL, and Metal's front end folds (x * 65536) >> 16 to x, so the whole word becomes the mantissa and the lightmap / probe UV's x comes out 65536x too large (in-shader 13184 where y gives 0.20). The probe lookup then lands on bright texels and the pixel shader's SH ambient term reads about 8 instead of 0.001. The extraction's shl is gone from the shader's IR without -fwrapv and present with it; the stock mesh shader rebuilt with -fwrapv and loaded as a loose library renders the wall uniformly dark (before / after below). 64 of the game's 42,848 translated shaders use the idiom.

#1844 adds -fwrapv to the Mac and iOS compiles (new chain). A full rebuild with it is running here now; I'll report on the Kokiri patch once it is done.

After (mesh shader rebuilt with -fwrapv):

[williamstonef]

iPad Pro 11" M5 (iPad17,1, 12 GB, iPadOS 27.2): measured data, Ocarina of Time v131072

Setup. Engine pin f25fa956 + NX2Apple patches (shader chain 1d619d0f520b), built with NX2Apple's script (--shell-framework) from the NUP. Kokiri Forest, from an autosave. Handheld, MetalFX for the present and for the game's upscaler, Frame buffering on, Metal 4 pipelines on, guest memory in files (DM_GUEST_FILE=1) unless noted. All figures come from the host log (timing, power, memory lines, 300-frame windows). No Instruments traces.

1. TBDR on this iPad (Encode thread off in both runs; fresh launch each, same save, Link standing still, the first 5,400 frames after the title discarded; 24 and 25 windows):

	DM_NVN_TBDR unset	DM_NVN_TBDR=1
fps	48.5	52.5
frame p50 / p95	19.5 / 26.8 ms	18.0 / 24.3 ms
GPU p50	18.4 ms	14.7 ms
NVN submit p50	10.8 ms	8.2 ms

A slightly different cut (fixed frames >= 5,400) gives 48.9 -> 51.8 fps and GPU 18.4 -> 15.6 ms, so about +3 to +4 fps. No crash. The README says TBDR did not help on the M4 Pro; here it did. One eyes-on observation, not measured: with TBDR on there is a water transparency glitch. In every water area I have looked at, the water is not transparent and the bottom is not visible. I did not compare screenshots.

2. Encode thread on iOS (#1664). 39,600 frames (~20 min) on f25fa956, no abort, with TBDR on: 53.1 fps, frame p50 17.2 / p95 20.5 ms, GPU 13.1 ms, NVN submit 5.0 ms (9.2 ms with the encode thread off and TBDR on, a different run). Thermal state reached 2 during that run.

3. Memory limit. Without DM_GUEST_FILE: jetsam ended the app at 8.04 GiB (lifetime max 8.15), reason vm-pageshortage, frontmost, at the title screen. The host's last resident line, 46 s earlier, read 3,775 MB, so that figure under-reads what iOS counts. An earlier 20-minute run (host resident 10.6 GB, engine 88c7e03) ended in a kernel panic memorystatus: all vm-pageshortage victims exhausted. With DM_GUEST_FILE=1: host resident 4.3 GB (13 min), 5.3 GB (20 min), and at most 7.0 GB after a 25-min prepare; no kill so far.

4. Cold vs warm caches (same app, same settings, same scene):

	warm	after the app container was emptied
fps	58.0	30.2
frame p95	20.0 ms	69.5 ms
GPU p50 / submit p50	13.2 / 5.9 ms	13.6 / 5.3 ms
task/mesh pipelines built in the background	68, 0.6 s in total	94, 239 s in total

GPU and submit time did not change; the lost frames are compile time. This fits #1368's description (mesh pipelines are not in the prewarm lists).

5. Time before the first frame, on iOS.

Prewarm lists	Time	Bases compiled
9,713 lines (1,192 render), warm archive	1.05 s	0
14,568 lines (3,331 render), empty archive	223.6 s	3,009
14,615 lines (3,377 render), full archive	251.3 s	0 (3,051 from the archive)
#1732 manifest, 40,145 lines (24,806 shaders, 11,100 render, 2,134 compute), first start	1,542.7 s	7,015, 2,105 MDB_MAP_FULL
same, second start (13-file archive)	1,526.0 s	7,863, 1,465 MDB_MAP_FULL

That reproduces #1735 on an iPad (engine without eed7c429). Prepare time did not follow the number of compiled bases (251 s with 0 compiled, 224 s with 3,009), so the cost seems to be loading shader libraries and building specializations; I did not test #1825 / #1840. On the second start the thermal state went to 1 at 1,404 s and to 2 at 2,479 s after launch.

6. Panics. Besides the one above, a GFX DATA ABORT ... far=0 - agx_scheduler(2) kernel panic ~28 s after pressing Load on an autosave (engine 88c7e03, before f25fa956). I have not seen it again on f25fa956 over several save loads, which proves nothing.

Not measured: resolution on this iPad (no controlled A/B), Dynamic resolution, the 1.3.0 engine, the grass pre-pass. I can run any A/B or DM_NVN_GPUPROF if useful.

Drafted with Claude; reviewed by me.

[padsiwono]

The 1536x864 main pass that main calls full resolution is the top of the game's own level list, and that list is data, so the game can render above it without a host change.

The file is Pack/Bootup.Nin_Ounce_NVN2.pack.zs → System/GameGraphics/ModelSceneExtensionParam/Default.game__gfx__ModelSceneExtensionParam.bgyml (BYML v7 with hash-keyed maps, the same file Tears of the Kingdom keeps its ResolutionParam in). In v131072:

• DynamicResolutionParam/LevelPaletteDict has two palettes: 10 levels 0.75, 0.711, 0.667, 0.622, 0.583, 0.533, 0.5, 0.444, 0.372, 0.333 (IncreaseCounterMax 30, LoadThreshold 95, ReduceCounterMax 2) and 8 levels 0.75 … 0.444 (LoadThreshold 108, ReduceCounterMax 3). Both start at 0.75, so the game never renders the base size, on the Switch 2 either.
• ResolutionParam has a 1920x1080 and two 2560x1440 entries (base = final).

With every level set to 1.0 in a copy of that file, loaded through DM_ROMFS_OVERLAY, the main pass renders at 2048x1152. On an M5 Max (macOS 26.6.2), headless on the same route, GPU p50 went from 7.51 ms at 0.75 to 8.68 ms at 1.0. That fits what you measured on the M4 Pro: on a Mac the frame hardly follows the resolution.

In a 13.6-minute play session at 2048x1152 (Kakariko, its graveyard and Death Mountain Trail, HDR on) it held 100 % in every window, GPU p50 12.2 ms, frame p95 19.1 ms (median of the 300-frame windows). I haven't separated GPU from CPU in the slower windows; the longest stalls were first-sight mesh pipeline builds.

Scaling the levels instead (×1.333, so the top is 1.0 and the rest keep their ratios) lets the game's controller move between 100 % and about 71 % under load (headless, with a larger output from our own render-size change), which might suit the climbing controller in 7de11da: as it stands its top is the game's 0.75.

This was on our own build (main of 7 November plus later fixes from main), not current main. The overlay is made with a small BYML editor adapted from our Tears of the Kingdom work; happy to share it.

[padsiwono]

Follow-up: the game also crops its frame to 80 %, and that part is code, not data. With the crop removed, the 3D renders at the full output size and nothing is scaled.

In v131072's main the crop is computed at main+0x6505b4–0x6505d8: the base size ([entry+0x64]) times a float pair per mode and level, 0.8 here, read from [obj+0xf0 + mode×0x208 + level×0x60 + (signed byte [obj+0x54a])×8 + 0x1e4 / +0x1e8]. Each axis is clamped to a size from main+0x112c5e0 and stored through the setter at main+0x1087f40; main+0x650538 reads it back for nvnWindowSetCrop (called from main+0x11e7b74). The 0.8 is not in ModelSceneExtensionParam. I found the writer by watching the crop object's size field: one store, 3024x1702 → 2419x1361, pc main+0x1087f64, lr main+0x650604.

Removing it takes two instructions: the two fmul become nop, main+0x6505c4 1e200840 → d503201f and main+0x6505d4 1e200820 → d503201f. nvnWindowSetCrop then gets 0,0 3024x1702 instead of 2419x1361, and the scene and its depth target are 3024x1702. With every level at 1.0 (the overlay above), the 3D is at the output size.

Measured on an M5 Max, macOS 26.6.2, with our host (main 270b6fd5 plus our Ocarina of Time changes, including an output-size change of our own, so 3024x1702 is not a size main offers). Castle Town market save, Link standing, HDR, headless, one run each, the 300 frames after the scene settled:

	crop, 3D 2419x1361	crop removed, 3D 3024x1702
frame p50 / p95	16.67 / 17.64 ms	19.47 / 21.29 ms
GPU p50 / p95	13.36 / 14.20 ms	17.34 / 18.37 ms

1.56 times the pixels cost 1.3 times the GPU here, and at this size the market no longer holds 60 on this Mac. Framing and field of view are the same, and the pause menu has the same layout, text and item art.

Not tested: main's own output (2560x1440 would render at 2560x1440 instead of 2048x1152), play beyond the pause menu (text boxes, HUD), other areas, handheld.

On main this would fit as a patch in the game's patches.py behind a launcher setting. Happy to open a PR if you want it.

Castle Town market, frame 2600, output 3024x1702. A: the crop, 3D at 2419x1361, scaled up; B: crop removed. The two runs are not in step, so the NPCs stand differently and their speech bubbles show only in A; text boxes are among the things not tested yet.

The same frames at 1:1 (800x450; A scaled up for the picture with Lanczos, since the capture is the 2419x1361 crop itself):

[williamstonef]

iPad Pro 11" M5 (12 GB): history of my Ocarina of Time sessions, in one table

Follow-up to my earlier comment. Same iPad, always Kokiri Forest, from autosaves of the same playthrough. Every row is one session, from the host log: timing / power / memory lines, 300-frame windows, counted from the frame the save has loaded. fps is frames per wall-clock second (the power line), so it includes the stalls. "Stalls" are frames of 1 s or more, almost all of them pipeline builds.

#	Engine	Setup	min	fps	frame p95 ms	GPU ms	submit ms	mem max GB	stalls >= 1 s
1	1.2.3 (88c7e03)	Docked, Encode on, Metal 4 off	9	35.4	30.3	27.7	0.3	4.8	3 (48 s)
2	1.2.3 (88c7e03)	Handheld, Metal 4 on, Encode off	8	38.1	32.8	23.8	5.6	7.3	16 (50 s)
3	1.2.3 (88c7e03)	same, 20 min; ended in a kernel panic (memory)	20	39.9	34.5	19.7	12.4	10.6	45 (155 s)
4	1.2.3 (88c7e03)	loading a save from the title; jetsam (proc-thrashing) at 8.26 GiB	8	32.8	34.6	28.4	1.4	9.5	1 (2 s)
5	1.2.3 (88c7e03)	DM_GUEST_FILE=1	7	26.6	59.0	20.2	16.9	5.4	20 (76 s)
6	f25fa95	new engine, Encode off, TBDR off	5	49.2	26.7	18.3	10.7	4.6	2 (4 s)
7	f25fa95	DM_NVN_TBDR=1	5	51.3	25.4	16.1	9.8	4.8	2 (6 s)
8	f25fa95	Encode on + TBDR, 13 min	13	53.2	20.7	13.4	5.1	5.3	29 (87 s)
9	f25fa95 + #1543 + #1497	render scale 0.5, Encode on, TBDR off, warm caches (includes Hyrule Field)	16	57.4	20.0	12.6	5.3	5.1	18 (127 s)
10	f25fa95 + #1543 + #1497	same settings, caches emptied	5	30.2	69.5	13.6	5.3	5.0	6 (28 s)
11	f25fa95 + #1543 + #1497	TBDR + #1732 manifest (40,145 lines), render scale 0.5	20	53.1	24.2	14.8	8.0	7.0	27 (195 s)

What can and cannot be read from it

• Only rows 6 and 7 are a controlled A/B (same build, same save, fresh launch each, one switch). Every other step changes several things at once (engine, Handheld, Encode thread, caches, render scale), so the table shows the path, not what each switch is worth.
• GPU time per frame fell from 27.7 ms (row 1, Docked) to 12.6 ms (row 9), but frame time is also set by the caches: rows 9 and 10 have the same settings and almost the same GPU (12.6 / 13.6 ms) and submit time (5.3 / 5.3 ms), and differ by 57.4 vs 30.2 fps and p95 20.0 vs 69.5 ms. Row 10 is the same app right after its caches (pipeline archive and Metal's) were emptied.
• Rows 3 and 4 ended at or over the 8 GiB per-process limit (a kernel panic and a jetsam kill); from row 5 on, guest memory is in files (DM_GUEST_FILE=1) and the host-reported resident figure stays at or under 7.0 GB.
• Row 11 is the games/0400486024DEC000: update Ocarina of Time pso_manifest.tsv, shader_seen.tsv and notes #1732 manifest on a cold archive on an engine without eed7c429: 2,105 and then 1,465 MDB_MAP_FULL errors on the two starts (the numbers are in my earlier comment).

Not measured on this iPad: resolution (no controlled A/B), Dynamic resolution, and anything built on the 1.3.0 engine. I will add the next build as a new row.

Drafted with Claude from the host logs.

[bmjask]

Hello, I have a 120fps build on my branch

#1663

I've added a ton of performance fixes, I'm getting 90+fps on my m4 max in kakariko / square.

I'm getting 120fps in most other locations, however.

Should I make a stacked PR ?

[Know-It-All-Brother]

I will confirm this should go to main.

You may like #1894 :3

[Know-It-All-Brother]

I got nerd sniped by nxnative! I'm an edge inference expert so I thought I'd start taking a crack at some of the networks. I was able to fully reimplement the NDVE movie network as Metal kernels: #1894. M1 Pro goes from 36.3 → 10.8 ms of GPU a frame on the title.

Found a lane race in the translated HMMA shaders along the way: #1893

What's working well for me (M5 Max 64 GB, macOS 27.0.1, v131072):

• main e3ad825 + Ocarina of Time: experimental Mac playability and performance checkpoint #1663 Profile key gpu_ts_hz, the tick rate of raw counter reports (default 614.4 MHz, unchanged): Ocarina of Time reads them at 31.25 MHz, so its wind and leaves moved 19.66 times too fast #1714 The host reports the task / mesh draws it dropped while their pipeline was still being built #1720 games/0400486024DEC000: update Ocarina of Time pso_manifest.tsv, shader_seen.tsv and notes #1732 DM_NVN_GPUPROF_GEOM: split each render pass's GPU time into geometry and pixels #1751 Raw global memory loads through one pointer resolve its pool once, and the pool table search is a loop: Ocarina of Time's colour-pass mesh shaders stop looking up the pool table per word (chain 723dc3a3e17d) #1752 Compute threads that exit before a barrier keep arriving at it, and every shader loop has a cap (GPU hang defences), chain change #1774 Shader build: a simdgroup-scope fence after each threadgroup store (warp-synchronous scans), chain change #1775 Host GPU-hang defences: a watchdog that names a stuck submission's shaders, clamped GPU-written dispatch counts, a SyncWait timeout line #1777 Compare-and-swap loops in helper functions retry with the value the exchange returned #1796 A 64- or 128-bit raw global load that misses every pool reads 0 in all its words again, as the word loads it replaced did: the vector helper returned 0 for word 0 but never wrote gmem_vec, so words 1..3 read the previous load's words or undefined data; a vector load's follower scan takes only its own LDG's words (chain 6ed0e0a1fc9f) #1813 The pipeline archive bundled with the app is mapped too, not read into memory #1817 A texel fetch past the edge of a sampled 2D texture reads the nearest texel inside: Ocarina of Time's ground mist no longer has a hard straight edge at the picture's margins #1824 Avoid repeated archive scans for missing pipeline bases #1825: all 42,847 shaders build, pipeline prebuild 0 failed
  • Compute threads that exit before a barrier keep arriving at it, and every shader loop has a cap (GPU hang defences), chain change #1774 needs the one-statement-per-line fix reefbluesky found or all 520 task/mesh shaders fail. left suggestions on it
• dynamic res off, game upscaler → MetalFX to 2560x1440
• title: locked 60, 0 missed frames over 60 s (Metal HUD), ~14 ms GPU
• gameplay mostly 60 with some dips, no GPU hangs so far. haven't run a proper route yet

Happy to look at NDSR next if nobody's on it.

[ElevateAIAgencysv]

M5 (10-core GPU), 16 GB: Kakariko Village and Hyrule Field drop to 12–20 fps and slow motion while GPU memory sits near 14 GB; the rest of the game holds 60

Setup. Mac17,2, Apple M5 (10-core GPU), 16 GB, macOS 27.2, built-in Retina display. main 9b69d27 (shader chain 39417042cfe8), built from the NUP with the NSP script (--platform mac-app). Handheld (DM_OPMODE=0), DM_UPSCALER=metalfx, DM_PRESENT_UPSCALE=metalfx, fullscreen, DM_GUEST_FILE=1. The NUP's 3 damaged textures are replaced by substitutes in RomFS.img (same files as #1530). Game Mode is off (LSSupportsGameMode false), swap stays near 0, other apps about 1.5 GB. All figures are from the host log (timing, frame step, memory, residency set lines, 300-frame windows).

Session	Dynamic resolution	Where	Result	GPU p50	GPU allocated, max	Residency added, max per 300 frames
A (~52 min)	on	Goron City, Death Mountain Trail and Crater, Dodongo's Cavern	60 fps held in every window, game speed 1.000	median 9.1 ms	14.1 GB	2.8 GB
B	on	Great Fairy above the trail → owl → Kakariko	12–16 fps, game speed 0.63–0.66, held 3 vsyncs most of the time	16–31 ms	13.4 GB	4.1 GB
C	off	Kakariko, Hyrule Field, back to Kakariko	frame p50 33–80 ms, game speed 0.58–0.95	17–34 ms	14.8 GB	7.5 GB

What I see in B and C:

• The pause menu and the heart-piece get-item screen run smooth at the same spot in Kakariko; back in the scene it is slow again.
• Kakariko frames like whole frame 90–100 ms, submit+present 66–79 ms, GPU 41–43 ms, with no new pipelines and few texture uploads in those frames.
• The residency set adds and removes thousands of allocations per 300 frames there (up to 4,161 added / 3,357 removed, 7.5 GB added with dynamic resolution off), while in session A it stays far lower.
• No MDB_MAP_FULL, no pipeline failures, no crash, exit 0. The only [RED] lines are the known 2–3 texture type mismatch … slot 7 … type 7 (issue Ocarina of Time (M5 Mac): blue holes in Kokiri Forest terrain, and full-screen horizontal stripes since 0c7e88b9 #1566).
• Separately: first-sight task / mesh pipelines take 1.8–5.4 s each on the submit thread (mtk_pso), one stall per new area; on a later run the same pipelines take 1–4 ms, so Metal's cache keeps them.

Is this the render-target churn listed as open under Memory in the game note, or GPU memory pressure on a 16 GB Mac (the same churn shows with dynamic resolution off)? I can run Kakariko again with any switches or extra logging you suggest.

[ElevateAIAgencysv]

Update: Kakariko is fixed on 7479adf (7f50222 and bf420fd); one new case of the frame step staying at 20 fps after a load spike

Same Mac (Mac17,2, M5 10-core GPU, 16 GB, macOS 27.2), same settings as above, DM_DYNAMIC_RESOLUTION=1, rebuilt on main 7479adf (chain 3df44a5a4e0c).

Long session, about 2 hours (Kakariko, Hyrule Field, and the Ruto / Jabu-Jabu part):

• 99 % of frames at hold 1 (60 fps), game speed 1.000 (±0.005) in 217 of 235 frame step lines, below 0.97 in 9 (lowest 0.85).
• GPU p50 10–11 ms in every quarter of the session; GPU p95 at most 20.3 ms.
• Kakariko and Hyrule Field no longer drop: the residency churn from the report above is gone. heaps made 38046, reused 96395 at the end.
• The first run built and captured the profile's 123 task / mesh pipelines in the prewarm; the next run took 232 from the archive, so first-sight mesh stalls are gone in places visited before.
• Late in the session (last ~20 minutes): stretched textures and missing floor patches, while GPU allocated rose from a 17.2–17.7 GB maximum in the first three quarters to 21.0 GB in the last. GPU time did not rise. The only [RED] lines were two texture type mismatch … slot 7 … type 7 (issue Ocarina of Time (M5 Mac): blue holes in Kokiri Forest terrain, and full-screen horizontal stripes since 0c7e88b9 #1566).

Short session the next morning, same build and settings (about 5 minutes):

Loaded a save in Castle Town, played there at 60 (GPU p50 about 10 ms), then walked out to Hyrule Field:

frame step: hold 1 → 3 vsyncs at present 15084 (3 late frames; cost 2.08, GPU 0.77; hold 1 held at 0.68)
frame step: hold 3 → 1 vsyncs at present 15125 (probe, GPU 2.90× its time at hold 1; cost 3.32, GPU 2.04)
frame step: hold 1 → 3 vsyncs at present 15252 (3 late frames; cost 6.30, GPU 3.66; hold 1 held at 0.68)
frame step: hold 3 → 1 vsyncs at present 15282 (probe, GPU 5.47× its time at hold 1; cost 6.34, GPU 2.96)
frame step: the probe failed at present 15315 (CPU 92% on performance cores)

• The step to hold 3 came from a load spike with the GPU under one vsync (GPU 0.77): about 450 texture first uploads in one frame, the residency set adding 3.5 GB in that 300-frame window, GPU allocated 17.3 → 18.4 GB, alias region mapped up to +5.25 GB, about 1.5 GB of guest-file page-ins.
• After it, frames ran at 90–130 ms with GPU 35–48 ms, and probes back to hold 1 failed with GPU 2.9–5.5× its hold-1 time. Game speed 0.70–0.84. It had not recovered after about a minute, when the game was closed.
• Power adapter connected, Low Power Mode off, Game Mode off, swap near 0, thermal -1 throughout.

It looks like the GPU clock case in "A probe's grace lasts while its frames keep getting cheaper", entered at hold 3 instead of 2. Happy to run it again with any switch you want compared.

[ElevateAIAgencysv]

Second case on main 8190693 (chain 0c434f1b68df): after entering a heavy area, the submit thread waits for the GPU about 4× as often per frame and stays at 15–20 fps until the game is restarted

Same Mac and settings as above (M5 10-core GPU, 16 GB, handheld, MetalFX, DM_GUEST_FILE=1, DM_DYNAMIC_RESOLUTION=1). Power adapter on, Low Power Mode off, Game Mode off, swap near 0. The pipeline archive was warm from the build.

Loaded in Hyrule Field: 60 fps, GPU p50 12–14 ms. Then walked into Castle Town Market. From the translation detail and timing lines (300-frame windows; "waits" is the change in GPU waits so far per frame):

Window	Submit CPU	Submit wall clock	GPU waits per frame	GPU p50	Frame p50
Hyrule Field (f4200–f9000)	5.9–6.3 ms	7.0–7.9 ms	4.0–5.9	12.3–13.7 ms	16.7 ms
Castle Town (f9600)	5.2 ms	42.6 ms	13.2	31.3 ms	84.8 ms
Castle Town (f10200)	5.6 ms	42.3 ms	18.5	25.2 ms	65.9 ms
Castle Town (f10800)	4.3 ms	48.3 ms	19.2	26.1 ms	70.4 ms

• Game speed fell to 0.66 (held 3 vsyncs in 1,059 of 1,800 frames). It did not recover in about 1.5–2 minutes, and the game was closed.
• At f9600: 19 bases compiled (5.6 s on the background threads), 280 hybrid draws, 3,469 residency additions (2.5 GB). After f10200 there were no new pipelines, and the waits stayed high.
• GPU allocated 17.2–17.5 GB, resident about 5.8 GB. The indirect-count masking line shows waits for the GPU-written count 0 throughout, so the waits come from elsewhere.
• No [RED] lines other than the known slot 7 mismatch.

The same pattern happened the day before on 7479adf, walking from Castle Town into Hyrule Field: 20 waits per frame, wall clock 67.8 ms, also right after an area load. A two-hour session on that build never showed it, and both times a restart made the area play at 60 again. So it looks intermittent and tied to the first heavy area load after a start, not to the new build.

Is there a log switch that names what each wait is for? I can reproduce it with one if you tell me which.

[winkmix-png]

M6 MacBook Pro, 32 GB, macOS 27.2: Kokiri Forest holds 60, Castle Town 20–25 fps and looks CPU-bound

Setup. Built from main as of 2026-11-08 with build_ipa_from_nsp.py --platform mac-app, from the NUP. v131072. MetalFX for both the present and the game's upscaler (the game's own upscaler looked soft and ran slower here). Encode thread off. Game files and app on an external exFAT SSD (Samsung T5).

Frame rate:

Area	fps	GPU per frame	Frame time
Kokiri Forest	locked 60	~12 ms	16.7 ms
Castle Town	20–25	~15.5 ms	40–49 ms

In Castle Town the GPU finishes well inside a 60 Hz frame, but frames still take 40–49 ms. So it looks bound by the CPU or the translation step rather than the GPU, and I'd expect lowering the resolution not to help there. It matches what @sleepyeyedjoe described for the 30 fps stretches (frame step near 100% on the performance cores). I can run DM_NVN_GPUPROF or scripts/bench.py in Castle Town if a profile from an M6 would help.

exFAT saves, in case it helps someone: with the game on an exFAT drive, saves failed because exFAT can't do the game's save transactions. Pointing DM_SAVE at a folder on the internal APFS disk (~/Library/Application Support/nxnative/0400486024DEC000/save) fixed it.

[todd866]

Another data point for the 10-core M5, from a build a couple of days behind main (host e9c2f7e, shader chain dc815e0f38ca), 32 GB, MetalFX, handheld, dynamic resolution left to the game.

A 27-minute stretch of heavy scenes averaged about 37 fps. Windows with no shader compilation in them still had a GPU median around 34 ms, so this is not a cold pipeline cache. That matches the same-day 10-core M5 report above: dungeons can hold 60, and Kakariko / Hyrule Field stay in the 17–34 ms band.

In the settled heavy section the attachment traffic sat at about 710 MB loaded and 710 MB stored per frame. Every 300 frames the host still created roughly 5,000 Metal textures (2–6 GB) and committed the residency set for 200–600 ms, while the set itself stayed near 12,000 allocations and about 15 GB. By frame 72,900 storage takeover had run 131,003 times, shared textures were about 1,200 (about 140 of them unreferenced), and idle retirement was still 0. The create-and-replace cost is real and it is about 1–2 ms a frame, not the rest of the GPU time.

Two host findings from this tree, in case they are not on main yet:

• DM_TEX_RELEASE=2 rebuilt textures from zeroed pages after area changes (magenta and black surfaces, reproduced after the Saria / Fairy Ocarina cutscene). Level 1 keeps the CPU-visible copies and those surfaces stayed intact.
• Opening the ocarina aborted in SDK 22.2 nn::hid::IsOunceVibrationDeviceMounted with ResultNotRegistered (0xe15). A two-symbol HLE that reports no Ounce vibration device lets it open. IsOunceVibrationDeviceMounted and the hid::detail variant.

DM_NVN_DISPATCH_OVERLAP was off for the run above. This tree's own Kokiri measurement is 27.1 to 24.9 ms a frame with it on (M4 Pro). It is on for the next run here.

[Amost-]

Hello, to help with debugging and checking how different locations render/work, I created some save checkpoints. Below the details, hope it helps:

Test saves for v131072: the same spots for everyone (10 checkpoints, child and adult, plus 9 more spots)

Many reports here and on #1326 come from each player's own autosaves, so the scenes differ from Mac to Mac. These saves (zip, 11.6 MB) load Link at fixed places instead:

• Checkpoints: Hyrule Field and Zora's Domain as child and as adult (each pair at the same spot); Death Mountain as child (in front of the Goron on Death Mountain Trail, @lunks's spot: A talks to it) and as adult (the Bolero of Fire pad in the crater, in the Goron Tunic); and the adult warp pads of the other four temples (Minuet, Serenade, Nocturne, Requiem).
• More spots: Kakariko's gate, Castle Town's square, the Lost Woods (and the Goron City tunnel that loads them when you walk in), inside Dodongo's Cavern, Gohma's room before and after the fight, Link's porch in Kokiri Forest, and the Hyrule Castle vines (@ichampnick's hang spot: hold up to climb).
• Use one: quit the game, then open --env "DM_SAVE=/path/to/ocarina-test-saves-v131072/<folder>/save" "/path/to/<your app>.app" and pick Continue (the launch environment wins over env.txt, so your own save is untouched). Or copy a folder's contents into your save folder after backing yours up. Turn Cloud saves off first, or the sync may trade your cloud save and the test save (the newer one replaces the older, a backup is kept). The README in the zip lists each save's stage, position and facing.
• Adult saves are a child save made adult: the Master Sword and Hylian Shield equipped, the Ocarina of Time and the six warp songs (Sheik's lessons marked as seen), but the adult story flags as in the child save and 4 hearts. Fine for testing scenes, not a real late-game save.
• Checked: each one loaded headless on my M2 Pro (main 028a2631 + local picks, docked, 0.93 render scale), with a screenshot. It held 30 fps (33 ms frames) in most places, with 40–50 ms frames in Castle Town and about 50 ms in the frozen Zora's Domain.
• How: where a save loads is three fields of progress.sav: the stage Sequence_CurrentBanc, the position PlayerStatus.SavePos and the facing PlayerStatus.SavePosRadY. The file is a table of {MurmurHash3(field name), value} pairs with no checksum, and the game searches a safe spot near the saved position on arrival. The tool that makes such saves from any save is in Ocarina of Time: test saves that load Link at a chosen place (scripts/ocarina_test_save.py) #2002.

The save-format work builds on Traiver's OOT-SW2SE save editor (MIT): its parser and field catalog showed the file layout and many field names, and its catalog's hashes let me confirm the hash function.

[JoaoArturCosta]

MacBook Pro 14" (M3 Pro, 11-core CPU, 14-core GPU), 36 GB, macOS 27.0.1, Xcode 27.2

Build: main c68706f9, chain 0c434f1b68df, Mac app via build_ipa_from_nsp.py --platform mac-app from the NUP, all 42,848 shaders, bundled pipeline archive applegpu_g15s (3.5 GB). v131072, Docked.

Benchmark: scripts/bench.py on the kokiri-deck route, starting from my own Kokiri Forest autosave, --fps 60 so the GPU's headroom shows. Settings: DM_UPSCALER=metalfx, dynamic resolution off (full 1536x864), Metal 4 pipelines on, DM_FOCUS_KEEP=1. Small window, no HUD. One warm pipeline archive shared by all runs. Three runs per variant, run in turn, medians (range), 45 s window. 0 [RED] lines in all 9 runs. The shots match apart from camera position.

Variant	fps	GPU p50	GPU p95	Frames over 25 ms
Docked (default)	32 (30–42.8)	17.66 (15.75–17.78) ms	22.13 ms	1154 (698–1227)
DM_NVN_TBDR=1	31.7 (31.6–45.1)	17.85 (15.88–17.97) ms	22.37 ms	1209
Handheld (DM_OPMODE=0)	48.8 (44–50.6)	15.03 (14.04–15.12) ms	18.26 ms	473

• 60 does not hold at full resolution on a 14-core M3 Pro: the GPU sits just above the 16.7 ms budget in Kokiri Forest. In my play sessions in heavier areas it was 17–21 ms p50.
• Handheld takes 2.6 ms (15 %) off the GPU but lands around 49 fps, which is uneven. A locked 30 with DM_FPS_HOLD=on is the steadier choice on this Mac. At 30, Docked already holds: GPU p95 22 ms here, at most about 31 ms in play.
• DM_NVN_TBDR=1 gave no gain here, as on the M6 in the game note, unlike the M5 iPad above.

Metal 4 pipelines off (DM_PSO_MTL4=0) cost stutter here. I had it off on the earlier advice that it shortens the prepare step. Over a 15-minute session at a locked 30 on 2a6f9b32 (chain 3510ea767778), the log showed 200–5,137 ms of on the submit thread pipelines per 300 frames as new areas loaded, with frames of 190–867 ms. With it on, the benchmark runs above built 90–175 ms of pipelines on the submit thread per whole run. These are different scenes, so this is not a controlled A/B. On a Mac whose GPU matches the bundled archive, I'd keep it on.

First start with Metal 4 pipelines on: the bundled archive was found (holds the bases of every shader pair (14 of 14 probed bases found)), but the prepare step still compiled 887 bases (29 from the archive) in 61 s. That is the open item in the game note: the render bases the manifest's R rows need are not in the offline archive. With the per-user archive kept, the next start prepared in 0.46 s.

A 75 Hz external screen (2560x1080, it also offers 60 Hz): with DM_PRESENT_LINK=0 in env.txt, the screen stayed at 75 Hz in fullscreen. At 30 fps, frames then stayed up 2 and 3 refreshes in turn (on-screen p99 40 ms). From reading link_sync in nvn_mtl.m: with the link at 0, front_screen_match is never called, so the fullscreen switch to 60 Hz does not happen. Without that line, the switch should take place. I have not yet confirmed that in a log.

[LeonardoBai12]

M1 Pro (16 GB), macOS 26.6.1: two long sessions with no GPU hang, and the vanishing ground fixed by #1920

Builds: current main plus #1774 (the shader loop guard) and #1377 / #1379 (AV1 movies). Session 1 was on main 2a6f9b3, session 2 on main 04811a6. All 42,848 shaders built. Handheld, 60 fps target, dynamic resolution off, Metal 4 pipelines off.

	Session 1 (2a6f9b3)	Session 2 (04811a6, with #1920)
Length	83 min	131 min
GPU hangs	0	0
fps per 600-frame window, p50 / p10	30.9 / 26.5	23.7 / 20.6
GPU allocated	11–14 GB	about 16 GB most of the time, peak 20.9 GB
Ground in Hyrule Field	stopped drawing for moments, late in the session	never vanished
Where	Dodongo's Cavern, Goron City, Death Mountain, Kakariko, Hyrule Field, Gerudo Valley	Zora's Domain, Lake Hylia, Gerudo Valley, graveyard, Jabu-Jabu's Belly

• The lower fps in session 2 is not a clean comparison, because the places were different. Its long frames had more GPU time (median 33 ms against 26 ms), so I think the water areas are heavier rather than the build being slower. I plan to replay one of session 1's places on the new build to check.
• Gerudo Valley: in session 1 a texture covered the area above the river and hid everything below it. In session 2 I went through Gerudo Valley again and saw no wrong texture, there or anywhere else.
• Frame pacing: in both sessions the game felt mostly smooth, with occasional short stutters.

[JDSavvy]

Data point from an M4 Pro MacBook Pro (20-core GPU, 24 GB, macOS 27.0.1, AC power): main 04811a6 (chain 39eb2f781d07) plus eight open PRs stacked as host code only, no new shader image.

Stacked on 04811a6: #1927, #2085, #2077, #2079, #1888, #2065 (host part), #1901, #1951. All apply cleanly; #1901 needs a 3-way apply for nvn_mtl.m.

ABBA against an older build (from before the move to main), Kokiri Forest route from a fixed save with scripted inputs, docked, window, 190 s per run, stats from 60–120 s:

build	fps	frame p50 / p99 ms	GPU p50 ms	frames >25 ms (20–180 s)
main + PRs	59.1 / 59.4	16.67 / 18.79, 16.67 / 18.55	14.25 / 14.17	22 / 24
plain main (earlier series)	59.4 / 58.2	16.67 / 18.47, 16.67 / 33.35	15.16 / 15.31	137 / 290

Most of the gain over plain main on this route is #1927: plain main dropped to 30 fps for stretches of the walk (one run never came back to 60 in a later series). #1901 accounts for about -0.8 ms of GPU time. No crashes, no GPU hangs, no [RED] lines in any run.

One warning: #2065's host part without an archive rebuilt by its pso_offline.py makes the bundled archive fail its probe (this system does not take its bases (14 of 15 probed bases found)), so the device builds all 34,649 pairs in the background. First frame came after 36–37 s instead of 18–20 s, and the first scene change had two frames of 254 / 281 ms. Rebuild the archive with that PR, or leave it out.

[aksmith-git]

Pushed a fix to pub/oot-ipad-m5 (4b2d73c8) for black foliage with the grass compute pre-pass (DM_NVN_MESH_CPRE=1).

Symptom: in the Deku Tree (1F), ferns and vine leaves intermittently render as black silhouettes and flip back to green as the camera moves.

Cause: when the pre-pass draws get a render pass of their own (the next pass has other targets, about once a frame in the Deku Tree), cpre_flush encoded them without udc_note_draw. The colour attachment then counted as loaded and unwritten, so DM_NVN_UNWRITTEN_DC gave it a DontCare store and the foliage was discarded.

Fix: cpre_flush notes each draw's writes the same way draw_ does (6 lines). Checked on the Mac and on an iPad Pro M5 with the same save: no black foliage. t_nvn 264/0, t_unwritten_store 66/0.

If you are on an older build of the branch, DM_NVN_UNWRITTEN_DC=0 also works around it, at a small cost in store bandwidth.

[dpadr]

M2 profile, as asked: MacBook Air M2 (8-core CPU, 10-core GPU, 16 GB, fanless), macOS 27.0

Ocarina of Time v131072, handheld at 1440x810 (dynamic resolution off), MetalFX, frame-rate hold off. The route walks Kokiri Forest.

Build: our local OoT build on base 7f50222, plus #1833 (shader_raw_direct), #1902 / #1897 and the movie PRs #1377 / #1379. Frame p50 is 29.45 ms (32.72 ms without #1833).

Measured with DM_NVN_GPUPROF, Metal System Traces and a small standalone Metal benchmark. Nothing here is from main as of today, so the shares may shift there.

Where a frame goes on the M2

	per frame
GPU span	~26.6–28 ms
GPU busy (trace)	~67–70 %
Encoder time (sum of encoder durations)	~20 ms
GPU idle	~8–10 ms
CPU busy (replay thread)	~15 ms

By pass shape (66 frames, GPUPROF):

• the 1440x810 RG11B10+D32 scene family is ~10.8 ms (54 %), of which the main pass is ~5.5 ms and the foliage mesh pass ~2.7 ms;
• compute is ~3.7 ms (18 %);
• the 720x405 lighting target is ~0.9 ms;
• transparents and effects are ~1.2 ms;
• post is ~0.6 ms;
• shadows are ~0.5 ms.

So the M2 is GPU-bound at full size, unlike the M4 Pro.

What the idle is (trace, ~9.9 ms a frame)

• Pass seams: ~4–5.5 ms. These are short gaps between encoders, ~157 encoders a frame.
• Late first commit: ~3–3.7 ms. The next frame's work is translated only after the end-of-frame event releases it (details in Ocarina / M4 Pro: event-gated replay delays next-frame translation; experimental encode-ahead results #2096).
• Command-buffer boundaries: ~0.8 ms.
• Cross-queue waits: ~0.5 ms.
• Present: 0.

Two hardware facts that may help everyone

From a ~100-line standalone Metal 3 benchmark on this M2:

• A render-pass boundary costs a fixed 43–47 µs, whatever the attachment size (64² to 1440x810), format or load/store actions. A compute encoder between passes adds ~20 µs. Across ~157 encoders that's ~7 ms a frame, more than the visible idle, because ramp-down and ramp-up don't show as idle.
• The M2 overlaps independent work by ~40 %, for fragment beside vertex, fragment beside compute, and compute beside fragment. In the host, though, every compute, copy and upload encoder opens with barrierAfterQueueStages:All and ends with an All→All barrier. So in traces only 225 of 68k GPU intervals overlap another channel.

Results on this Mac

• Raw global memory at Metal's own GPU addresses, opt-in per game (shader_raw_direct, Ocarina of Time) #1833 (raw loads at Metal's addresses): −3.27 ms frame p50 (comment there).
• Draw passes overlap again while image heaps are alive and after fragment-only game barriers #1970 (relaxed draw-pass barriers): active, ~26k barriers saved per 300 frames, but no frame change; comment there. Overlap only rose from 0.02 to 0.39 ms a frame, because compute and copy encoders keep their full barriers.
• Mesh stage: after Raw global memory at Metal's own GPU addresses, opt-in per game (shader_raw_direct, Ocarina of Time) #1833, doubling the mesh shader body adds only ~50 µs to the 720x405 mesh pass.
• Main-pass fragment shaders: knocking out the three largest saves only ~1.4 ms in total.
• Depth prepass: ~1.1 ms.
• GPU clock: it holds its Maximum perf state 98–99 % of the time whenever the CPU keeps it fed. It drops only when starved (one run at 33 % busy spent 38 % of its time at Minimum).
• Thermal: on the fanless Air the CPU falls to ~2 GHz after about 5 back-to-back runs. Whole A/B blocks are lost to that unless runs whose clock dropped are excluded.

Encoder boundaries in one frame (LKG, frame 2400): 142 encoders

Source	Count
Game render-target changes	73
Compute↔render alternations the game's order allows reordering	43
Host-inserted (10 of them the culling-hang split)	~24

A dependency graph of the frame schedules in 109 encoders (−23 %). Its critical path is 79 % of node time.

What this suggests for the M2

The biggest levers here look structural, not per-shader:

1. feed the GPU at frame start: the Ocarina / M4 Pro: event-gated replay delays next-frame translation; experimental encode-ahead results #2096 release, early commit, frame lag (DM_NVN_FRAME_LAG=<k>: a frame starts once the GPU has completed the work presented k presents before the newest (off by default) #2073);
2. barriers scoped to what compute and render really share, so the ~40 % overlap is used;
3. fewer encoders.

We're testing these next on this machine and will add results.