[Know-It-All-Brother]

What and why

The translator's HMMA helper accumulates on simdgroup_half8x8, which rounds the f16 accumulator after every product. Ampere's HMMA with an f16 accumulator adds 8 products to C and rounds once. So an HMMA.16816 is k 0–7 rounded to f16, then k 8–15 on top of that, rounded again.

I checked this against the 5,000 recorded A100 fp16 cases in Khattak and Mikaitis's MATLAB-tensor-core toolbox (model_validation/A100/fp16, C rounded to f16 as HMMA's accumulator is):

Model	Cases matching the A100
Per-product f16 rounding (main)	2,420 of 5,000
8 products summed in f32, rounded once	5,000 of 5,000

The change: HMMA_BODY accumulates each block of 8 into a simdgroup_float8x8 from the half operands and a half C (simdgroup_multiply_accumulate takes mixed types), then rounds it to f16 per lane through thread_elements().

It's a chain change, so every Switch 2 game's HMMA shaders rebuild: agl's NDSR upscaler and the NDVE movie networks among them.

Reference entry: HMMA rounds its f16 accumulator once per 8 products.

Checked

• Tests:

  • test_dmsh_hmma.py with DMSH_TEST_XLATE: 12 tests pass. Its GPU runs now compare every output bit for bit with the two-block reference instead of within a tolerance. main's body fails that on 21,461 outputs.
  • docs_symptoms.py --check and test_docs pass.
  • lint_python.py --against origin/main: no new findings in the changed files. The 3 new pyright findings are in dm/tools/unity_bucket.py, from main.

• Games run, with title ID and version: none in-game yet. Ocarina of Time's (0400486024DEC000, v131072) NDSR shaders ran isolated on synthetic inputs, and their outputs change in 1.5% (fa3033d05c0a41d6) and 0.8% (d99fbeeb53d38847) of halves.

• Devices and OS: MacBook Pro M5 Max (64 GB), macOS 27.0.1, Xcode 27.2.

• Measurements, isolated, best of 200 runs on the M5 Max:

  Shader	main	This PR
  NDSR fa3033d05c0a41d6	0.276 ms	0.316 ms
  NDSR d99fbeeb53d38847	0.393 ms	0.446 ms

  Rounding through a multiply-accumulate with zero A and B gives the same bits but took about twice as long, so I round per lane.

• Not checked: regress.py check on the other Switch 2 games, in-game picture and frame time, iOS.

Related

• This doesn't fix the lane race in the translated NDVE network (NDVE movie network: the translated HMMA shaders give different results from run to run #1893). That's a separate bug.
• Run nn::ndve's movie network as host Metal kernels behind DM_NDVE_ENGINE=1 #1894's engine copies this helper's arithmetic. If this merges first, I'll update the engine to round the same way.
• Ampere HMMA through Metal 4 tensor operations on Apple10 GPUs: Donkey Kong Bananza's own upscaler goes from 3.82 to 3.00 ms a frame on an M5 #1334 runs HMMA through Metal 4 matmul2d on Apple10 GPUs. That path would need the same once-per-8 rounding.

Docs

• The finding is in docs/reference/shaders.md ("HMMA rounds its f16 accumulator once per 8 products"), and python3 dm/tools/docs_symptoms.py was run
• A chain tool changed: dmsh_chain.py --header was run and dm/gfx/nvn/dmsh_chain.h is part of the change
• No game files, keys or data extracted from a game in the change or the attachments

[fxd0h]

Measured the same thing on real Ampere on Nov 7, a Jetson Orin (sm_87): mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 through inline PTX, 3 seeds × 4,096 warps, inputs uniform in [-1, 1].

main's body matched the Orin bit for bit on 34.9% of outputs, and 26.6% for two chained HMMAs. Summing each block of 8 in float and rounding once matched 99.97% in both cases. The 0.03% left were near-cancellations around zero, which looked like my numpy model rather than the hardware. So the Orin agrees with your A100 numbers.

[Know-It-All-Brother]

@fxd0h I have Orin Nanos, that's a great idea to validate that way.
I wonder if I could get it running on one?

[fxd0h]

yeah it's a tiny cuda kernel, calls the mma through inline ptx with random halves and dumps the bits, then I just diff against metal on the mac: https://gist.github.com/fxd0h/ee97de9e608d5d93932f7855a73973d4

[seedds]

Landed as 356a713, with you as co-author. Thanks, and thanks to fxd0h for the Orin run; its numbers are in the reference entry.

What changed on top of the PR:

• The host's NDVE network rounds the same way. Run nn::ndve's movie network as host Metal kernels behind DM_NDVE_ENGINE=1 #1894 landed before this, and test_ndve_engine.py checks that its kernels match the helper bit for bit, so ndve_engine.metal's hmma() now sums each block of 8 in f32 and rounds once too. It rounds into a new matrix: writing the rounded halves into the accumulator arrays' own thread_elements(), as HMMA_BODY does, made its layers 15 to 50 times slower on an M6. With a new matrix, the seven layer kernels take 2.94 ms against 2.85. In the translated helper both forms cost the same, so HMMA_BODY is yours unchanged.
• Rebased onto the current chain (now 0c434f1b68df), with the code comment about the zero multiply-accumulate reworded, and a line added to Ocarina of Time's note.

Checked on an M6: test_dmsh_hmma.py passes, and main's body differs in 21,461 outputs there, the same count as on your M5 Max. test_ndve_engine.py, test_docs and test_nsp_chain pass too. No in-game run yet.