Repository navigation
feat(hipblaslt): add SwizzleA and swizzleB for subtile MXFP4 - #13282
Merged
Merged
Conversation
The optimizations that follow all apply to one narrow family of kernels: fp4 inputs with a bf16 destination, built through UseSubtileImpl on gfx950. Everything else in the library has to come out byte-identical, so the question "is this kernel in scope?" gets asked from the solution resolver, the kernel namer, and a dozen emit sites. Giving that question a single answer -- isMxf4SubtilePath -- keeps those callers from each spelling out their own type test and drifting apart. The narrower helpers it is built from are exposed too, because the tile-shape predicates (plsinBlockSchedTile, plsinEarlyStoreTile) select subsets of the same family for individual optimizations. These helpers are deliberately tolerant of partially-built solution state: assignPostLoopStoreInNll runs before the ProblemType is fully populated, so the predicate has to return False rather than raise on a half-filled dict. The accompanying tests pin that, along with the exact tile shapes that do and do not qualify. Co-authored-by: Cursor <cursoragent@cursor.com>
PostLoopStoreInNll fuses the D-tile buffer_store into the No-Load-Loop, so accumulator pairs that finish early are stored while the final-K-step MFMAs for later pairs are still running. The parameter defaults on and assignPostLoopStoreInNll turns it off for everything that cannot support it, rather than the other way around, so a tuning file does not have to opt in per solution. The eligibility test routes its destination-type check through plsinSubtileTypes instead of spelling one out, which keeps it from accepting kernels the emit side will later decline to build -- a mismatch there allocates a flag SGPR and renumbers every register after it for no reason. Naming needs a matching exception. Required parameters are appended to every kernel name, so naming this one unconditionally would put PLSIN0 into the name of every kernel in the library and change its hash. It is dropped for kernels that cannot set it. No kernel emits anything differently yet; this commit only resolves and names the parameter. The golden updates are that new key appearing. Co-authored-by: Cursor <cursoragent@cursor.com>
Fusing the store into the No-Load-Loop only pays off if the loop hands it work early, which takes three changes to the subtile scheduler. Block scheduling groups the MFMAs so that a whole accumulator tile retires at once instead of dribbling out, giving the fused store a complete tile to write while later tiles are still computing. It is restricted to 256x256 macro tiles, where there are enough MFMAs in flight to hide the store; at smaller tiles the grouping costs more in occupancy than it returns. The preloop global read is now unconditional. It used to be conditional on scheduling state that the block-scheduled path invalidates, and reasoning about when it did and did not run was the source of most of the scheduler's complexity here. The MFMA-to-ACC read latency is modelled from a new arch capability rather than assumed. gfx950 needs 17 cycles between an MFMA and a read of its accumulator, and the store weaving has to respect that or it reads a result that has not landed. A D accumulator tile that spills into the arch VGPR pool stays live from the MFMA loop through the epilogue, so anything that writes those registers corrupts one 16x16 output block per wave, silently. The scheduler records the spilled ranges and assertNotSpilledDTile turns that into a generation failure instead of a wrong answer. The rocisa side gains the latency capability plus TextBlock handling in macro inlining, since CustomSchedule annotates the main loop with comments that were being dropped when the macro was expanded. Co-authored-by: Cursor <cursoragent@cursor.com>
This is the store itself: the D-tile buffer_store moves into the No-Load-Loop, a flag SGPR records which workgroup owns the fused store so the post-loop body can skip the work it already did, and the epilogue kernarg load learns to skip its reload for that owner. The skip target is shared with the existing GSU skip, since both mean "this kernel does not store here", but it keeps develop's label name whenever GSU is the only condition reaching it -- the name is cosmetic and reusing it leaves every out-of-scope kernel byte-identical. The StreamK changes are scoping rather than new behavior. The partials store was biasing an absolute VGPR index by a ValuC base, which makes the workspace store read the wrong accumulators; correcting it is gated on the MXF4 subtile path because every other subtile kernel ships against the uncorrected form. Co-authored-by: Cursor <cursoragent@cursor.com>
…ions Nothing in the unit suite covered this feature surface, so a scope leak only showed up by generating the whole gfx950 library and diffing it -- an expensive way to find out that a type test was one predicate too wide. The config carries two problem groups that differ only in DestDataType, one bf16 and one fp32, at two macro tiles. That gives each test an in-scope case and a control that should look exactly like develop, so a leak fails a test instead of surviving to the library diff. The tile pair does the same for block scheduling, which applies at 256x256 but not at 128x128. The store-repack assertion matches the permlane16 that swaps dwords rather than any permlane16, because develop already emits one for swizzling and matching that would pass no matter what this branch did. Co-authored-by: Cursor <cursoragent@cursor.com>
_selectPlsinFusedStorePolicy takes a different branch once a macro tile exceeds 256: last-K MFMAs move into the store gaps and the K=0 A/B tiles are lent out so store temps can reuse those holes. Every PLSIN test so far sits at or below 256x256, so that branch shipped untested. MT320x256 exercises it, and the characterization test pins the resolution (PostLoopStoreInNll on, MacroTile0 320, MacroTile1 256) so a future change to the policy cannot silently reclassify these kernels as small-tile. The mxfp8 MT320x256 and MT256x320 regression configs already covered the geometry; what was missing was an MXF4 case that actually resolves PLSIN.
0754223 dropped the batch-tail workgroup barrier when the gfx950 permlane16 store path was active, reasoning that "the permlane16 store has no LDS traffic, so its single-DU path keeps the read drain but does not need the following workgroup barrier". b2da475 then widened that to every single-DU subtile path. The barrier was not only protecting this wave's own LDS traffic: it also kept a sibling wave's bias/SAV LDS reads from overlapping this wave's stores. The retained s_waitcnt dscnt(0) is a per-wave counter and cannot order anything across waves, so dropping the barrier dropped a cross-wave guarantee with nothing standing in for it. The result was a race in the store path for single-DU subtile kernels with bias or ScaleAlphaVec active. It surfaced as wrong data at correct addresses, confined to partial-M edge tiles where the skewed wave timing opens the window: MT256x320 MIWT8_10 at 12384x4096x4096 miscompared on 78% of runs, with MT128x384, MT128x192 and MT256x192 also affected. Restoring the barrier for every path that drains bias/SAV takes that shape to 0 failures over 1700 verifications spanning 71 macro-tile configurations, and leaves the emitted barrier count identical to develop's. Measured cost on 6144x32768x4096 is 1.3% for MT256x256 PLSIN1; 0754223 had claimed 0.4% for the removal. Co-authored-by: Cursor <cursoragent@cursor.com>
f0c5ded restored this barrier on every store batch, which is more than correctness needs and cost 1.3-1.7% on 6144x32768x4096. Every LDS op a store batch issues is a load: bias and ScaleAlphaVec are staged into LDS before the store loop, never inside it. Readers cannot race readers, so the batches do not need fencing from each other. What does need fencing is the end of the store against whatever rewrites that LDS next -- under StreamK the same workgroup returns to the mainloop and restages A/B while a sibling wave may still be reading bias/SAV -- and one barrier on the last batch covers that. Multi-DU keeps the per-batch placement: it emits its stores after the barrier rather than before, so the two are not interchangeable there. An assert on the non-barrier batches holds the "loads only" premise, so a future change that introduces a ds_write into the store fails the build instead of silently reintroducing the race. Verified on gfx950: the 71-config all-solutions sweep at m=12384 over 5 passes is 1700 results with worst norm_error exactly 0, and the fp4 / swizzleA / swapAB suites are clean. The 7 remaining failures are the pre-existing swizzle-B k<=128 miscompares, which fail identically without this change. Measured back at parity with the unsynchronized baseline: MT256x320 PLSIN0 4331.8 vs 4331.7, MT256x256 PLSIN1 4532.0 vs 4549.9. Co-authored-by: Cursor <cursoragent@cursor.com>
Three knobs on the mxfp4-bf16 subtile store's fused arm, all default off, all measured neutral on the eleven perf shapes. They are kept because the placement they buy is real and the assembly is now where the FlyDSL reference says it should be; what they show is that store placement is not what bounds this kernel. TENSILE_PLSIN_TAIL_PARTM splits the tail partition grid along M. An N split floors at four partitions once the MX scale granularity forbids a partition narrower than two MFMA tiles, so M is the only axis left. The store enumerates its elements N-group-outermost, so an M split interleaves the partitions rather than giving each a contiguous run; the cut still works because a boundary marker may only open a later stage, but the stages come out lopsided, so buildSubtileFusedStore sorts the element list into partition order. An element's D address comes from its tuple while the accumulator it drains comes from its position, so the acc-read queue is permuted in lockstep -- without that every accumulator pairs with the wrong address. Both are scoped to the list this one fused store consumes. 2x4 is the floor: 4x4 fails vgpr tile assignment. TENSILE_PLSIN_DRAIN_EARLY places each drain unit inside its own partition at the bar where its accumulators go final, instead of deferring whole stages to the next partition where every bar is vacuous. First store moves from MFMA 72 to 28 of 256, cover 184 to 228 against a reference of 232, trailing stores 6 to 3. TENSILE_PLSIN_TAIL_DROP_LAST_SYNC drops the one barrier build_nll keeps in the four-deep body. That barrier separates the body's final LDS read from the next macro tile's global reads, but since ee8f6c3 the fused arm ends with the store's own cross-wave barrier, which sits a whole store later and separates the same two things. ATT confirms the removal: eight barriers become seven, its 22604 cycles do not resurface elsewhere, total stall falls 2.85%. Verified on the eleven perf shapes (11/11, worst norm_error 0), the ten verify suites (212 results, only the known pre-existing STB failures), and five passes of the all-solutions sweep at m=12384 (71 of 71 configurations clean). Performance against the baseline e6d15de is 1.000x geomean, 0.992x to 1.015x. Co-authored-by: Cursor <cursoragent@cursor.com>
The three knobs landed off so they could be measured. They are correct and neutral, and the placement they buy is what the FlyDSL reference describes, so they become the default for the tiles block scheduling is already scoped to: MXF4 in, bf16 out, Weave store, MacroTile exactly 256x256. Every other geometry keeps its shipped schedule, because none of them has been traced. Checked by generating the eleven kernels of the F4BS STA Equality logic three ways -- new defaults, all three knobs forced off, all three forced on. Only MT256x256 MIWT8_8 moves, at both PLSIN0 and PLSIN1, and it lands byte-identical to the forced-on build. MT256x192, MT256x320, MT320x256 and MT256x256 MIWT4_16 are byte-identical to the forced-off build. A plain build of the full library verifies 11/11 on the perf shapes with worst norm_error 0. Co-authored-by: Cursor <cursoragent@cursor.com>
The tail entry waited vmcnt(0) because a scale local-read was placed two sub-k-iterations ahead of its consumer, putting it in front of the only barrier that makes the other unroll half's LDS writes visible. Nothing else could cover that half, so the entry had to drain every outstanding load before the body could start. Clamp the scale read to one sub-k-iteration of lookahead in the tail sub-scheduler only, which leaves room for a second wait/barrier pair, and emit that pair in front of the first read that touches the in-flight half. The entry wait then only has to make the resident half visible, so it drops to one DepthU batch of loads. On MT256x256 fp4 the entry goes from vmcnt(0), charged 354 cycles per wave, to vmcnt(18) charged 28; the second wait costs 73. Tail vmcnt stall falls 54%, from 91.7k cycles to 42.4k. Throughput is unchanged at 1.001x geomean over 11 shapes: the arm is MFMA-bound at ~54% v_mfma occupancy, so the recovered cycles land back on the critical path. Kept on because a bounded wait is the correct shape for a half the barrier already covers. Scoped to the block-scheduled tiles, so only three MT256x256 MIWT8_8 PLSIN1 kernels change; mainloops are byte-identical. Verified over 1912 results: 11-shape suite, a five-pass all-solutions sweep at m=12384 covering 71 macrotile configurations, and 10 verify suites. The only 7 failures are tiny-K STB cases that reproduce with identical norm_error on the pre-change library. TENSILE_PLSIN_TAIL_SPLIT_WAIT=0 restores the single vmcnt(0). Co-authored-by: Cursor <cursoragent@cursor.com>
The five picked commits carry TENSILE_PLSIN_* env overrides, which this branch removed. Fold each back to the default it shipped with: the finer tail grid at 4 M tiles per partition, early drain on with reserve 2 and no extra spacing, the last-sync drop on, and the split entry wait on. Scoping is unchanged -- plsinStagingEligible still holds the grid and the sync drop to the tiles they were measured on.
The picked commit narrowed the barrier for every subtile kernel. The loads-only premise behind it was read off the MXF4 store epilogue and measured there, so scope it the same way the rest of this branch is scoped and leave other subtile kernels on develop's per-batch barrier.
_plsinWeaveStat is debug instrumentation this branch removed with the env gates; it was default-off upstream and emits nothing. The scoping test pinned the tail at four partitions, which was the N-only split -- the 2x4 grid makes it eight.
Dropping the scope here added lgkmcnt waits to 552 kernels outside the MXF4 subtile path -- F4SS, F4HS, F8 and BBS -- which this branch holds byte-identical to develop. The two forms agree for in-scope kernels, so the restore costs the race fix nothing: single-DU MXF4 still takes the per-element route, and the barrier it needs is emitted at the store.
plsinBlockSchedTile scoped on tile and operand types alone, so a kernel whose PostLoopStoreInNll had been disabled -- StreamK atomic, a non-zero StoreRemapVectorWidth, wave32, MultipleBuffer accumulation -- could still be built block-scheduled. It then emitted the NGLL/NLL arm split for a fused store that does not exist. No shipped kernel reaches that corner today: all three MT256x256 F4BS solutions resolve PLSIN on. PLSIN-on codegen is byte-identical with this check, so the gate now states the dependency the staged store already had.
The mock carried tile and operand types only, so it described a kernel with no PostLoopStoreInNll key at all -- which the block-scheduling gate now reads. Give it the parameter, defaulting on as every real MT256x256 fp4 kernel resolves it, and test the off case directly.
Removing the TENSILE_PLSIN_* env overrides left two comments pointing at things a reader cannot reach: the split entry wait still described a knob that overrides its bound, and both sites cited a FINDINGS document that does not exist in the tree. The bounds they describe are now derived and fixed, so say only that. Co-authored-by: Cursor <cursoragent@cursor.com>
The block-scheduled tail, its staged store and the finer store grid are all reached through plsinStagingEligible, while the store barrier and the per-element wait accounting are scoped on isMxf4SubtilePath directly. Nothing checks that the first is contained in the second, so a tile predicate that was loosened off the MXF4 path would build the block schedule for kernels none of this was measured on, and the only sign of it would be a library diff. Sweep the type, subtile, tile and PLSIN axes and assert the containment holds at every point. Dropping the type check from plsinEarlyStoreTile fails it on an fp4-in/fp32-out kernel, which is the leak being guarded. Co-authored-by: Cursor <cursoragent@cursor.com>
Develop moved StreamK behind the TileProcessingStrategy/ExecutionPolicy abstraction and renamed its constant-parking machinery from StreamK* to Persistent*, so the resolutions worth recording are: * The solution-time StreamK check became isStreamK(state); the codegen-time ones became isStreamK(kernel) and isPersistentDataParallel(kernel), since develop no longer carries a "StreamK" key in either dict. * Develop's #12615 narrowed subtileBarrierDrains to multi-DU for every subtile kernel, which is what this branch's needsCrossWaveBarrier was doing for the MXF4 path alone. The indirection is now tautological in the wait accounting, so it is dropped there; the last-batch barrier at the emission site, which develop did not change, is kept. * The post-loop store-init is emitted once. This branch relocates it after the fused-store dedup guard, and develop added skipPhantomTileStore next to the original call site; keeping both emitted the SRD init twice and collided the per-channel labels. * Persistent constants are parked in VGPRs on gfx950 under PostLoopStoreInNll. The fused epilogue peaks at 104 SGPRs against a cap of 102 on tiles above 256, and parking the kernarg run frees a contiguous block to cover it.
Develop added subtile fp4 support for the NN/NT/TT layouts, along with the GRTag_TLU1 strided-K pointer update they need and 1559 lines of new coverage in subtile_mxfp4.yaml. Those kernels satisfy isMxf4SubtilePath as it stood (UseSubtileImpl plus fp4 in / bf16 out), so merging develop silently opted them into optimizations that were only ever built and measured on TN. The pre-loop SRD hold and the fused post-loop store reason about the TN global-read shape, which those kernels do not have, and 170 of the new non-TN cases miscompared. Add the layout to the gate, which is where the rest of the scoping already lives, so the non-TN kernels keep develop's schedule exactly. subtile_mxfp4 then matches the develop baseline at 9937 passed / 0 failed on gfx950. Also give _emitGRPtrUpdate_TLU1 the dispatcher's holdOnLastIter parameter. It was added to the TLU=0 implementation only, so any TLU=1 kernel aborted codegen with a TypeError. It is accepted and deliberately not honoured, for the same reason the gate is now TN-only. test_mxf4_subtile_scope covers all four layouts, so the restriction is pinned down rather than implicit. Its fixture previously omitted the transpose keys entirely.
This preserved Macro::addComment0 bodies when a macro was expanded, so that CustomSchedule's mfmaIndex annotations survived into the assembly. It broke rocisa's test_prnd_generator: PRND_GENERATOR's macro body opens with a comment naming the macro, so the inlined output started carrying the string the test asserts is absent. The annotations are inert and nothing here needs them. The subtile path builds no macros at all, so the only bodies this ever affected belong to CustomSchedule and PRND_GENERATOR, neither of which is part of this work. Dropping it keeps the change confined to the MXF4 subtile path, leaves the develop-owned test untouched, and restores develop's behaviour of discarding those comments. No effect on the shipped library: the 12 gfx950 kernels carrying mfmaIndex comments get them from Module::addComment0 rather than from macro expansion, and that count is the same on develop.
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
Codecov Report❌ Patch coverage is Additional details and impacted files@@ Coverage Diff @@
## develop #13282 +/- ##
===========================================
- Coverage 76.11% 72.48% -3.63%
===========================================
Files 2886 2948 +62
Lines 483963 490629 +6666
Branches 73003 71998 -1005
===========================================
- Hits 368338 355608 -12730
- Misses 90392 110164 +19772
+ Partials 25233 24857 -376
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
GlobalWriteBatch imported plsinEarlyStoreTile and never called it, in the commit that added the import or in any since. Remove it; every other name on that line is used. The docstring was stale in a way that misread as a limit on PLSIN itself. It opened by explaining that the predicate exists so several epilogue knobs could be scoped alike, but those knobs were inlined at their shipped values earlier in this branch, leaving plsinBlockSchedTile as the only caller -- which narrows the bound to the MT256x256 equality, so it decides nothing on its own. Say that, say why the two predicates stay separate rather than collapsing into one constant, and say the part that was actually missing: MT>256x256 does get PLSIN, through largeTileLend. What it cannot take is an *early* store, because it lends the K=0 operand registers a store moved earlier would still be reading. Emission is unchanged: the designed scoping pair hashes identically before and after, both problem groups. Co-authored-by: Cursor <cursoragent@cursor.com>
liu-yiqian
force-pushed
the
users/liu-yiqian/mxf4-swizzleb
branch
from
October 8, 2026 14:39
9081f10 to
91ef49d
Compare
The TextBlock handling was removed from the C++ macro_inline pass but not from its Python mirror in the stinkytofu adaptor, leaving the two implementations disagreeing about whether a macro-body comment survives inlining. Remove it there too, for the reason it was removed from C++: the subtile path builds no Macros at all, so this only ever reshaped CustomSchedule and PRND_GENERATOR, both out of scope for this work. mfma.hpp's added #include <optional> was hygiene, not a requirement -- the file's existing std::optional members already resolve through a transitive include, and rocisa builds clean without it. It and a stray blank line were the only rocisa changes outside hardware_caps.hpp; revert both so the rocisa diff is exactly the MfmaToAccReadLatency capability, which is load-bearing. Both files are now byte-identical to develop. Emission is unchanged: the designed scoping pair hashes the same before and after, and rocisa's suite passes 130 including test_prnd_generator, the test this had broken in CI. Co-authored-by: Cursor <cursoragent@cursor.com>
…liu-yiqian/mxf4-swizzleb
Same defect the MT320x256 weave config had: findConfigs turns a declared mark into a skip only on exact string equality with a detected architecture, and gfx1250-strict is its own architecture name rather than a qualifier on gfx1250. Both configs declared skip-gfx1250 alone, so both would have run on the strict lane, where these gfx950-only subtile kernels cannot build. With this, all 53 configs under Tests/common/gemm/gfx950 declare the pair.
The pre-shuffled global read enumerates one K-subtile at a time, but three
places scaled by the whole DepthU window instead. Those agree only when
localSubtileGrid[1] is 1, which for FP4 means DepthU 256. At DepthU 512 the
window splits in two and each wave read every other 128-byte swizzle
sub-block, so the MFMA multiplied rows against the wrong K positions and
produced wrong results rather than failing.
- the per-row step becomes subIterKBytes, the K-subtile extent, instead of
the whole DepthU window
- the aligned 16-row padding correction uses the same K-subtile extent
- a K-subtile step is scaled by the rows in a swizzle block, since one K
step consumes a K-extent on every row
Guard the boundary at both layers so an unsupported DepthU cannot come back
silently. Deriving a swizzled solution past SWIZZLE_MAX_DEPTHU is rejected
with a swizzle-aware reason; until now DepthU 1024 was only turned away by
the LDS budget, which says nothing about why. The emitter additionally
refuses a K-window step that does not fit the 12-bit MUBUF immediate, which
rocisa does not range-check and would otherwise emit verbatim.
All eight shipped swizzle kernels stay byte-identical, and 96 validations
pass on gfx950 across both operands, DepthU 256 and 512, tail and odd sizes,
and StreamK 0 and 3.
Three defects made the host pre-swizzled path read the wrong operand data. The global-read stride rounded K up to swizzleK, which covers a whole wave, but the host pads K to a whole operand before reshaping and that padded K is the 16-row block stride. For any K landing between the two, every block after the first was addressed at the wrong offset. Round to the host's granularity instead. A K shorter than one DepthU window makes the row padding correction negative. It was computed in elements and narrowed to bytes by shifting the product, which is a logical shift of a two's-complement value and turned a small negative step into roughly 2**31. Narrow the stride before the multiply instead. StreamK's partial-tile start offset stepped by DepthU in unswizzled terms. A pre-shuffled block interleaves 16 rows into one K step, which the unroll loop already accounts for, so a workgroup whose range started mid-tile read 16x too close to the tile start. Verified on gfx950: the 15 sizes from the reported tuning run go from 9/15 to 15/15 and a 26-point K sweep from 13/26 to 26/26, while the existing 96-validation swizzle sweep, both shipped swizzle configs and the unswizzled control stay green. Regenerating gfx950/Origami changes exactly the 8 pre-shuffled kernels and leaves the other 61373 byte-identical.
Block_32_UE8M0_32_8_EXT scales go through preSwizzleScalesGFX950, which lays them out as a 32x8 grid: the MN extent pads to 32 and the K-block extent to 8. The bench sized that buffer with the granularity of swizzle_mx_scale instead, 128 / MXBlock = 4 K blocks and no MN padding, so whenever K was not a multiple of 256, or M or N not a multiple of 32, the allocation came up a whole group short and the swizzle's tail was truncated. The kernels then read the missing group as zeros. Every F4BS solution failed identically, with a relative norm error of sqrt(4 / (K / 32)) when only K was unaligned: 0.279 at K=2944, 0.519 at K=384. Verified on gfx950: 229/229 heuristic solutions pass at 4096x32768x2944 with swizzled A, where all 229 failed at 0.2708 before. K=384, 640, 896, 1152, 1408 and 2944 and M or N of 1000 and 1008 all pass now; the already-aligned sizes are unchanged.
✅ 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. |
Name the instruction the path actually uses instead of an external project, so the comments describe the code in front of the reader.
…lsin-blksched-preloop-v2 # Conflicts: # projects/hipblaslt/tensilelite/tensilelite/KernelWriterAssembly.py # projects/hipblaslt/tensilelite/tensilelite/SolutionStructs/Naming.py # projects/hipblaslt/tensilelite/tensilelite/Tests/unit/characterization/SolutionClass/__snapshots__/test_solution_class_char.ambr # projects/hipblaslt/tensilelite/tensilelite/Tests/unit/test_subtile_bias_sav_waits.py
liu-yiqian
marked this pull request as ready for review
October 21, 2026 16:42
Restore per-file coverage above the ratchet baseline for the three Subtile files this PR touches: InstructionEmitter.py 94.92 -> 97.56 (baseline 97.35) LogicalScheduler.py 91.35 -> 93.80 (baseline 93.46) SubtileLREmit.py 92.31 (baseline 89.12) Adds unit coverage for the loop-control setup hoist, LR tile set counting, tail tile release, the tail entry wait bounds, the multi-partition GR wait counts, and the schedule dump printers.
…liu-yiqian/mxf4-swizzleb Resolve testing_matmul.hpp in favor of develop's blockScaleBufferSize. Develop landed a helper that computes the same pre-swizzled scale extent this branch was computing inline: for Block_32_UE8M0_32_8_EXT it routes to scaleBufferSize, which pads the K-block count to 8 and the MN extent to 32, matching mxPreSwizzledScaleSize. The non-swizzled padding path matches as well, so the inline block and its helper are dropped.
SWIZZLE_MAX_DEPTHU bounds the subtile pre-shuffled global read, whose K step rides in the MUBUF immediate. The derivation gate asked it of every SwizzleTensorA/B solution instead, including the classic swizzle path, which builds its step differently and has shipped well above it: gfx942 F8 swizzle-A logic carries 191 solutions at DepthU 1024 and 2048. A rejection raised during library-logic parsing rather than skipping the solution, so TensileCreateLibrary aborted and the whole gfx942 build failed on the first one. Scope the check to isMxf4SubtilePath, which is also what decides whether the subtile emitter runs, so derivation and emission agree.
Two of the pre-shuffled global-read changes keyed off SwizzleTensorA/B alone, so they also moved codegen for the classic swizzle kernels gfx942 ships. All 514 swizzle tests across the six gfx94X shards failed with unspecified launch failure; the same shards pass on the base PR. They had never run before: the gfx942 build died on the DepthU ceiling until that was scoped too. graAddresses aligned the A/B stride to swizzleLayoutK, the granularity the host pads to when it lays out a pre-shuffled tensor. That is right for the subtile path but not for the classic one, and the two values are far apart -- 8 against 32 for gfx942 f16, 16 against 64 for f8 -- so the SRD addressed outside the buffer. Align to swizzleLayoutK only on the subtile path. StreamK._preShuffleRows scaled a mid-tile range start by the rows in a block for any swizzled tensor. Only the subtile global read walks blocks that way; the classic path steps one row per K window, so the scale moved those workgroups off their own range. Return 1 off the subtile path. The remaining two arithmetic changes -- the fractional bpe in swizzleGeometry and the fractional MX read ratio in KernelWriter -- are unreachable rather than unscoped: both agree with the old expression everywhere the old one did not divide by zero, checked exhaustively over the operand ranges. The r9 StreamK test asserted the unscoped behaviour, so it now builds a subtile kernel and a companion pins the classic path at 1. 11290 unit tests pass and all 802 snapshots are unchanged.
nakajee
reviewed
Oct 27, 2026
Base automatically changed from
users/liu-yiqian/plsin-blksched-preloop-v2
to
develop
October 27, 2026 14:21
…xf4-swizzleb # Conflicts: # projects/hipblaslt/tensilelite/tensilelite/Components/Subtile/LogicalScheduler.py # projects/hipblaslt/tensilelite/tensilelite/Components/Subtile/SubtileGREmit.py # projects/hipblaslt/tensilelite/tensilelite/KernelWriterAssembly.py # projects/hipblaslt/tensilelite/tensilelite/SolutionStructs/Naming.py # projects/hipblaslt/tensilelite/tensilelite/SolutionStructs/Solution.py # projects/hipblaslt/tensilelite/tensilelite/Tests/unit/characterization/SolutionClass/__snapshots__/test_solution_class_char.ambr # projects/hipblaslt/tensilelite/tensilelite/Tests/unit/test_SubtileBasedLogicalScheduler.py # projects/hipblaslt/tensilelite/tensilelite/Tests/unit/test_assign_post_loop_store_in_nll.py # projects/hipblaslt/tensilelite/tensilelite/Tests/unit/test_mxf4_subtile_scope.py # projects/hipblaslt/tensilelite/tensilelite/Tests/unit/test_plsin_srd_hold.py # projects/hipblaslt/tensilelite/tensilelite/Tests/unit/test_vgpr_overflow_force_generate.py
… gate Review feedback. SWIZZLE_MAX_DEPTHU = 512 was the MXF4 instance of a limit that belongs to the 12-bit DTL immediate, so derive it from the rows in a swizzle block and the element size instead. isSubtileSwizzledTensor names the question the pre-swizzled global read actually asks. Named off SwizzleTensorA/B so it cannot be read as the MX scale swizzle, which has its own layout and its own read. It delegates to isMxf4SubtilePath: MXFP8 and BF16 are out of scope here.
swizzleMaxDepthU read as though it covered the MX scale swizzle too. It does not: the scale blocks have their own layout and their own read.
Swept in by accident; it is a build artifact of a local library regeneration and holds absolute paths into my checkout.
liu-yiqian
enabled stacked auto-merge (squash)
October 27, 2026 20:55
nakajee
approved these changes
Oct 28, 2026
nakajee
left a comment
nakajee
left a comment
Contributor
There was a problem hiding this comment.
Thanks for your update.
I do not have further comments.
Please make sure no CI regression due to this change.
Contributor
TheRock Submodule Bump ActivityNewest first
|
1 task done
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-4083
Motivation
SwizzleTensorAandSwizzleTensorBmean the host has already interleaved the operand into the MFMA-friendly 16-row layout, so the kernel does not pay for the rotation. This enables both of them on the subtile MXFP4 path, which previously supported neither.Technical Details
1. Reading the pre-shuffled layout
TileInfogrows anisPreShuffledflag and the emitters branch on it. A pre-shuffled tensor interleaves one K tile across all rows of an MFMA tile, so: the GR pointer advance steps 16 row fragments perDepthU; GR offsets index byrowId * depthUBytesrather thanrowId * stride; the colId swizzle skips pre-shuffled tensors, so the operand keeps the host's layout while the other one still gets the rotation; the aligned 16-row HBM padding is folded back in; the SRD limit and tile start scale by the 16-row block stride; a StreamK range that starts mid-tile scales its start by the same row count; and the MX scale buffer is sized the way the layout pads it.Where nothing is pre-shuffled the two colIds are equal, so B stays derived from A by copy rather than recomputed — that is what keeps kernels off this path byte-identical.
2. Keeping it off the architectures it was not built for
gfx942 ships 35 logic files (30 swizzle-A, 5 swizzle-B) that set the same
ProblemTypeflags for the classicDirectToVgprswizzle, which is a different host layout with a different K step. Three changes here first read those flags and so leaked onto that path; all three now askisSubtileSwizzledTensorinstead:graAddressesstride alignment. The subtile path pads to the pre-shuffled block stride (swizzleLayoutK), the classic path to the per-wave block (swizzleK). They differ wheneverswizzleKexceeds one operand — 32 against 8 for gfx942 f16, 64 against 16 for f8.DepthUceiling. 382 of those gfx942 solutions shipDepthU1024, above the subtile ceiling.Guards written inside
Components/Subtile/are implicitly scoped because only that module constructs aTileInfo; these three live outside it, which is why they needed an explicit gate.3. The
DepthUceiling is derived, not hard-codedA K step covers every row of a swizzle block, so the byte step is
rows * DepthU * bpeand has to fit the 12-bit DTL immediate.swizzleTensorMaxDepthU(rowsPerBlock, bpe)returns that limit: 512 for MXF4's half-byte element over 16 rows, 256 for FP8, 128 for BF16. It is enforced both when deriving a solution and when emitting one, and the rejection message now prints the real limit for the data type at hand.4. Gate naming
isSubtileSwizzledTensoris named off theSwizzleTensorA/SwizzleTensorBparameters it reads, to distinguish it from the MX scale swizzle, which has its own layout and its own read. It currently delegates toisMxf4SubtilePath: nothing about reading a pre-swizzled operand is MXF4-specific, but FP8 and BF16 are untuned and untested here and out of scope for this ticket. Widening it later is one edit in one place rather than three. The store and schedule gates stay onisMxf4SubtilePath, which is genuinely MXF4-in/bf16-out specific.5. Logic files
Swizzle is a
ProblemTypeproperty rather than a solution parameter, so it needs its own logic files rather than extra entries in the existing ones. Six are added, from gfx950 tuning runs, mirroring theProblemTypeof the shipped F4BS library withSwizzleTensorAorSwizzleTensorBset — one STA and one STB in each ofOrigami/,Origami_nta4/andOrigami_ntb4/, 76 solutions each. Equality, GridBased and Range are untouched.The library grows by exactly 152 kernels: 76 STA and 76 STB, since the three NonTemporal variants resolve to the same kernels.
Test Plan / Test Result
tensilelite/Tests/unit)subtile_mxfp4_swizzlea.yamlandsubtile_mxfp4_swizzleb.yamlThe regeneration is the load-bearing one for scope: every kernel the library built before this change is byte-identical afterwards, so the only new code any existing problem can reach is none.
Neither
SwizzleTensorAnorSwizzleTensorBhad unit coverage on the subtile path before this — the existing tests all drive the classicDirectToVgprpath, which never builds a subtileTileInfo, soisPreShuffledwas False everywhere and every arm it guards shipped untested. The new designed gfx950 config has problem groups differing only inSwizzleTensorAorSwizzleTensorB, pinning the pre-shuffled code against an otherwise identical control. It deliberately sets noDirectToVgpr, which is what pins the subtile path's exemption from that requirement: if the exemption regresses, the config stops producing solutions and the test fails rather than silently passing.The A/B asymmetry gets its own test, since a check written per kernel rather than per tensor passes a markers-present test and is still wrong.
test_r9_preshuffled_char.pyadditionally pins the three scoping gates from section 2 against their classic-swizzle counterparts, the derivedDepthUceiling and its boundary, and the equivalence of the two gate helpers so they cannot silently drift.test_subunit_division.pycovers two division fixes directly, including the operands that used to raise.Submission Checklist