Donkey Kong Bananza on iPhone: border corrections out of the main pixel libraries, NDSR replacements, pacing and fence fixes by givethesourceplox · Pull Request #1073 · seedds/nxnative

givethesourceplox · 2026-11-03T11:45:09Z

What and why

Donkey Kong Bananza (040074A01BF12000, v458752) reaches gameplay on an iPhone 15 Pro Max, but the GPU is the limit there: 46–48 ms a gameplay frame at about half render scale, and the phone soon throttles. This PR has the general fixes found while profiling it. Game-specific patches (widescreen, 40 fps, render-scale floor) are not part of it.

Border colour corrections out of the main pixel libraries (builds on #1032, co-authored by @AdyaGMD). border_emu's dm_bs branch took registers in every pixel where it never ran. A pixel shader's main library is now built without it (border_plain), and the correcting source goes into a side library <key>_bb, taken only while one of the shader's samplers has a twin (pso_key.logic bit 1). Three fixes on top of #1032:

<key>_b.metal keeps the correcting source, so a cache hit makes _bb, and the alpha test and logic libraries, with the correction. With the stripped source there, a cache hit made none of them.
Binding a sampler does not mark the pipeline dirty. The draw now compares the border bit and fetches the pipeline again when it changes. g_ntwin skips that check while no registered sampler has a twin.
Blending and the dual-source entry point test logic bit 0 only. With Pixel shaders without the border-colour correction in their main library: it spilled registers in every pixel where it never ran #1032 on main as it is, the border bit turned both off. t_missing_src1 has a new check that fails without this fix.

If this lands, #1032 can be closed.

NDSR replacements, opt-in (DM_NVN_NDSR=metalfx|bilinear). Only the reviewed 2x kernels listed in the profile's ndsr_spatial.tsv are replaced; anything else runs the game's shader. metalfx runs MetalFX spatial HDR on the Metal 3 queue, ordered with the Metal 4 queue by shared events. bilinear is one bilinear draw in the replay's own command stream. On the phone, NDSR's single dispatch takes 15.5–16.7 ms of the frame. Reference: Reviewed spatial upscalers on MetalFX.

Fixes

Present snapshots: the acquire waits for the display link's present-interval ticks, as without snapshots. With only the ring slot wait, DK at the 30 fps cap ran 43–65 fps of game time (fast-forward) and dropped about 110 of every 300 frames. After: 30.0 fps, 300 of 300 frames on screen (Mac).
Queued future fences: a guest wait on a fence the submit worker has published but not yet committed now takes the shared-event path, instead of waiting for CPU feedback (perf-order-cpu, t_perf_order).
Segmented submissions count draws from mtl_sub_begin. Before, a submission drawing into the target it inherited left its first pass's draws out of the count (t_reorder_chain --submission-overlap-only).
A mip-chain pass whose sampler pins one level samples that level alone (fb_pinned_level). DK's bloom upsample read a sharper level, so the stage intro's neon letters came out fat and filled.
HMMA with an F16 accumulator sums in FP32 (simdgroup_float8x8); test_dmsh_hmma.py adds an HDR-sized input case.
iOS: the prepared project and the unsigned IPA ask for Extended Virtual Addressing when the guest reservation is over 6 GiB (gen_entitlements.py). Without it, Play aborted before graphics with DK's 9.75 GiB block. DM_CPU picks the -mcpu target (apple-a17 for iPhone 15 Pro). DmGamepad.m builds with the iOS 27 SDK.

The shader chain changes (dmsh_chain.h regenerated).

Checked

Tests (M1 Pro, macOS 26, Xcode 26.5): make -C dm/gfx/tests GAME=040074A01BF12000 check_t_border, check_t_ndsr, perf-order-cpu, check_t_perf_order, check_t_reorder_chain, check_t_logic, check_t_alpha, check_t_missing_src1, check_t_integer_blend, check_t_link_ticks, and check_t_present_snap (01007EF00011E000); test_dmsh_border.py, test_dmsh_hmma.py, test_ipa_unsigned.py, test_docs.py; lint_python.py (no new finding in the files changed here). The full make -C dm/gfx/tests GAME=040074A01BF12000 run passes.
Build note: main's nvn_mtl.m uses MTL4FXTemporalScaler's *ContentOffset* properties, which the 26.5 SDK does not have. I commented those 4 lines out locally to build and test; they are not changed in this PR.
Mac M1 Pro, measured on our fork's build with the same changes (plus DK's widescreen patch, which this PR does not include). DK on the Banana Mine save, same route, 30 fps cap, DM_RENDER_SCALE=0.5, frames 85–118 s, one run at a time:

	GPU ms a frame	G-buffer pass ms	submit CPU ms
