Run nn::ndve's movie network as host Metal kernels behind DM_NDVE_ENGINE=1 by Know-It-All-Brother · Pull Request #1894 · seedds/nxnative

Know-It-All-Brother · 2026-11-09T01:46:52Z

What and why

On M1-class Macs, Ocarina of Time's title screen and movies run at about 24 fps, because 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 Metal kernels. On an M1 Pro the network takes 10.8 ms instead of 36.3 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.
Same arithmetic as main. The kernels use the same HMMA steps as dmsh_hmma16816_f16, so the output is bit-identical to main's wherever main's output is stable.
Deterministic. main's translated network has a lane race (NDVE movie network: the translated HMMA shaders give different results from run to run #1893), and the engine doesn't.

Fixes #1375. Reference entry: The NDVE movie network by the host.

Two findings for the translator

Lane race (NDVE movie network: the translated HMMA shaders give different results from run to run #1893). 10 of the translated network's 12 dispatches give different values from run to run, and three runs of one frame differ in 42–46% of pixels, by up to 502 of 1,023 codes.
Rounding. main's helper rounds the f16 accumulator after every product, as simdgroup_half8x8 does 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: that matches 5,000 of 5,000 cases, against 2,420 for per-product rounding. The engine follows main for now, so the test catches any drift. I'll propose rounding once in the translator and the engine together in a separate PR.

Checked

Tests:
- python3 dm/gfx/tests/test_ndve_engine.py passes on an M1 Pro and an M5 Max. It runs all 10 kernels at the network's sizes on random inputs and weights, bit-for-bit against 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.
- make -C dm/host/tests GAME=0400486024DEC000 run passes on an M5 Max.
- make -C dm/gfx/tests GAME=0400486024DEC000 run passes on an M5 Max, except t_counter_clock's check "feedback conversion uses the resumed guest epoch". It's intermittent on unmodified main too: it failed in 1 of 3 runs there and in 0 of 3 on this branch.
- lint_c.py tidy --against origin/main with LLVM 23.1.4: 0 file/check pairs above the baseline.
Games run, with title ID and version: The Legend of Zelda: Ocarina of Time, 0400486024DEC000, v131072.
- Against main's translated shaders, on 2 title-movie frames and 1 random picture, the engine matches every output that main computes the same way on every run.
- On an M5 Max, I captured 60 s of the title screen with the Metal HUD for each path, at 2048x1152 upscaled to 2560x1440 with MetalFX. Both held 60 fps.
Devices and OS: MacBook Pro M1 Pro (16-core GPU, 16 GB), macOS 26.5.2; MacBook Pro M5 Max (64 GB), macOS 27.0.1.
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 didn't run regress.py check.

Measurements:

Mac	Measurement	Game's network	Engine
M1 Pro	Network GPU time per frame, standalone	36.3 ms	10.8 ms
M5 Max	Title screen, whole-frame GPU time, 60 s	14.3 ms	13.8 ms
M5 Max	Title screen fps, 60 s	60	60

The gain is largest on M1-class GPUs, where the network is the bottleneck. The M5 Max runs the game's network fast enough to hold 60 fps either way.

Not checked: the iOS compile of ndve_engine.metal, the DM_NVN_ENC_THREAD path, and Macs other than these two.

Scope

The change is 853 lines, 361 of them the test. It's a second path for work the translator already does, so CONTRIBUTING asks to agree on scope first; I'm happy to settle that in #1375. #1331 skips the network instead, which is cheaper, but the picture isn't the game's. The two can coexist.

Docs

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

…NDVE_ENGINE=1, off by default): Ocarina of Time's title-movie network takes 10.8 ms of GPU instead of 36.3 on an M1 Pro with 16 GPU cores, standalone, with the translated shaders' bits wherever their lane race leaves them stable (issue seedds#1375) The title and the movies run 14 compute dispatches a frame: the colour conversion 49dff010, eleven HMMA.16816 convolutions of a U-Net over the picture packed 2x2 at 960x540, the last convolution added to the picture (5a7d69b0) and a 4/3 resampler with a small network of its own (4f41ccf2). On an M1 Pro the translated network costs about 38 ms of GPU a frame, so the title runs at about 24 fps (24.8 ms on an M1 Max in the issue). The translated HMMAs still take B through the SIMD-group's scratch and the weights' LDSM and MOVM through shuffles. nvn_mtl.m's dispatch now hands those dispatches to nvn_ndve.inc when DM_NDVE_ENGINE is 1 (or ab with DM_AB). It recognizes the network by its 14 shader keys in order and its weight table (each dispatch's weight and bias buffers at the network's offsets from the first's), then the picture and image sizes and formats and the resampler's constants; anything else leaves the game's dispatches running, logged once, and a mismatch after some dispatches were skipped turns the host's network off for the rest of the run. It skips the first eleven, runs the whole network at 5a7d69b0 into the image that dispatch writes, from the picture it reads and the weights the game bound, and runs its own resampler at 4f41ccf2. The kernels are dm/gfx/msl/ndve_engine.metal, built into the host with its other shaders (host_msl.py), so nothing compiles at run time. They do the translated shaders' arithmetic in their order: each HMMA as dmsh_build.py's HMMA_BODY does it (simdgroup_half8x8, simdgroup_multiply for a chain's first call, then multiply_accumulate over input channels 0-7 and 8-15), the chain over the taps and then the 16-channel blocks, the epilogue's D + bias in f32 with ReLU and one rounding, the picture's sum in the game's order under safe math. The translated network has a lane race on main: run twice on the same input, 10 of its 12 dispatches give different values in some outputs (5a7d69b0 in 13.5 % of its pixels over 12 runs); end to end, three runs on one frame differ from each other in 42-46 % of the pixels, by up to 502 of 1,023 codes. The host's network gives at every output a value some run gives, and at outputs that never change, that value. The race's cause in the translation is open. The matrix hardware rounds simdgroup_half8x8's accumulator after every product (M1 Pro, 40,000 of 40,000 random 8-product cases), and an A100's HMMA with an F16 accumulator rounds the sum of 8 products once (5,000 of 5,000 cases recorded on an A100 in Khattak and Mikaitis's tensor core models, against 2,420 for the per-product rounding). So the translated helper, and the host's network with it, differ from the Switch in the low bits of most HMMAs; shaders.md says what an exact helper would do. The host's network follows the helper, and its test fails when they differ. Checked: dm/gfx/tests/test_ndve_engine.py on the M1 Pro: every kernel at the network's sizes on random inputs and weights, bit for bit against references that call dmsh_hmma16816_f16 as hmma_msl rewrites it, in PTX's fragment layout with TensorPass's chain flags (all identical; with an f32 accumulator per HMMA instead, 6-25 % of each HMMA kernel's outputs differ), and the kernel names the host binds. Main's 13 translated dispatches (chain c096b1d1dc53) standalone on two title-movie frames and a random picture against the host's: equal at every output that is the same in 6 or 12 runs, and every changing output equals the host's in some run; the resampler identical. make -C dm/host/tests run; make -C dm/gfx/tests GAME=0400486024DEC000 run (t_present_crop crashes in its MetalFX scaler resize check, as on main on this Mac); dm/tools/tests (the same five failures as on main here: test_env_switches, test_nsp_scan, test_tree_tar, test_triage_report, test_usage). Not run: the game itself with the switch on (frame rate, GPU span, the picture compared in the game), other Macs, the iPhone.

…default: Ocarina of Time's title movie takes 6.77 ms of GPU a frame on an M4 Pro instead of 12.78, the output the game's network gives (issue #1375, PR #1894) The title and the opening movies run nn::ndve's model Movie/Ndve/Ndve-Nvn2-vid_seq5.dat (a Rouleau archive: per network its 9 compute shaders and 0x3c320 bytes of weights, no layer graph): the colour conversion, eleven HMMA.16816 convolutions of a U-Net over the picture packed 2x2 at 960x540, the last one added to the picture, and a 4/3 resampler. Translated, the network is the game's NVIDIA code, which loads each layer's weights before its tile loop and holds up to 144 of them in registers through it; on Apple GPUs that costs 36 ms of GPU a frame on an M1 Pro (24 fps), 18.9 ms of compute on an M4 with 8 cores. No general lowering measured does better: B as simdgroup matrices once before the loop, B loaded from its storage buffer at every HMMA (built in TensorPass, bit-exact, slower), B in threadgroup slots (shaders.md). nvn_mtl.m's dispatch hands the network to nvn_ndve.inc (PR #1894's engine, reworked), which runs it as dm/gfx/msl/ndve_engine.metal's kernels in one encoder: the same arithmetic as the translated shaders in their order. It is bound to the library's model, not to a game: the keys are the model's shaders' hashes and the weights' offsets follow from its layers (each layer's weights then its biases, padded to 256 bytes, to the archive's 0x3c320), checked against the offsets Ocarina of Time binds. The picture the last layer adds must be the one the first read; anything else leaves the game's dispatches running. DM_NDVE_ENGINE=0 runs the game's; check runs both and logs the texels that differ. Checked: test_ndve_engine.py on an M4 Pro (bit for bit against the translated HMMA helper, as on the PR author's M1 Pro and M5 Max), test_docs, test_env_switches; the title-movie route on an M4 Pro, three runs each, GPU p50 12.78 -> 6.77 ms, both at 60 fps; DM_NDVE_ENGINE=check over 1,500 frames, with and without DM_NVN_ENC_THREAD: the resampler identical, the network in all but 10 texels, each 1/1023 (the translated shaders' fast math; 5a7d69b0 built without it gives other bits in a few texels); the engine's Metal compiles for iOS. Not run: an M1-class Mac or the iPhone in the game, regress.py (the hook acts only on the model's keys, which only Ocarina of Time runs). Issue #1893's run-to-run differences did not reproduce on an M4 Pro with chain dedc9391c30b or 39417042cfe8. Co-authored-by: Know-It-All-Brother <354264050+Know-It-All-Brother@users.noreply.github.com>

seedds · 2026-11-09T11:37:49Z

Thank you. This landed in c66a36a, on by default, with you as co-author. The kernels, the hook and the test are yours; what changed:

Bound to the library's model, not to the game. nn::ndve is linked into main (nn::ndve::RouleauArchiveIO), and Ndve-Nvn2-vid_seq5.dat is a Rouleau archive: each network holds its 9 shaders and 0x3c320 bytes of weights, but no layer graph. So the keys stay the model's identity, and the weight offsets are now computed from the layer shapes (each layer's weights then its biases, padded to 256 bytes) instead of two tables. A new test checks the rule against the offsets the game binds.
One more check: the picture 5a7d69b0 adds must be the one the first layer read.
On by default: DM_NDVE_ENGINE=0 runs the game's dispatches, ab still alternates, and the new check mode runs both side by side, comparing the outputs texel by texel.
Docs: the reference entry drops the lane-race claims. They didn't reproduce here (see NDVE movie network: the translated HMMA shaders give different results from run to run #1893).

Measured on an M4 Pro (20 GPU cores), title-movie route, three runs each: the frame's GPU time went from 12.78 to 6.77 ms (p50), both at 60 fps. In check mode over 1,500 frames, the resampler was identical in every frame and the network in all but 10 texels, each off by 1/1023. The likely cause is the translated shaders' Metal fast math (5a7d69b0 built without it gives other bits too).

Before landing this, I tried to fix it in the translator instead: B as matrices, B loaded from its storage buffer at each HMMA (built and bit-exact), and B in threadgroup slots. None beat the current lowering on all the heavy layers. The numbers are in shaders.md, "The NDVE movie network by the host".

Not measured yet: the game on an M1-class Mac. A run of the title movie on your M1 Pro with the current main, plus the GPU time line, would be very welcome.

…es, instead of after every product: the translated helper and the host's NDVE network sum each block of 8 in f32 and round it to f16, which matches 5,000 of 5,000 recorded A100 cases (per-product rounding 2,420) and 99.97 % of a Jetson Orin's outputs (34.9 %); chain 0c434f1b68df (PR #1897) HMMA_BODY accumulated on simdgroup_half8x8, whose multiply-accumulate rounds the f16 accumulator after every product. Each block now accumulates into a simdgroup_float8x8 from the half operands and a half C and is rounded to f16 per lane through thread_elements(). test_dmsh_hmma.py's GPU runs compare every output bit for bit with the two-block reference (main's body differs in 21,461 outputs on an M5 Max and on an M6). ndve_engine.metal's hmma() (#1894, bit for bit with the helper) rounds the same way, into a new matrix: written into its accumulator arrays' own thread_elements(), its layers ran 15 to 50 times slower on an M6. On the M6 the network's seven layer kernels take 2.94 ms against 2.85; isolated on an M5 Max, NDSR's fa3033d05c0a41d6 takes 0.316 ms against 0.276. Not measured in a game. Co-authored-by: Know-It-All-Brother <354264050+Know-It-All-Brother@users.noreply.github.com>

…default: Ocarina of Time's title movie takes 6.77 ms of GPU a frame on an M4 Pro instead of 12.78, the output the game's network gives (issue seedds#1375, PR seedds#1894) The title and the opening movies run nn::ndve's model Movie/Ndve/Ndve-Nvn2-vid_seq5.dat (a Rouleau archive: per network its 9 compute shaders and 0x3c320 bytes of weights, no layer graph): the colour conversion, eleven HMMA.16816 convolutions of a U-Net over the picture packed 2x2 at 960x540, the last one added to the picture, and a 4/3 resampler. Translated, the network is the game's NVIDIA code, which loads each layer's weights before its tile loop and holds up to 144 of them in registers through it; on Apple GPUs that costs 36 ms of GPU a frame on an M1 Pro (24 fps), 18.9 ms of compute on an M4 with 8 cores. No general lowering measured does better: B as simdgroup matrices once before the loop, B loaded from its storage buffer at every HMMA (built in TensorPass, bit-exact, slower), B in threadgroup slots (shaders.md). nvn_mtl.m's dispatch hands the network to nvn_ndve.inc (PR seedds#1894's engine, reworked), which runs it as dm/gfx/msl/ndve_engine.metal's kernels in one encoder: the same arithmetic as the translated shaders in their order. It is bound to the library's model, not to a game: the keys are the model's shaders' hashes and the weights' offsets follow from its layers (each layer's weights then its biases, padded to 256 bytes, to the archive's 0x3c320), checked against the offsets Ocarina of Time binds. The picture the last layer adds must be the one the first read; anything else leaves the game's dispatches running. DM_NDVE_ENGINE=0 runs the game's; check runs both and logs the texels that differ. Checked: test_ndve_engine.py on an M4 Pro (bit for bit against the translated HMMA helper, as on the PR author's M1 Pro and M5 Max), test_docs, test_env_switches; the title-movie route on an M4 Pro, three runs each, GPU p50 12.78 -> 6.77 ms, both at 60 fps; DM_NDVE_ENGINE=check over 1,500 frames, with and without DM_NVN_ENC_THREAD: the resampler identical, the network in all but 10 texels, each 1/1023 (the translated shaders' fast math; 5a7d69b0 built without it gives other bits in a few texels); the engine's Metal compiles for iOS. Not run: an M1-class Mac or the iPhone in the game, regress.py (the hook acts only on the model's keys, which only Ocarina of Time runs). Issue seedds#1893's run-to-run differences did not reproduce on an M4 Pro with chain dedc9391c30b or 39417042cfe8. Co-authored-by: Know-It-All-Brother <354264050+Know-It-All-Brother@users.noreply.github.com>

…es, instead of after every product: the translated helper and the host's NDVE network sum each block of 8 in f32 and round it to f16, which matches 5,000 of 5,000 recorded A100 cases (per-product rounding 2,420) and 99.97 % of a Jetson Orin's outputs (34.9 %); chain 0c434f1b68df (PR seedds#1897) HMMA_BODY accumulated on simdgroup_half8x8, whose multiply-accumulate rounds the f16 accumulator after every product. Each block now accumulates into a simdgroup_float8x8 from the half operands and a half C and is rounded to f16 per lane through thread_elements(). test_dmsh_hmma.py's GPU runs compare every output bit for bit with the two-block reference (main's body differs in 21,461 outputs on an M5 Max and on an M6). ndve_engine.metal's hmma() (seedds#1894, bit for bit with the helper) rounds the same way, into a new matrix: written into its accumulator arrays' own thread_elements(), its layers ran 15 to 50 times slower on an M6. On the M6 the network's seven layer kernels take 2.94 ms against 2.85; isolated on an M5 Max, NDSR's fa3033d05c0a41d6 takes 0.316 ms against 0.276. Not measured in a game. Co-authored-by: Know-It-All-Brother <354264050+Know-It-All-Brother@users.noreply.github.com>

Know-It-All-Brother mentioned this pull request Nov 9, 2026

NDVE movie network: the translated HMMA shaders give different results from run to run #1893

Closed

1 task done

github-actions Bot added nvn NVN on Metal and its tests (dm/gfx/nvn, dm/gfx/tests) docs Documentation game: 0400486024DEC000 A game labels Nov 9, 2026

This was referenced Nov 9, 2026

Ocarina of Time on an M1 Max: the movie library's NDVE network takes 35 ms of GPU a frame, so movies and the title screen run at ~28 fps #1375

Open

HMMA with an f16 accumulator rounds once per 8 products, as Ampere does (chain change) #1897

Closed

seedds closed this Nov 9, 2026

seedds mentioned this pull request Nov 9, 2026

Scale decoded NDVE movies on Metal and skip matched reconstruction #1331

Closed

3 tasks done

Navigation Menu

Repository navigation

Run nn::ndve's movie network as host Metal kernels behind DM_NDVE_ENGINE=1 - #1894

Run nn::ndve's movie network as host Metal kernels behind DM_NDVE_ENGINE=1#1894
Know-It-All-Brother wants to merge 1 commit into
seedds:mainfrom
Know-It-All-Brother:ndve-engine

Know-It-All-Brother commented Nov 9, 2026

Uh oh!

seedds commented Nov 9, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

Know-It-All-Brother commented Nov 9, 2026

What and why

Two findings for the translator

Checked

Scope

Docs

Uh oh!

seedds commented Nov 9, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants