[AdyaGMD]

What and why

Donkey Kong Bananza upscales its frame with agl's NDSR network (agl_ndsr.pck): one compute dispatch between the render-resolution passes and the output-resolution ones, which then draw the HUD over it. On an M1 in handheld mode (960x540 → 1920x1080) it is the largest single item of a heavy frame, 3.5–5 ms (DM_NVN_GPUPROF), because its HMMA products are emulated and its shared-memory matrix loads move their row addresses between lanes with shuffles (41 % of its time). MetalFX's spatial scaler does the same job, same sizes, in about 1 ms.

A new launcher option: Game's upscaler. The menu before the game starts gets a Game's upscaler row with Game's own (the default) and MetalFX. It appears only for games whose profile names their upscaler shaders (the new profile key metalfx_upscaler; today only Bananza, its 24 NDSR compute variants), so other games' menus do not change. It is a choice and not a default because the picture is then Apple's scaler rather than the game's trained network. It is not the same as the existing Upscaling row: that one scales the finished picture, HUD included, to the screen; this one replaces the game's own step inside the frame. Both can be on. docs/settings.md explains it for players; the variable is DM_UPSCALER=game|metalfx.

How it works (up_mfx_encode in nvn_mtl.m). A dispatch of a named shader is replaced by MetalFX's Metal 4 spatial scaler in HDR mode, from the input rectangle to the output rectangle the shader would use (its push constants c[24]: origin and inclusive last texel of each). Bananza's colour is R11G11B10F, which MetalFX does not take, so the input rectangle is drawn into an RGBA16F scaler texture and the result drawn into the output rectangle (copy_t2t_draw); RGBA16F textures with whole rectangles go straight through. MetalFX is encoded through rec_cb_block (new in nvn_rec.h), so on the encoding thread (iOS's default) it gets the real command buffer. The game's shader still runs for an output not larger than the input, an indirect dispatch, an output texture that is not a render target, or a size MetalFX refuses.

Credit: the rectangle decoding, HDR mode and the encoding-thread handling follow @givethesourceplox's NDSR replacement in #1073 (credited as co-author). This PR differs from that one in being a player setting in the launcher, covering every NDSR variant and size ratio, and encoding on the Metal 4 queue instead of a Metal 3 queue synchronized by events.

Checked

• Tests: make -C dm/gfx/tests GAME=040074A01BF12000 check_t_upscaler_mfx: a named shader's 960x540 → 1920x1080 dispatch is drawn by MetalFX; a cropped R11G11B10F input with a highlight above 1.0 into an offset output rectangle, every pixel inside and outside checked; an output no larger than the input and an unnamed shader keep the shader; nothing changes without the switch; also on the encoding thread. dm/tools/tests/test_game_profile.py (the key's values) and test_settings.py (the row shows for Bananza with "Game's own" as default, not for NieR:Automata); test_docs.py.

• Games run, with title ID and version: Donkey Kong Bananza 040074A01BF12000 v458752, handheld, the same heavy area, on a build with this change, the border-colour change (Pixel shaders without the border-colour correction in their main library: it spilled registers in every pixel where it never ran #1032) and mesh draw runs (a separate PR):

  Game's upscaler	GPU ms a frame	frame ms	fps
  Game's own (NDSR)	17.37	19.44	51.4
  MetalFX	15.32	16.6	60

  No visual difference noticed between the two in play.

  The launcher's new row (shown only for Bananza):

  

  The same heavy area with the frame-rate counter, the game's own upscaler (NDSR):

  

  With MetalFX:

  Memory was flat over a long MetalFX run (resident +15 MB between frames 8,100 and 11,100; one texture usage upgrade in the whole run), after an earlier build that upgraded the output texture every frame was fixed by keeping the shader for a non-render-target output.

• Devices and OS: MacBook Pro (13-inch, M1, 8-core GPU), macOS 26.6.2, Xcode 27.2.

• scripts/regress.py: not run. The shader chain does not change; the host change only acts for a game with metalfx_upscaler and DM_UPSCALER=metalfx.

• Not checked: the iPhone. One open note (in the reference entry): a standalone MetalFX test on the M1 intermittently produced invalid pixels at small sizes, and one handheld run showed one wrong HDR pixel; not seen in play since.

• Known Metal-side quirk: MetalFX's Metal 4 scaler stops under Metal validation (MTL_DEBUG_LAYER=1) on a barrier-stage setting it does not expose (_outputTextureBarrierStages not set), so that test runs without validation; Donkey Kong Bananza on iPhone: border corrections out of the main pixel libraries, NDSR replacements, pacing and fence fixes #1073 reports the same.

Docs

• The finding is in its docs/reference/rendering.md entry ("A game's own upscaler drawn by MetalFX"), the game note and docs/settings.md; python3 dm/tools/docs_symptoms.py was run
• A chain tool changed (no)
• No game files, keys or data extracted from a game in the change or the attachments

🤖 Generated with Claude Code

[seedds]

Landed in 85468b6, with you and givethesourceplox as co-authors. Thank you for this. The MetalFX path, the RGBA16F conversion, the fallbacks and rec_cb_block are yours as written. Changes on top:

• No profile key. The game's upscaler shaders now come from the shader map: upscaler_keys in dm/games/game.py takes the compute shaders found in agl_ndsr.pck. For Bananza that gives exactly your 24 keys. Splatoon Raiders ships the same network (20 shaders) and now gets the row too; the swap hasn't been run on it. The header macros are now DM_GAME_UPSCALER_KEYS / DM_GAME_UPSCALER_N.
• DLSS through rec_cb_block as well. dlss_eval encoded MetalFX on the recording's stand-in command buffer. With DM_NVN_ENC_THREAD=1 it stopped the host (DmRecCB has no stand-in for -useResidencySet:). It now encodes through your block, with the scaler's settings made inside it. The new test check_t_dlss_mfx runs it with the thread off and on.
• check_t_upscaler_mfx also runs with the encoding thread on.
• Docs: env_switches.tsv order (its test failed), a settings.tsv trailing tab, apostrophes and the make -C command in the reference entry, and a link to the open LDSM/MOVM item in "Ampere HMMA through simdgroup_matrix". The network's own cost on the default path is still there. I also dropped the "sharpens edges" line from settings.md, since it wasn't measured.

Not checked: the iPhone, and Raiders with MetalFX.