before (2 runs)	15.56 / 15.37	2.37 / 2.40	10.21 / 10.14
border side libraries (2 runs)	13.97 / 14.11	1.58 / 1.66	10.31 / 10.19
+ `DM_NVN_NDSR=bilinear`	12.22 (upscaler 2.78 → 0.08)		10.49
+ `DM_NVN_NDSR=metalfx`	12.67		10.54

At 100 s the output frames with and without the border change differ by a mean 1.12/255; two runs of the same build differ by 1.65/255. Bilinear differs from NDSR by 1.07/255 and is visibly softer; MetalFX differs by 1.28/255 and looks close to NDSR.

iPhone 15 Pro Max, iOS 27, on the fork's build: with Extended Virtual Addressing the game reaches gameplay, and with the snapshot pacing fix game time is back to 1.0× real time. The border and NDSR GPU savings are not measured on the phone yet.
Not run: scripts/regress.py check --games. regress.py impact says the change reaches every NVN game (nvn_mtl.m, hle_nvn.c), and DK's full shader build (22,852 shaders, 12,897 with a _bb side library) was done on the fork's chain with the same dmsh_build.py changes, not on this branch's chain 92c201edd640.

Docs

The findings are in docs/reference/{rendering,frame-rate,shaders,memory}.md, docs/development.md, docs/env_switches.tsv and DK's README; docs_symptoms.py was run
dmsh_chain.py --header was run and dmsh_chain.h is part of the change
No game files, keys or data extracted from a game

…el libraries, NDSR replacements, pacing and fence fixes Border colour corrections (builds on seedds#1032): a pixel shader's main library is built without border_emu's dm_bs branch, which spilled registers in every pixel where it never ran; the correcting source goes into a side library <key>_bb, taken only while one of the shader's samplers has a twin (pso_key.logic bit 1). On top of seedds#1032: - <key>_b.metal keeps the correcting source, so a cache hit makes _bb, and the alpha test and logic libraries, with the correction (with the stripped source there it made none); - binding a sampler does not mark the pipeline dirty, so the draw compares the border bit and fetches the pipeline again when it changes; g_ntwin skips that look while no registered sampler has a twin; - blending and the dual-source entry point test logic bit 0 only: the border bit turned both off (t_missing_src1 fails without this). Mac M1 Pro, Banana Mine, 30 fps cap, scale 0.5: GPU 15.4-15.6 -> 14.0-14.1 ms a frame, the G-buffer pass 2.4 -> 1.6 ms, the output frame unchanged. NDSR replacements (DM_NVN_NDSR, off by default, reviewed 2x kernels from the profile's ndsr_spatial.tsv only): metalfx runs MetalFX spatial HDR on the Metal 3 queue ordered by shared events; bilinear is one bilinear draw in the replay's own command stream. On an iPhone 15 Pro Max NDSR's dispatch takes 15.5-16.7 ms of a 46-48 ms GPU frame; on the M1 Pro 2.78 ms -> 0.08 ms (bilinear), whole frame 14.12 -> 12.22 ms (bilinear) / 12.67 ms (metalfx). No phone measurement yet. Fixes: - Present snapshots: the acquire waits for the display link's present-interval ticks; the ring slot wait alone let DK at the 30 fps cap run 43-65 fps of game time and drop ~110 of 300 frames. - Queued future fences: guest waits on a fence the submit worker published but has not committed yet take the shared-event path instead of waiting for CPU feedback (t_perf_order). - Segmented submissions count draws from mtl_sub_begin, so a submission drawing into the target it inherited no longer leaves its first pass's draws out (t_reorder_chain). - A mip-chain pass whose sampler pins one level samples that level alone (fb_pinned_level): DK's bloom upsample read a sharper level, the stage intro's neon letters came out fat and filled. - HMMA with an F16 accumulator sums in FP32 (simdgroup_float8x8), as the hardware does. - iOS: the prepared project and unsigned IPA ask for Extended Virtual Addressing when the guest reservation is over 6 GiB (gen_entitlements.py); DM_CPU picks the -mcpu target (apple-a17 for iPhone 15 Pro); DmGamepad builds with the iOS 27 SDK. Co-authored-by: Adya <114347520+AdyaGMD@users.noreply.github.com>

…he level its sampler pins, and a submission's segment count starts with it (from #1073) - Present snapshots: with the display link ticking, the acquire waits for the present interval's ticks like any present; only a stalled link falls back to the texture's slot time. The ring slot wait alone did not pace: the link drops all but the newest two frames and a dropped frame frees its slot, so Donkey Kong Bananza at the 30 fps cap ran 43-65 fps of game time and dropped about 110 of every 300 frames (measured by the PR's author). t_link_ticks checks the acquire with snapshots on. - Same-pass feedback in mip chains: a sampler whose LOD clamp pins one level (fb_pinned_level) gets a view of that level alone, on either side of the attachment. The view of the levels below clamped an upsample's read of level n + 1 to n - 1, a sharper level (DK's bloom: fat, filled neon letters), and an upsample into the view's base level read the attachment itself. t_feedback mip_chain_case. - Segmented submissions: the draw count starts in mtl_sub_begin, not at the first sub_cut, which only a target change reaches; draws into an inherited target were left out. t_reorder_chain inherited_target. Regress (M4 Pro): BotW PASS; TotK screenshots differ only because the baseline run lost window focus three times and fell behind (38.5 fps, frame 4241 vs 5294; 0 changed MSL, 0 new RED). Co-authored-by: givethesourceplox <272928635+givethesourceplox@users.noreply.github.com>

