Skip to content

Revert "fix(hipblaslt): correct MI300A kernel selection in the WGMXCC predicate and KD-tree search" - #13479

Merged
arjun-raj-kuppala merged 1 commit into
developfrom
revert-13293-users/nelumala/hipBLASLt/mi300a-wgmxcc-predicate
Oct 12, 2026
Merged

arjun-raj-kuppala merged 1 commit into
developfrom
revert-13293-users/nelumala/hipBLASLt/mi300a-wgmxcc-predicate

Conversation

@TriveniTadapaneni

@TriveniTadapaneni TriveniTadapaneni commented Oct 12, 2026 •

Copy link
Copy Markdown
Contributor

@therock-pr-bot

therock-pr-bot Bot commented Oct 12, 2026 •

Copy link
Copy Markdown

⏳ Waiting on Required Checks

⏳ Required checks have not all concluded yet. This bot run is exiting without failing so those checks can report independently.

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🔎 pre-commit ⏳ Pending ⏳ Still running…
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —

⏳ Required checks are still running or have not started yet.
This PR is not ready for review until they complete.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

@therock-pr-bot

therock-pr-bot Bot commented Oct 12, 2026 •

Copy link
Copy Markdown

⏳ Waiting for required checks to start or finish:

  • ⏳ pre-commit

@TriveniTadapaneni
TriveniTadapaneni marked this pull request as ready for review October 12, 2026 07:06
@TriveniTadapaneni
TriveniTadapaneni requested a review from a team as a code owner October 12, 2026 07:06
@arjun-raj-kuppala
arjun-raj-kuppala self-requested a review October 12, 2026 07:33

@arjun-raj-kuppala arjun-raj-kuppala left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Approved

@arjun-raj-kuppala
arjun-raj-kuppala merged commit 01f6f9c into develop Oct 12, 2026
348 of 375 checks passed
@arjun-raj-kuppala
arjun-raj-kuppala deleted the revert-13293-users/nelumala/hipBLASLt/mi300a-wgmxcc-predicate branch October 12, 2026 07:39
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest 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.
@TriveniTadapaneni TriveniTadapaneni added the bump-pr-revert reverted pr which were created to unblock bump pr in throck label Oct 15, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants