Repository navigation
IMMA (8-bit integer mma) translates (chain change) - #2056
Open
Know-It-All-Brother wants to merge 3 commits into
Open
Know-It-All-Brother wants to merge 3 commits into
Know-It-All-Brother wants to merge 3 commits into
Conversation
…re_orin.py runs about 190 operations, each the GPU compiler's own machine code for one PTX operation, through the translator, spirv-cross, the chain's MSL rewrites and Metal on this Mac's GPU, and compares every output word with what a Jetson Orin (sm_87) computed from the same 1,024 input rows. The translator tests checked the MSL a few encodings give, not the values; this one finds the FSETP.NE, SHFL lane, F2FP, rounding-mode and signed-zero differences the tracking issue lists, which KNOWN names until each fix takes its entry out. dm/gfx/tools/ampere_orin/capture.py records the fixture on any Linux machine with an Ampere GPU and NVIDIA's driver, with no toolkit or compiler on either side
…ules with its functions, so a rewrite that uses one (struct in PR seedds#1902's hfma_msl) runs as it does in a build instead of raising NameError
…bit accumulator) is translated and gives the Ampere GPU's results bit for bit, where the translator used to stop with "Ampere instruction 237". It had HMMA but not its integer twin. InstEmitAmpereTensor.cs now translates IMMA per thread as it does HMMA: each lane fetches the bytes of A and B it needs from the lanes that hold them with subgroup shuffles, in PTX's layouts for m16n8k32, m16n8k16 and m8n8k16, and sums the products onto C in 32-bit integers that wrap, or clamp with .SAT, with signed or unsigned bytes and the shape read where Mesa NAK's OpImma puts them; the 4-bit shapes still stop the translator. test_dmsh_ampere_orin.py's imma_16832_s8 went from not translating to 0 of 4,096 output words differing and leaves KNOWN; the other shapes, unsigned bytes and .SAT share the code but aren't in the fixture; chain a10a4bd270b8
github-actions
Bot
added
nvn
Know-It-All-Brother
added a commit
to Know-It-All-Brother/nxnative
that referenced
this pull request
Nov 9, 2026
…h's translator rebuilt from its own source; chain 8bee24f0d0e2
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Part of the Ampere accuracy series (#2040). It depends on #2041 for the test: the commits before this PR's own are #2041's and drop out once it lands, so this PR's own change is the last commit.
What and why
IMMA is HMMA's 8-bit integer twin (
mma.sync.m16n8k32.s32.s8.s8.s32, opcode 0x237). The translator had no emitter for it, so a shader with one failed to translate.AImmaworks per lane, like the translator's HMMA. Each lane gathers the A rows and B columns it needs from the lanes that hold them, with subgroup shuffles in PTX's documented fragment layout, then unpacks the bytes (signed or unsigned per the encoding) and sums the products in int32 onto C with wraparound (.SAT clamps that add). The products can't overflow, so the result is exact integer arithmetic and bit-exact. The fields come from Mesa NAK's IMMA: the shape from bit 75 and bits 85–86 (m8n8k16, m16n8k16, m16n8k32), A signed at bit 76, B signed at bit 78, and .SAT at bit 82. The 4-bit shapes still fail the shader.An m16n8k32 IMMA costs each lane about 32 shuffles and 128 byte extracts, multiplies and adds. It's a heavy instruction, as it is on the GPU.
Reference entry: "IMMA, the 8-bit integer matrix multiply" in
docs/reference/shaders.md.How to verify
imma_16832_s8failed to translate on main and now matches the Orin in all 4,096 output words, so it comes out ofKNOWN. Not checked against the GPU: the m16n8k16 and m8n8k16 shapes, unsigned bytes and .SAT, which share the code but aren't in the fixture.Checked
test_dmsh_ampere_orin.py,test_dmsh_hmma.py; every othertest_dmsh_*gives the same result as on main (test_dmsh_ampere_mem.py,test_dmsh_meshstage.pyandtest_dmsh_pairs.pyfail on main the same way);test_nsp_chain.py,test_xlate_anycpu.py,test_docs.py.tool_lock.py --update-lockanddmsh_chain.py --headerwere run (chain a10a4bd270b8).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. Ocarina of Time's build already translates every shader, so it can't contain this instruction.Docs
docs/reference/shaders.mdhas the entry, andpython3 dm/tools/docs_symptoms.pywas rundmsh_chain.py --headerwas run, anddmsh_chain.h, the translator dlls and the lock are part of the changelint_python.pyshows no new findings.