Fix 2 kernel bugs by mengzcai · Pull Request #8406 · ROCm/rocm-libraries

mengzcai · 2026-06-14T16:16:22Z

Motivation

Fix kernel bugs:

Fix TDM tail local-read using wrong LDS buffer (bf6_tdm_gfx1250, tdm_gfx1250,...)
-- key config: 1LDSBuffer=0
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.

--> the bug already fixed in the merged PR [hipBLASLt] Fix SIA0+PGR2+PLR0: drain dscnt before barrier sync #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.

Technical Details

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).
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)

Modified GroupedGemm logic yamls from TDM+PGR2+PLR0 to PLR1

Test Plan

Tox test for all tests.

Test Result

tox passed on ffm&B0

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

codecov-commenter · 2026-06-14T16:54:10Z

Codecov Report

✅ All modified and coverable lines are covered by tests.

❌ Your project status has failed because the head coverage (77.89%) 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    #8406   +/-   ##
========================================
  Coverage    71.43%   71.43%           
========================================
  Files         2612     2612           
  Lines       407002   407002           
  Branches     60775    60775           
========================================
+ Hits        290731   290733    +2     
+ Misses       94967    94966    -1     
+ Partials     21304    21303    -1

Flag	Coverage Δ		*Carryforward flag
TensileLite	`76.82% <100.00%> (+0.01%)`	⬆️
hipBLAS	`90.81% <ø> (ø)`		Carriedforward from 0531052
hipBLASLt	`41.33% <ø> (-0.02%)`	⬇️
hipCUB	`82.68% <ø> (ø)`		Carriedforward from 0531052
hipDNN	`86.71% <ø> (ø)`		Carriedforward from 0531052
hipFFT	`50.79% <ø> (ø)`		Carriedforward from 0531052
hipRAND	`76.12% <ø> (ø)`		Carriedforward from 0531052
hipSOLVER	`69.18% <ø> (ø)`		Carriedforward from 0531052
hipSPARSE	`86.55% <ø> (ø)`		Carriedforward from 0531052
rocBLAS	`48.08% <ø> (ø)`		Carriedforward from 0531052
rocFFT	`47.22% <ø> (ø)`		Carriedforward from 0531052
rocRAND	`57.07% <ø> (ø)`		Carriedforward from 0531052
rocSOLVER	`77.89% <ø> (ø)`		Carriedforward from 0531052
rocSPARSE	`72.37% <ø> (ø)`		Carriedforward from 0531052
rocThrust	`91.34% <ø> (ø)`		Carriedforward from 0531052

*This pull request uses carry forward flags. Click here to find out more.

Files with missing lines	Coverage Δ
...blaslt/tensilelite/Tensile/Components/LocalRead.py	`54.72% <100.00%> (-0.12%)`	⬇️
...ects/hipblaslt/tensilelite/Tensile/KernelWriter.py	`70.74% <100.00%> (+0.01%)`	⬆️
...blaslt/tensilelite/Tensile/KernelWriterAssembly.py	`69.55% <100.00%> (+0.04%)`	⬆️

... and 2 files with indirect coverage changes

🚀 New features to boost your workflow:

❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

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). Repro/affected tests: TDM kernels with a tail e.g. bf6_tdm_gfx1250, tdm_gfx1250, ...

1. Set LDS>=64KB and 1LDSB and SIA0 to reproduce the issue 2. Fix test yaml: f64_gfx1250 (using LocalReadVALU)

Copilot

Pull request overview

This PR addresses three gfx1250 kernel correctness issues around TDM tail local-reads, LDS addressing for LocalReadVALU when LDS offsets exceed 64KB, and a known-bad TDM prefetch configuration (PGR=2/PLR=0) by rejecting it.

Changes:

Adjust tail-loop local-read offset reset logic to avoid forcing buffer-0 for TDM tails unless required (e.g., wider local-read / StreamK).
Fix LocalReadVALU LDS source-address selection so reads past 64KB use the correct LocalReadAddr+N VGPR segment.
Update test/logic YAMLs to exercise the bug scenario and to move affected solutions off PLR=0 (and reject the unsupported combo in solution derivation).

