Repository navigation
[tensile] Fix SIA0 PGR2 global-read placement - #8417
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.90%) 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 #8417 +/- ##
===========================================
- Coverage 65.62% 65.54% -0.08%
===========================================
Files 2610 2613 +3
Lines 405606 406914 +1308
Branches 60526 60726 +200
===========================================
+ Hits 266159 266708 +549
- Misses 119412 120093 +681
- Partials 20035 20113 +78
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
mengzcai
approved these changes
Jun 16, 2026
Contributor
Author
|
I will do more test for gfx1250 to make sure it works correctly. |
Defer SIA0 non-TDM PGR2 global reads until after the local-write iteration consumes the previous vgprG2L values, and keep tail local-read offsets reset before tail MACs. This prevents next-prefetch buffer loads from clobbering data before it is stored to LDS.
cycheng
force-pushed
the
users/cycheng2/fix-f32_gfx1250.yaml
branch
from
June 18, 2026 08:29
8020b2b to
312c91d
Compare
Avoid clearing the tail local-read Blk offset for SIA0 TDM kernels, whose tail descriptors already select the correct LDS bank. This preserves the non-TDM SIA0 tail reset fix without regressing TDM tail reads.
jichangjichang
approved these changes
Jun 23, 2026
2 of 3 tasks
davidd-amd
added a commit
that referenced
this pull request
Jun 24, 2026
…avior Update the SIA0 tail-reset golden to develop's current emission (tail_lr_reset_a/b = True), so the characterization is green on develop and guards against future regressions of the SIA0 PGR2 non-TDM placement. The snapshot was first captured against #8417's parent (markers False) to prove the guard fails-before / passes-after the fix; this commit pins the post-fix behavior that should hold going forward. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
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.
Defer SIA0 non-TDM PGR2 global reads until after the local-write iteration consumes the previous vgprG2L values, and keep tail local-read offsets reset before tail MACs.
This prevents next-prefetch buffer loads from clobbering data before it is stored to LDS.
Test Result
(1) Codegen test: asm code (before v.s. after)
3062 kernels, 69 yamls
difference only on:
(2) In runtime test, the following are failed but not related to this PR.
spmm_f8hs.yaml, spmm_f8_ml.yaml, spmm_f8.yaml, fp16_gfx1250.yaml, mxf8_gfx1250.yaml, mxf8f4ss_gfx1250.yaml, subtile_bf16_gfx1250_cluster.yaml
Submission Checklist