seedds · 2026-11-03T17:08:22Z

Thanks for this. Three of the fixes landed in 56e4577, with you credited as co-author:

Present snapshots: the acquire waits for the display link's ticks while the link is running, as you found. t_link_ticks now checks the acquire with snapshots on.
fb_pinned_level: an upsample pass samples the level its sampler pins, on either side of the attachment. That also covers rendering into the view's base level, which got no view before. I left out the DM_NVN_FB_PIN switch (a getenv on every draw) and took the sampler ID from the draw's own texture handle. New test: t_feedback mip_chain_case.
Segment draw count: the count starts in mtl_sub_begin. The test is a smaller inherited_target case in t_reorder_chain.

Regression run against main on an M4 Pro: BotW passes. TotK was flagged only because the run on main lost window focus and fell behind; both builds' pictures look right.

The rest is still open here:

NDSR: superseded by A game's own upscaler drawn by MetalFX, the player's choice in the launcher: Donkey Kong Bananza's NDSR network 3.5-5 ms → about 1 ms a frame on an M1 #1086 (DM_UPSCALER=metalfx, which takes the upscaler keys from the shader map), so ndsr_spatial.tsv and DM_NVN_NDSR won't land. If MetalFX turns out too slow on the phone, a bilinear option could become a third DM_UPSCALER value.
Fence wait, HMMA FP32 accumulator, Extended Virtual Addressing entitlement: these look right. Before landing, I want to check the HMMA claim against the translator's own body.
DM_CPU: the README's minimum is the iPhone 15 Pro, so the default should become apple-a17 rather than a new knob.
Border side libraries (Pixel shaders without the border-colour correction in their main library: it spilled registers in every pixel where it never ran #1032 plus your three fixes): the design fits the _ba / _bl pattern. It still needs a chain build and a regression run on a game that uses twin samplers (TotK), and a look at the extra pipeline variants that existing prewarm lists don't have yet.
csh_dispatch_ahead: not described in the PR. Its op list copies state_op's, so both should share one function, and it needs its own measurement.
DmGamepad.m guard, CFBundleName: the project needs the iOS 27 SDK (Xcode 27.x), so the guard isn't needed. CFBundleName isn't explained.

Also: the tests and measurements were run with Xcode 26.5, but dmsh_chain.lock requires the Xcode 27 Metal compiler. Please rebuild the chain with 27.x when you rebase.

…ing them: lint-c's ratchet passes again (seven nullability findings from #1086, #1084 and #1073, whose own runs were cancelled)

…he level its sampler pins, and a submission's segment count starts with it (from seedds#1073) - Present snapshots: with the display link ticking, the acquire waits for the present interval's ticks like any present; only a stalled link falls back to the texture's slot time. The ring slot wait alone did not pace: the link drops all but the newest two frames and a dropped frame frees its slot, so Donkey Kong Bananza at the 30 fps cap ran 43-65 fps of game time and dropped about 110 of every 300 frames (measured by the PR's author). t_link_ticks checks the acquire with snapshots on. - Same-pass feedback in mip chains: a sampler whose LOD clamp pins one level (fb_pinned_level) gets a view of that level alone, on either side of the attachment. The view of the levels below clamped an upsample's read of level n + 1 to n - 1, a sharper level (DK's bloom: fat, filled neon letters), and an upsample into the view's base level read the attachment itself. t_feedback mip_chain_case. - Segmented submissions: the draw count starts in mtl_sub_begin, not at the first sub_cut, which only a target change reaches; draws into an inherited target were left out. t_reorder_chain inherited_target. Regress (M4 Pro): BotW PASS; TotK screenshots differ only because the baseline run lost window focus three times and fell behind (38.5 fps, frame 4241 vs 5294; 0 changed MSL, 0 new RED). Co-authored-by: givethesourceplox <272928635+givethesourceplox@users.noreply.github.com>