Reviewed changes

Copilot reviewed 12 out of 12 changed files in this pull request and generated 2 comments.

Show a summary per file

File	Description
projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/tdm_gfx1250.yaml	Adds an “odd DU + tail” exact size to exercise the TDM tail behavior.
projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f64_gfx1250.yaml	Forces `1LDSBuffer=0` (and SIA0) to reproduce/cover the LocalReadVALU >=64KB addressing case.
projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py	Rejects the known-bad TDMInst=3 + PGR=2 + PLR=0 configuration.
projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py	Extends tail TDM descriptor reset conditions to align with tail LR offset reset policy (incl. StreamK).
projects/hipblaslt/tensilelite/Tensile/KernelWriter.py	Refines tail local-read reset gating for TDM tails and preserves wider-LR state across recalculation.
projects/hipblaslt/tensilelite/Tensile/Components/LocalRead.py	Fixes `LocalReadVALU` to use the correct LDS base VGPR when offsets exceed 64KB.
projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx1250/GridBased/gfx1250_Cijk_Alik_Bjlk_HSS_BH_Bias_GG_HA_S_SAV_UserArgs.yaml	Moves solution from PLR=0 to PLR=1 and updates solution/kernel name strings accordingly.
projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx1250/GridBased/gfx1250_Cijk_Alik_Bjlk_HHS_BH_Bias_GG_HA_S_SAV_UserArgs.yaml	Same: PLR=0 → PLR=1 update.
projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx1250/GridBased/gfx1250_Cijk_Ailk_Bljk_HSS_BH_Bias_GG_HA_S_SAV_UserArgs.yaml	Same: PLR=0 → PLR=1 update.
projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx1250/GridBased/gfx1250_Cijk_Ailk_Bljk_HHS_BH_Bias_GG_HA_S_SAV_UserArgs.yaml	Same: PLR=0 → PLR=1 update.
projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx1250/GridBased/gfx1250_Cijk_Ailk_Bjlk_HSS_BH_Bias_GG_HA_S_SAV_UserArgs.yaml	Same: PLR=0 → PLR=1 update.
projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx1250/GridBased/gfx1250_Cijk_Ailk_Bjlk_HHS_BH_Bias_GG_HA_S_SAV_UserArgs.yaml	Same: PLR=0 → PLR=1 update.

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

## Summary This change fixes tensor-DMA bounds contracts exposed while reducing the stock gfx1250 TensileLite TDM tail failure tracked by #9935. It is independent of the VMID correction in #9775 at the source level. Both changes are needed to advance the original workload through the currently known rocJITsu model gaps. ## Related producer-side evidence - ROCm/rocm-libraries#7976 introduced TDM iterate mode and the `tdm_gfx1250.yaml` configuration lineage that exposed this gap. - ROCm/rocm-libraries#8406 contains the same affected configuration after iterate mode was enabled and reports passing FFM and B0 validation. This is file-level evidence rather than a published per-solution trace. - ROCm/rocm-libraries#9410 added iterate-mode edge cases with waves that have zero valid rows specifically to exercise descriptor bounds clamping, and reports passing FFM validation. ## Failure before this fix Iterate mode advanced the starting global address for each repeated tile, but rocJITsu continued to apply bounds using tile-local coordinates. A final row outside the tensor was therefore copied into LDS instead of being zero-filled. An active tensor extent of zero was treated as unbounded. The complete dense or gather tile could therefore be copied instead of being masked. ## Fixes ### Include iterate-mode address advancement in tensor bounds The model decomposes each advancing iteration's linear global-address increment into an absolute tensor-coordinate origin, then applies one bounds predicate to that origin plus the tile-local coordinate. This handles sub-row advances, row advances, and rank-3 plane advances through the same path. Overlapping-stride layouts are rejected only when multiple iterations actually advance the global origin; repeated-origin descriptors do not require inverse coordinate decomposition. ### Treat zero active tensor extents as empty For an active dense or gather tile dimension, tensor extent zero masks the complete dimension. Loads zero-fill the affected LDS elements and stores suppress the corresponding global writes. Optional null descriptor groups retain their architectural zero-filled value. A rank-3 tile with null D2 therefore has a zero third tensor extent: it masks the transfer without aborting execution, still completes an enabled barrier, and releases TENSORCNT normally. ### Avoid unnecessary origin decomposition Non-iterating descriptors use an all-zero origin directly. This avoids inverse stride decomposition, including accepted rank-5 layouts whose large default strides can wrap. ## Test coverage Focused decoded-instruction regressions verify: - partial final-row load zero-fill and corresponding store suppression; - active zero-length dense and gather extents; - null-D2 rank-3 load masking and store suppression; - completion-barrier arrival and TENSORCNT release for null D2; - sub-row, row, and rank-3 plane iteration origins; - rejection of advancing overlapping-stride layouts; - acceptance of repeated-origin overlapping-stride layouts; and - safe non-iterating rank-5 execution. The complete `Gfx1250ExecutionTest.TensorDma*` test set and changed-file pre-commit checks pass locally. ## Workload status This PR intentionally does not claim to fix the complete TensileLite reproducer. It removes independently testable descriptor-model discrepancies and gets the workload closer to the unresolved tail-loop semantic question. ## Issue tracking Progress toward #9935.

