[zephleggett]

What and why

On M1-class Macs, Ocarina of Time's title screen and movies run at about 24 fps. The movie library's NDVE network takes 36 ms of GPU a frame (#1375).

This PR adds an opt-in engine that runs the same network as 10 hand-written Metal kernels. On an M1 Pro it takes 10.8 ms.

• Off by default. DM_NDVE_ENGINE=1 turns it on.
• Fail-safe. The hook recognizes the network by its 14 shader keys in order, its weight layout, and its image sizes and formats. On any mismatch, the game's own dispatches run.
• Main's bits. The engine uses the same HMMA operations as main's dmsh_hmma16816_f16, so it reproduces main's output.

Closes #1375.

Exactness

• test_ndve_engine.py: all 10 kernels at the network's sizes, on random inputs and weights. They are bit-identical to references built from main's HMMA helper. An f32-accumulator variant fails 6–25 % of each HMMA kernel's outputs, so the test can tell the difference.
• Main's translated shaders: 2 title-movie frames and 1 random picture. The engine matches at every output that main computes the same way on every run. The remaining outputs vary between main's own runs; see the first finding below.

Two findings for the translator

1. Lane race. Main's translated network is not deterministic.

   • 10 of its 12 dispatches give different values from run to run.
   • Three end-to-end runs of one frame differ in 42–46 % of pixels, by up to 502 of 1,023 codes.
   • The engine has no race.

2. Rounding. Main's helper rounds the f16 accumulator after every product, which is how simdgroup_half8x8 behaves on Apple GPUs. Recorded A100 data (Khattak and Mikaitis's tensor core models) shows that HMMA with an F16 accumulator rounds the sum of 8 products once:

   • rounding the sum once matches 5,000 of 5,000 cases;
   • per-product rounding matches 2,420.

   The engine follows main for now. If the helper changes, the test fails until the engine is updated.

Checked

• Device: M1 Pro (16-core GPU, 16 GB), macOS 26.5.2, Xcode 26.6.

• Tests:

  • test_ndve_engine.py passes.
  • make -C dm/host/tests run passes.
  • make -C dm/gfx/tests GAME=0400486024DEC000 run passes except t_present_crop, which crashes the same way on unmodified main on this Mac.
  • dm/tools/tests fails the same tests as unmodified main.

• Speed: measured standalone on one title-movie frame.

  Path	GPU per frame
  Translated network	36.3 ms (sum of best-of-10 per dispatch)
  Engine	10.8 ms (best of 19)

• Reach: regress.py impact lists BotW, Scarlet, TotK and Hades only because nvn_mtl.m's dispatch changed. With the switch unset, the hook returns on its first line. I did not run regress.py check.

• Not checked on this branch:

  • in-game fps and picture (on a downstream tree with the engine, the title held 60 fps at 12.5 ms of GPU);
  • the iOS compile of ndve_engine.metal;
  • the DM_NVN_ENC_THREAD path;
  • other Macs.

Scope

The change is 853 lines, 361 of them the test. It is a second path for work the translator already does, so CONTRIBUTING asks to agree on scope first. If you would rather settle that in #1375, I'll move it there.

#1331 skips the network instead: it is cheaper, but the picture is not the game's. The two can coexist.

Docs

• docs/reference/shaders.md: new entry "The NDVE movie network by the host". python3 dm/tools/docs_symptoms.py was run.
• The game note and docs/env_switches.tsv are updated.
• No game files, keys or extracted data. The weights come from the game's buffers at run time, and the test generates its own.