[beecrepaldi-afk]

What and why

On an M5, Donkey Kong Bananza's own upscaler (agl's NDSR, DM_UPSCALER=game) still takes about a quarter of the GPU frame docked (3.1-4.7 ms of 12-14 ms), with its HMMAs already through simdgroup_matrix (Ampere HMMA through simdgroup_matrix). Apple10 GPUs (M5) have neural accelerators in each core, which Metal 4 reaches only through tensor operations (MetalPerformancePrimitives).

dmsh_build.py now builds a second library for each shader with HMMA.16816, <key>_bt (hmma_tensor_msl): the same scratch block, with the product as one matmul2d 16x8x16 of one SIMD-group over threadgroup tensors. The reflection names it (hmma_tensor), and shader_load takes it on an Apple10 GPU only; other GPUs and iOS keep <key>_b, and DM_NVN_HMMA_TENSOR=0 turns it off. A tensor library that fails to build or load leaves <key>_b.

Reference entry: Ampere HMMA through Metal 4 tensor operations on Apple10 GPUs.

Checked

• Tests: dm/gfx/tests/test_dmsh_hmma.py (6 tests, including the tensor body on the GPU against the CPU reference of mma.m16n8k16), dm/tools/tests/test_docs.py, dm/tools/tests/test_nsp_chain.py: pass.
• Measurements, MacBook M5 16 GB, macOS 27.0.1, Xcode 27.2:

  • microbenchmark of the helper's block round trip: 1.57 ns a SIMD-group per HMMA through matmul2d against 1.84 through simdgroup_matrix, identical results;

  • Donkey Kong Bananza (040074A01BF12000, v458752), docked, DM_UPSCALER=game, DM_NVN_GPUPROF=every:300, two runs of about 2 minutes on one route, the 24 NDSR libraries swapped by hand on chain 7a8a6262ad3a:

    Frame	simdgroup_matrix	tensor
    1500	3.08 ms	2.38 ms
    1800	3.76 ms	2.60 ms
    2100	4.68 ms	3.64 ms
    2400	3.80 ms	2.59 ms
    2700	3.79 ms	3.81 ms
    mean	3.82 ms (29 % of GPU time)	3.00 ms (25 %)

    Five comparable frames only; macOS clocks the GPU to the frame, so the share of the frame's GPU time is the comparison that clock changes do not move.

  • this branch's host and chain ce59435185e4 on the same game: HMMA: Metal 4 tensor operations (<key>_bt libraries, Apple10 GPU), about 2 minutes at 60 fps (p50 16.6 ms), no new [RED] line.

• Regression (regress.py chain): Donkey Kong Bananza 23638/23638 → 23638/23638, 24 changed (reflection and variants only, no MSL); Mario Kart World (0400C3F00006E000, v0) 11561/11561 → 11561/11561, 30 changed (its NDSR upscaler, likewise). These are the only two games on this Mac, so no other game was run.
• Not checked: another Apple10 GPU (M5 Pro / Max), iOS (A19 is Apple10 too, but the build makes no _bt for iOS yet), Mario Kart World in game, and whether the pipeline prebuild / prewarm build the _bt functions or leave them to first use.

Docs

• The finding is in its game note or its docs/reference/<subsystem>.md entry, and python3 dm/tools/docs_symptoms.py was run (the symptom index)
• 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

🤖 Generated with Claude Code

[Know-It-All-Brother]

Thanks for this! I checked it against fxd0h's Jetson Orin HMMA recordings (gist) on an M5 Max (Apple10, Metal 4), and it rounds differently from the Orin.

This branch merged main before #1897, which made the f16 accumulator sum each block of 8 products in f32 and round once to f16, as Ampere does. matmul2d with a half destination sums all 16 products and rounds once (99.96 % of outputs match that model), which is closer than per-product rounding but not the hardware's. The tests here compare within a tolerance, so they couldn't see it; main's test now checks every bit.

Body (outputs equal to the Orin's, 3 seeds)	One call	Two chained	ns per HMMA
This PR, simdgroup_matrix	34.87 %	26.57 %	0.30
This PR, matmul2d (_bt)	72.25 %	62.91 %	0.245
main, simdgroup_matrix (#1897)	99.97 %	99.97 %	0.458
matmul2d, once per 8	99.98 %	99.97 %	0.40

The last row is what I'd suggest. Since matmul2d rounds its half destination once per run, two runs over one k-block each give #1897's rounding: load C into a half cooperative tensor, split B into two copies with the other block's k zeroed (the second copy can go in C's place once C is in the tensor), then run, run, store. That's still about 13 % faster than main's body. One caveat: "a half destination rounds once per run" is what I measured on this M5 Max, not something Apple documents, so it's worth a test that checks it bit for bit. An f32 cooperative destination rounded between two 16x8x8 runs also matches, but it was about 7x slower.

The branch also conflicts with main: main's HMMA_BODY now has chain slots and A loads from shared memory, and HMMA_SIMDGROUP slices on simdgroup_half8x8 a00, which is gone, so dmsh_build.py would fail at import after a rebase.

To rerun, use the gist's metal_run.py with HMMA_NX set to your checkout. The Orin suite in #2041 covers main's body; it doesn't build the _bt library, so it can't check this path until the tensor body sits on main's. I'm happy to share my runner for the _bt body.