Repository navigation
fix(tensilelite): AmaxD inter-wave reduction LDS write-after-read race - #15040
Merged
mengzcai merged 4 commits intoNov 6, 2026
Merged
Conversation
Every step of the inter-wave tree reused LDS slots from 0, and its one barrier did not order step i's loads before step i+1's stores, so a partial max could be overwritten before it was read and AmaxD came out low. Store wave w's value to slot w and load slot w + Offset: each slot is written once per reduction, so no extra barrier is needed. JIRA: AIHPBLAS-5448
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
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. |
added 3 commits
November 4, 2026 10:19
…Groups Take the wave counts from makeValidWorkGroups() and the target's WavefrontSize instead of a hardcoded list, emit the reduction for each, and check its start Offset. Only power-of-two counts are tested: the tree halves Offset each iteration.
The inter-wave amax reduction halves the wave count each step, so with e.g. 3 or 6 waves some partial maxima are never reduced. No shipped AmaxD solution uses such a count (1, 2 or 4 waves); reject it at solution derivation instead of producing a wrong amax.
Every group used MIWaveGroup [1, 1], a one-wave workgroup, so the inter-wave amax reduction never ran a step. Add [2, 1] and [2, 2]; the second step, where the LDS slot reuse raced, only exists at 4 waves.
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project check has failed because the head coverage (53.94%) 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 #15040 +/- ##
===========================================
+ Coverage 77.43% 77.64% +0.21%
===========================================
Files 2902 2902
Lines 489842 490184 +342
Branches 74142 74226 +84
===========================================
+ Hits 379301 380580 +1279
+ Misses 85573 84660 -913
+ Partials 24968 24944 -24
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
AndySu12
approved these changes
Nov 5, 2026
jichangjichang
approved these changes
Nov 5, 2026
Contributor
Author
|
The only FFM failure is fixed by "fix(tensilelite): turn off USO in gsuasb_fp32_atomic_gfx1250" (AIHPBLAS-5447). |
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.
JIRA ID : AIHPBLAS-5448
Motivation
matmul_fp8_dst_fp8_amaxD_1250fails intermittently on gfx1250 HW CI: the GPU AmaxD is low by a fixed amount per shape (TT 768x4096x768 bf8 x f8: CPU 211540, GPU 207480).Root cause: in
amax_inter_wave_reduction, each tree step has upper waves store to slotw - Offset, a barrier, then lower waves load slotw. Every step reuses slots from 0, and nothing orders step i's loads before step i+1's stores. With 4 waves, w1 can store max(w1, w3) to slot 0 in step 2 before w0 has read w2's value from it in step 1, so w2's partial max is lost. The code is old and shared by all archs; it surfaced after #13300 made gfx1250 ship AmaxD kernels.Technical Details
wstores to slotw; lower wavewloads slotw + Offset. Each wave stores once per reduction, so no slot is reused and no barrier is added. Slot 0 stays free for the broadcast; LDS reservation is unchanged.OutputAmaxDwith a non-power-of-twoNumWaves: the tree halves the wave count, so e.g. 3 waves would drop w2. Shipped AmaxD solutions use 1, 2 or 4 waves.Test Plan
On gfx1250 (MI450X B0), develop d1dd5a8 vs this branch: hipblaslt-test
*amaxD_1250*, the tox YAML, and the new unit tests. To make the race deterministic, a debug-only build delays wave 0 withs_sleepbefore its inter-wave load.Test Result
*amaxD_1250*x200*amaxD_1250*+ forced delayamax_gfx1250.yaml+ forced delayamax_gfx1250.yamlgfx942 shares the codegen and is covered by the wave64 unit test only.
Submission Checklist