Repository navigation
[hipBLASLt] Fix SIA0+PGR2+PLR0: drain dscnt before barrier sync - #8204
Merged
Merged
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
Codecov Report❌ Patch coverage is
❌ Your project status has failed because the head coverage (77.83%) 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 #8204 +/- ##
========================================
Coverage 61.54% 61.54%
========================================
Files 2095 2095
Lines 361375 361377 +2
Branches 54697 54698 +1
========================================
+ Hits 222389 222396 +7
+ Misses 120140 120135 -5
Partials 18846 18846
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
chicyang
force-pushed
the
users/chicyang/fix_SIA0_PGR2_PLR0
branch
2 times, most recently
from
June 9, 2026 11:16
84b8e1f to
0884567
Compare
Contributor
There was a problem hiding this comment.
Pull request overview
Fixes incorrect results in specific gfx12 TDM GEMM kernels by ensuring cross-wave LDS buffer swap synchronization does not occur while LDS local reads are still in flight (PLR=0 case).
Changes:
- Add an explicit
s_waitcnt dscnt=0(via_wait(..., skipLocalRead=0)) immediately before the cross-wave TDM swap_syncThreadsbarrier whennumItersPLR == 0. - Add a new Tensile YAML benchmark case to cover the problematic configuration (SIA0 + PGR2 + PLR0) on gfx1250.
Reviewed changes
Copilot reviewed 2 out of 2 changed files in this pull request and generated 1 comment.
| File | Description |
|---|---|
| projects/hipblaslt/tensilelite/Tensile/KernelWriter.py | Inserts a local-read drain (dscnt=0) before the cross-wave TDM swap barrier for the PLR=0 path to prevent buffer reuse hazards. |
| projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/tdm_gfx1250.yaml | Adds a targeted benchmark problem group intended to exercise the fixed synchronization scenario on gfx1250. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
jichangjichang
approved these changes
Jun 10, 2026
chicyang
force-pushed
the
users/chicyang/fix_SIA0_PGR2_PLR0
branch
2 times, most recently
from
June 12, 2026 09:58
3e896bb to
00aaddc
Compare
mengzcai
added a commit
that referenced
this pull request
Jun 25, 2026
## Motivation
Fix kernel bugs:
1. Fix **TDM tail** local-read using wrong LDS buffer (bf6_tdm_gfx1250,
tdm_gfx1250,...)
-- key config: **1LDSBuffer=0**
2. Fix LocalReadVALU overwriting correct LDS>=64KB read address
(f64_gfx1250.yaml)
-- key config: **1LDSBuffer=0 + EPS=True** (both SIA0/SIA1)
The bug requires DeviceLDS > 64KB (so that num > 0 becomes possible) on
the double-buffer path.
Only gfx950 and gfx1250 satisfy this. On gfx950 the non-MI test datatype
is bf16, which derives EPS = False → hasMultipleBuffer = False; the only
way bf16 can make a single operand region (A or B) reach ≥ 64KB is
DepthU = 512, but such kernels exceed the 256-VGPR limit and fail to
generate. Hence no buildable num > 0 kernel exists on gfx950.
<img width="810" height="225" alt="image"
src="https://github.com/user-attachments/assets/e4765089-30ff-4262-8389-149fc32ad39c"
/>
3. --> the bug already fixed in the merged PR
#8204
Fix **TDM PGR=2/PLR=0**: force ds waitcnt before LR->TDM barrier to fix
write-after-read :
The "wait current LR finish for next GR(TDM)" barrier had no preceding
dscnt wait, and PGR>=2 made syncThreads skip the force waitcnt0, so the
next trip's tensor_load_to_lds could overwrite the LDS buffer still
being local-read.
<!-- Explain the purpose of this PR and the goals it aims to achieve.
-->
## Technical Details
1. Fix TDM tail local-read using wrong LDS buffer
Before: TDM + double-buffer tail forced the tail LR offsets back to
buffer0, but the TDM write pointer may stay at buffer1.
(after an odd number of main-loop iters)
After: No longer forces buffer 0 in the tail, unless wider local read
needs the offset recomputed `(numReadsIterCoalesced>1)`.
2. Fix the overwrite bug when using function `LocalReadVALU`
ds_read_b64 destVgpr, srcAddr offset:paramList[0]
..........................└─dst─┘ └─src─┘ └────ds────┘
srcAddr: N*64KB
ds = DSModifiers ; immediate offset
**Before (develop):**
```
num = paramList[0] // 65536
paramList[0] = paramList[0] - num * 65536 # paramList[0] is now reduced mod 65536, so it is always < 65536
srcVgpr = vgpr("LocalReadAddr%s+%d"%(tc,num)) # (A) correct high-part split, but stored in srcVgpr -> never used (**dead**)
if numOffsets == 1:
addrIdx = paramList[0] // 65536 # (B) paramList[0] is already < 65536, so addrIdx is ALWAYS 0
srcAddr = vgpr("LocalReadAddr%s+%u"%(tc, addrIdx)) # -> always LocalReadAddr+0 (high 64KB part dropped: BUG)
paramList[0] -= addrIdx * 65536 # -> subtract 0, a no-op
ds = DSModifiers(na=1, offset=paramList[0])
```
**After (PR):**
```
# The ds_read const offset field is only 16-bit. For LDS >= 64KB, split the byte
# offset into a high part (selects the base-address VGPR, each pre-advanced by
# 64KB) and a low part (the 16-bit offset immediate).
num = paramList[0] // maxLDSConstOffset # high part: which 64KB base register
paramList[0] = paramList[0] - num * maxLDSConstOffset # low part: fits the 16-bit offset field
srcAddr = vgpr("LocalReadAddr%s+%d"%(tc, num)) # base VGPR already advanced by num * 64KB
if numOffsets == 1: //-->f64 pass from here
# e.g. ds_read_b64
ds = DSModifiers(na=1, offset=paramList[0])
if numOffsets == 2:
# combine two reads
# e.g. ds_read2_b32 : two chunks share one base with two offsets.;one base &offset0 & offset1;
ds = DSModifiers(na=2, offset0=paramList[0], offset1=paramList[1])
srcAddr = vgpr("LocalReadAddr%s"%tc)
```
3. Modified GroupedGemm logic yamls from TDM+PGR2+PLR0 to PLR1
<!-- Explain the changes along with any relevant GitHub links. -->
## Test Plan
Tox test for all tests.
<!-- Explain any relevant testing done to verify this PR. -->
## Test Result
tox passed on ffm&B0
<!-- Briefly summarize test outcomes. -->
## Submission Checklist
- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
---------
Co-authored-by: mengzcai <mengzcai@amd.com>
Co-authored-by: test <test@test>
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
SIA0+PGR2+PLR0 TDM kernels gave wrong results because the cross-wave TDM buffer swap barrier didn't wait for in-flight local reads to drain
Technical Details
Insert a dscnt=0 wait before the TDM swap _syncThreads barrier when numItersPLR==0, ensuring all local reads finish before the buffer is reused.
Test Plan
Add a case
Test Result
Pass
tox test (those failed case are known issues: sk no solution)
Submission Checklist