…ing them: lint-c's ratchet passes again (seven nullability findings from seedds#1086, seedds#1084 and seedds#1073, whose own runs were cancelled)

…sampler with a twin: every pixel shader with border-colour corrections declares the function constant dm_border_off (DMSH_BORDER_FC, 216) and tests it before the mask bit in dm_bs, and the host specializes each function with it true when no sampler bound to the draw's stages has a border twin (border_need, pso_key.border), so the branch and the registers it kept live across every sample are gone from the pipeline instead of skipped per pixel. The mask branch made Apple's compiler keep the sample's operands for the correction that never ran, which spilled registers in shaders with many samples. A pipeline builder takes an explicit off flag from the draw's key (pso_build, pso_build_ex, pso_build_dev, pso_desc4, pbase_key, cps_new_b, the mesh and capture pipelines, the prewarm and pso_offline.py's archives), and the specialized function is found there (fn_off, fn4_of_b, fn_dev_b); descriptors hold only real Metal functions. A draw compares its stages' need with the bound pipeline's (prog_border), so a sampler with a twin bound after the pipeline was picked gets the correcting pipeline. DM_NVN_BORDER_KEEP=1 keeps the corrections in every pipeline and =ab alternates by DM_AB phases. Ocarina of Time standing on the forest deck, in-process A/B on an Apple M6: GPU p50 10.66-10.88 ms without the corrections, 11.01-11.26 ms with them, frame time 16.67 ms both. Compiled for the M1's GPU (applegpu_g13g, 300 of its pixel shaders): mean code 16,232 -> 14,531 bytes, per-thread scratch lower in 77 shaders, higher in 11, the worst 240-352 bytes -> 0. AdyaGMD's measurement of the side-library form on an M1, Splatoon Raiders: 31.6 -> 19.8 ms a frame; not reproduced here. t_border (113 checks), the 136 gfx suites, test_dmsh_border, test_pso_offline, test_docs pass; regress.py check on Bastion and NieR:Automata: 60 fps, every shader builds, no new red line. The side library of PR #1203 is not taken: it compiles every pixel shader with corrections twice. Not run: Tears of the Kingdom (the one known game that binds a twin; its build is not on hand), the iPhone, an M1. The function constant is FeReuss's design and code (PR #1929); the diagnosis of the spill and the M1 numbers are AdyaGMD's (PRs #1203, #1032); the need for a per-draw choice is givethesourceplox's finding (PR #1073) Front-ends: the shader chain version changes; shaders are rebuilt. Co-authored-by: FeReuss <102842515+FeReuss@users.noreply.github.com>

seedds · 2026-11-11T06:06:07Z

The border-correction part of this PR is covered on main by 511a0a2: a function constant compiles the corrections out of every pipeline that binds no sampler with a twin, and a draw checks its stages' need against the bound pipeline's, which is the per-draw choice this PR showed was needed (credited in the commit message). The side libraries are not taken. The other parts (NDSR replacements, pacing, fences) are not touched by that commit and stay open here. Not run: Donkey Kong Bananza on an iPhone; a retest of the border part on main would be welcome.

AdyaGMD mentioned this pull request Nov 3, 2026

A game's own upscaler drawn by MetalFX, the player's choice in the launcher: Donkey Kong Bananza's NDSR network 3.5-5 ms → about 1 ms a frame on an M1 #1086

Closed

2 of 3 tasks

seedds added a commit that referenced this pull request Nov 3, 2026

Tests check the Metal buffers, bitmaps and events they make before us…

ad7af9c

…ing them: lint-c's ratchet passes again (seven nullability findings from #1086, #1084 and #1073, whose own runs were cancelled)

AdyaGMD mentioned this pull request Nov 4, 2026

Pixel shaders without the border-colour correction in their main library, on the shared side-library compile #1203

Closed

givethesourceplox mentioned this pull request Nov 5, 2026

Recheck GPU feedback after an unsignaled fence event #1409

Open

3 tasks done

Navigation Menu

Repository navigation

Donkey Kong Bananza on iPhone: border corrections out of the main pixel libraries, NDSR replacements, pacing and fence fixes - #1073

Donkey Kong Bananza on iPhone: border corrections out of the main pixel libraries, NDSR replacements, pacing and fence fixes#1073
givethesourceplox wants to merge 1 commit into
seedds:mainfrom
givethesourceplox:dk-iphone-gpu-pacing

givethesourceplox commented Nov 3, 2026

Uh oh!

seedds commented Nov 3, 2026

Uh oh!

seedds commented Nov 11, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

givethesourceplox commented Nov 3, 2026

What and why

Checked

Docs

Uh oh!

seedds commented Nov 3, 2026

Uh oh!

seedds commented Nov 11, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants