Repository navigation
fix(hipblaslt): correct MI300A kernel selection in the WGMXCC predicate and KD-tree search - #13293
Merged
NaveenElumalaiAMD merged 5 commits intoOct 9, 2026
Conversation
…MappingXCCCheck On MI300A (gfx942, 228 CUs) almost every solution in the generic gfx942 library logic is rejected before ranking, so the heuristic chooses from a tiny pool and collapses onto a handful of macro-tiles. WorkgroupMappingXCCCheck keys on the pair (WorkGroupMappingXCC, WorkGroupMappingXCCGroup). A group of -1 resolves to the device CU count, and 228 = 2^2 * 3 * 19 is not divisible by 8, so every solution carrying XCC=8 with XCCGroup=-1 fails the divisibility test. That is 150,576 of the 195,438 solutions in the generic gfx942 logic (77.1% overall, 99.7% of the GridBased tables). The rejection is spurious. Those solutions are never launched at their nominal XCC: isFallbackForHW() marks any bare-Processor (generic) solution as a CU fallback on a part where isStandardCU() is false, and the launch sites in tensile_host.cpp then call setWGMXCC(isCUFallback ? 1 : 0), demoting XCC to 1 before dispatch. The predicate runs during findTopSolutions(), before setFallbackStatus() has been called, so fallbackStatus() is still false and the solution is filtered at an XCC it will never run with. It is an ordering gap, not a hardware constraint. Force XCC to 1 in the predicate when the device is MI300A, matching what the launch path already does. Every kernel selected after this change still reports WGMXCC1, so no illegal XCC reaches a launch. Scope is deliberately narrowed to MI300A rather than to !isStandardCU(). The CU-partitioned modes (20/38/64/80/152 CU) have the identical gap and are intentionally left alone here. No change is needed in ValidWorkGroupMappingXCC.py or SolutionValidation.hpp: both derive the CU count from a *_Ncu directory name and skip the generic gfx942 directory entirely, so the build-time rule and this runtime rule do not overlap.
The grid-based KD-tree index in PropertyMatching is two-dimensional, keyed
on (M, N) only. A cell therefore carries no batch dimension and may hold no
row at all for the problem's batch count: querying b == 1 against a cell
whose rows are all batched makes lower_bound return begin(), takes the
"b is 1" branch, and leaves bStart == bEnd == begin().
The code then ran on that empty range:
auto lower = std::lower_bound(bStart, bEnd, k, compK);
if(lower == bEnd)
lower--;
lower == bEnd == begin(), so lower-- steps before the first element and the
following comparison dereferences it -- an out-of-bounds read whose result
then decides which solution is returned. The Debug branch above it had the
same shape, dereferencing bEnd unconditionally even when bEnd == end().
Fixes:
- skip the cell when bStart == bEnd. Its solutions are all for a different
batch size, so there is nothing to return from it; the remaining cells,
or the binary search below, resolve the problem instead.
- print "<end>" rather than dereferencing bEnd when bEnd == end().
- rewrite "auto prev = lower == bStart ? lower : lower--" as an explicit
copy plus a guarded decrement. This is behaviour-preserving -- the
post-decrement already yielded the old iterator -- but the guard now
reads as a guard rather than as a side effect buried in a ternary.
The predicate change in the preceding commit widens the MI300A candidate
pool, which makes more of these cells reachable in practice, and the
selection runs used to validate it enabled the KD index explicitly via
TENSILE_GRIDBASED_KDTREE=1. Both changes were exercised together.
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. |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #13293 +/- ##
===========================================
+ Coverage 71.94% 71.98% +0.04%
===========================================
Files 2866 2872 +6
Lines 473058 474664 +1606
Branches 69884 70226 +342
===========================================
+ Hits 340319 341683 +1364
- Misses 109141 109288 +147
- Partials 23598 23693 +95
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
NaveenElumalaiAMD
deleted the
users/nelumala/hipBLASLt/mi300a-wgmxcc-predicate
branch
October 9, 2026 15:42
arjun-raj-kuppala
pushed a commit
that referenced
this pull request
Oct 12, 2026
… predicate and KD-tree search" (#13479) Reverts #13293 Fix issue :#13480 test run passed: https://github.com/ROCm/TheRock/actions/runs/38203742146/job/114696224272
Contributor
TheRock Submodule Bump ActivityNewest first
|
NaveenElumalaiAMD
added a commit
that referenced
this pull request
Oct 14, 2026
…te and KD-tree search- #13293 (#13586) ## Motivation MI300A is locked out of nearly all the generic gfx942 (MI300X) kernels: **150,639 of 195,438** solutions (77.1%; 99.7% of GridBased) are rejected before ranking at 228 CUs, so SGEMM NT collapses onto a handful of tiles and fp8/bf8 fall back to a single tile per transpose combo. This PR lets MI300A use the full MI300X kernel set, and re-lands #13293 (reverted by #13479) with **only the code changes** — both `clients/tests/data` YAML commits are dropped. ## Technical Details Generic gfx942 solutions carry `WorkGroupMappingXCCGroup: -1`, so `WorkgroupMappingXCCCheck` substitutes `cuCount` (228) and rejects them because `228 % XCC != 0`. They are CU-fallbacks that already launch at `WGMXCC=1` (`setWGMXCC(isCUFallback ? 1 : 0)`), but the predicate runs before `setFallbackStatus()` marks them — `isMI300A()` closes that gap by evaluating the check at `XCC=1`. Net **2 files, +59/−6**, provably inert off MI300A (`fallbackXCC()` reduces to the original `fallbackStatus()` when `isMI300A()` is false); also guards the KD-tree batch search against empty cells (`PropertyMatching.hpp`). ## Test Plan Built against `develop` @ `b5c1c01faf3` on MI300A (gfx942, 228 CU) and ran `hipblaslt-test` plus CPU-verified `hipblaslt-bench` SGEMM NT at 8192³, 4096³ and 2048×8192×8192. The predicate change was previously swept over 3189 SGEMM NT shapes against #12609 with the KD-tree held equal between arms. ## Test Result Admitted SGEMM NT solutions on MI300A go from **97 of 8192 (1.2%) to all 8192**, and the distinct kernels actually selected across the 3189-shape sweep rise **44 → 756** — no longer collapsing onto `MT128x128x32` for a third of shapes. Selection matches #12609 on **3188/3189** shapes, delivering that PR's 353 MB of 228-CU tuning logic as code rather than data; `hipblaslt-test` **57/57 passed** and all three shapes verify against CPU (8192³ at 88.2 TFLOPS, `norm_error` 1.36527e-06). The 97→8192 and 44→756 figures come from the earlier canonical launch-site dump and 3189-shape sweep, not from this session's build — my attempted differential re-check here didn't discriminate, as I noted. The correctness numbers (57/57, 88.2 TFLOPS, norm_error) are from this session.
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.
Motivation
On MI300A (gfx942, 228 CUs)
WorkgroupMappingXCCCheckrejects 150,639 of 195,438 generic gfx942 solutions (77.1%; 99.7% ofGridBased) because228 % 8 != 0, starving the heuristic across every data type — fp8/bf8 collapse to a single macro-tile per transpose combination. A second, independent bug lets the grid-based KD-tree search read out of bounds when a cell holds no row for the problem's batch count.Technical Details
The rejection is spurious: generic gfx942 solutions are CU-fallbacks on MI300A and are already launched at
WGMXCC=1bysetWGMXCC(isCUFallback ? 1 : 0), but the predicate runs duringfindTopSolutions()beforesetFallbackStatus(), so it filters them at an XCC they never run with — fixed by forcingXCC = 1whenisMI300A(). Separately,b == 1against an all-batched KD cell leavesbStart == bEnd ==begin(), after whichif(lower == bEnd) lower--;steps beforebegin()and is dereferenced; that cell is now skipped (same fix as #12609, byte-identical).Test Plan
Compared per-shape kernel selection (
TENSILE_DB=0x8000) on 4x MI300A against #12609's tree over all 3,189 shapes of the f32 NT sweep, re-running every disagreement withTENSILE_GRIDBASED_KDTREE=1to isolate the KD index from the predicate. Added a 120-dispatch probe over f32 / fp16 / bf16 / fp8 / bf8 x 3 transpose combinations x 8 shapes to confirm the fix is not fp32-only.Test Result
3,188 of 3,189 shapes select the identical kernel to #12609 once the KD-tree is held equal (2,743 as-run; the 446 raw diffs were all KD), every selected name ends
WGMXCC1, and the reachable pool grows from 44 solutions / 25 tiles to 792 / 437.Submission Checklist