[Know-It-All-Brother]

Part of the Ampere accuracy series (#2040). It depends on #2041 for the test and builds on #2061 (the f32-accumulator model): the commits before this PR's own are #2041's and #2061's and drop out once those land, so this PR's own change is the last commit.

What and why

HMMA with bf16 sources (mma.sync m16n8k16 .f32.bf16.bf16.f32) and with tf32 sources (m16n8k8 .f32.tf32.tf32.f32) failed to translate: AHmma only took f16 sources.

This adds both source types and routes them through #2061's tensor-core model, which matches the Orin exactly for these too:

• bf16 uses the f16 m16n8k16 fragment layout. Each element is its bits moved up 16. Blocks of 8 products, as for f16.
• tf32 has its own m16n8k8 layout, one element per register (A: (g, t), (g+8, t), (g, t+4), (g+8, t+4); B: k = t and t+4 in column g). The tensor core ignores the low 13 bits of each tf32 input, and with them kept, no output matched. Blocks of 4 products.

To do this, AHmmaF32 now takes the source type: it sets the block size, the exponent width and how each element is unpacked. bf16 and tf32 need an f32 accumulator, tf32 needs m16n8k8, and the compute shaders' simdgroup_matrix route still takes f16 only.

One change touches #2061's code: the model's final right shift is clamped to 31. A sum below f32's denormals from an 8-bit-exponent source could ask for a shift of 32 or more, which is undefined in MSL. For f16 sources the results are unchanged, at the cost of one min per block.

These are heavy instructions: about 1,200 MSL statements for a tf32 HMMA and 1,900 for a bf16 one. In Ocarina of Time's shader code, every HMMA but 2 has f16 sources and an f16 accumulator, and that path doesn't change.

Reference entry: "HMMA with bf16 and tf32 sources" in docs/reference/shaders.md, which links the f32-accumulator entry.

How to verify

DMSH_TEST_XLATE=dm/gfx/tools/dmsh_xlate_bin/dmsh_xlate AMPERE_ORIN_ONLY=tensor python3 dm/gfx/tests/test_dmsh_ampere_orin.py

Op (outputs of 4,096)	main	with #2061 alone	this PR
hmma_16816_bf16	failed to translate	failed to translate	0
hmma_1688_tf32	failed to translate	failed to translate	0
hmma_16816_f32	2,693	0	0

Both come out of KNOWN. Before applying #2061's model, a plain FMA chain for these left 1,334 and 2,189 outputs differing. The fixture's bf16 and tf32 inputs all have ordinary exponents and every output is normal. So these aren't checked against the GPU: product overflow, sums below f32's denormals, denormal / Inf / NaN sources, and a tf32 NaN whose mantissa bits are all in the ignored low 13 (it reads as Inf).

Checked

• Tests: test_dmsh_ampere_orin.py, test_dmsh_hmma.py; every other test_dmsh_* gives the same result as on main (test_dmsh_ampere_mem.py, test_dmsh_meshstage.py and test_dmsh_pairs.py fail on main the same way); test_nsp_chain.py, test_xlate_anycpu.py, test_docs.py. lint_python.py shows no new findings.
• Translator: the dlls are rebuilt from this tree with .NET 10; tool_lock.py --update-lock and dmsh_chain.py --header were run (chain b0cbe916d032).
• Devices and OS: MacBook Pro M5 Max, macOS 27.
• regress.py: not run. My only local Switch 2 build's shader package comes from an older shader split than main's map, so main's chain can't build it.
• Games run: none.

Docs

• docs/reference/shaders.md has the entry, and python3 dm/tools/docs_symptoms.py was run
• A chain tool changed: dmsh_chain.py --header was run, and dmsh_chain.h, the translator dlls and the lock are part of the change
• No game files, keys or data extracted from a game