Skip to content

fix(hipblaslt): omit reserved wave32 directive on gfx125x - #13365

Merged
newling merged 1 commit into
developfrom
users/newling/AIHPBLAS-5123-wave32-descriptor
Oct 10, 2026
Merged

newling merged 1 commit into
developfrom
users/newling/AIHPBLAS-5123-wave32-descriptor

Conversation

@newling

@newling newling commented Oct 8, 2026 •

Copy link
Copy Markdown
Contributor

JIRA ID : AIHPBLAS-5123
https://amd-hub.atlassian.net/browse/AIHPBLAS-5123

My understanding is that LLVM commit b452df1 makes hipBLASLt's gfx125x kernels fail to assemble. It adds a check that correctly rejects a directive we're emitting. This is holding up a downstream ROCm compiler update: picking up the LLVM change would break the hipBLASLt build in rocm-libraries testing. This PR aims to get those builds passing so that update can go ahead.

gfx125x always execute threads in groups of 32 (wave32). The fix in this PR removes the directive that explicitly requests that group size, and makes sure the kernel generators and assembly readers still treat these kernels as wave32 when the directive is missing.

Focused Python tests pass, StinkyTofu assembly tests also pass. Local checks confirm that the assembly still preserves wave32.

An agent has also tested with the updated compiler before and after this change, and reproduced a failure before and success after.

@therock-pr-bot

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

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🔎 pre-commit ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —
🤖 therock-pr-bot ✅ Pass —

🎉 All checks passed! This PR is ready for review.

📖 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 8, 2026 •

Copy link
Copy Markdown

🎉 All checks passed! This PR is ready for review.

@newling
newling marked this pull request as ready for review October 8, 2026 22:41
@newling
newling requested review from a team as code owners October 8, 2026 22:41
@newling
newling enabled auto-merge (squash) October 8, 2026 23:46
@codecov

codecov Bot commented Oct 9, 2026 •

Copy link
Copy Markdown

Codecov Report

✅ All modified and coverable lines are covered by tests.

Additional details and impacted files
@@           Coverage Diff            @@
##           develop   #13365   +/-   ##
========================================
  Coverage    71.89%   71.89%           
========================================
  Files         2876     2876           
  Lines       475549   475549           
  Branches     70268    70268           
========================================
+ Hits        341879   341887    +8     
+ Misses      109892   109888    -4     
+ Partials     23778    23774    -4     
Flag Coverage Δ *Carryforward flag
TensileLite-CPP 49.67% <ø> (ø)
TensileLite-Unit 83.44% <ø> (ø)
hipBLAS 90.62% <ø> (ø) Carriedforward from 23eb21a
hipBLASLt 42.75% <ø> (+0.03%) ⬆️
hipCUB 82.68% <ø> (ø) Carriedforward from 23eb21a
hipDNN 86.80% <ø> (ø) Carriedforward from 23eb21a
hipFFT 42.58% <ø> (ø) Carriedforward from 23eb21a
hipRAND 76.12% <ø> (ø) Carriedforward from 23eb21a
hipSOLVER 68.96% <ø> (ø) Carriedforward from 23eb21a
hipSPARSE 86.99% <ø> (ø) Carriedforward from 23eb21a
rocBLAS 48.17% <ø> (ø) Carriedforward from 23eb21a
rocFFT 46.51% <ø> (ø) Carriedforward from 23eb21a
rocRAND 57.60% <ø> (ø) Carriedforward from 23eb21a
rocSOLVER 76.83% <ø> (ø) Carriedforward from 23eb21a
rocSPARSE 74.07% <ø> (ø) Carriedforward from 23eb21a
rocThrust 93.15% <ø> (ø) Carriedforward from 23eb21a

*This pull request uses carry forward flags. Click here to find out more.
see 1 file with indirect coverage changes

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
  • 📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

@newling
newling force-pushed the users/newling/AIHPBLAS-5123-wave32-descriptor branch from e11acf2 to 1fe40e9 Compare October 9, 2026 17:16
@newling
newling force-pushed the users/newling/AIHPBLAS-5123-wave32-descriptor branch from 1fe40e9 to caccb28 Compare October 10, 2026 00:10
@newling
newling merged commit 01eccd6 into develop Oct 10, 2026
148 of 153 checks passed
@newling
newling deleted the users/newling/AIHPBLAS-5123-wave32-descriptor branch October 10, 2026 03:41
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

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