## Summary This change fixes tensor-DMA bounds contracts exposed while reducing the stock gfx1250 TensileLite TDM tail failure tracked by #9935. It is independent of the VMID correction in #9775 at the source level. Both changes are needed to advance the original workload through the currently known rocJITsu model gaps. ## Related producer-side evidence - ROCm/rocm-libraries#7976 introduced TDM iterate mode and the `tdm_gfx1250.yaml` configuration lineage that exposed this gap. - ROCm/rocm-libraries#8406 contains the same affected configuration after iterate mode was enabled and reports passing FFM and B0 validation. This is file-level evidence rather than a published per-solution trace. - ROCm/rocm-libraries#9410 added iterate-mode edge cases with waves that have zero valid rows specifically to exercise descriptor bounds clamping, and reports passing FFM validation. ## Failure before this fix Iterate mode advanced the starting global address for each repeated tile, but rocJITsu continued to apply bounds using tile-local coordinates. A final row outside the tensor was therefore copied into LDS instead of being zero-filled. An active tensor extent of zero was treated as unbounded. The complete dense or gather tile could therefore be copied instead of being masked. ## Fixes ### Include iterate-mode address advancement in tensor bounds The model decomposes each advancing iteration's linear global-address increment into an absolute tensor-coordinate origin, then applies one bounds predicate to that origin plus the tile-local coordinate. This handles sub-row advances, row advances, and rank-3 plane advances through the same path. Overlapping-stride layouts are rejected only when multiple iterations actually advance the global origin; repeated-origin descriptors do not require inverse coordinate decomposition. ### Treat zero active tensor extents as empty For an active dense or gather tile dimension, tensor extent zero masks the complete dimension. Loads zero-fill the affected LDS elements and stores suppress the corresponding global writes. Optional null descriptor groups retain their architectural zero-filled value. A rank-3 tile with null D2 therefore has a zero third tensor extent: it masks the transfer without aborting execution, still completes an enabled barrier, and releases TENSORCNT normally. ### Avoid unnecessary origin decomposition Non-iterating descriptors use an all-zero origin directly. This avoids inverse stride decomposition, including accepted rank-5 layouts whose large default strides can wrap. ## Test coverage Focused decoded-instruction regressions verify: - partial final-row load zero-fill and corresponding store suppression; - active zero-length dense and gather extents; - null-D2 rank-3 load masking and store suppression; - completion-barrier arrival and TENSORCNT release for null D2; - sub-row, row, and rank-3 plane iteration origins; - rejection of advancing overlapping-stride layouts; - acceptance of repeated-origin overlapping-stride layouts; and - safe non-iterating rank-5 execution. The complete `Gfx1250ExecutionTest.TensorDma*` test set and changed-file pre-commit checks pass locally. ## Workload status This PR intentionally does not claim to fix the complete TensileLite reproducer. It removes independently testable descriptor-model discrepancies and gets the workload closer to the unresolved tail-loop semantic question. ## Issue tracking Progress toward #9935. (cherry picked from commit a659172)

