feat: add non-TN support for gfx950 mxfp4 with subtile impl - #11369
Open
WorldofKerry wants to merge 147 commits into
Open
WorldofKerry wants to merge 147 commits into
WorldofKerry wants to merge 147 commits into
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
WorldofKerry
changed the title
WorldofKerry
commented
Aug 31, 2026
WorldofKerry
changed the title
WorldofKerry
changed the title
WorldofKerry
force-pushed
the
gfx950-transposed
branch
3 times, most recently
from
September 15, 2026 22:27
3ddf2e6 to
1015150
Compare
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
The emitted transpose-load mnemonic is invalid and non-TN client references consume swizzled scales as though they were canonical.
Get a fresh assessment by requesting another Copilot review.
Pull request overview
Extends gfx950 MXFP4 subtile GEMM support from TN to NT, NN, and TT layouts.
Changes:
- Adds TLU=1 geometry planning, LDS layouts, cooperative reads, and transpose loads.
- Updates scale pre-swizzling and StreamK addressing.
- Adds broad unit, characterization, and gfx950 correctness coverage.
Review assessment: The test breadth is strong, but unsupported ISA mnemonics and incorrect non-TN client reference scaling are blocking. The PR description also names the wrong test YAML and lacks reproducible test/performance evidence and work-item tracking.
Suggested PR Description Updates
- Replace
subtile_mxfp4_tlu1.yamlwithsubtile_mxfp4.yaml. - Add the exact gfx950 test commands and device results.
- Include baseline versus updated MT256 performance measurements.
- Link the corresponding work item and document the cross-architecture instruction risk.
File summaries
| File | Description |
|---|---|
projects/hipblaslt/clients/common/src/mxDataGen.cpp |
Adds orientation-aware gfx950 scale pre-swizzling. |
projects/hipblaslt/tensilelite/client/src/DataInitialization.cpp |
Allocates padded scale staging storage. |
projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py |
Updates swizzled-scale StreamK offsets. |
projects/hipblaslt/tensilelite/Tensile/KernelWriter.py |
Applies centralized subtile LDS layouts. |
projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py |
Selects and validates TLU=1 stack geometry. |
projects/hipblaslt/tensilelite/Tensile/Components/Subtile/Kernel.py |
Integrates cooperative TLU=1 planning. |
projects/hipblaslt/tensilelite/Tensile/Components/Subtile/SubtileGeometry.py |
Adds cooperative-read geometry planning. |
projects/hipblaslt/tensilelite/Tensile/Components/Subtile/SubtileGREmit.py |
Implements TLU=1 global-read emission. |
projects/hipblaslt/tensilelite/Tensile/Components/Subtile/SubtileLREmit.py |
Implements transposed LDS reads. |
projects/hipblaslt/tensilelite/Tensile/Components/Subtile/SubtileLdsLayout.py |
Centralizes operand and scale LDS allocation. |
projects/hipblaslt/tensilelite/Tensile/Components/Subtile/SubtileTLUSwizzle.py |
Adds bank-conflict-avoiding layouts. |
projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx950/subtile_mxfp4.yaml |
Expands gfx950 layout and geometry coverage. |
projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_gr_coop_spread.py |
Tests cooperative-read planning. |
projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_subtile_stack_for_tile.py |
Tests stack selection and fallback. |
projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/test_r5_subtile_kern_char.py |
Characterizes implemented DTL dispatch. |
projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/mem.hpp |
Defines the transpose-load instruction. |
projects/hipblaslt/tensilelite/rocisa_stinkytofu_adaptor/rocisa_stinkytofu_adaptor/instruction.py |
Mirrors the instruction in the Python backend. |
Review details
- Files reviewed: 20/20 changed files
- Comments generated: 3
- Review effort level: Balanced
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
The characterization golden update lacks the repository-required ADR documenting the intended behavior change.
Get a fresh assessment by requesting another Copilot review.
Review details
- Files reviewed: 27/27 changed files
- Comments generated: 1
- Review effort level: Balanced
Add the TLU=1 (NT) path for UseSubtileImpl kernels, where both operands are free-dim contiguous in global memory. Global reads become a per-lane K-ramp with a single wide buffer_load spanning the free dim, and transpose local reads use ds_load_b64_tr_b4. Also drop the StinkyTofu MatrixInstructionNaming pre-check (module, validator wiring, and its tests) that this path supersedes. Adds nt_fp4_bringup.yaml (gfx950 NT MXFP4 subtile bring-up).
The TLU=1 (NT) subtile global-read pointer update advanced the SRD by DepthU * bpe bytes per main-loop iteration. For NT the free dim is unit-stride and the K/unroll dim is strided, so one DepthU K-window is DepthU * strideK * bpe bytes. The missing strideK factor only cancels when the main loop runs once (K == DepthU), so any K > DepthU produced incorrect results. Scale the increment by the runtime K stride in the TLU=1 branch, keeping a unit-stride fast path for the non-transposed operand. Rename nt_fp4_bringup.yaml to f4_nt.yaml and sweep sizes covering single/multi K-iteration and single/multi workgroup.
The NT (TLU=1) fp4 subtile path wrote each operand free-dim contiguous, K-major, into LDS. The ds_read_b64_tr_b4 transpose reads then hit a 2-way LDS bank conflict: the 32 lanes of each read half map onto only 32 of the 64 banks (see format.md analysis). Apply a chunk-index XOR swizzle (chunk[6] ^= chunk[5]) plus an 8-byte pad inserted between load-blocks. The XOR is an involution, so the GR write (which global K-row each lane fetches) and the LR read (which chunk each lane addresses) apply the same transform and the LDS image round-trips A exactly. The transform is centralized in SubtileTLUSwizzle and keyed by stack size (subtileShape[0]); only the 2x1 fp4 stack is wired, all others keep baseline addressing. Measured on gfx950 at 4096x4096x4096: SQ_LDS_BANK_CONFLICT drops from 8388608 to 0 and SQ_LDS_IDX_ACTIVE halves, with validation still PASSED.
The NT (TLU=1) fp4 subtile path only worked at MT 32x32 (a single 32x128 subtile strip per operand). Larger macro tiles crashed because the GR emit indexed its soffset register list with a scheduler tileId in MMA-tile units, overrunning the per-subtile-row list, and neither the GR write base nor the LR transpose-read base applied a per-strip LDS offset. Map the scheduler tileId to (subtileRow, mTileWithinStrip): the GR write and LR read now place strip subtileRow at subtileRow*stripStride in LDS and select the M-tile within the strip. Fill the per-subtile-row global-M soffset for the TLU path (free dim is unit-stride, so stride is subtileM*bpe). Centralize the pad-aware stripStride in SubtileTLUSwizzle so GR write, LR read, and the LDS size computation agree, and fold the swizzle pad into the LDS footprint. Single-wave macro tiles from 32x32 up to 128x128 (2x2 and 4x4 subtile grids, including non-power-of-2 like 96x96) now build and validate, with LDS scaling accordingly (8KB -> 48KB). Multi-wave (MIWaveGroup>1) partition offsets remain a separate follow-up. f4_nt.yaml gains a second problem block exercising MT 64x64 and 128x128.
Extend the NT (TLU=1) fp4 subtile path to MIWaveGroup>1 so wg2x2 (4 waves) and single-axis wg2x1/1x2 macro tiles work. Waves partition the free dim (M for A via MIWaveGroup[0], N for B via MIWaveGroup[1]) into disjoint subtile strips. Three coupled per-wave offsets are added, all keyed on the wave's axis id (waveId%mWaves for A, waveId//mWaves for B): the GR global read starts at the wave's first strip (axisId*localSub0*subtileM elements, unit-stride free dim); the GR DTL LDS write base and the LR transpose-read base target the wave's region of the full-macro-tile LDS image (axisId*localSub0*stripStride bytes). Also fix the cooperative-load math: bytesPerLoad folds in numWaves for TLU=0 (waves share a strip's K), but TLU waves partition free-dim strips and load their own K independently. Using numWaves there halved numGRPerSubtile and undersized the soffset register list, so drive both the load-ratio and localGRGranularity from a single-wave grLoadWaves for TLU. Derive the swizzle strip pad from DepthU/wavesize so it is identical across all register-allocation snapshots. Verified on gfx950: wg2x1, wg1x2, and wg2x2 from MT64 up to MT256x256 build and validate, single-wave regression stays green, and the LDS bank conflict counter stays 0 for the multi-wave kernel. f4_nt.yaml gains a wg2x2 (4-wave) MatrixInstruction variant.
NT (TLU=1) subtile GEMM stacked one MFMA K-window per DepthU. Extend it to DepthU that spans multiple K-windows (e.g. DepthU 256 = 2 windows of instK 128) so each strip holds several K-windows in LDS. Global read: for NT the free (M/N) dim is unit-stride and K is strided, so stepping to K-window sId1 must advance the GLOBAL address by sId1*instK*strideK*bpe bytes. The prior code encoded this as a DTL offset12 immediate, which only shifts the LDS m0 write and left window sId1>0 fetching window 0's global rows. Move the K-window step onto soffset (scaled by the runtime K stride) and drop the offset12 trick on the TLU=1 path. Local read: transpose read adds sId1 * globalSubtileGrid[0] * stripStride to reach the window's LDS strip, matching the GR m0 placement. Swizzle pad: a strip spans exactly one K-window, so swizzlePadPerStrip counts load-blocks per K-window (instK/waveSize) instead of per full DepthU; this keeps stripStride correct when DepthU > instK. f4_nt.yaml gains DepthU 256 coverage across MT32x32, MT64x64, and MT128x128 wg2x2. Bank conflict stays 0 (verified via rocprofv3 PMC).
Exercise DepthU 256 combined with the combined-subtile macro tiles up to MT256x256 wg2x2. The 256x256x256 tile fits gfx950 resources (144KB LDS of 160KB, 496 VGPR) and keeps the LDS bank conflict at 0 (verified via rocprofv3 PMC). f4_nt.yaml now covers 26 problems across DepthU 128/256 and MT 32x32 through 256x256.
The span is roundUp(ceil(K/mxBlock), 8) * 32 on both layouts, but only the TLU=1 branch computed it. The TLU=0 branch shifted Strides<tc>+0 left by 5, which reaches the same number only because that register happens to hold paddedKBlocks when the free dim is slow and the host did pad. That is an assumption about the host layout the kernel cannot check, and the register it reads is the same one three separate attempts at a StreamK bug reached for by mistake. Derive it from K on both. Verified the two forms agree at K = 256, 288, 1024 and 4096 before making the change, and subtile_mxfp8 -- nine TN groups, the layout that previously took the shift -- still reports 107 passed / 0 failed. Costs four SALU in the prologue, once per kernel. The padding and the grouping are now named, SWIZZLE_GROUP_ROWS and K_BLOCK_PAD, rather than appearing as 5, 7, 3 and 8 across five lines where they are easy to mistake for each other. Both numbers come from preSwizzleScalesGFX950, so the calculation is swizzle-specific by construction. Making that explicit is what lets the next commit refuse NoSwizzle rather than leave it silently wrong. 7408 unit tests, 764 snapshots unchanged, subtile_mxfp4 9953 / 0 / 131. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The subtile scale group span is derived from the swizzled layout -- the K-block padding and the 32-row grouping are both properties of preSwizzleScalesGFX950 -- so a canonical scale buffer is read with the wrong stride. Nothing rejected the combination, so it produced faulting or wrong-answer kernels instead of refusing. Measured on gfx950, MXScaleFormat=NoSwizzle with UseSubtileImpl=1: TN K=256, StreamK 0 and 3 FAILED, 5 illegal memory accesses NN aborts during kernel generation NT K=256, 288 passes, the one case the padding happens to suit StreamK=0 failing rules out the partial-tile K stride fixed earlier in this branch; this is the scale addressing itself. Nothing selects it: MXScaleFormat=Auto resolves to HostPreSwizzle on gfx950 and InMemorySwizzle under TDM, and all 30 MXScaleFormat settings in the gfx950 suites are 1. So refuse it, rather than implement a canonical-layout scale path nobody asked for. The TN probe that produced 5 illegal accesses now yields 0 valid solutions and 0 faults, with the reason printed. state.get for UseSubtileImpl because this runs during MX format derivation, before the key is necessarily present -- test_MXScaleLayoutDerivation calls it with a minimal state, and absent means not subtile. 7408 unit tests, 764 snapshots unchanged; subtile_mxfp8 107 / 0 / 9 and subtile_mxfp4 9953 / 0 / 131, neither of which touches NoSwizzle. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…tile K unit numSubIterK answered scales with a post-hoc clamp bolted onto the side of a per-operand rule. Scales are a per-operand property, so answer them where every other property is answered and let the max do the rest. Equivalent for all 64 input combinations: subIterKFor returns only 1 or 2, so the old clamp yielded 2 whenever either operand had scales, which is what max(2, other) now yields. 764 codegen snapshots unchanged. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
AB_B16_TLU1 is selected for bf16/half whenever TLU=1, but nothing gives the
free dim the element multiple its 16B chunk needs, the way the fp4 branch sets
AssertFree{0,1}ElementMultiple to 32. An NN bf16 UseSubtileImpl config therefore
cleared SolutionStructs (300 of 372 solutions) and then died in kernelBodySubtile
on a bare AssertionError rather than failing cleanly.
Reject the combination, matching the NoSwizzle precedent. Enabling it instead
would generate kernels on an untested path and needs its own coverage.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
generateData swizzled scaleBytes in place, but the reference path indexes scales canonically: getAlignedFloat uses scale_id = kBlock * M + m. The early return covers transposed A and non-transposed B, so the layouts that fall through -- the TLU=1 ones this branch enables -- would build their reference from unrelated scale entries and reject a correct GPU result. Swizzle a copy for the device buffer and leave scaleBytes in generator order. The bytes written to the device are unchanged in every case, so configurations covered today (all preswizzled gtests are transA T / transB N, which return early) are unaffected. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
ldsRowBankSize was recomputed inline at three sites in SubtileGREmit, and the row-major split into loads and bank rows, with its two asserts, was duplicated when _graTileAssignment_rowMajorSingle was factored out for NN/TT. Give both a name and call them. Also spell out the writer attribute each LDS region publishes instead of building it with a format string, so the names the emit paths read are greppable to their source, and drop SubtileLdsLayout.offsetOf, which has no callers. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…unner The TLU=1 work is exercised only by the gfx950 common suites, which are filtered out on a runner without that GPU, so the coverage job saw SubtileGREmit drop to 52% and SubtileLREmit to 63% as the files roughly doubled in size. Add one curated NN logic file to the codegen characterization set, which drives the real emitter with no GPU. NN is the shape that pairs a TLU=1 operand with a row-major one, so a single file reaches both the col_scatter stack and the row-major pairing path: SubtileGREmit 52.76 -> 73.18%, SubtileLREmit 63.58 -> 91.22%, the latter above its pre-branch baseline. Measured against the alternatives before settling on one file: NT covers TLU=1 on both sides but never the pairing path, TT mirrors NN, and a StreamK variant was byte-identical in coverage to plain NT. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
… paths The one NN logic file left SubtileGREmit at 73.18%, under its 85.92 floor in coverage-baseline.json, which the ratchet enforces with a 1pp tolerance. Add unit tests driving the TLU=1 global-read offset helpers directly. Which split a geometry turns on is chosen by Solution.py from the tile shape, so a curated logic file reaches one combination at a time; calling the helpers lets each branch be stated as a case. They assert on the emitted stream, the returned flags and register balance, not on address values, which the gfx950 suites cover. Add the short-stack NT file as well: NN selects a 16x1 stack and so always takes the col_scatter path, leaving the XOR-swizzle branch of _graTileAssignment_tlu unreached by any curated file. SubtileGREmit 73.18 -> 86.12%, SubtileLREmit 91.22 -> 94.96%, both now above their floors without relying on the tolerance. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
findConfigs walks the whole Tests tree, so the characterization data files are offered to configMarks, which guards against them by testing for a top-level sequence. The LibraryLogic step emits a mapping today, so the two TLU=1 files captured from it passed that guard and crashed on doc["BenchmarkProblems"], taking out collection for test_config, test_config_build and test_config_run. Convert both to the sequence form every other per-arch data file uses, and test for the BenchmarkProblems key rather than the document shape so a future capture in either form is skipped instead of crashing. The conversion is lossless: rawLibraryLogic returns an equal tuple before and after, and the emit digests are unchanged, so the goldens need no update. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
SubtileGeometry gained the strided-subtile mapping on this branch, which no unit test reached, dropping the file to 83.55% against its 86.89 floor. subtileForMmaTile is pure geometry, so test it directly: the contiguous case, the zero-stride fallback, the strided case where a subtile's rows interleave with its neighbours', and a self-consistency sweep that every row lands in a subtile containing it. The interleaved case is the one worth pinning, since getting it wrong still returns a plausible contiguous answer. SubtileGeometry 83.55 -> 90.60%, above its floor. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The gfx950 emit golden for MX_Subtile_TLU1_NN.yaml was generated locally, and CI computes a different kernel basename for the same logic file. Running --snapshot-update locally is a no-op, since it can only re-pin the local value, so take the digest CI reports. The divergence itself is unexplained: CUCount, the LdsPad solution keys and the HasLDSTrB64B4 assembler probe were each tested and ruled out, and the basename is stable across runs on either side. Only this file diverges; the twelve pre-existing gfx950 files and the sibling NT_Stack2 file agree. This unblocks the coverage lane, which runs the characterization suite first and so never reached coverage.json while the golden mismatched. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
col_scatter had no test of its own. Its derivation ran only incidentally, via the K-slot tests, and its emit side only through one characterization file, so a wrong thread-bit placement would not fail to build -- it would build a kernel that reads the wrong LDS bytes. Test the properties the layout exists to provide rather than a table of numbers: bit 3 carries col_group[gdBit] (the bank-pair separation), the thread bits are a collision-free permutation, every lane decodes to a distinct (m_chunk, col_group), and the K step is a non-zero per-lane constant -- the degenerate case the N=32 guard exists for. The stack-16 step is anchored to the offset:256 the shipped kernels actually emit. Also drop the 'verified against the bank model' claims: that verification is not checked in, so the comments now say what the code does, not what was measured. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Syrupy reports only the first differing entry, so pinning NN surfaced the same divergence on the sibling NT_Stack2 file. Pin it the same way. The cause looks to be the toolchain, not the test: the harness builds its capability map by probing the local assembler (makeIsaInfoMap(SUPPORTED_ISA, cxx)), and CI runs clang 24.0.0 against clang 22.0.0git here. Python is ruled out -- 3.12.3 and 3.14.6 agree locally, while CI's 3.10.8 and 3.12.10 agree with each other, so the split is by machine. The twelve pre-existing gfx950 goldens match on both toolchains, so only the fp4 subtile kernels are sensitive to whichever capability differs. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Comments and docstrings that narrated the change rather than the code they sit on, plus function specs longer than the bodies they document. Two docstrings explained themselves against a previous implementation: the s_waitcnt builder claimed to replace "the old comment-marker scheme", and the LDS token back-edge map called the identity carry "the classic behaviour". Both now state the property that holds. Two subtile_mxfp4 configs justified their K values by a crash they once reproduced; they now name the path they cover. The s_waitcnt docstring also carried an em dash. Trimmed eight docstrings on functions this branch adds where the spec outweighed or duplicated the body -- _b128ChunkTiling restated a decomposition its caller already comments, and _allocLROffsetRegs_tlu spent eight lines on a four-line body. The offset formulas on the large emitters stay: those bodies emit instructions rather than perform the arithmetic, so the equation is the only statement of the result. Comments only. No executable line changes, verified by filtering the diff for assignments, statements and braces. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Matching a mnemonic string was the wrong fix for the CI failure. The root cause is that VAddU32 has two spellings: native rocisa renders it v_add_u32 on gfx9, while rocisa_stinkytofu_adaptor/instruction.py:1379 defines it as v_add_nc_u32 unconditionally. Which one appears depends on the active backend, so any string assertion is a coin flip that happened to land differently on CI. Check isinstance against the instruction class instead, the convention already used in test_r6_subtile3_char.py, and apply it to every mnemonic assertion in the file rather than only the two that failed. Text matching is kept for the one operand check, since register names are stable across backends. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The scale GR thread split divides Serial by the scale K-subtile count. That count is lrSubtileShape[0] * mmaTileSize * (DepthU // instK) / loadWidth, and every factor except DepthU // instK is a power of two, so a DepthU that is merely a multiple of 2 * MatrixInstK is not enough: DepthU 768 over instK 128 gives 6, where int(math.log2(6)) truncates to 2 and the mask 6-1 is not a modulo. Reject that at solution validation instead of widening the split to cover it. The emit goes back to a shift and a mask, dropping the reciprocal-multiply fallback and its two scratch VGPRs. Emitted assembly is unchanged for every DepthU that still validates: the gfx950 characterization digest is identical with and without this change. subtile_mxfp4 loses its DepthU 768 group and drops 768 from the DepthU fork, since both are now rejected. No other MX config in Tensile/Tests uses a non-power-of-two DepthU. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Replaces the two curated characterization logic files with unit tests that drive the same emitters directly. The goldens for those files were not reproducible: they depend on the assembler version (clang 22 here, 24 on CI), so their digests had to be pinned to CI's values and were red for anyone on a different ROCm. CUCount, the LdsPad solution keys, the HasLDSTrB64B4 probe and the Python version were each tested and ruled out; the cause is still unknown. A golden nobody can regenerate does not do the job a golden exists for, so drop both files and restore the .ambr untouched. The tests build a real TileInfo from the shipped geometry table and stub only the writer, so the subtile shape, swizzle selection and register allocation are genuine. The stack height picks the layout, which is how one fixture reaches both: 2x1 takes the chunk-index XOR, 16x1 column-scatter, and a shared strip reroutes even the short stack. They assert branch selection, register allocation and that temporaries are returned -- not assembled output, which the gfx950 hardware suites establish. SubtileGREmit 67.90 -> 86.12%, SubtileLREmit 64.33 -> 92.36%, SubtileGeometry 90.60%, all above their floors without relying on the ratchet tolerance. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The coverage gate measured SubtileLREmit at 86.67% on CI against 87.08 with tolerance, while the same commit measures 92.36% here. CI runs a much smaller selection -- its last report showed 139 skipped and 143 deselected against 29 skipped locally -- so the coverage other test files contribute on a developer box is not there, and this file has to stand on its own. Cover _lraTileAssignment_fp8_legacy, which pairs both operands off the writer's tile state. That lifts this file's standalone LREmit coverage from 50.41 to 58.37%, 41 lines that no longer depend on another test happening to run. Also stub the LDS start offsets applyLdsLayout publishes, which that path reads. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Two docstrings belonging to pre-existing code were trimmed while tidying the subtile comments. Neither function's code is touched by this branch, so the edits drop context we did not write and cannot vouch for: _make_swaitcnt lost the note that the current scheme replaced comment-marker plus assembly-text post-processing, which left a combined s_waitcnt that waitReconstruction could not rebuild and perturbed scheduling. _ldsTokenBackEdgeMap lost that an empty map is the classic behaviour the binary 0<->1 toggle already amounts to, and that _nextLdsToken is the same primitive codegen swaps with. Audited the rest of the branch for the same thing: no other function has unchanged code with an edited docstring, no Python file is a comment-only change, and the two remaining comment-only hunks are a comment above code this branch does change and a newly added one. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The file had no header at the develop base; the block was added when the TLU=1 groups were folded in and grew from there. Most of it restated what the config already shows: which groups cover which layout, and that MT is 16 * stack * waveGroup when the MacroTiles are in the fork parameters. The copyright/SPDX lines go too: test yamls here do not carry one (1 of the 48 pre-existing files in this directory does, 41 of 414 across common), and this file did not have one before this branch. Parsed document is unchanged -- 130 groups, identical when loaded. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Solution.py carried 178 lines of subtile layout rules -- strip heights, wave sharing, the GR K-partition bug check -- interleaved with the rest of solution validation. Move them to SolutionStructs/Validators/Subtile.py, next to the MatrixInstruction, MXScaleFormat and WorkGroup validators, and import the three entry points Solution.py actually calls. Validators is the right home rather than Components/Subtile: it is the established place for Solution-level validation (MXScaleFormat.py says so in its own docstring), and it imports cleanly, where Components/Subtile at Solution.py module scope deadlocks the package load and needs lazy imports. The three externally called helpers lose their leading underscore to match the validateX convention; the rest stay private to the new module. Two unit tests import the moved names and follow. No behaviour change: 764 codegen snapshots unchanged, 7497 unit tests pass, and subtile_mxfp4 is 9937 passed / 0 failed on gfx950. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The guard was widened to skip mapping-format library logic yamls, because the two curated files added to data/gfx950 were captured in the form the LibraryLogic step emits today and crashed collection on doc[BenchmarkProblems]. Those files were later converted, then dropped entirely when the subtile coverage moved to unit tests, so nothing in the tree needs the wider guard: every logic yaml under Tests parses as a sequence, which the original check catches. Leaves the file untouched by this branch. test_config, test_config_build and test_config_run each collect 414 tests. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The rule rejected any non-power-of-two DepthU for a subtile MX solution, and took TN's DepthU 768 out of subtile_mxfp4.yaml with it. The requirement comes from the scale GR thread split, which divides Serial by numThreadsPerGroup with a shift and a mask; that count's other factors are powers of two only when an operand is free-dim contiguous, so the requirement lands on DepthU for NT, NN and TT and not for TN. Gate it on TLUA or TLUB, which is exactly non-TN (TLUA = !transA, TLUB = transB), and put TN's DepthU list back to its develop form so this branch does not touch TN config at all. All 12 base TN groups are now present unchanged. DepthU 768 still generates nothing on TN, there and on develop alike: the pre-existing LDS pad rule rejects it (LdsBlockSizePerPadA 1024 % 768 x 0.5 != 0), verified by running the base yaml against the base commit. That is unrelated to the scale split and is not this branch's to change. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
WorldofKerry
force-pushed
the
gfx950-transposed
branch
from
September 18, 2026 20:14
76e20b0 to
e72f887
Compare
WorldofKerry
marked this pull request as ready for review
September 18, 2026 20:15
nakajee
reviewed
Sep 18, 2026
| : 0; | ||
| size_t const scaleCols = static_cast<size_t>(sizes[1]); | ||
|
|
||
| bool const kIsRows = (isMatrixA && isTranspose) || (!isMatrixA && !isTranspose); |
Contributor
There was a problem hiding this comment.
Why can't we put these changes in case MXScaleLayout::GFX950:?
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.
JIRA ID AIHPBLAS-4765
Motivation
Support non-TN mxfp4 GEMM for gfx950.
Technical Details
ds_read_b64_tr_b4to transpose out of LDSTest Plan
Include test yamls:
Tensile/Tests/common/gemm/gfx950/subtile_mxfp4.yamlfor coverageTest Result
Tensile/Tests/common/gemm/gfx950/subtile_mxfp4.yamlpassesMT256x256x256on 4k by 4k by 32k problem.