Repository navigation
feat(stinkytofu): issue a barrier once its dependency threshold is met - #14734
Merged
Merged
Conversation
The threshold already includes the last WMMA a feeding ds_load depends on and that load's drain. Holding a met barrier for another tail of WMMAs skipped it and let a later barrier be selected.
hcman2
changed the title
hcman2
changed the title
✅ 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. ❌ Your project check has failed because the head coverage (51.10%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage. Additional details and impacted files@@ Coverage Diff @@
## develop #14734 +/- ##
===========================================
+ Coverage 77.34% 77.36% +0.01%
===========================================
Files 2893 2893
Lines 486262 486371 +109
Branches 73502 73525 +23
===========================================
+ Hits 376094 376238 +144
+ Misses 85342 85314 -28
+ Partials 24826 24819 -7
*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/preserve_barrier
branch
from
October 30, 2026 08:21
f580d9b to
80c08a9
Compare
hcman2
force-pushed
the
users/hcman2/preserve_barrier
branch
from
October 30, 2026 08:33
80c08a9 to
57bf1b3
Compare
Contributor
Author
|
hipblaslt test passed on B0 |
Contributor
Author
|
tox passed except sk |
hcman2
force-pushed
the
users/hcman2/preserve_barrier
branch
from
October 30, 2026 10:01
e838579 to
042115c
Compare
KKyang
approved these changes
Nov 2, 2026
Contributor
TheRock Submodule Bump ActivityNewest first
|
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
JIRA ID: https://amd-hub.atlassian.net/browse/AIHPBLAS-5417
Summary
When a barrier is in both the after and before threshold maps, averaging those thresholds can no longer place it earlier than the parent-WMMA drain. depThreshold (the last WMMA a feeding ds_load depends on, plus the WAR gate and that load's drain) is a floor. Even when the overlap-after placement is later than that drain, the averaged result stays at or after depThreshold.
A barrier issues as soon as its WMMA threshold is met. WmmaBarrierCoverCycles and the extra WMMA tail (default 4) are removed. The threshold already includes the last dependent WMMA and the load's drain; holding the barrier for another tail skipped it and let a later barrier be selected. --wmma-barrier-cover-cycles and the WmmaBarrierCoverCycles module option are still accepted, but they only warn and are ignored.
When an after threshold is earlier than the before threshold and the claim windows do not overlap, the same-token tensor_load still sits between the two barriers: barrierAfter → tensor_load → barrierBefore. A tensor_load is the next LDS producer, so def-use does not place it after the barrier; that order is now an explicit constraint. Only pairs whose windows actually overlap become overlap candidates and are eligible to merge.
Technical Details
Test Plan
Test Result
Submission Checklist