[sheepyd]

What and why

This continues the open item of Ampere HMMA through simdgroup_matrix: it follows LDSM → HMMA → HMMA chains for the accumulator and for A. B and MOVM are not covered.

Symptom. Ocarina of Time's title and opening movies run the NDVE network. It is 11 compute shaders in cs_movie.bin with 2,217 HMMA.16816.F16, and they take 27 ms of GPU a frame on an M4 MacBook Air with 8 GPU cores. The title screen runs at 35 fps. The OoT profile is in #1299; I ran it on main plus #1299, #1300 and #1312.

Cause. Each dmsh_hmma16816_f16 call took every operand through the SIMD-group's scratch block. The lanes stored A, B and C, simdgroup_load read them back, and D returned to the registers the same way. That is three SIMD-group barriers per call. In these networks:

• 1,890 of the HMMAs continue an accumulation whose D nothing else reads (chains of 9 to 36).
• Four calls in a row share A.
• A comes from an LDSM whose rows form a strided tile in shared memory, and the LDSM's emulation costs four shuffles per lane.

Measured on the 4 frames from 1200 to 1203, with wrong pictures:

Variant	Compute per frame
Baseline	26.9 ms
Without the helper's body	15.4 ms
Without the helper's body or the shuffles	11.0 ms

Change.

• TensorPass.Prepare works on the optimised IR of each block and appends flags to every HMMA, as a constant argument after C:
  • An accumulator slot for a chain (8 per function). D stays in the slot and becomes the next call's C (FlagCSlot, FlagDSlot).
  • A that is the previous call's (FlagAReuse).
  • C = 0 (FlagCZero).
  • A that is four matrices of one LDSM in the same block, with no shared memory write or barrier in between (FlagADirect; FlagATrans for .MT88). The call also gets the LDSM's row address, each matrix's offset from lane 0's row, and the row stride.
• FlagADirect needs the row address to be U + f(lane), where U is the same in every lane and f is known for all 32 lanes. A divergence analysis over the function decides U:
  • The lane is divergent, and so are the components of the thread index that a subgroup spans.
  • A loop header's phis are uniform when every exit of the loop branches on a uniform condition.
  • Phis where paths meet are divergent.
  • f is then evaluated for each lane; the rows must be 16-byte aligned and strided.
• In a compute shader, an LDSM becomes one MatrixLoad per matrix, lowered to the old per-lane code after Prepare. There is one destination each because Optimizer.RemoveNode treats an operation as dead once one of its destinations loses its last use.
• The translator's own body still reads only the registers, so each flag is a promise about values that body computes anyway.
• dmsh_build.py's hmma_msl:
  • declares the 8 slots and A's four matrices in each function that calls the helper, and passes them by reference, so the inlined helper keeps them in registers;
  • passes the caller's shared memory array, which is the one threadgroup array besides the scratch. With no such array or several, FlagADirect is ignored and A goes through the block;
  • with FlagADirect, loads A by simdgroup_load at lane 0's row (simd_broadcast) plus the matrix's offset. The LDSM's per-lane code is then unused and the Metal compiler drops it.
• After this, a call in the middle of a chain stores only B.
• DMSH_TENSOR_TRACE=1 on a dmsh_xlate job prints why an HMMA's A is not loaded where its LDSM read it.

The new reference entry is HMMA chains keep their matrices in the SIMD-group. This is a chain change.

Size: about 1,180 added lines. 250 of them are the test. TensorPass.cs is 888 lines, of which about 400 are code; the rest are blank lines, braces and the 75 lines of the uniformity whitelist. That is over the roughly 1,000 lines in CONTRIBUTING.md, so tell me if you would rather have it split, for example the divergence analysis and FlagADirect as a second PR.

Checked

• Tests:

  • All 45 dm/gfx/tests/test_*.py files pass with DMSH_TEST_XLATE (385 tests together with test_nsp_chain, test_docs and test_xlate_anycpu).
  • test_dmsh_hmma.py now runs four flagged calls on the GPU against a CPU reference: A from shared memory, plain and transposed; A kept; C = 0; two accumulator slots. Each call gets the registers the translator's body would read, including a stale A where A comes from shared memory. The largest error is 0.014 for two chained HMMAs with an f16 accumulator between them. The single-call check still has a largest error of 0.009.
  • With DMSH_TEST_XLATE, test_dmsh_hmma.py also checks the flags the translator gives two chained HMMAs on an LDSM's matrices.
  • make -C dm/gfx/tests was not run.

