Commit 8b911e1
authored
fix(stinkytofu): fix potential bug when calculating overlap (#12169)
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
<!-- Briefly summarize test outcomes. -->
## Submission Checklist
- [ ] Look over the contributing guidelines at
https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.1 parent bc358d8 commit 8b911e1
2 files changed
Lines changed: 164 additions & 186 deletions
0 commit comments