feat(stinkytofu): Add RemoveDscntPass for WMMA Loops - #10389
Open
hcman2 wants to merge 7 commits into
Open
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
shared: stinkytofu
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 #10389 +/- ##
===========================================
- Coverage 69.69% 68.52% -1.18%
===========================================
Files 2754 2754
Lines 453338 460598 +7260
Branches 66803 73250 +6447
===========================================
- Hits 315952 315585 -367
- Misses 117083 123179 +6096
- Partials 20303 21834 +1531
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
hcman2
force-pushed
the
users/hcman2/remove_dscnt_pass
branch
3 times, most recently
from
August 7, 2026 01:57
e4c9300 to
98024c4
Compare
Contributor
Author
|
For the default conservative setting, 6 dscnt are removed for compute bound F8 256x256 MT |
hcman2
force-pushed
the
users/hcman2/remove_dscnt_pass
branch
from
August 7, 2026 02:10
98024c4 to
d4e288b
Compare
Contributor
Author
|
tox passed wo sk cases on gfx1250 |
hcman2
changed the title
KKyang
changed the title
KKyang
approved these changes
Aug 7, 2026
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
2 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
JIRA ID: https://amd-hub.atlassian.net/browse/AIHPBLAS-4230
RemoveDscntPass is motivated by the gap between conservative s_wait_dscnt insertion and the actual dataflow/timing behavior of LDS reads in WMMA-heavy kernels.
In practice, upstream wait insertion tends to be safety-first: it often keeps dscnt values higher than necessary, and may emit multiple waits that become redundant once instruction scheduling, overlap, and issue timing are considered together. This is especially common in long loop bodies with dense ds_load streams and interleaved v_wmma instructions.
The goal of this pass is to tighten or eliminate unnecessary LDS wait pressure while preserving correctness:
Track in-flight LDS loads (ds_load) over a continuous scan.
Model when waits are still required versus when they are over-conservative.
Reduce dscnt counts to the minimum safe value when possible.
Remove waits that are provably redundant and leave a traceable comment marker.
By doing this, the pass improves instruction-level overlap (less over-waiting), reduces synchronization overhead in hot loops, and keeps the generated assembly both safer-than-manual tuning and more efficient than purely conservative wait insertion.
Test Plan
Check some performances
Stinkytofu ctest
Test Result
tox passed wo sk cases on b0
stinkytofu ci passed
Submission Checklist