Repository navigation
fix(hipblaslt): build without rocRoller when HIPBLASLT_ENABLE_ROCROLLER is off - #14086
Conversation
✅ 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. |
Codecov Report❌ Patch coverage is Additional details and impacted files@@ Coverage Diff @@
## develop #14086 +/- ##
========================================
Coverage 77.66% 77.66%
========================================
Files 2903 2904 +1
Lines 490366 490404 +38
Branches 74218 74237 +19
========================================
+ Hits 380819 380870 +51
+ Misses 84597 84572 -25
- Partials 24950 24962 +12
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
jaopaulolc
changed the title
AlexBrownAMD
left a comment
AlexBrownAMD
left a comment
There was a problem hiding this comment.
Tested on a local repo with no rocroller checkout and it works. Thanks!
28ef88b to
e290401
Compare
1d0cb23 to
e6175fd
Compare
peymanr
left a comment
peymanr
left a comment
There was a problem hiding this comment.
Thanks, I reran the PR Pundit agent and it did not find anything new.
PR Pundit review(testing PR Pundit Code Review agent. Not all comments may be relevant. Please reach out to @peymanr for suggestions and feedback.)
1. 🟠 Should fix: Ninja deps fallback passes object paths that never match targets
The Ninja fallback in Who is affected: Developers who run the check by hand on a Ninja-generated build tree from anywhere other than the build root; the Makefile-based CI workflow never reaches this path. Verified by running it. 2. 🔵 Nit: build.ninja search loop can hang on relative depdir
The walk-up loop in Who is affected: Anyone running the check script with a relative directory that has no Verified by running it. 3. 🔵 Nit: CI guard never exercises the rocRoller-off build it claims to protect
The new Who is affected: Only maintainers who manually dispatch the drop build against a ref older than this PR, plus the general lack of CI coverage for 4. 🔵 Nit: Unsupported-format lookup throws through C API instead of empty list
The new comments in Who is affected: Callers whose problem takes the rocRoller path ( Verified by running it. 5. 🔵 Nit: Diagnostics mention non-called function and print bare integer types
The rocRoller rejection messages in Who is affected: Users reading stderr or assert diagnostics when hipBLASLt rejects a rocRoller problem for a non-F32 scale or accumulation type (log readability only). 6. 🔵 Nit: Integer-comparison comment contradicted by direct enum use elsewhere
Comments say Who is affected: Maintainers reading Verified by running it. 7. 🔵 Nit: Changed files not clang-formatted, adding unrelated diff noise
The PR lengthens three already over-limit lines in Who is affected: Code readers and anyone later running clang-format on these files; there is no functional impact. Verified by running it. 8. 🔵 Nit: No test covers new type-translation and cache-key mapping layer
There are no unit tests for the new mapping helpers Who is affected: Only future edits to the scale-format helpers or the custom-kernel table could cause a mismatch, which would make the preloaded mxfp4 kernels miss the cache; no current configuration is affected. Related past reviews: #372 on 9. 🔵 Nit: SingleScale mode silently removed and genScaleModeString misnamed
Who is affected: Readers and maintainers of the rocRoller gemm code; there is no runtime effect. Verified by running it. 10. 🔵 Nit: Long historical comments narrate past tree state and go stale
Some comments added by this PR describe past history or are wrong. In Who is affected: Maintainers reading the rocRoller ranker sources; there is no runtime impact. Verified by running it. Related past reviews: #7717 on 11. 🔵 Nit: Unused rocroller_interop.hpp include pulls in rocRoller headers
The PR adds Who is affected: Maintainers only, in builds with rocRoller enabled; there is no runtime or build-time effect. Verified by running it. Shared from a pilot of PR Pundit, an agentic code reviewer that builds and tests the change and reads the project's past reviews. Reviewed at |
e3c0774 to
8b8a122
Compare
|
Addressed the review of
|
e15c463 to
b056b96
Compare
…eaders The tile ranker in hipblaslt-rocroller-selection is built in every configuration, including HIPBLASLT_ENABLE_ROCROLLER=OFF, and reached shared/rocroller/lib/include for two enumerations: rocRoller::DataType and rocRoller::Operations::ScaleMode. It calls no rocRoller function and links no rocRoller symbol, but those two includes pulled in 80 headers and a FetchContent clone of libfmt, and made the target unbuildable in a checkout that omits shared/rocroller. Declare the twelve DataType and two ScaleMode enumerators the ranker names in rocroller_enums.hpp, generated from the single list in rocroller_enum_list.hpp. kernel_type.hpp and analytical_utils.hpp pick rocRoller's headers only when HIPBLASLT_HAVE_ROCROLLER_HEADERS is defined, which is true of hipblaslt-rocroller alone. rocroller_enum_conformance.cpp expands the same list into static assertions against rocRoller's headers, so the two declarations of KernelType that meet in a rocRoller-enabled build cannot disagree. Drop the rocRoller include directory, the libfmt FetchContent and link, and the now-unused ROCROLLER_USE_HIP definition from the ranker target.
… plainly Review feedback on the previous commit. No behaviour change: the ranker's instruction stream for chooseSolutionIndexParameters, genKernelType, hipDataType_to_rocRoller_type and rankRocRollerSelectedTiles is unchanged. Drop the HIPBLASLT_HAVE_ROCROLLER_HEADERS definition and read the existing HIPBLASLT_USE_ROCROLLER instead. It already discriminates the two targets: it is private to hipblaslt-rocroller, whose sources are the only ones built against rocRoller's headers, and it is absent from the ranker target. kernel_type.hpp and analytical_utils.hpp are internal headers that nothing outside library/src/amd_detail/rocblaslt/src includes, so the definition being public on hipblaslt does not reach a translation unit that would then look for rocRoller headers it has no path to. Using the hipBLASLt declarations in every configuration is not possible: solution_cache.hpp includes both gemm.hpp, which includes rocRoller's headers, and kernel_type.hpp, so one translation unit would define rocRoller::DataType twice. Replace the X-macro list and its two expansions with plain declarations in rocroller_enums.hpp and explicit static assertions in rocroller_enum_conformance.cpp, one per enumerator. The list no longer keeps the two in step by construction, so each file carries a comment pointing at the other. All 14 values and both underlying types were read back out of rocRoller's headers at run time before being written down.
rocroller_enums.hpp reuses rocRoller's namespace and type names, so no translation unit includes both declarations.
HIPBLASLT_USE_ROCROLLER already selects the mirror header. The Windows uint shims no longer include anything that needs them.
… headers The compiler dependency files for the ranker sources are the check.
The ranker copied rocRoller's DataType and ScaleMode so a build with HIPBLASLT_ENABLE_ROCROLLER=OFF could compile without rocRoller. Use hipDataType and RocblasltContractionProblem::ScalingFormat instead, and translate to rocRoller only in rocroller_interop.hpp, which the ranker does not include.
The comments named the conversion they replaced. They now say what normalizeOperandType and hip2RocRollerType do, and why the comparison is an integer compare.
hip2RocRollerType throws on a format it cannot name, and the public catch turns that into an internal error. The reject path now prints the hipDataType and returns invalid_value. The ranker dependency check only flags rocRoller and fmt headers, and reads ninja -t deps when a Ninja build left no Makefile .d files.
The ninja dependency fallback asked for objects by a path ninja does not use as a target, and the ranker comments described a throw as an empty tile list.
b056b96 to
7696adb
Compare
Auto-merge was disabled for another pull request in the stack
…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.
TheRock Submodule Bump ActivityNewest first
|
…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.
JIRA ID
Motivation
With
HIPBLASLT_ENABLE_ROCROLLERoff, a checkout may omit the rocRoller sources. The tile ranker is compiled in both builds. It included rocRoller's headers only to name a number format and a scale layout, so a checkout without those sources failed to compile, and the build also fetched libfmt, the formatting library those headers use.Technical Details
A kernel's operands are now described with the names hipBLASLt already uses. An element format is a
hipDataType. A scale layout isRocblasltContractionProblem::ScalingFormat.ScaleType::formatstarts asScalingFormat::None, which means the operand is not scaled.KernelTypeandScaleTypeare defined inkernel_type.hpp. That header includes hipBLASLt's own type headers and does not include rocRoller.hip2RocRollerType, inrocroller_interop.hpp, converts ahipDataTypeinto rocRoller's format name. Only three files include it:gemm.cpp,rocroller_host.cpp, andruntime_args_selection.cpp. Those files are built intohipblaslt-rocroller, the library created only when rocRoller generation is on. A scale layout is not converted. The generator callsisBlockScaling()and reads the format's own block sizes.The ranker library
hipblaslt-rocroller-selectionno longer adds rocRoller's include directory and no longer links libfmt. Its sources arerocroller_selected.cppandsolution_selection.cpp.normalizeOperandTypemaps both 8-bit E4M3 spellings ontoHIP_R_8F_E4M3, and both E5M2 spellings ontoHIP_R_8F_E5M2. A format no ranked kernel handles becomesHIPBLASLT_DATATYPE_INVALID, which matches no tile.Each tile is still emitted in both non-temporal forms.
solution_selection.cppsetsnon_temporal_a_availableandnon_temporal_b_availableon all three tile configs, becauserank_configsreads those flags.The package-build workflow runs
check_rocroller_selection_deps.shon the ranker's compiler dependency files. The script fails when a file namesrocRoller/orfmt/. The in-tree directoryrocroller/does not matchrocRoller/.This pull request is based on
develop.Test Plan
Host compilation only. Both builds used
-DGPU_TARGETS=gfx950and a Release configuration, on commit3685ce7d7c4, beforedevelopwas merged into this branch.HIPBLASLT_ENABLE_ROCROLLER=OFFin a checkout that does not containshared/rocroller. Linklibhipblaslt.soandlibhipblaslt-rocroller-selection. Runninja -t depsonsolution_selection.cpp.oandrocroller_selected.cpp.o, then runcheck_rocroller_selection_deps.shon the ranker's dependency directory.-DCMAKE_POLICY_VERSION_MINIMUM=3.5, because rocRoller's bundled yaml-cpp declares a CMake minimum that CMake 4 rejects. Build the targethipblaslt-rocroller.Test Result
Completed on
3685ce7d7c4, before thedevelopmerge. Not repeated on the current head. The only commit after749e9f2f1bdrewords comments.libhipblaslt.soandlibhipblaslt-rocroller-selection.ninja -t depsfor the two ranker objects named norocRoller/header and nofmt/header.check_rocroller_selection_deps.shfailed because ninja wrote no.dfiles.ninja hipblaslt-rocrollerexited 0.parameter_selection.cpp,gemm.cpp, androcroller_host.cppare inlibhipblaslt-rocroller.a, andlibrocroller.solinked.Automated GitHub checks on the current head are pending. No kernel-selection comparison was run on this branch.
Submission Checklist
Risk level
Medium (3/5). The ranker compares formats using
hipDataTypevalues, which differ from rocRoller's numbering, so the tile it keeps can change.hip2RocRollerTypestops the run when a format has no translation, rather than generating a different format. The public functions hipBLASLt exposes to callers are unchanged.