Repository navigation
feat(hipblaslt): ship rocRoller MX kernels and route the ext API to them - #14045
Merged
jaopaulolc merged 39 commits intoOct 27, 2026
Merged
jaopaulolc merged 39 commits into
jaopaulolc merged 39 commits into
Conversation
…lect them from production logic Adds 39 rocRoller custom kernels for the three pure MX operand pairs -- MXFP8xMXFP8, MXBF8xMXBF8, MXFP4xMXFP4 -- with non-pre-swizzled scales, completing the diagonal that #12965 left out. Adds production GridBased logic for all 129 checked-in rocRoller kernels. They were previously reachable only from benchmark configs, so a default hipBLASLt build could not select any of them. LibraryLogic.py gains the branches that emit exact-size entries for custom kernels; a unit test pins tie-breaking between an exact custom kernel and a generated solution of the same size. The kernels declare `gfx950:sramecc+` and require the preceding target id change to assemble. Batched MX stays out of scope: no kernel declares a scale batch stride, so batch_count > 1 returns no solution through this path.
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
jaopaulolc
changed the title
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## users/jolabega/kfa-rr-mx-2-target-id #14045 +/- ##
=======================================================================
Coverage ? 75.95%
=======================================================================
Files ? 2876
Lines ? 477024
Branches ? 71576
=======================================================================
Hits ? 362300
Misses ? 90031
Partials ? 24693
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
…ga/kfa-rr-mx-3-pure-mx-kernels
✅ 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. |
…ga/kfa-rr-mx-3-pure-mx-kernels
…ga/kfa-rr-mx-3-pure-mx-kernels
1 task done
…ga/kfa-rr-mx-3-pure-mx-kernels
A kernel built for MXScaleFormat HostPreSwizzle indexes the MX scale tensor
through the gfx950 e8m0_shuffle permutation; one built for NoSwizzle or
InMemorySwizzle reads a compact row-major E8M0 array. The two buffers have
identical extents and the same element type, so nothing in the problem
separated them and a mismatch was silently wrong answers, not a failed launch.
That is what the 128 checked-in RR_GEMM custom kernels hit. They read compact
scales, but the gfx950 MX shards they were registered into also serve callers
passing Block_32_UE8M0_32_8_EXT, and the heuristic preferred them: norm_error
0.48-0.88 on every shape, against 1.5e-06 for the Cijk control.
hipBLASLt already knows the layout -- scaleAType distinguishes Block_32_UE8M0
from Block_32_UE8M0_32_8_EXT, which is what useRocRoller() reads today -- it
just stopped at ConstructTensileProblem, where both formats collapse onto the
same setHostMXScaleA/B call. So:
- ContractionProblemGemm carries mxScaleHostPreSwizzle, set by hipBLASLt from
the scaling format and by the TensileLite client from --mx-scale-format.
It also keys the problem hash and comparison, since selection depends on it.
- A new MXScaleHostPreSwizzle ProblemPredicate, emitted per MX solution from
its MXScaleFormat. Per solution rather than from ProblemType.predicates(),
because MXScaleFormat is a Solution-level knob: the file-level ProblemType
does not carry it and a library-level predicate built from it would read
NoSwizzle for every file and reject whole HostPreSwizzle libraries.
- InMemorySwizzle maps to false: that kernel permutes the scales itself and
still reads the compact buffer, which is how DataInitialization treats it.
No logic file changes: the 75 custom-kernel YAMLs carry MXScaleFormat only in
DefaultSolution and only as Auto, which already resolves to NoSwizzle -- the
layout those kernels actually read.
Measured on gfx950 (MI355X). At scaleA=1001 the nine known-bad shapes now pick
a HostPreSwizzle Cijk and verify clean; across 299 pre-swizzled problems all
251 pre-existing Cijk selections are unchanged, the 48 RR_GEMM selections
become Cijk, and no problem loses a selection. At scaleA=3 the RR_GEMM kernels
stay selectable with byte-identical selections before and after.
1 task done
…tured for A rocRoller custom kernel is a pre-assembled object captured at one tile. It carries no generated edge handling in M and no partial-k-tile drain, so it is correct only where M is a whole number of its MacroTile0 and K a whole number of its DepthU. The generated logic declared neither, and a GridBased lookup hands the kernel neighbouring sizes it never captured. The result is wrong answers that are also nondeterministic, so they do not reproduce reliably and a single clean run proves nothing: at 1824x512x2048 f8->bf16 the same record gave norm_error 0.3505 on seven of eight runs and 0.078 on the eighth, and at 392x1408x2816 it alternated between 0.0123 and 0.0043. On bf16 output the bad runs can sit near the 3.5e-03 rounding floor and read as clean, which is how this survived review. Measured on gfx950 across the FP8, BF8 and FP4 families, heuristic forced: every size with M % MacroTile0 == 0 and K % DepthU == 0 is correct and stable (20 of 20 in the M/N/K sweep); break either and it is wrong. N is not constrained -- N = 160, 192 and 320 against a 128-wide tile all verify clean, so the kernels do handle the N edge, and rejecting those sizes would make the kernels unreachable for no reason. The rocRoller-off path already guards all three in getRocRollerSelectedSolutions, which is why that route picks a smaller tile at M=640 and verifies at 1.6e-06 where the heuristic does not. No new predicate type: this reuses Free0SizeMultiple and BoundSizeMultiple, the predicates AssertFree0ElementMultiple and AssertSummationElementMultiple already produce. Those assertions exist on all 75 logic files but are hardcoded to 1, so they emit nothing. Scoped by the RR_GEMM_ prefix, which is how matchesRocRollerSelectedKernel identifies the same family; the AITER and Tensile-generated custom kernels are deliberately untouched, having no equivalent measurement. At scaleA=3 with the heuristic forced, the five shapes in verify_rr.yaml that satisfy the rule keep the same kernel and the same norm_error; the four that violate it -- two visibly wrong at 0.45, two masked near the bf16 floor -- now select no RR_GEMM at all. The production path is unchanged: those sizes go to the rocRoller JIT, which computes them correctly.
…ga/kfa-rr-mx-3-pure-mx-kernels
jaopaulolc
marked this pull request as ready for review
October 22, 2026 18:02
caio96
reviewed
Oct 22, 2026
| # direction. The rocRoller-off path (getRocRollerSelectedSolutions) guards | ||
| # all three; that is stricter than correctness requires. | ||
| customKernel = state.get('CustomKernel') if hasattr(state, 'get') else None | ||
| if customKernel and str(customKernel.get('name', '')).startswith('RR_GEMM_'): |
Contributor
There was a problem hiding this comment.
RR_GEMM_ kernels can be selected for batch_count > 1 and return wrong results for batches 1..N-1.
The RR kernels are 2-D only: no batch strides, and workgroup_id_z is unused. They still launch grid: [..., Batch], so every z-slice recomputes batch 0. This block emits no BatchSizeEqual. Two paths reach these kernels for batched problems: HIPBLASLT_USE_ROCROLLER=0, and an explicit or reused algo.
HIPBLASLT_USE_ROCROLLER=0 hipblaslt-bench --transA T --transB N -m 512 -n 512 -k 512 \
--a_type f8_r --b_type f8_r --c_type f16_r --d_type f16_r --compute_type f32_r \
--scaleA 3 --scaleB 3 --batch_count 4 -v
On MI355X, with the same RR_GEMM_..._256x128x128 kernel, norm_error is 2e-4 at batch 1, 1.0 at batch 2 and 3.0 at batch 4.
Fix: add BatchSizeEqual(1) here, regenerate the gfx950_id75a3 logic, and add a batched MX test.
Contributor
Author
There was a problem hiding this comment.
Added in 51a28abcbe4: the RR_GEMM_ block now emits BatchSizeEqual(index=0, value=1), outside the macrotile-length guard so a kernel with a degenerate tile still declares it. Confirmed from the checked-in logic that these kernels take no batch strides and launch grid: [TilesX, TilesY, Batch], exactly as you describe.
On regenerating the gfx950_id75a3 logic: not needed for this predicate, and I checked rather than assumed. The predicate list each RR_GEMM_ solution carries is derived at library-build time by CompoundPredicates in Contractions.py, so the two new predicates take effect without touching the YAMLs — the gfx950_id75a3 corpus is byte-identical between 440325fc409 and this branch while both were added, and a msgpack dump of a library built at the older commit carries neither on any of its 128 RR_GEMM_ solutions. One correction to that, though: the YAMLs do bake CustomKernelSizePredicate, a six-entry size window on 38 solutions across 12 files, which Contractions.py reads back. So the general claim "the logic files carry no predicates" is wrong, and anything that needs a tuned rather than a derived value does belong there.
Measured on MI355X (gfx950, PCI 0x75a3), f8 x f8 -> f32 at 512x512x512 with scaleA=3 scaleB=3, and with norm_check set in the YAML record so a numerical failure aborts before the timing row:
- Before (
440325fc409),batch_count=4with the Tensile heuristic forced (HIPBLASLT_USE_ROCROLLER=0) selectedRR_GEMM_TN_FP8_FP8_Float_Float_Float_SA_BE8M0_32_SB_BE8M0_32_WGT_256x128x128_UR_2and failed verification on 3 of 3 runs. - After, the same problem reports "NO solution found" on 3 of 3 runs.
batch_count=1on the same shape is unchanged — same kernel,norm_error1.32878e-06 before and after — so the predicate does not cost the unbatched case.
On the batched MX test you asked for: there is now Python coverage that the predicate is emitted and is scoped to RR_GEMM_ (test_batch_is_pinned_to_one, test_batch_and_buffer_limits_are_rocroller_only in test_rocroller_custom_kernel_tile_predicate.py), but no GPU-level regression test in a shared lane — the measurement above was run by hand.
| if tileM and tileM > 1: | ||
| rv += [cls('Free0SizeMultiple', index=0, value=tileM)] | ||
| if depthU and depthU > 1: | ||
| rv += [cls('BoundSizeMultiple', index=-1, value=depthU)] |
Contributor
There was a problem hiding this comment.
RR_GEMM_ kernels silently zero their output once any tensor reaches 4 GiB.
The RR kernels size each buffer descriptor over the whole tensor with a 32-bit num_records, and the per-lane offset is 32-bit too. At 4 GiB the output comes back all zeros; above that, writes also wrap and overwrite the start of D. BufferLoad/StoreOffsetLimitCheck only bound the per-tile extent, so nothing rejects these problems.
On MI355X, f8×f8→f32 (--scaleA 3 --scaleB 3):
- 32768×32768×256 gives norm_error 1, while N=32640 on the same kernel gives 7.7e-5.
- 1024×1024×256 with
--ldc/--ldd 1048576also gives 1, so padding counts.
The same happens with A or B over 4 GiB, and with f4 or mixed f8×f4.
Fix: add an RR-only predicate requiring every whole-tensor extent (A, B, C, D, scales, padding included) to stay under 2^32 bytes, like exceedsBufferAddressingLimit on the rocRoller-on path. The padded-ldd case is a cheap regression test.
Contributor
Author
There was a problem hiding this comment.
Added in 51a28abcbe4 as a new BufferAddressingLimitCheck predicate (ContractionProblemPredicates.hpp plus serialization, emitted from the RR_GEMM_ block), the selection-time counterpart of RocRollerGemmKernel::exceedsBufferAddressingLimit. It covers A, B, C, D and the two MX scale tensors.
One adjustment to what you proposed. I used max(stride[i] * size[i]) rather than the tight allocated extent, because the tight extent would have accepted your own padded repro: M=N=1024, K=256 with ldc/ldd 1048576 has a tight extent of 4290777088 bytes (3.996 GiB, admitted) but a stride product of exactly 4294967296 (rejected), and you measured that shape as wrong. Your other two shapes agree under either formula. So the stride product is the one that matches the hardware, consistent with num_records being programmed from the stride product rather than from the tightest allocation that holds the elements. All three of your shapes are now regression tests in Predicates_test.cpp; the file passes 16/16 locally.
1 task done
…stom kernels The RR_GEMM_ kernels are 2-D: no batch strides, workgroup_id_z unread, but a grid of [TilesX, TilesY, Batch], so batches 1..N-1 recomputed batch 0. They also size one buffer descriptor over each whole tensor with a 32-bit num_records, so a tensor at 4 GiB wrapped and the kernel zeroed its output instead of failing to launch. Declare BatchSizeEqual(1) and a new BufferAddressingLimitCheck, the selection-time counterpart of RocRollerGemmKernel::exceedsBufferAddressingLimit. The extent is the stride product, not the tight allocated extent, so leading-dimension padding counts: a 1024x1024x256 problem with ldc/ldd 1048576 is wrong on hardware and only the stride product catches it.
… kernels The 39 kernels this branch adds carry rocRoller's codegen narration in whole-line `//` comments -- register allocation, expression lowering, liveness. The assembler discards them. Same treatment as the kernels already on develop (#14206), applied with the tool that lands there. These 39 were captured without a `.kernel_graph` dump, so comments are all there is to remove. Verified: assembling each kernel at its `.amdgcn_target` before and after gives a byte-identical code object; `custom.config` and `.amdgcn_target` compare equal line for line. 1,084,926 -> 180,903 lines, 54.4 -> 6.1 MiB.
The rocRoller tile ranker built its hardware model from HIP device 0 unconditionally, so on a multi-GPU or partitioned system it could rank against a device the caller was not using. Ask for the current device instead, and cache the model per device: get_hardware_for_device queries HIP, which is far too expensive to repeat for every ranked problem. Taking the hardware as an argument rather than querying for it inside the ranker is also what makes the ranker callable without a device at all. Name the problem fields the ranker actually reads as a RankerProblem, so the ranker body no longer needs a whole RocblasltContractionProblem to run. The existing entry point is unchanged; it builds one and calls through. No behaviour change on a single-GPU system.
…a problem Block scaling routes a problem to rocRoller. rocRoller then declined some of them with an error, so the caller got a hard failure for a GEMM the generated library could have run. A GPT-OSS-120B MoE workload hits this on 88 problems: every one is batched, and batch_count != 1 returned invalid_value. Route instead of fail, and do it generally rather than for batching alone: - rocRollerSupportsProblem() is the one definition of what rocRoller declines outright - a bias, a batch, a scale type other than f32, non-f32 accumulate. useRocRoller() consults it, so a declined problem never enters the rocRoller branch at any call site and reaches the Tensile heuristic the normal way. rankRocRollerSelectedTiles and getRocRollerBestSolutions share the predicate instead of restating it. - When rocRoller accepts a problem but no kernel fits it, getBestSolutions and getAllSolutions now continue into the Tensile heuristic instead of returning an empty list the caller reports as not implemented. That second fallback means a problem useRocRoller() accepts can come back with a Tensile algo, so the algo, not the problem, now decides the backend on the launch and support paths. rocRoller solution indices have bit 31 set, which getSolutionsFromIndex already relies on. A caller-supplied algo is never replaced; only a call that picked no algo may fall through, and only on not_implemented, so a launch failure still surfaces. The rocroller_predicate tests asserted the old contract - that a declined problem reports "NO solution found!". Both cases are now served by the fallback, so they assert that instead. Their norm_check is off: the generated MX kernels they land on fail it with HIPBLASLT_USE_ROCROLLER=0, without any rocRoller routing, which is a TensileLite defect these tests must not be coupled to. (cherry picked from commit 4c915202cd978a5223b6e67f0c70e0a59ffd2d00)
The rocRoller-off path refused a ranked tile unless all three of M, N and K divided it. Only two of those are real: M against the macro tile and K against the depth unroll. The kernel edges the free dimension, so N need not divide the tile width -- N of 160, 192 and 320 against a 128-wide tile were measured on this project and verify. The spurious N term only ever refused tiles that would have worked. With nothing left to match, the problem came back with no rocRoller kernel at all. Covers the gate with a non-power-of-two N, which no existing case exercised: every shape in the rocRoller suite divided its tile in N by luck.
rocRollerSupportsProblem arrived from a branch that predates this stack's removal of rocRoller's headers from hipblaslt-rocroller-selection, so it named rocRoller::DataType and would not compile here. Ask the same question in the types the rest of this target uses -- hipDataType, via normalizeOperandType -- which is what lets a checkout without shared/rocroller build it. Same predicate, same answers.
A first name hit is wrong once two kernels share a type, tile, StreamK flag, and workgroup mapping.
A 192-wide macrotile hoists ceil(size / tile) as its own kernarg, and every other RR_* value is 32-bit.
A 192-wide macrotile emits a libdivide tree instead of a shift, and that tree had no CustomArgSemantic.
…cates The size, batch and scale-layout conditions bounding a rocRoller custom kernel are derived at library-build time from the solution's state, not written in the logic files. Nothing checked that the derivation actually fired for a kernel in the tree: a missing macrotile, a name that stopped matching the RR_GEMM_ prefix, or a solution dropped by the skip-on-missing-config path would ship an unbounded kernel with every existing test still passing. Add two checks from opposite ends. A corpus test parses the shipped logic the way a build does, converts each solution with the same call, and asserts all 128 rocRoller kernels carry their five predicates with values tracking their own tile; a checked-in manifest pins the per-kernel result so a derivation change is a reviewable diff. A TensileLogic --check-all invariant checks each macrotile is usable and agrees with the _WGT_ token in the kernel's own name, guarding the input the derivation reads.
These captured kernels declare gfx950:sramecc+, the target the library build loads.
GFLOPS is rocprof MinNs, so a shared size keeps the fastest kernel and every captured kernel stays in the logic.
GridBased only returns a kernel that owns an ExactLogic point. The measured tile was already taken, so each of the 446 keeps the next free N and drops the N bound that would reject that point.
The saved ON selection uses one 128x64x128 StreamK tile on 24 of 375 problems, twelve kernels. The symbol omits _WGM_ because StreamK forces workgroup mapping off.
…sed point GridBased only returns a kernel that owns an ExactLogic point. The 128x64 tile key was already taken, so each kernel keeps the next free N and no size pin. The speed is TotalFlops/MinNs at that tile.
The rocRoller-off path rejected that kernel on 4096x4096x1024 because its size window required M, N, and K to stay at 256x256x128, then fell through to the 32x32x64 kernel. Drop the window so N stays free and the same kernel the on path selects can pass.
…nels Those kernels were the saved ON choice at 4096x4096x1024, but SizeLessThan 257 kept OFF on a smaller tile. N stays free because the window is gone and the multiple checks remain.
OFF had no assembly for these tiles, so it could not launch the kernel the measured problems picked.
N stays unbounded, and none of these kernels gained a SizeLessThan 257 pin. The exact-row efficiency is 0 because the rank was not measured.
…oller kernels BoundSizeMultiple already requires K to be a multiple of the depth. SizeLessThan 129 rejected K of 512 and 1024.
Those SizeLessThan caps rejected legal tile multiples, so OFF never launched the kernel ON runs. Paired rocprof mins stay within 2%.
Those SizeLessThan windows rejected the checked-in tiles, so the OFF library fell through to a slower Cijk kernel.
… kernels The measured and StreamK assembly files, and the predicate manifest, had no license comment. They now start with the two-line header.
jaopaulolc
changed the title
Contributor
Author
Contributor
Author
1 task
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-5263
Motivation
hipBLASLt picks a kernel for each matrix-multiply problem by consulting heuristic selection files, called library logic, that TensileLite generates at build time. For the MX ("microscaling") number formats, where each run of 32 values along the summation axis shares one scale factor, two gaps prevented a stock build from ever running the hand-written rocRoller kernels on the gfx950 architecture (the MI350 series). First, the kernels for the three pure MX operand pairs — MXFP8 with MXFP8, MXBF8 with MXBF8, and MXFP4 with MXFP4 — did not exist: #12965 added the mixed pairs and excluded the diagonal. Second, all of the checked-in rocRoller kernels were reachable only from benchmark configurations, never from production library logic, so a default build could not select any of them.
Technical Details
Stacked on #14044. The 39 new kernels declare
.amdgcn_target "amdgcn-amd-amdhsa--gfx950:sramecc+"and cannot assemble without the target-id change in that pull request.This adds 39 rocRoller custom kernels covering the three pure MX operand pairs with non-pre-swizzled scales, completing the diagonal #12965 left out.
It also generates production library logic for all 129 checked-in rocRoller kernels.
preferredExactCustomKernelinLibraryLogic.pygains two new rungs on its selection ladder: a large-MN rung without workgroup mapping, which picks the big tile when it divides the problem, and a small-MN rung with workgroup mapping at a 32x32x128 tile. Those are the two pure-MX tile families the existing ladders could not reach. Each of the two synthetic exact rows is gated on the corresponding kernel being present in the pool, so the F6 and mixed-MX libraries come out unchanged.The generated logic lands under the
gfx950_id75a3chip-id overlay rather than the basegfx950/directory, because a base-directory entry would fail theValidChipIdlint. The logic declares library typeGridBased, matching what both the configurations and the generated files already declare.Batched MX remains out of scope. An earlier revision of this description claimed it was already unreachable because no kernel declares a scale batch stride; that was wrong, and the measurement is below. It is made unreachable here by an explicit predicate.
MX scale layout is now a selection predicate
Registering these kernels through ordinary library logic exposed a gap in what the heuristic knows. The rocRoller kernels read MX block scales as a compact row-major
E8M0array (MXScaleFormat: NoSwizzle). The gfx950 MX problem shards they join also serve callers passingHIPBLASLT_MATMUL_MATRIX_SCALE_BLK32_UE8M0_32_8_EXT, whose scale buffer is already permuted by the AITERe8m0_shuffle. The two buffers have identical extents and the same element type, so no predicate separated them and the heuristic preferred the rocRoller kernel:norm_error0.48-0.88 against 1.5e-06 for theCijk_*control on the same problem.hipBLASLt already carried the distinction --
scaleATypeseparatesBlock_32_UE8M0fromBlock_32_UE8M0_32_8_EXT, which is whatuseRocRoller()reads -- but it stopped atConstructTensileProblem, where both formats collapse onto the samesetHostMXScaleA/Bcall. This pull request forwards it:ContractionProblemGemmgainsmxScaleHostPreSwizzle, set by hipBLASLt from the scaling format and by the TensileLite client from--mx-scale-format. It also keys the problem hash and comparison, since selection depends on it.MXScaleHostPreSwizzleProblemPredicate, emitted per MX solution from itsMXScaleFormat. Per solution rather than fromProblemType.predicates(), becauseMXScaleFormatis a Solution-level knob: the file-levelProblemTypedoes not carry it, and a library-level predicate built from thatProblemTypewould readNoSwizzlefor every file and reject wholeHostPreSwizzlelibraries.InMemorySwizzlemaps to false. That kernel permutes the scales itself and still reads the compact buffer, which is howDataInitializationtreats it.No logic file changes: the 75 custom-kernel files carry
MXScaleFormatonly insideDefaultSolutionand only asAuto, which already resolves toNoSwizzle-- the layout those kernels actually read.Custom kernels now declare the sizes they were captured for
Registering the kernels through ordinary library logic exposed a second, independent gap, of the same kind: over-broad registration on size rather than layout. A rocRoller custom kernel is a pre-assembled object captured at one tile. It carries no generated edge handling in M and no partial-k-tile drain, so it is correct only where M is a whole number of its
MacroTile0and K a whole number of itsDepthU. The generated logic declared neither, and aGridBasedlookup hands the kernel neighbouring sizes it never captured.The resulting wrong answers are nondeterministic, so a single clean run proves nothing: at 1824x512x2048 f8 -> bf16 the same record gave
norm_error0.3505 on seven of eight runs and 0.078 on the eighth. On bf16 output a bad run can sit near the 3.5e-03 rounding floor and read as clean, which is how this survived review.Measured on gfx950 across the FP8, BF8 and FP4 families with the heuristic forced: every size with
M % MacroTile0 == 0andK % DepthU == 0is correct and stable (20 of 20 in an M/N/K sweep); break either and it is wrong. N is deliberately not constrained -- N = 160, 192 and 320 against a 128-wide tile all verify clean, so the kernels do handle the N edge, and rejecting those sizes would make the kernels unreachable for no reason.The size constraint needs no new predicate type: it reuses
Free0SizeMultipleandBoundSizeMultiple, whichAssertFree0ElementMultipleandAssertSummationElementMultiplealready produce. Those assertions are present on all 75 logic files but hardcoded to 1, so they emit nothing. The constraint is scoped by theRR_GEMM_prefix, which is howmatchesRocRollerSelectedKernelidentifies the same family; the AITER and Tensile-generated custom kernels are untouched, having no equivalent measurement.The rocRoller-off path already guards all three dimensions in
getRocRollerSelectedSolutions, which is why that route picks a smaller tile at M=640 and verifies at 1.6e-06 where the heuristic did not.Batch count and addressable extent are now selection predicates
Registering the kernels through ordinary library logic exposed two further gaps of the same kind: a kernel reachable for problems it was never built to handle.
A rocRoller custom kernel is two-dimensional. It takes no batch strides and never reads the third component of its workgroup index, yet it is launched over a grid of
[TilesX, TilesY, Batch], so every slice along the batch axis recomputes batch zero. Measured on MI355X with the Tensile heuristic forced (HIPBLASLT_USE_ROCROLLER=0), f8 x f8 -> f32 at 512x512x512 withbatch_count=4selected a 256x128x128 rocRoller kernel and failed its numerical check on every run, while the same shape atbatch_count=1selected the same kernel and verified at 1.32878e-06. The kernels now declareBatchSizeEqual(1), which is what makes the earlier statement true.Each kernel also reaches a tensor through a single buffer descriptor whose byte-count field is 32 bits wide, and it programs that field once over the whole tensor rather than once per tile. A tensor of 4 GiB or more wraps the field, after which loads read back zero and stores are discarded, so the kernel runs to completion and returns a wrong answer rather than failing. The existing
BufferLoadOffsetLimitCheckandBufferStoreOffsetLimitCheckbound one tile's extent and do not catch this. A newBufferAddressingLimitCheckpredicate — the one new predicate type in this pull request — bounds every tensor the kernel addresses, meaning A, B, C, D and the two MX scale tensors, at 4 GiB. It is the selection-time counterpart ofRocRollerGemmKernel::exceedsBufferAddressingLimiton the rocRoller-enabled path.The extent it compares is the largest
stride x sizeproduct rather than the tightest allocation that holds the elements, because the kernel programs the field from the stride product, so padding in a leading dimension counts in full. The distinction is not academic. At 1024x1024x256 withldcandlddof 1048576 the stride product is exactly 4 GiB and the problem returns wrong answers, while the tightest allocation is 3.996 GiB and would have been admitted.The same pull request also routes block-scaled GEMMs on the object API through rocRoller, and it ships the measured MX kernels plus the StreamK kernels. The object API stores the rocBLASLt problem when the GEMM is created and uses that problem for the same routing decision as the C API. If rocRoller returns no kernel, selection uses the TensileLite heuristic. A StreamK kernel is loaded as a custom kernel and launched on the grid rocRoller computed. The off path matches a checked-in kernel by name, including its unroll and scale-block size, and keeps it only when that kernel's predicates accept the problem. N does not have to divide the tile. Upper bounds on M and K that rejected a legal multiple of the tile are absent on the kernels where that bound was hiding the tile the on path selects.
Test Plan
ValidateMetadata --strictover the full custom-kernel set, to check that every kernel's metadata is self-consistent.TensileLogic --check-allover the gfx950 logic corpus, to check that no solution in the generated logic is rejected at load.test_exact_custom_kernel_ties.py,test_CustomKernels_buffer_oob.py, andtest_rocroller_wgm.py, covering tie-breaking between exact custom-kernel rows, out-of-bounds buffer accesses, and workgroup-mapping selection.HIPBLASLT_ENABLE_ROCROLLER=OFFagainst one with itON, over 12 problem shapes crossed with 3 input type pairs and 3 destination types. For each combination, record which kernel each build selects and compare the numerical output byte for byte.Predicates.BufferAddressingLimitCheck_*inPredicates_test.cpp, covering a tensor exactly at 4 GiB, one a tile narrower that must stay selectable, and one pushed over the limit by leading-dimension padding alone.test_batch_is_pinned_to_one,test_buffer_addressing_limit_is_declaredandtest_batch_and_buffer_limits_are_rocroller_only, checking that both predicates are emitted for a rocRoller custom kernel and for no other family.batch_count4 and 1. The numerical check must be requested in the test record itself: a--verifyflag on the command line does not reach a record supplied by--yaml, and a record that fails its check exits before printing its timing row, so a silent abort is the failure signal rather than a number.Test Result
Completed evidence:
ValidateMetadata --strict: 257 kernels, 0 errors, 0 warnings.TensileLogic --check-all: all 84,499 solutions kept, 0 rejects.OFFbuild as in theONbuild, and on those 90 the numerical results are byte-identical between the two builds.Scale-layout predicate, measured on MI355X (gfx950, PCI 0x75a3):
scaleA=1001now select aHostPreSwizzleCijk_*and verify clean:norm_error1.52e-06, 0, 0, 0, 1.61e-06, 0, 0, 6.40e-05, 5.82e-05, against 0.48-0.88 before.Cijk_*selections are unchanged, the 48RR_GEMM_*selections becomeCijk_*, and no problem loses a selection. NoRR_GEMM_*is selected atscaleA=1001any more.scaleA=3with the Tensile heuristic forced (HIPBLASLT_USE_ROCROLLER=0): of the nine shapes, the five that satisfy the size rule keep the sameRR_GEMM_*kernel and the samenorm_error; the four that violate it now select noRR_GEMM_*at all. Two of those four were visibly wrong before (0.449663 and 0.450189) and two were masked near the bf16 rounding floor but are demonstrably nondeterministic.scaleA=3is unchanged end to end:useRocRoller()routes compact-scale MX to the rocRoller JIT, and all nine shapes select the same JIT kernels and return the same numbers as before.HIPBLASLT_ENABLE_ROCROLLER=OFFbuild,scaleA=3is byte-identical before and after -- that route matches kernels by name and already carries its own divisibility guard -- andscaleA=1001gets the same nine cleanCijk_*picks as theONbuild.ValidateMetadata --strict: 287 kernels, 0 errors, 0 warnings, unchanged.tensilelite/Tests/unit: the failure node-ID set is unchanged (9, all intest_arch_steppings.pyand pre-existing); 11 new tests pass.Predicates.MXScaleHostPreSwizzle*gtests pass, and both the C++ and the Python guards fail when the predicate is mutated to always accept.Size predicate, same build: with N=128, K=128, f8 -> f32 the heuristic previously selected a 256-wide tile at M=640 and 896 and returned
norm_error0.45 and 0.38. Those sizes are now rejected, and M=256, 512, 768 and 1024 still select the same kernel and return the same numbers (1.49843e-06, 1.58511e-06, 1.63033e-06, 1.66026e-06).Batch predicate, measured on MI355X against an
HIPBLASLT_ENABLE_ROCROLLER=OFFbuild with the Tensile heuristic forced, f8 x f8 -> f32 at 512x512x512: before the changebatch_count=4selectedRR_GEMM_TN_FP8_FP8_Float_Float_Float_SA_BE8M0_32_SB_BE8M0_32_WGT_256x128x128_UR_2and failed its numerical check on 3 of 3 runs; after it, the same problem reports no solution on 3 of 3 runs.batch_count=1on that shape is unchanged, selecting the same kernel atnorm_error1.32878e-06 before and after, so the predicate costs the unbatched case nothing.Addressable-extent predicate: the three
Predicates_test.cppcases are new and the file passes 16 of 16. The padded case is the one that discriminates the chosen formula — its stride product is exactly 4 GiB and its tightest allocation is 3.996 GiB, so a predicate written against the allocation would have admitted a shape that returns wrong answers on hardware. The 4 GiB behaviour itself was reported rather than reproduced here; the predicate reproduces the reported boundary, not the wrong answer.Pending: automated GitHub checks have not reported on this branch. The selection sweep was run locally on MI355X and does not run in a shared CI lane today.
Submission Checklist
Risk level
Medium. This makes 129 kernels selectable from a default build that previously could not reach any of them, so kernel choice changes for gfx950 MX problems that now land on one of the two new ladder rungs. The blast radius is bounded by the overlay and the gating: the logic is confined to the
gfx950_id75a3chip-id overlay, and each new exact row is emitted only when its kernel is in the pool, which leaves the F6 and mixed-MX libraries byte-unchanged. 90 of 108 combinations in that earlier sweep selected the same kernel as the rocRoller-enabled build. The 18 StreamK rows were measured before this pull request included the StreamK kernels, and that sweep has not been repeated.The scale-layout predicate narrows selection rather than widening it, and the measured blast radius is the 48 problems that were returning wrong answers: 251 of 251 pre-existing
HostPreSwizzleselections are unchanged and no problem loses a solution.All four predicates narrow selection rather than widening it. The size predicate removes sizes the kernels were never captured for; in the measured set the production path already served those sizes from the rocRoller JIT, so no problem loses a working kernel. The batch predicate removes a case that was returning wrong answers, and leaves
batch_count=1byte-identical. The addressable-extent predicate removes shapes at or past 4 GiB, which is the only predicate here that rejects on a bound rather than on a layout or a multiple; it is the one to revisit if a future rocRoller kernel programs its descriptors per tile rather than once over the whole tensor.The folded kernel and routing commits were not re-run as a suite after they were moved onto this branch. Automated GitHub checks are pending.