mengzcai requested a review from a team as a code owner June 14, 2026 16:16

mengzcai force-pushed the users/mengzcai/fix_waitcnt0 branch from 31966e6 to 3cf854b Compare June 14, 2026 16:16

github-actions Bot added project: hipblaslt project: hipsparselt ci:hipsparselt-fast labels Jun 14, 2026

assistant-librarian Bot added the organization: ROCm label Jun 14, 2026

mengzcai force-pushed the users/mengzcai/fix_waitcnt0 branch from 3cf854b to 17f96fb Compare June 16, 2026 08:21

mengzcai mentioned this pull request Jun 16, 2026

[tox] Modify test yamls #8409

Closed

1 task done

mengzcai changed the title ~~[Fix] TDM PGR=2/PLR=0: force ds waitcnt before LR->TDM barrier~~ Fix TDM PGR=2/PLR=0; Fix kernel bugs under SIA0 Jun 16, 2026

mengzcai requested a review from jichangjichang June 16, 2026 08:27

mengzcai changed the title ~~Fix TDM PGR=2/PLR=0; Fix kernel bugs under SIA0~~ Fix 3 kernel bugs Jun 17, 2026

mengzcai force-pushed the users/mengzcai/fix_waitcnt0 branch 8 times, most recently from 4676650 to e5754d6 Compare June 23, 2026 13:23

mengzcai added 2 commits June 23, 2026 20:36

Fix LocalReadVALU overwriting correct LDS>=64KB read address

5d1d6dc

1. Set LDS>=64KB and 1LDSB and SIA0 to reproduce the issue 2. Fix test yaml: f64_gfx1250 (using LocalReadVALU)

mengzcai force-pushed the users/mengzcai/fix_waitcnt0 branch 2 times, most recently from 401898e to 41887aa Compare June 24, 2026 04:10

jichangjichang requested a review from Copilot June 24, 2026 05:31

Copilot started reviewing on behalf of jichangjichang June 24, 2026 05:32

View session

Copilot AI reviewed Jun 24, 2026

View reviewed changes

Comment thread projects/hipblaslt/tensilelite/Tensile/Components/LocalRead.py Resolved

Comment thread projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py Resolved

mengzcai force-pushed the users/mengzcai/fix_waitcnt0 branch from 41887aa to 4db9b4d Compare June 24, 2026 06:44

Modified GroupedGemm logic yamls PLR0->PLR1

65166e3

mengzcai force-pushed the users/mengzcai/fix_waitcnt0 branch from 4db9b4d to 65166e3 Compare June 24, 2026 08:04

mengzcai changed the title ~~Fix 3 kernel bugs~~ Fix 2 kernel bugs Jun 24, 2026

jichangjichang approved these changes Jun 25, 2026

View reviewed changes

mengzcai merged commit 8bbd00a into ROCm:develop Jun 25, 2026
59 checks passed

mengzcai deleted the users/mengzcai/fix_waitcnt0 branch June 25, 2026 06:24

newling mentioned this pull request Aug 10, 2026

fix(rocjitsu): honor TDM tensor bounds across iterations ROCm/rocm-systems#9936

Merged

tony-davis mentioned this pull request Oct 30, 2026

test: adds tensilelite sia0 pgr2 characterization testing #8788

Open

2 of 3 tasks

Navigation Menu

Repository navigation

Fix 2 kernel bugs - #8406

Fix 2 kernel bugs#8406
mengzcai merged 3 commits into
ROCm:developfrom
mengzcai:users/mengzcai/fix_waitcnt0

mengzcai commented Jun 14, 2026 •

edited

Loading

Uh oh!

codecov-commenter commented Jun 14, 2026 •

edited

Loading

Uh oh!

Copilot AI left a comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

Conversation

mengzcai commented Jun 14, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Uh oh!

codecov-commenter commented Jun 14, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

mengzcai commented Jun 14, 2026 •

edited

Loading

codecov-commenter commented Jun 14, 2026 •

edited

Loading