fix(stinkytofu): fix potential bug when calculating overlap - #12169
Merged
Merged
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. |
Contributor
Author
|
tox passed for non streamk cases on a0 machine except unsupported feature(gl2 prefetch) |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #12169 +/- ##
========================================
Coverage 70.29% 70.29%
========================================
Files 2812 2812
Lines 462735 462769 +34
Branches 68113 68122 +9
========================================
+ Hits 325241 325271 +30
- Misses 113939 113943 +4
Partials 23555 23555
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
KKyang
approved these changes
Sep 16, 2026
Contributor
Author
|
[----------] Global test environment tear-down |
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
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.
JIRA ID: https://amd-hub.atlassian.net/browse/AIHPBLAS-4714
Summary
Fix CDNA5 Layer 2 exclusive after/before overlap detection: claim the full unclamped after demand (wmmaWindowsNeeded + latencyWmmaBudget) instead of the promote-clamped issue window.
Previously, after barriers whose base threshold exceeded issuedCount could miss overlap with earlier before-barriers, so they stayed pinned at the region end.
Clean up Layer 2 structs/helpers without changing scheduling policy beyond the claim-window fix; update unit tests accordingly.
Motivation
Layer 2 decides whether exclusive after/before barrier groups fight for the same WMMA windows. The after claim window used overlapWmmaWindow, which is capped to adjustedAfterThreshold / issuedCount for promote timing. When base after demand was large (issue windows + drain latency) but the region had few WMMAs, the claimed interval shrank and failed to overlap a nearby before group — so no reconcile ran and the after barrier stayed at the end.
Technical Details
Overlap claim (behavior change): after exclusive groups now claim wmmaWindowsNeeded + latencyWmmaBudget (unclamped). Promote threshold still clamps to issuedCount.
Proportional split demand: after still uses issue-only wmmaWindowsNeeded; before uses its claim window.
Cleanup: remove dead overlapWmmaWindow; rename summary fields (claimWindow / splitNeeded / pendingThreshold); split after/before group builders; fold pending thresholds and descendants into the group; write setGroupThreshold once after optional tensor-load spacing.
Tests: rebuild the “separate windows” fixture so lastOverlap > wmmaWindowsNeeded still yields non-overlap under the new claim; fold end-to-end merge checks into that test; drop duplicated/dummy smoke tests.
Test plan
unit_tests --gtest_filter='CDNA5ReadyQueueTest.:DAGSchedulerPassTest.:InFlightQueue*'
DAGSchedulerPassTest.Layer2*
hipblaslt / tensilelite kernel spot-check (e.g. mxf4 barrier placement) as needed
Test Result
Submission Checklist