• Games run:

  • The Legend of Zelda: Ocarina of Time (0400486024DEC000, v131072 NUP, base v0) on the Mac, on a branch of main plus [WIP] The Legend of Zelda: Ocarina of Time (Switch 2, 0400486024DEC000, v0, SDK 22.2.8) runs on the Mac through the title movie into Kokiri Forest #1299, Ampere F2IP and a task shader's own payload read: Ocarina of Time's untranslated shaders translate, all but one that is data #1300 and Zelda: Ocarina of Time — Restore missing mesh and texture effects on Mac #1312.
  • Only the 11 NDVE shaders were rebuilt with this chain; the rest of the library is from the previous chain. The dmsh_hmma16816_f16 body and the flags are identical to what this branch builds from main.
  • Measured on the title screen. Screenshots of it show the same picture as before, with no artefacts.
  • I did not check the opening movie with this build, and did not play into gameplay.

• Devices and OS: MacBook Air M4 (8 GPU cores, 16 GB), macOS 27. Not run on an iPhone.

• Regression check: scripts/regress.py check --games was not run. The only local game with this engine is OoT, and Donkey Kong Bananza (the other HMMA user) is not available here. Everything this change does is limited to compute shaders that contain an LDSM or an HMMA. Their MSL changes in two ways:

  • each LDSM's per-lane code is now emitted after optimisation;
  • the helper's signature and body are new.
    Bananza's upscaler (7629d857a0a93833) should be the first thing to check.

• Measurements: median over 321 profiled frames (frames 1000 to 2600, DM_NVN_GPUPROF=every:5, DM_NVN_DISPATCH_OWN=1) of the frame's compute on the title screen:

  	Compute per frame	Title screen
  Before	26.65 ms (quartiles 26.62 to 26.72)	35.5 fps
  After	18.85 ms (quartiles 18.82 to 18.89)	48.7 fps

  Of the 2,217 HMMAs, 540 load A where its LDSM read it and 1,249 keep the previous call's A. Of the other 428, 427 have a shared memory store between the LDSM and the HMMA. Most of those are in d15370a382630d4e.

• What is left is B's trip through the block. B is the weights: LDSM.MT88 reads them before the tile loop, they are kept in registers, and the shared memory is reused for activations. With B loaded from shared memory as well (a fake address, measured only), the frame took 12.6 ms. Keeping B as matrices in registers took it to 17.4 ms, but slowed 0a56002cdc7e1fce (4 SIMD-groups, 144 matrices) by 13 %, so it is not part of this PR. Removing the barrier at the end of a mid-chain call changed nothing (18.84 against 18.88 ms).

Docs

• The finding is in its docs/reference/<subsystem>.md entry, and python3 dm/tools/docs_symptoms.py was run (the symptom index). It is not in a game note yet, because OoT's note comes with [WIP] The Legend of Zelda: Ocarina of Time (Switch 2, 0400486024DEC000, v0, SDK 22.2.8) runs on the Mac through the title movie into Kokiri Forest #1299.
• 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

[seedds]

Landed as bd9f422, with you as co-author. Thank you, this is a clean piece of work: flags that are promises about values the translator's own body computes anyway kept every fallback correct.

What changed on the way in:

• Rebased onto main's chain (now 23a1ea1db092): the translator rebuilt, the lock and header updated, gather_offsets kept beside your hmma_msl.
• The divergence analysis also treats a loop header's phis as divergent when the loop has several back edges and any branch inside it depends on the lane: lanes that took different back edges bring the header different values. It changes none of Ocarina of Time's 71 HMMA shaders (traced with DMSH_TENSOR_TRACE).
• The reference entry links OoT's note, and the note has the new figures.

Checked here on an M6 (12 GPU cores): the title screen's compute 14.78 → 11.41 ms a frame and the GPU span 15.63 → 11.87 ms, main against this change. The title and opening movie are clean on both. OoT's 42,854 shaders build as on main. Donkey Kong Bananza's 23,638 shaders build and its opening and first play render cleanly. agl's NDSR upscaler (in both games) takes the chain flags, but none of its A loads directly: its LDSM rows are addresses read from a table in shared memory, so A goes through the scratch there. That's noted in the entry.

[jonpage0]

On Ocarina of Time's gameplay upscaler (agl_ndsr_vr.pck) this rewrite draws regular dotted rows across the whole 3D scene on an M3 Max, from the first frame in Kokiri Forest; the title and opening movie are fine, which is what the commit checked. Bisected on one app down to the 89 network shaders rebuilt with hmma_msl returning the MSL unchanged (clean) — details, screenshots and what was excluded in #1547.