[hipblaslt] Manual revert KRingShift#7443
Merged
Merged
Conversation
Remove the BAddrInterleave (B address interleave) and KRingShift (K ring-shift) codegen features introduced by commit e5bbd57 and its hard dependencies (8afbeff, f4dee28, 24212f9). Removed across 13 files (1052 lines deleted): - Python codegen: KRingShift tail-loop branching, initBInterleaveG(), initKRingShift(), interleaved tile offsets, store SRD/stride scaling - C++ predicates: Free1SizeDivByValueLowbitGT1, KRingShiftTailWrapOnly - Hardware caps: vL1DCacheLineBytes - YAML: disabled BAddrInterleave/KRingShift in equality kernel entries, deleted kringshift.yaml test file Kept: container.hpp macroSlash bug fix (standalone, unrelated). Soft dependency YAML files (18 tuning commits with false/0 defaults) left as-is since the parser tolerates unknown keys. Co-authored-by: Cursor <cursoragent@cursor.com>
…iles Remove the four now-defunct solution parameter keys from all 189 library logic YAML files (gfx950, gfx1200, gfx1250): - BAddrInterleave: false - KRingShift: false - AssertFree1DivByMT1LowbitGT1: 0 - AssertKRingShiftTailWrapOnly: 0 These keys were left at their default/disabled values by tuning commits and are no longer recognized after the feature codegen was removed. Co-authored-by: Cursor <cursoragent@cursor.com>
nakajee
reviewed
May 14, 2026
yenong-amd
force-pushed
the
users/yenong-amd/revert_kringshift
branch
from
8ce64a0 to
7a6f866
Compare
AlexBrownAMD
approved these changes
May 14, 2026
nakajee
approved these changes
May 15, 2026
Contributor
nakajee
left a comment
nakajee
left a comment
There was a problem hiding this comment.
I can still see the change from 7443.diff.
Looks good.
This was referenced May 15, 2026
eidenyoshida
pushed a commit
that referenced
this pull request
May 15, 2026
…hift, AssertKRingShiftTailWrapOnly, BAddrInterleave (#7513) ## Summary Downstream **hipSPARSELt** / **TensileCreateLibrary** builds broke after **[#7443](#7443 (*Manual revert KRingShift*) because **TensileLite** stopped understanding several **Assert\*** / **KRingShift** fields that were **still present** in checked-in **Tensile logic YAML**. This PR **removes those stale keywords** from the affected **hipBLASLt** and **hipSPARSELt** logic files so parsing matches the reverted **Python** behavior. ## Root cause (what broke) - **[#7443](#7443 removed handlers in `projects/hipblaslt/tensilelite/Tensile/Contractions.py` for keys such as **`AssertFree1DivByMT1LowbitGT1`**, **`AssertKRingShiftTailWrapOnly`**, and related **KRingShift** / **BAddrInterleave** plumbing. - Shipped **YAML** under **`projects/hipsparselt/.../Tensile/Logic/...`** (e.g. **gfx950** / **gfx942**) and **`projects/hipblaslt/.../Tensile/Logic/...`** (e.g. **gfx1250** GridBased) still contained those fields. - **TensileCreateLibrary** then hit **`RuntimeError: Unknown assertion key: AssertFree1DivByMT1LowbitGT1`** (and would be vulnerable to the next unknown key if only partially cleaned). So the regression was **YAML ↔ parser skew** after the revert, not the drop Docker image or venv. ## What this PR does - **hipSPARSELt:** strips **`AssertFree1DivByMT1LowbitGT1`**, **`KRingShift`**, **`AssertKRingShiftTailWrapOnly`**, **`BAddrInterleave`** from the **Equality** logic YAMLs under **`aquavanjaram/gfx942`** and **`gfx950`** (the paths exercised by **gfx942;gfx950** builds). - **hipBLASLt:** same cleanup on the **8** **`gfx1250` GridBased** logic YAMLs that still carried those keys. No change to kernel math here—this aligns **on-disk logic** with the **post-#7443** **TensileLite** parser. ## How we know it’s fixed - **Math CI:** **[hipblaslt-drop-build #117](https://math-ci.amd.com/job/image-builder/job/hipblaslt-drop-build/117/)** completed **SUCCESS** with a **`ROCM_LIBS_REF`** that includes this fix, exercising **hipBLASLt** + **hipSPARSELt** with **gfx942;gfx950** (full drop-style lane). - **This PR’s presubmit:** touches `projects/hipsparselt/**` (and **hipBLASLt** logic), so **TheRock** runs the **hipSPARSELt** / **sparselt** workstream for this change—not only **blas**—which directly exercises **TensileCreateLibrary** on the updated YAML. - **PR tag:** includes **`[project:hipsparselt]`** so **hipSPARSELt** is **built and tested** on this PR alongside **hipBLASLt**, not only after merge. ## CI coverage context (why #7443 slipped) - **[#7519](#7519 — documents that **hipBLASLt-only** PRs previously did **not** reliably pull **hipSPARSELt** into the **TheRock** matrix, so **#7443** could merge green while still breaking downstream **hipSPARSELt** / drop builds. - **[#7514](#7514 — follow-up to wire **`projects/hipblaslt/**`** so **`sparselt`** runs when **hipBLASLt** changes, closing that gap long-term (orthogonal to this YAML fix, but related). ## Checklist - [ ] Confirm no remaining **`AssertFree1DivByMT1LowbitGT1`** / stray **KRingShift** keys in the touched trees (`git grep` on the branch). - [ ] After merge, spot-check **TheRock** + optional **hipblaslt-drop-build** on **`develop`**. --------- Co-authored-by: Cursor <cursoragent@cursor.com>
assistant-librarian Bot
pushed a commit
to ROCm/hipBLASLt
that referenced
this pull request
May 15, 2026
[hipblaslt][hipsparselt] Removed AssertFree1DivByMT1LowbitGT1, KRingShift, AssertKRingShiftTailWrapOnly, BAddrInterleave (#7513) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Summary Downstream **hipSPARSELt** / **TensileCreateLibrary** builds broke after **[#7443](ROCm/rocm-libraries#7443 (*Manual revert KRingShift*) because **TensileLite** stopped understanding several **Assert\*** / **KRingShift** fields that were **still present** in checked-in **Tensile logic YAML**. This PR **removes those stale keywords** from the affected **hipBLASLt** and **hipSPARSELt** logic files so parsing matches the reverted **Python** behavior. ## Root cause (what broke) - **[#7443](ROCm/rocm-libraries#7443 removed handlers in `projects/hipblaslt/tensilelite/Tensile/Contractions.py` for keys such as **`AssertFree1DivByMT1LowbitGT1`**, **`AssertKRingShiftTailWrapOnly`**, and related **KRingShift** / **BAddrInterleave** plumbing. - Shipped **YAML** under **`projects/hipsparselt/.../Tensile/Logic/...`** (e.g. **gfx950** / **gfx942**) and **`projects/hipblaslt/.../Tensile/Logic/...`** (e.g. **gfx1250** GridBased) still contained those fields. - **TensileCreateLibrary** then hit **`RuntimeError: Unknown assertion key: AssertFree1DivByMT1LowbitGT1`** (and would be vulnerable to the next unknown key if only partially cleaned). So the regression was **YAML ↔ parser skew** after the revert, not the drop Docker image or venv. ## What this PR does - **hipSPARSELt:** strips **`AssertFree1DivByMT1LowbitGT1`**, **`KRingShift`**, **`AssertKRingShiftTailWrapOnly`**, **`BAddrInterleave`** from the **Equality** logic YAMLs under **`aquavanjaram/gfx942`** and **`gfx950`** (the paths exercised by **gfx942;gfx950** builds). - **hipBLASLt:** same cleanup on the **8** **`gfx1250` GridBased** logic YAMLs that still carried those keys. No change to kernel math here—this aligns **on-disk logic** with the **post-#7443** **TensileLite** parser. ## How we know it’s fixed - **Math CI:** **[hipblaslt-drop-build #117](https://math-ci.amd.com/job/image-builder/job/hipblaslt-drop-build/117/)** completed **SUCCESS** with a **`ROCM_LIBS_REF`** that includes this fix, exercising **hipBLASLt** + **hipSPARSELt** with **gfx942;gfx950** (full drop-style lane). - **This PR’s presubmit:** touches `projects/hipsparselt/**` (and **hipBLASLt** logic), so **TheRock** runs the **hipSPARSELt** / **sparselt** workstream for this change—not only **blas**—which directly exercises **TensileCreateLibrary** on the updated YAML. - **PR tag:** includes **`[project:hipsparselt]`** so **hipSPARSELt** is **built and tested** on this PR alongside **hipBLASLt**, not only after merge. ## CI coverage context (why #7443 slipped) - **[#7519](ROCm/rocm-libraries#7519 — documents that **hipBLASLt-only** PRs previously did **not** reliably pull **hipSPARSELt** into the **TheRock** matrix, so **#7443** could merge green while still breaking downstream **hipSPARSELt** / drop builds. - **[#7514](ROCm/rocm-libraries#7514 — follow-up to wire **`projects/hipblaslt/**`** so **`sparselt`** runs when **hipBLASLt** changes, closing that gap long-term (orthogonal to this YAML fix, but related). ## Checklist - [ ] Confirm no remaining **`AssertFree1DivByMT1LowbitGT1`** / stray **KRingShift** keys in the touched trees (`git grep` on the branch). - [ ] After merge, spot-check **TheRock** + optional **hipblaslt-drop-build** on **`develop`**.
aledudek
pushed a commit
that referenced
this pull request
May 20, 2026
## Motivation Revert KRingShift changes. ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> [----------] Global test environment tear-down [==========] 31685 tests from 12 test suites ran. (1506665 ms total) [ PASSED ] 31685 tests. hipBLASLt version: 100400 hipBLASLt git version: 38d55f3 <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Cursor <cursoragent@cursor.com>
aledudek
pushed a commit
that referenced
this pull request
May 20, 2026
…hift, AssertKRingShiftTailWrapOnly, BAddrInterleave (#7513) ## Summary Downstream **hipSPARSELt** / **TensileCreateLibrary** builds broke after **[#7443](#7443 (*Manual revert KRingShift*) because **TensileLite** stopped understanding several **Assert\*** / **KRingShift** fields that were **still present** in checked-in **Tensile logic YAML**. This PR **removes those stale keywords** from the affected **hipBLASLt** and **hipSPARSELt** logic files so parsing matches the reverted **Python** behavior. ## Root cause (what broke) - **[#7443](#7443 removed handlers in `projects/hipblaslt/tensilelite/Tensile/Contractions.py` for keys such as **`AssertFree1DivByMT1LowbitGT1`**, **`AssertKRingShiftTailWrapOnly`**, and related **KRingShift** / **BAddrInterleave** plumbing. - Shipped **YAML** under **`projects/hipsparselt/.../Tensile/Logic/...`** (e.g. **gfx950** / **gfx942**) and **`projects/hipblaslt/.../Tensile/Logic/...`** (e.g. **gfx1250** GridBased) still contained those fields. - **TensileCreateLibrary** then hit **`RuntimeError: Unknown assertion key: AssertFree1DivByMT1LowbitGT1`** (and would be vulnerable to the next unknown key if only partially cleaned). So the regression was **YAML ↔ parser skew** after the revert, not the drop Docker image or venv. ## What this PR does - **hipSPARSELt:** strips **`AssertFree1DivByMT1LowbitGT1`**, **`KRingShift`**, **`AssertKRingShiftTailWrapOnly`**, **`BAddrInterleave`** from the **Equality** logic YAMLs under **`aquavanjaram/gfx942`** and **`gfx950`** (the paths exercised by **gfx942;gfx950** builds). - **hipBLASLt:** same cleanup on the **8** **`gfx1250` GridBased** logic YAMLs that still carried those keys. No change to kernel math here—this aligns **on-disk logic** with the **post-#7443** **TensileLite** parser. ## How we know it’s fixed - **Math CI:** **[hipblaslt-drop-build #117](https://math-ci.amd.com/job/image-builder/job/hipblaslt-drop-build/117/)** completed **SUCCESS** with a **`ROCM_LIBS_REF`** that includes this fix, exercising **hipBLASLt** + **hipSPARSELt** with **gfx942;gfx950** (full drop-style lane). - **This PR’s presubmit:** touches `projects/hipsparselt/**` (and **hipBLASLt** logic), so **TheRock** runs the **hipSPARSELt** / **sparselt** workstream for this change—not only **blas**—which directly exercises **TensileCreateLibrary** on the updated YAML. - **PR tag:** includes **`[project:hipsparselt]`** so **hipSPARSELt** is **built and tested** on this PR alongside **hipBLASLt**, not only after merge. ## CI coverage context (why #7443 slipped) - **[#7519](#7519 — documents that **hipBLASLt-only** PRs previously did **not** reliably pull **hipSPARSELt** into the **TheRock** matrix, so **#7443** could merge green while still breaking downstream **hipSPARSELt** / drop builds. - **[#7514](#7514 — follow-up to wire **`projects/hipblaslt/**`** so **`sparselt`** runs when **hipBLASLt** changes, closing that gap long-term (orthogonal to this YAML fix, but related). ## Checklist - [ ] Confirm no remaining **`AssertFree1DivByMT1LowbitGT1`** / stray **KRingShift** keys in the touched trees (`git grep` on the branch). - [ ] After merge, spot-check **TheRock** + optional **hipblaslt-drop-build** on **`develop`**. --------- Co-authored-by: Cursor <cursoragent@cursor.com>
bnemanich
added a commit
that referenced
this pull request
May 20, 2026
CI hipBLASLt build on PR #7636 fails with: RuntimeError: Unknown assertion key: AssertFree1DivByMT1LowbitGT1 at `Contractions.py:ProblemPredicate.FromOriginalKeyPair` while `parseLibraryLogicData` walks a gfx950 logic YAML. Root cause: PR #7443 ("manual revert KRingShift") removed the parser handlers for `AssertFree1DivByMT1LowbitGT1` / `AssertKRingShiftTailWrapOnly`, and PR #7513 cleaned the matching keys out of the hipBLASLt / hipSPARSELt library logic YAML files. Both reverts are in our branch's history, but origin/develop has new gfx950 logic YAMLs (e.g. `gfx950/gfx950_128cu/Equality/ gfx950_Cijk_Alik_Bljk_HHS_BH_BiasSH_HAS_SAV_UserArgs.yaml`, added by #7125 "update cucount") that were generated *before* the revert pair and therefore still carry the deprecated keys. When CI merges our branch with origin/develop the merged tree contains a YAML the parser can no longer read. Fix: silently ignore the two deprecated keys in `ProblemPredicate.FromOriginalKeyPair`. Matches the spirit of the revert + cleanup pair without requiring follow-up YAML cleanup in develop, and is forward-compatible against any future stale YAML that slips through the same gap. Genuinely unknown `Assert*` keys still raise loudly. New unit tests in `test_Contractions_deprecated_asserts.py` pin the silent-ignore behavior for the two specific keys, verify that unrecognized `Assert*` keys still raise, and check that the existing known `AssertFree0/1ElementMultiple` + `AssertSummationElementMultiple` predicate path is unaffected. Validation: - new tests: 7 passed - full unit suite: 828 passed, 5 skipped, 2 xfailed (was 821 passed) - end-to-end smoke: drove `FromOriginalKeyPair` on a representative stale-key entry set; no `RuntimeError` raised; recognized keys still produce real predicates. Co-authored-by: Cursor <cursoragent@cursor.com>
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
4 participants
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
Revert KRingShift changes.
Technical Details
Test Plan
[----------] Global test environment tear-down
[==========] 31685 tests from 12 test suites ran. (1506665 ms total)
[ PASSED ] 31685 tests.
hipBLASLt version: 100400
hipBLASLt git version: 38d55f3
Submission Checklist