Repository navigation
feat(tensilelite): add gfx950 pure F6 custom kernels - #12964
Merged
Merged
Conversation
3 tasks done
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. |
jaopaulolc
added
the
rocm:gemm algos
3 tasks done
Codecov Report❌ Patch coverage is Additional details and impacted files@@ Coverage Diff @@
## develop #12964 +/- ##
===========================================
- Coverage 76.02% 76.02% -0.01%
===========================================
Files 2867 2867
Lines 474745 474806 +61
Branches 71078 71092 +14
===========================================
+ Hits 360920 360935 +15
- Misses 89314 89351 +37
- Partials 24511 24520 +9
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
jaopaulolc
force-pushed
the
users/jolabega/kfa-rr-pure-f6
branch
from
October 6, 2026 20:07
872e957 to
7ab6832
Compare
AlexBrownAMD
approved these changes
Oct 6, 2026
jaopaulolc
force-pushed
the
users/jolabega/kfa-rr-pure-f6
branch
10 times, most recently
from
October 14, 2026 22:24
c4dcaae to
4ffc4ae
Compare
Check in the selector-chosen FP6 and BF6 Gemm-From-Anywhere kernels so those GEMMs can ship as Tensile custom kernels.
rocRoller kernels declare StaggerU: 0 and do not implement the in-loop wrap. Counting and assembling them made this test fail for every kernel added to the stack.
New files take the two-line SPDX header from the project conventions.
Prebuilt custom kernels do not carry EnableMatrixInstruction or the derived MI keys. test_config rejected every rocRoller config for that.
The coverage ratchet failed because the new early return was untested.
gfx950 F6/B6 rocRoller custom-kernel tests abort in tensilelite-client
for two reasons.
Pure F6/B6 problems have Half or BFloat16 destinations, and the client
reference has no case for those destination types. Add them.
Mixed F6/B6 problems (B6F6_B6{B,S,H}S and F6B6_F6{B,S,H}S) still abort.
Custom kernels report their compute inputs as (DataType, DataType), that
is (B6,B6) or (F6,F6), but the client aliases expect (A,B). Set the
custom-kernel compute inputs from the A and B types in Contractions.py.
The same logic is on #12981 as 32a0e97, and the aliases here assume
it, so it belongs on this PR.
Verified with test_config[custom_rr_f6.yaml] against a prebuilt client.
jaopaulolc
force-pushed
the
users/jolabega/kfa-rr-pure-f6
branch
from
October 16, 2026 15:02
cc6ada9 to
43ac1d7
Compare
jaopaulolc
added a commit
that referenced
this pull request
Oct 17, 2026
JIRA ID : AIHPBLAS-5026 ## Motivation Pure F6 rocRoller kernels are checked in on the parent branch. The other half of the request is the mixed MX set: A and B drawn from MXFP8, BF8, MXFP6, BF6, and MXFP4, excluding pure FP8, pure BF8, and pure FP4. ## Technical Details Adds 54 gfx950 TN kernels for those mixed pairs, again the kernel hipBLASLt's rocRoller selector returned for the in-tree MX client problems. Assembly is under `Tensile/CustomKernels/rocroller/`. `Tensile/Tests/custom/custom_rr_mx.yaml` benchmarks them on gfx950 only. The census moves from 141 kernels / 81 declared-zero StaggerU to 195 / 135. The new names are on `NO_BUFFER_OOB_KERNELS`. Not included: 20 workgroup-mapped kernels from `matmul_mx_all_kernels` (f6×f6 to f16, and f8×f4 to f16). Their kernargs carry launch-time `MagicMultiple`, `MagicShiftAndSign`, `Subtract`, and `WGM` values that Tensile custom-arg semantics do not compute, so a config would not launch them. Those belong in a follow-up PR. Production logic does not reference these kernels. rocRoller stays a dependency. ## Test Plan - [x] `python -m Tensile.ValidateMetadata --strict` - [x] `pytest` on `test_custom_kernel_staggeru_metadata.py` and `test_CustomKernels_buffer_oob.py` ## Test Result `ValidateMetadata --strict`: 195 kernels, 0 errors, 0 warnings. Stagger census and buffer-bounds tests: 822 passed. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low. Additive custom-kernel sources and a gfx950-only Tensile test, stacked on the pure-F6 PR. No production logic change. Rollback is reverting this PR. Depends on #12964. Merge that one first.
Contributor
TheRock Submodule Bump ActivityNewest first
|
jaopaulolc
added a commit
that referenced
this pull request
Oct 19, 2026
JIRA ID : AIHPBLAS-5026 ## Motivation The mixed-MX branch left out 20 gfx950 rocRoller kernels because their kernargs include launch-time MagicMultiple, MagicShiftAndSign, Subtract, and WGM values. Tensile custom-arg semantics could not compute them, so a custom.config would not launch the kernel. This PR teaches custom-kernel launch to evaluate those values the way rocRoller does, and checks the kernels in. ## Technical Details The eleven hoisted arguments are one closed formula, the same on all 20 kernels. Tile counts are int32 ceil-divs of M (`Tensor_0_size_0`) and N (`Tensor_2_size_1`) by the power-of-two macrotile. hipBLASLt passes WGM = 2 (`DEFAULT_WGM`). The quotients are rocRoller's signed libdivide branchfree division, and the other slots are MagicMultiple / MagicShiftAndSign of WGM, numTilesN, the tail, and the main block. Products wrap as int32. Divisor 0 uses rocRoller's special case (`INT32_MAX/2`, shift-and-sign 0), not a raw libdivide call and not Tensile `magicNumber()`. rocRoller is not linked into TensileLite. The same arithmetic lives in `RocRollerWgmKernargs.hpp` for launch and `Tensile/RocRollerWgm.py` for the unit test. New `CustomArgSemantic` values are appended at the end of the macro, so existing ordinals stay put. Only custom-kernel launch reads them. `Tensile/Tests/custom/custom_rr_wgm.yaml` covers the 20 kernels on gfx950 only. Assembly is under `Tensile/CustomKernels/rocroller/`, with the `_gfx950-sramecc+` suffix dropped. The rocRoller `.kernel_graph` debug section is omitted so metadata parsing stays tractable; instructions and `.args` expressions are unchanged. The census moves from 195 kernels / 135 declared-zero StaggerU to 215 / 155. The new names are on `NO_BUFFER_OOB_KERNELS`. This PR stacks on #12965 which stacks on #12964. rocRoller stays a dependency. Production logic is unchanged. ## Test Plan - [x] `python -m Tensile.ValidateMetadata --strict` - [x] `python -m pytest Tensile/Tests/unit/test_rocroller_wgm.py Tensile/Tests/unit/test_custom_kernel_staggeru_metadata.py Tensile/Tests/unit/test_CustomKernels_buffer_oob.py Tensile/Tests/unit/test_CustomKernelMetadata.py -q --tb=line` ## Test Result From the worktree `projects/hipblaslt/tensilelite`, with `PYTHONPATH` set to the existing rocisa build and then that tensilelite tree: `python -m Tensile.ValidateMetadata --strict` ``` Summary: 215 kernel(s), 0 error(s), 0 warning(s) ``` `python -m pytest Tensile/Tests/unit/test_rocroller_wgm.py Tensile/Tests/unit/test_custom_kernel_staggeru_metadata.py Tensile/Tests/unit/test_CustomKernels_buffer_oob.py Tensile/Tests/unit/test_CustomKernelMetadata.py -q --tb=line` ``` 1041 passed in 175.63s (0:02:55) ``` The stagger census matches 215 kernels and 155 declared-zero StaggerU. The new unit test checks libdivide branchfree magic for signed and unsigned divisors, including the rocRoller 0 and 1 edges, and checks the C++ launch header against the same values. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low. Additive custom-kernel sources, a gfx950-only Tensile test, and custom-kernel launch evaluation for the new semantics. Generated-kernel ABI is unchanged. Rollback is reverting this PR. Depends on #12965. Merge that one first.
1 task done
1 task done
tony-davis
added a commit
that referenced
this pull request
Oct 22, 2026
… custom.config where possible (#14058) ISSUE ID : #14049 ## Motivation TensileLite ships hand-written GPU kernels, called custom kernels, as assembly files. Each file carries a block of YAML metadata describing the kernel's arguments and how Tensile may use it. The 69 rocRoller kernels (rocRoller is a GPU kernel generator) added by #12964 and #12965 also embed rocRoller's internal kernel graph in that block: about 13,000 lines each, against about 175 for the other kernels. Tensile parsed the block with PyYAML's pure-Python loader, which takes over a second per rocRoller kernel, and the unit tests read every shipped kernel more than once. TheRock's `Test tensilelite` job on gfx94X (MI300-series GPUs) went from 11 to 14 minutes to 17 to 21 minutes, and quick-tier runs, which have a 15-minute limit, now time out (#14049). PyYAML's C loader makes each parse faster but still parses everything. Across the 248 shipped kernels the metadata YAML totals 70.2 MB, while the `custom.config` entry that most callers read is 0.34 MB (0.48%). The rest, mostly the rocRoller kernel graph, is parsed and then discarded. ## Technical Details - `CustomKernels._readEmbeddedYaml` parses with PyYAML's C loader (`CSafeLoader`, built on the libyaml library) and falls back to the pure-Python `SafeLoader` when PyYAML was built without libyaml. `LibraryIO.py` already loads YAML this way; the loader is declared locally because `LibraryIO` imports `CustomKernels`, so importing it from there would be circular. - `readCustomKernelConfig` parses only the top-level `custom.config` entry, which `_customConfigYaml` extracts as text: from the `custom.config:` line up to the next line that starts in column 0 with anything other than a comment. The text ends with a newline, so a block scalar that ends the entry keeps its final newline, as it does in a whole-block parse. - The whole block is parsed instead, as before, in three cases: `custom.config` is not found in that form (for example a quoted key); `custom.config:` appears at the start of more than one line, since PyYAML keeps the last duplicate and the same line can sit inside a multi-line quoted scalar; or the extracted text does not parse, as with a flow collection closed in column 0 or an alias to an anchor defined outside the entry. A real error inside `custom.config` still raises `RuntimeError("Failed to parse YAML for custom kernel ...")`. - `getCustomKernelConfig` reads `custom.config` that way and parses the whole block only for a kernel with no explicit `CustomKernel` block, because only then does it need `amdhsa.kernels` to infer the argument list. 109 of the 248 shipped kernels have the block, including all 90 rocRoller kernels, so those never parse their roughly 1 MB kernel graph. The other 139 (119 Tensile kernels and 20 gfx1151 kernels) are parsed once, as before. - Behavior change: `readCustomKernelConfig` no longer reports malformed YAML outside `custom.config`, so `validate_all` and the default-root CLI test no longer catch a broken `amdhsa.kernels` section. Parsing the whole block still reports it. - `Tests/unit/test_CustomKernelMetadata.py` gains 17 test functions (26 test cases). They check that `CSafeLoader` is used when available; that both loaders return the same document for a large rocRoller kernel; that the whole block is parsed once for a kernel with no `CustomKernel` block; that only the `custom.config` entry is parsed; that comments, blank lines, neighbouring keys, flow style and a block scalar ending the entry are read correctly; that the duplicate-key, quoted-key and unparseable-slice cases take the whole-block path; that the slice equals a whole-block parse for shipped kernels under 600 KB; and that every shipped kernel takes the slice, which holds only `custom.config`. - `tensilelite/TESTING.md` gains a Known Risks row: TheRock's quick tier runs the same unit tests as the standard tier under a 15-minute limit, and the job's time grows with each custom kernel added. No kernel, tuning logic, or C++ code changes. ## Test Plan - Compare `getCustomKernelConfig` output for every shipped kernel between `develop` and this branch, with an empty and a non-empty `InternalSupportParams`. - Time `readCustomKernelConfig` and `getCustomKernelConfig` over every shipped kernel, collection of `rocisa/test` plus `tensilelite/Tests/unit`, and `test_validatemetadata_main_default_root_uses_shipped_kernels`. - Run `test_CustomKernelMetadata.py` and the other test files that read custom-kernel config. - Check the scanner tests by temporarily breaking the scanner and confirming a test fails. - In TheRock Multi-Arch CI, compare this PR's gfx94X `Test tensilelite` collection and `test_custom_kernel_cli.py` times against `develop`. ## Test Result Local runs, on a shared 96-core Linux machine with rocisa built from this branch. Times are for a fresh process over all 248 kernels. The first two columns are single runs of the collection and CLI test, measured before this branch was rebased onto `develop` (on `d1947ecc545`, 42 commits earlier); the `CustomKernels.py` code and the kernel files are the same in both bases. | | develop (no change) | libyaml only (first commit) | this PR | |---|---|---|---| | `getCustomKernelConfig` over all kernels | 155.6 s | 28.9–29.2 s | 1.43–1.45 s | | `readCustomKernelConfig` over all kernels | 155.3 s | 29.4–29.5 s | 0.85–0.89 s | | Collecting `rocisa/test` + `tensilelite/Tests/unit` | 242 s | 134 s | 56–57 s | | `test_validatemetadata_main_default_root_uses_shipped_kernels` | 158 s | 31.5 s | 0.56 s | - `getCustomKernelConfig` output matches `develop` (`745b66c929a`) for all 248 kernels with two `InternalSupportParams` inputs: 496 results, 0 mismatches, 0 errors. - `test_CustomKernelMetadata.py`: 130 passed. The metadata file together with the other custom-kernel tests (CLI, library IO, persistent arguments, host and occupancy tests, and the characterization tests that read custom kernels): 673 passed, 3 skipped. - Scanner checks: breaking the scanner in five ways (dropping the `(\s|$)` boundary, an unescaped `.` in the key pattern, `.search` instead of `.match`, returning `None`, not stopping at the end of the entry) makes at least one test fail each time. - `test_custom_kernel_staggeru_metadata.py` did not pass locally: its tests fail when the local assembler rejects gfx1151 `.amdhsa_*` directives. This is an environment limit, not a parsing failure; collection of that file works. - In TheRock Multi-Arch CI, on the first commit of this PR: gfx94X `Test tensilelite` ([job 119006629002](https://github.com/ROCm/rocm-libraries/actions/runs/39498719943/job/119006629002)) passed 11,307 tests with 147 skipped, 2 expected failures and 1 unexpected pass, on the standard tier. Against `develop`'s first gfx94X job after the rocRoller kernels landed ([job 117503989326](https://github.com/ROCm/rocm-libraries/actions/runs/39064938907/job/117503989326)), collection went from 285 s to 177 s, `test_custom_kernel_cli.py` from 169 s to 32 s, and the pytest stage from 17.5 to 14.0 minutes. That run was under the standard tier's 30-minute limit, so it does not show the quick tier fitting under 15. - Automated GitHub checks on the current head are pending. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (2/5): a host-side Python change to how custom-kernel metadata is read; no GPU, kernel, tuning or C++ code changes, so no device coverage is needed. Output is identical to `develop` on all 248 shipped kernels, and the code falls back to the previous loader, and to a whole-block parse whenever it cannot isolate or parse `custom.config`. The one loss is that a malformed `amdhsa.kernels` section in a kernel with a `CustomKernel` block is no longer reported. --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Alex Vasile <48962821+Alex-Vasile@users.noreply.github.com>
1 task done
jaopaulolc
added a commit
that referenced
this pull request
Nov 4, 2026
…me matching (#14043) JIRA ID : AIHPBLAS-5261 ## Motivation Two independent defects on gfx950 MX return wrong numbers: the row pitch hipBLASLt reports for the scale tensor does not match the layout the kernel was compiled against, and the test that picks a kernel by inspecting its name accepts kernels that do not fit the problem. Measured over 108 gfx950 MX problem and type combinations, 36 produced numerically wrong results, some containing `inf`. ## Technical Details This pull request is stacked on #14086. The first fix concerns the scale row pitch. `ContractionProblemGemm::setMXScaleA/B` rounds the scale tensor's K-block dimension up with `RoundUpToMultiple(CeilDivide(K, mxBlock), 8)` on gfx950. `generateCustomCall` passed that padded value as the kernel arguments `StrideScaleA0` and `StrideScaleB0`, but the rocRoller kernels are built against the unpadded `ceil(K/32)`. The two values are equal only when `K % 256 == 0`. Otherwise every scale row after the first resolves to the wrong bytes, and the kernel reads past the end of the caller's scale buffer. `generateCustomCall` now takes both strides from the buffer the caller passed. `customKernelMXScaleAIsPacked()` and `customKernelMXScaleBIsPacked()` name that caller, one operand at a time, because hipBLASLt takes the A and B scale formats independently. When an operand's flag is set, its row stride is the unpadded `ceil(K/mxBlock)`. When it is not, the descriptor stride is already the one the kernel must use. The padding is deliberately not removed globally. The gfx950 MX kernels that TensileLite generates overwrite this pitch in their own prologue from the problem's `SizesSum` field, so the kernel argument is dead for them and the padded descriptor must stay for the rest of the library. The change is confined to `generateCustomCall`, which runs only for externally captured custom kernels. The batch stride `StrideScale{A,B}1` follows the same caller buffer: a packed caller gets the product of the unpadded scale sizes, and an unpacked caller keeps the descriptor stride. No shipped rocRoller kernel reads that batch stride, and those kernels run a single batch. This also repairs the FP6 custom kernels added by #12964 and #12965, which reach the same code path. hipBLASLt declares each operand packed unless its format is `Block_32_UE8M0_32_8_EXT`: that caller pre-swizzled its scales into the padded gfx950 layout, which the descriptor already describes. The TensileLite benchmark client leaves both flags false, because it allocates from the padded descriptor. The second fix concerns `matchesRocRollerSelectedKernel`, the predicate that decides whether a kernel's name describes the problem at hand. It has three defects, all reachable on the `HIPBLASLT_ENABLE_ROCROLLER=OFF` build: - The tile test was the bare substring `WGT_<m>x<n>x<k>_UR_`, which occurs inside both `_WGT_..._UR_` and `_SK_WGT_..._UR_`. A plain tile therefore matched a StreamK kernel — StreamK is a work-partitioning scheme that splits the summation axis across workgroups — and a StreamK tile matched a plain kernel. The test is now anchored as `_WGT_...` with an explicit `_SK` discriminator. - The prefix match stopped at `_Float_` while the tile match began at `WGT_`, so the scale-layout tokens between the two were never inspected. A pre-swizzled kernel, named with `PSTA_`, `PTA_`, `PSTB_`, or `PTB_`, could match a problem whose scales are not pre-swizzled. Those tokens are now rejected. - A StreamK kernel name never carries `_WGM_`, because `parameter_selection.cpp` forces `workgroupMappingDim = -1` for StreamK. The workgroup-mapping comparison now runs only when `workgroupMapping && !streamK`, and StreamK tiles are no longer rejected outright. `getRocRollerSelectedSolutions` also gains the workspace rejection that the rocRoller-enabled path already applies at `gemm.cpp:703-706`. Without it an `OFF` build hands back a StreamK kernel that an `ON` build would have skipped, because `_convertToHeuristicResultArray` reports `requiredWorkspaceSize` but never filters on it. ### Name matching is not enough on its own A kernel's name carries its tile shape and its operand types, and nothing else. Every other condition the kernel was captured under — which scale layout it reads, the size multiples and size windows it requires, how many batches it handles, how large a tensor it can address — lives in the kernel's selection predicates. An ordinary library lookup evaluates those; this hand-rolled search skipped them. A problem that block-scales only one of its two operands therefore matched a kernel built to read two scale tensors, and that kernel read a scale buffer the caller had never allocated. `getRocRollerSelectedSolutions` now requires a candidate's `hardwarePredicate` and `problemPredicate` to accept the problem before it takes a name match, and keeps searching when they do not. It tests hardware and problem only, as the grouped-gemm path in the same file does: `taskPredicate` needs a `Task`, which needs the workspace size that is tested immediately afterwards. This narrows what the `OFF` path selects. Where the kernel it used to match declares a size window that excludes the problem, it now picks a different one. That is the intended behaviour, but it has a cost, measured below: the only candidates on this path are other rocRoller kernels, so when the best-fitting tile is correctly rejected there may be nothing comparable left in the pool. The predicate moves to `rocroller_kernel_match.hpp`, whose only dependency is `<string>`, so it can be tested without a problem descriptor or a GPU. No longer rejecting StreamK outright changes which kernel an `OFF` build selects. That is intended — it is what lets an `OFF` build reach the StreamK kernels at all — but this pull request tests it only at the name-matching level. The selection-identity sweep that compares an `OFF` build against an `ON` build lands with the logic that makes these kernels selectable, later in this stack. ## Test Plan - `CustomKernelTest.ScaleStrideIsCompactNotPadded`, `CustomKernelTest.ScaleStrideIsCompactForEitherBoundDimension`, `CustomKernelTest.ScaleBatchStrideFollowsTheCallerBuffer`, `CustomKernelTest.ScaleStrideFollowsDescriptorWhenCallerDoesNotPack`, `CustomKernelTest.ScaleStridesIncludeTheInitialStrideForDataParallelArgs`, and `CustomKernelTest.ScaleStrideRoundsUpPartialBlocks`: assert that `generateCustomCall` emits the caller's scale stride. A packed caller with a `K` that is not a multiple of 256 gets `ceil(K/mxBlock)`, for both bound dimensions and for the batch stride. An unpacked caller gets the descriptor stride. - `HostUnitRocRollerKernelMatch.*`: five host-only tests covering the StreamK and plain tile discrimination, the pre-swizzle token rejection, and the workgroup-mapping comparison. Mutation-check them by running them against the pre-patch predicate; a test that passes against the old code does not discriminate. - `CustomKernelTest.ScaleStrideFollowsEachOperandsBuffer`: a packed A with a padded B, and the reverse, each get their own buffer's row and batch strides. - `HostUnitMXScaleSetup.smoke_PreSwizzledDeclarationIsPerOperand`: a pre-swizzled operand is declared unpacked, and declaring one operand does not change the other. - Numerical parity sweep on gfx950 over 108 MX problem and type combinations against a 1e-6 tolerance, comparing each result with and without this change. - A problem that block-scales one operand only — `scaleA 3`, `scaleB 0`, bf8 inputs, 1024x1024x1024 — on an `HIPBLASLT_ENABLE_ROCROLLER=OFF` build. It must report no solution rather than launching a kernel. - A selection sweep on the same build over 8 shapes that do block-scale both operands, recording which kernel each build picks and whether it verifies, to show the predicates narrow selection without rejecting problems the kernels do serve. Numerical checks must be requested in the test record itself, because a `--verify` flag on the command line does not reach a record supplied by `--yaml`. ## Test Result Completed evidence: - Both `CustomKernelTest` cases are new and pass. - All five `HostUnitRocRollerKernelMatch` tests are new and pass. Against the pre-patch predicate, 4 of the 5 fail, so they discriminate the change. - Parity sweep: 36 of the 108 combinations were numerically wrong before the change, with a normalized error as large as 6508 and `inf` in some cases. After the change 0 are wrong and none regressed; the 63 combinations that were already clean are bit-identical. Predicate evaluation, measured on MI355X (gfx950, PCI 0x75a3) against an `HIPBLASLT_ENABLE_ROCROLLER=OFF` build, with the numerical check set in the test record: - One-sided MX, bf8 x bf8 -> f32 at 1024x1024x1024 with `scaleA 3` and `scaleB 0`: before the change the run ends in `an illegal memory access was encountered (700)`; after it, the problem reports no solution, on 2 of 2 runs. - Selection sweep, f8 x f8 -> f32 with both operands block-scaled, over 8 shapes: every shape still receives a solution and every shape verifies, on both builds, so no problem the kernels serve is rejected. Five of the eight change kernel. For instance `RR_GEMM_..._WGT_32x32x64_UR_2` declares M and N between 64 and 128 and K exactly 128, and was being selected at 512x512x512. - The narrowing costs performance where the rejected tile was the only large one available. At 4096x4096x1024 the selection moves from `RR_GEMM_..._WGT_256x128x128_UR_2`, whose declared window ends at M = 2560, to `RR_GEMM_..._WGT_16x16x128_UR_2`, and the kernel time rises by a factor of about 3.6 over three repeats. That fallback, routing a problem the rocRoller path declines to ordinary Tensile selection, is #14497, later in this stack. It is not in this pull request. Pending: automated GitHub checks have not reported on this branch. The parity sweep was run locally on gfx950 hardware and does not run in a shared CI lane today. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Medium. Two pieces of device-visible behavior change. The scale stride passed to custom MX kernels on gfx950 changes value whenever `K % 256 != 0`, which is the defect being fixed; kernels generated by TensileLite are unaffected because they overwrite that argument. Separately, an `HIPBLASLT_ENABLE_ROCROLLER=OFF` build can now select StreamK kernels it previously rejected, so kernel choice on that path changes; the numerical consequence of that selection change is measured later in this stack rather than here. Evaluating the predicates narrows that path further, and the measured blast radius is kernel choice rather than correctness: across the 8 shapes swept, no problem lost a solution and none failed verification, while 5 changed kernel. The cost is performance, not accuracy, and it falls on shapes that no checked-in rocRoller kernel was captured for — previously served by a kernel used outside its declared window. The largest measured case is about 3.6x slower. That cliff is removed by the fallback in #14497, later in this stack, not by this pull request.
cycheng
pushed a commit
that referenced
this pull request
Nov 6, 2026
…me matching (#14043) JIRA ID : AIHPBLAS-5261 ## Motivation Two independent defects on gfx950 MX return wrong numbers: the row pitch hipBLASLt reports for the scale tensor does not match the layout the kernel was compiled against, and the test that picks a kernel by inspecting its name accepts kernels that do not fit the problem. Measured over 108 gfx950 MX problem and type combinations, 36 produced numerically wrong results, some containing `inf`. ## Technical Details This pull request is stacked on #14086. The first fix concerns the scale row pitch. `ContractionProblemGemm::setMXScaleA/B` rounds the scale tensor's K-block dimension up with `RoundUpToMultiple(CeilDivide(K, mxBlock), 8)` on gfx950. `generateCustomCall` passed that padded value as the kernel arguments `StrideScaleA0` and `StrideScaleB0`, but the rocRoller kernels are built against the unpadded `ceil(K/32)`. The two values are equal only when `K % 256 == 0`. Otherwise every scale row after the first resolves to the wrong bytes, and the kernel reads past the end of the caller's scale buffer. `generateCustomCall` now takes both strides from the buffer the caller passed. `customKernelMXScaleAIsPacked()` and `customKernelMXScaleBIsPacked()` name that caller, one operand at a time, because hipBLASLt takes the A and B scale formats independently. When an operand's flag is set, its row stride is the unpadded `ceil(K/mxBlock)`. When it is not, the descriptor stride is already the one the kernel must use. The padding is deliberately not removed globally. The gfx950 MX kernels that TensileLite generates overwrite this pitch in their own prologue from the problem's `SizesSum` field, so the kernel argument is dead for them and the padded descriptor must stay for the rest of the library. The change is confined to `generateCustomCall`, which runs only for externally captured custom kernels. The batch stride `StrideScale{A,B}1` follows the same caller buffer: a packed caller gets the product of the unpadded scale sizes, and an unpacked caller keeps the descriptor stride. No shipped rocRoller kernel reads that batch stride, and those kernels run a single batch. This also repairs the FP6 custom kernels added by #12964 and #12965, which reach the same code path. hipBLASLt declares each operand packed unless its format is `Block_32_UE8M0_32_8_EXT`: that caller pre-swizzled its scales into the padded gfx950 layout, which the descriptor already describes. The TensileLite benchmark client leaves both flags false, because it allocates from the padded descriptor. The second fix concerns `matchesRocRollerSelectedKernel`, the predicate that decides whether a kernel's name describes the problem at hand. It has three defects, all reachable on the `HIPBLASLT_ENABLE_ROCROLLER=OFF` build: - The tile test was the bare substring `WGT_<m>x<n>x<k>_UR_`, which occurs inside both `_WGT_..._UR_` and `_SK_WGT_..._UR_`. A plain tile therefore matched a StreamK kernel — StreamK is a work-partitioning scheme that splits the summation axis across workgroups — and a StreamK tile matched a plain kernel. The test is now anchored as `_WGT_...` with an explicit `_SK` discriminator. - The prefix match stopped at `_Float_` while the tile match began at `WGT_`, so the scale-layout tokens between the two were never inspected. A pre-swizzled kernel, named with `PSTA_`, `PTA_`, `PSTB_`, or `PTB_`, could match a problem whose scales are not pre-swizzled. Those tokens are now rejected. - A StreamK kernel name never carries `_WGM_`, because `parameter_selection.cpp` forces `workgroupMappingDim = -1` for StreamK. The workgroup-mapping comparison now runs only when `workgroupMapping && !streamK`, and StreamK tiles are no longer rejected outright. `getRocRollerSelectedSolutions` also gains the workspace rejection that the rocRoller-enabled path already applies at `gemm.cpp:703-706`. Without it an `OFF` build hands back a StreamK kernel that an `ON` build would have skipped, because `_convertToHeuristicResultArray` reports `requiredWorkspaceSize` but never filters on it. ### Name matching is not enough on its own A kernel's name carries its tile shape and its operand types, and nothing else. Every other condition the kernel was captured under — which scale layout it reads, the size multiples and size windows it requires, how many batches it handles, how large a tensor it can address — lives in the kernel's selection predicates. An ordinary library lookup evaluates those; this hand-rolled search skipped them. A problem that block-scales only one of its two operands therefore matched a kernel built to read two scale tensors, and that kernel read a scale buffer the caller had never allocated. `getRocRollerSelectedSolutions` now requires a candidate's `hardwarePredicate` and `problemPredicate` to accept the problem before it takes a name match, and keeps searching when they do not. It tests hardware and problem only, as the grouped-gemm path in the same file does: `taskPredicate` needs a `Task`, which needs the workspace size that is tested immediately afterwards. This narrows what the `OFF` path selects. Where the kernel it used to match declares a size window that excludes the problem, it now picks a different one. That is the intended behaviour, but it has a cost, measured below: the only candidates on this path are other rocRoller kernels, so when the best-fitting tile is correctly rejected there may be nothing comparable left in the pool. The predicate moves to `rocroller_kernel_match.hpp`, whose only dependency is `<string>`, so it can be tested without a problem descriptor or a GPU. No longer rejecting StreamK outright changes which kernel an `OFF` build selects. That is intended — it is what lets an `OFF` build reach the StreamK kernels at all — but this pull request tests it only at the name-matching level. The selection-identity sweep that compares an `OFF` build against an `ON` build lands with the logic that makes these kernels selectable, later in this stack. ## Test Plan - `CustomKernelTest.ScaleStrideIsCompactNotPadded`, `CustomKernelTest.ScaleStrideIsCompactForEitherBoundDimension`, `CustomKernelTest.ScaleBatchStrideFollowsTheCallerBuffer`, `CustomKernelTest.ScaleStrideFollowsDescriptorWhenCallerDoesNotPack`, `CustomKernelTest.ScaleStridesIncludeTheInitialStrideForDataParallelArgs`, and `CustomKernelTest.ScaleStrideRoundsUpPartialBlocks`: assert that `generateCustomCall` emits the caller's scale stride. A packed caller with a `K` that is not a multiple of 256 gets `ceil(K/mxBlock)`, for both bound dimensions and for the batch stride. An unpacked caller gets the descriptor stride. - `HostUnitRocRollerKernelMatch.*`: five host-only tests covering the StreamK and plain tile discrimination, the pre-swizzle token rejection, and the workgroup-mapping comparison. Mutation-check them by running them against the pre-patch predicate; a test that passes against the old code does not discriminate. - `CustomKernelTest.ScaleStrideFollowsEachOperandsBuffer`: a packed A with a padded B, and the reverse, each get their own buffer's row and batch strides. - `HostUnitMXScaleSetup.smoke_PreSwizzledDeclarationIsPerOperand`: a pre-swizzled operand is declared unpacked, and declaring one operand does not change the other. - Numerical parity sweep on gfx950 over 108 MX problem and type combinations against a 1e-6 tolerance, comparing each result with and without this change. - A problem that block-scales one operand only — `scaleA 3`, `scaleB 0`, bf8 inputs, 1024x1024x1024 — on an `HIPBLASLT_ENABLE_ROCROLLER=OFF` build. It must report no solution rather than launching a kernel. - A selection sweep on the same build over 8 shapes that do block-scale both operands, recording which kernel each build picks and whether it verifies, to show the predicates narrow selection without rejecting problems the kernels do serve. Numerical checks must be requested in the test record itself, because a `--verify` flag on the command line does not reach a record supplied by `--yaml`. ## Test Result Completed evidence: - Both `CustomKernelTest` cases are new and pass. - All five `HostUnitRocRollerKernelMatch` tests are new and pass. Against the pre-patch predicate, 4 of the 5 fail, so they discriminate the change. - Parity sweep: 36 of the 108 combinations were numerically wrong before the change, with a normalized error as large as 6508 and `inf` in some cases. After the change 0 are wrong and none regressed; the 63 combinations that were already clean are bit-identical. Predicate evaluation, measured on MI355X (gfx950, PCI 0x75a3) against an `HIPBLASLT_ENABLE_ROCROLLER=OFF` build, with the numerical check set in the test record: - One-sided MX, bf8 x bf8 -> f32 at 1024x1024x1024 with `scaleA 3` and `scaleB 0`: before the change the run ends in `an illegal memory access was encountered (700)`; after it, the problem reports no solution, on 2 of 2 runs. - Selection sweep, f8 x f8 -> f32 with both operands block-scaled, over 8 shapes: every shape still receives a solution and every shape verifies, on both builds, so no problem the kernels serve is rejected. Five of the eight change kernel. For instance `RR_GEMM_..._WGT_32x32x64_UR_2` declares M and N between 64 and 128 and K exactly 128, and was being selected at 512x512x512. - The narrowing costs performance where the rejected tile was the only large one available. At 4096x4096x1024 the selection moves from `RR_GEMM_..._WGT_256x128x128_UR_2`, whose declared window ends at M = 2560, to `RR_GEMM_..._WGT_16x16x128_UR_2`, and the kernel time rises by a factor of about 3.6 over three repeats. That fallback, routing a problem the rocRoller path declines to ordinary Tensile selection, is #14497, later in this stack. It is not in this pull request. Pending: automated GitHub checks have not reported on this branch. The parity sweep was run locally on gfx950 hardware and does not run in a shared CI lane today. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Medium. Two pieces of device-visible behavior change. The scale stride passed to custom MX kernels on gfx950 changes value whenever `K % 256 != 0`, which is the defect being fixed; kernels generated by TensileLite are unaffected because they overwrite that argument. Separately, an `HIPBLASLT_ENABLE_ROCROLLER=OFF` build can now select StreamK kernels it previously rejected, so kernel choice on that path changes; the numerical consequence of that selection change is measured later in this stack rather than here. Evaluating the predicates narrows that path further, and the measured blast radius is kernel choice rather than correctness: across the 8 shapes swept, no problem lost a solution and none failed verification, while 5 changed kernel. The cost is performance, not accuracy, and it falls on shapes that no checked-in rocRoller kernel was captured for — previously served by a kernel used outside its declared window. The largest measured case is about 3.6x slower. That cliff is removed by the fallback in #14497, later in this stack, not by this pull request.
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-5026
Motivation
hipBLASLt still depends on rocRoller to JIT pure F6 (FP6 and BF6) GEMMs on gfx950. These kernels are the first slice of checking that code in as Gemm-From-Anywhere custom kernels, the same way the existing rocRoller FP16 demo is shipped.
Technical Details
Adds 15 gfx950 TN kernels whose A and B types are both in {FP6, BF6}, captured from the kernel hipBLASLt's rocRoller selector returned for the in-tree MX client problems. Assembly lives under
Tensile/CustomKernels/rocroller/, withcustom.configembedded.Tensile/Tests/custom/custom_rr_f6.yamlbenchmarks them on gfx950 only.The custom-kernel census moves from 126 kernels / 66 declared-zero StaggerU to 141 / 81. The new names are on
NO_BUFFER_OOB_KERNELSbecause rocRoller does not use Tensile's BufferOOB store pattern.These kernels are not referenced from production logic, so hipBLASLt solution selection is unchanged. rocRoller remains a build dependency. Workgroup-mapped kernels whose kernargs Tensile cannot launch (
MagicMultiple,WGM, and related) are not in this stack.Test Plan
python -m Tensile.ValidateMetadata --stricton the combined pure-F6 plus mixed-MX treepytestontest_custom_kernel_staggeru_metadata.pyandtest_CustomKernels_buffer_oob.pyon that same treeTest Result
On the combined tree (this branch plus the stacked mixed-MX branch):
ValidateMetadata --strictreported 195 kernels, 0 errors, 0 warnings. The stagger census and buffer-bounds tests passed (822).On this branch alone (
872e9576985):python -m Tensile.ValidateMetadata --strict-> 141 kernels, 0 errors, 0 warningspython -m pytest Tensile/Tests/unit/test_custom_kernel_staggeru_metadata.py Tensile/Tests/unit/test_CustomKernels_buffer_oob.py-> 606 passed in 116.69sSubmission Checklist
Risk level
Low. Additive custom-kernel sources and a gfx950-only Tensile test. No production logic change. Rollback is reverting this PR.
Stacked PR: the mixed MXFP8/BF8, MXFP6/BF6, and MXFP4 kernels are the next PR, based on this branch.