Repository navigation
[tensilelite] Support TDM iterate mode - #7976
Merged
jichangjichang merged 3 commits intoJun 8, 2026
Merged
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
chicyang
force-pushed
the
users/chicyang_iterate_mode_large_VW
branch
from
June 2, 2026 15:21
a68cb70 to
00ac2b4
Compare
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 #7976 +/- ##
===========================================
- Coverage 61.66% 61.63% -0.03%
===========================================
Files 2091 2092 +1
Lines 359565 359961 +396
Branches 54389 54475 +86
===========================================
+ Hits 221701 221853 +152
- Misses 119084 119322 +238
- Partials 18780 18786 +6
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
hcman2
reviewed
Jun 3, 2026
chicyang
force-pushed
the
users/chicyang_iterate_mode_large_VW
branch
from
June 3, 2026 07:47
e2533e7 to
9653ddb
Compare
Contributor
There was a problem hiding this comment.
Pull request overview
This PR adds an opt-in (and auto) TDMIterateMode for TensileLite’s TDM path to bypass the pad_interval 1024B encoding limit by supplying padding via tensor_load_to_lds iterate mode (descriptor Group2/3), allowing larger VectorWidth/DepthU combinations to generate.
Changes:
- Introduces new solution parameter
TDMIterateMode(-1 auto / bitmask for A/B) with defaults + validation plumbing. - Updates TDM descriptor SGPR allocation and
tensor_load_to_ldsemission to optionally pass iterate operands (including wave-separated handling). - Expands gfx1250 TDM-related test YAMLs to cover iterate-mode permutations and larger
DepthU/problem sizes.
Reviewed changes
Copilot reviewed 9 out of 9 changed files in this pull request and generated 5 comments.
Show a summary per file
| File | Description |
|---|---|
| projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/tdm_gfx1250.yaml | Extends TDM test coverage for iterate-mode and broader VW/DU combinations. |
| projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf8_gfx1250.yaml | Adds iterate-mode parameter coverage for MXF8 cases. |
| projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf4_gfx1250.yaml | Adds iterate-mode parameter coverage and expands DepthU coverage for MXF4 cases. |
| projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py | Implements iterate-mode selection (auto/manual) and adjusts pad_interval validation behavior. |
| projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py | Adds descriptor Group2/3 handling and emits iterate-mode descriptor initialization. |
| projects/hipblaslt/tensilelite/Tensile/Components/TensorDataMover.py | Refactors data_size encoding helper and fixes iterate_enable bit setting/clearing. |
| projects/hipblaslt/tensilelite/Tensile/Common/ValidParameters.py | Registers TDMIterateMode as a valid parameter with documented values. |
| projects/hipblaslt/tensilelite/Tensile/Common/RequiredParameters.py | Adds TDMIterateMode to the minimal required parameter set. |
| projects/hipblaslt/tensilelite/Tensile/Common/GlobalParameters.py | Adds default TDMIterateMode: 0 to benchmark common defaults. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
jichangjichang
approved these changes
Jun 8, 2026
Contributor
|
The CI failures is not related to this PR. |
davidd-amd
pushed a commit
that referenced
this pull request
Jun 8, 2026
## Motivation
When a user-specified VectorWidth produces an LdsBlockSizePerPad that
exceeds the pad_interval encoding's 1024 B cap (e.g. large VW on large
DepthU), the kernel cannot be generated under the existing pad_interval
path. This PR adds an opt-in TDMIterateMode that delivers the LDS pad
through tensor_load_to_lds with iterate_enable=1 instead of the 3-bit
pad_interval field, allowing the requested VW to be honored.
<!-- Explain the purpose of this PR and the goals it aims to achieve.
-->
## Technical Details
New YAML parameter (bitmask, bit0=A, bit1=B)
- TDMIterateMode: [-1, 0, 1, 2, 3] # -1=auto(only enabling when LBSPP >
1024B), 0 = off (default), 3 = both A and B
**Descriptor layout**
The tensor_load_to_lds instruction takes four groups. Iterate-mode
produces the same LDS bytes as the pad_interval path, so read-side
codegen (lraFinalOffset, LocalRead, LraTileAssignment) needs no changes;
only the write-side descriptor setup differs:
- Group 0/1: base/LDS addr, data_size, dims; iterate-mode clears HW
padding via setPadding(0, 0) and overrides tile-row count via
setTensorTile1(N-rows-per-iter).
- Group 2 (4 SGPRs): iterate fields — iter_count, lds_addr_increment,
global_addr_increment, populated by _emitTdmIterateInit() (sub-byte
dtypes like fp4 get a dedicated byte-unit adjustment).
- Group 3: required by the instruction encoding but unused for 2D
descriptors → aliased onto Group 2's 4 SGPRs via RegSet, satisfying the
HW operand count at zero SGPR cost.
Both non-wave-separated and wave-separated paths are supported. WS
shares one issueLoad across even/odd waves; the non-iterate mode Group
2/3 SGPRs are zero-initialized.
Excluded: enableLDSTr{A,B} — ds_load_tr uses brute-force LDS layout
search that already keeps LBSPP within the pad_interval cap, and the
path operates at VW=1, so iterate-mode brings no benefit. Rejected at
Solution.py with a clear error.
## Test Plan
Extended Tensile/Tests/common/gemm/gfx12/{tdm,mxf4,mxf8}_gfx1250.yaml
with TDMIterateMode.
<!-- Explain any relevant testing done to verify this PR. -->
## Test Result
tox pass
<!-- Briefly summarize test outcomes. -->
## Submission Checklist
- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
archana-ramalingam
added a commit
that referenced
this pull request
Jun 11, 2026
These gfx1250 streamk tests fail with "No valid solutions found" due to a VectorWidth derivation regression in PR #7976 (TDM iterate mode). Mark as xfail until the kernel team resolves the VW halving issue. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
3 tasks
kuhar
pushed a commit
to ROCm/rocm-systems
that referenced
this pull request
Aug 18, 2026
## 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.
atgutier
pushed a commit
to ROCm/rocm-systems
that referenced
this pull request
Aug 19, 2026
## 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)
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
When a user-specified VectorWidth produces an LdsBlockSizePerPad that exceeds the pad_interval encoding's 1024 B cap (e.g. large VW on large DepthU), the kernel cannot be generated under the existing pad_interval path. This PR adds an opt-in TDMIterateMode that delivers the LDS pad through tensor_load_to_lds with iterate_enable=1 instead of the 3-bit pad_interval field, allowing the requested VW to be honored.
Technical Details
New YAML parameter (bitmask, bit0=A, bit1=B)
Descriptor layout
The tensor_load_to_lds instruction takes four groups. Iterate-mode produces the same LDS bytes as the pad_interval path, so read-side codegen (lraFinalOffset, LocalRead, LraTileAssignment) needs no changes; only the write-side descriptor setup differs:
Both non-wave-separated and wave-separated paths are supported. WS shares one issueLoad across even/odd waves; the non-iterate mode Group 2/3 SGPRs are zero-initialized.
Excluded: enableLDSTr{A,B} — ds_load_tr uses brute-force LDS layout search that already keeps LBSPP within the pad_interval cap, and the path operates at VW=1, so iterate-mode brings no benefit. Rejected at Solution.py with a clear error.
Test Plan
Extended Tensile/Tests/common/gemm/gfx12/{tdm,mxf4,mxf8}_gfx1250.yaml with TDMIterateMode.
Test Result
tox pass
Submission Checklist