Repository navigation
Conversation
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. |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #10871 +/- ##
========================================
Coverage 69.36% 69.36%
========================================
Files 2779 2779
Lines 458135 458133 -2
Branches 67569 67568 -1
========================================
Hits 317781 317781
+ Misses 117213 117211 -2
Partials 23141 23141
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
| # LBSPP = roundUp(DepthU * bpe * VectorWidth, 256). When full VW=8 would | ||
| # exceed the non-iterate cap and VectorWidth was halved to fit pad_interval, | ||
| # the non-iterate descriptor path does not match the LDS layout. | ||
| if vector_width < 8: |
There was a problem hiding this comment.
Why do we need this condition?
Shouldn't we just rely on full_vw_lbspp > TensorDataMoverLoad.MAX_NON_ITERATE_LBSPP_BYTES to decide if we need iterative mode or not?
There was a problem hiding this comment.
Agreed. Pushed a new patch
|
Please fix the typo (150) in the title. |
dsaffars
changed the title
| multiple = 256 | ||
| for tc in ["A", "B"]: | ||
| if state.get("_TDMIterateMode%s" % tc, False): | ||
| continue |
There was a problem hiding this comment.
I do not think we need to remove this condition?
Why did you remove it?
There was a problem hiding this comment.
This condition is about if not state.get("_inputVW%s_was_auto" % tc, False):
| if dwords > 0 and (dwords & (dwords - 1)) == 0: | ||
| pad_interval = int(math.log2(dwords)) - 1 | ||
| if pad_interval <= 7: | ||
| break # current VW fits pad_interval encoding |
There was a problem hiding this comment.
I do not think we need to remove this vw //= 2 logic.
This logic is just inconsistent with LdsBlockSizePerPadA/B != -1 case.
So, I think we just need to skip wv //=2 if LdsBlockSizePerPadA/B != -1
jichangjichang
left a comment
jichangjichang
left a comment
There was a problem hiding this comment.
Why did TDMIM0 VWB4 LBSPPB1024 cause illegal memory acess?
We need the root cause why TDMIM0 padding cause illegal memory access
There was a problem hiding this comment.
Pull request overview
Fixes gfx1250 TDM illegal global memory access for DepthU=128 cases where Tensile previously “solved” the 1024B pad_interval encoding limit by halving VectorWidth{A,B} (e.g., 8→4), by instead preferring TDM iterate mode selection via a shared helper.
Changes:
- Added
TensorDataMoverLoad.needsIterateModeForPad()to centralize iterate-vs-non-iterate selection for TDM padding constraints. - Updated
Solution.pyStage 1/2 TDM iterate-mode decision logic to use the helper and removed the Stage 2 VectorWidth-halving loop. - Added unit tests for the helper and adjusted a gfx1250 characterization YAML to keep the test case within SGPR budget.
Reviewed changes
Copilot reviewed 4 out of 4 changed files in this pull request and generated 3 comments.
| File | Description |
|---|---|
| projects/hipblaslt/tensilelite/Tensile/Components/TensorDataMover.py | Adds the new helper for deciding when to force TDM iterate mode for padding. |
| projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py | Uses the helper in Stage 1/2 and removes VectorWidth halving in favor of iterate mode. |
| projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_TensorDataMoverPad.py | New unit tests validating pad interval encoding and iterate-mode selection logic. |
| projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/_codegen/data/test_data/_designed/gfx1250/streamk_tdm_prefetchgl2.yaml | Adjusts DepthU to avoid SGPR budget overflow in the designed gfx1250 case. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
b78fc29 to
85fa7fe
Compare
| return False | ||
| if ldsBlockSizePerPad > TensorDataMoverLoad.MAX_NON_ITERATE_LBSPP_BYTES: | ||
| return True | ||
| full_vw_bytes = int(depth_u * bpe * 8) |
There was a problem hiding this comment.
I do not understand why we need this condition.
Why do we need to ignore VectorWidth setting and use 8 (= maximum of VectorWidth)?
| lbspp = roundUpToNearestMultiple(int(state["_DepthU%s" % tc] * bpe_tc * vw), 256) | ||
| else: | ||
| lbspp = lbspp_cfg | ||
| if TensorDataMoverLoad.needsIterateModeForPad(lbspp, state["_DepthU%s" % tc], bpe_tc): |
There was a problem hiding this comment.
My understanding is just to use original if condition if lbspp > 1024 here.
I think we just need to add the logic to use lbspp = lbspp_cfg if lbspp_cfg is not -1.
| lbspp = roundUpToNearestMultiple(int(depthU_tc * tmpBpe * vw), multiple) | ||
| # When TDMIterateMode is explicit (not -1), Stage 1 does not auto-select | ||
| # iterate; block VW halving and force iterate for the gfx1250 broken path. | ||
| if iterModeMask != -1 and TensorDataMoverLoad.needsIterateModeForPad(lbspp, depthU_tc, tmpBpe): |
There was a problem hiding this comment.
Shouldn't we simply reject the kernel if both TDMIterateMode and LdsBlockSizePerPad are not -1 and the setting is inconsistent?
|
Already addressed. |
JIRA ID: AIHPBLAS-4324
Fix gfx1250 TDM illegal memory access when
DepthU=128forces non-iterate TDM with halvedVectorWidthB=4. Prefer TDMIterateMode (iterate_enable) and keep auto-derivedVectorWidth=8instead of halving VW to fit the 1024 B pad_interval cap.Motivation
Tuning configs with
DepthU=128,TDMInst=3, and auto-derived vector widths hithipErrorIllegalAddress. The same macro-tile withDepthU=64passes.Root cause: when
LBSPP = roundUp(DepthU × bpe × VectorWidth, 256)exceeds 1024 B, Tensile Stage 2 halvedVectorWidthB(8→4) so the non-iteratepad_intervalfield could encode the padding. That produced kernels like:TDMIM0 VWB4 LBSPPB1024→ illegal global accessThe working workaround is explicit iterate mode:
TDMIM2 VWB8 LBSPPB2048→ passesThe non-iterate TDM descriptor path with reduced VectorWidth does not match the LDS layout on gfx1250. Halving VW to satisfy
pad_intervalis not a safe fallback for this case — iterate mode is.Technical Details
New helper:
TensorDataMoverLoad.needsIterateModeForPadCentralizes the iterate-vs-non-iterate decision in
TensorDataMover.py. ReturnsTruewhen non-iterate TDM must not be used:LBSPP == 0LdsBlockSizePerPad: 0) — no padding, do not force iterateLBSPP > 1024pad_intervalmax → 1024 B)roundUp256(int(DepthU × bpe × 8)) > 1024LBSPPB ≤ 1024The VW=8 reference LBSPP uses
int(depth_u * bpe * 8), matchingSolution.py's LBSPP formula. This is required for fractionalbpe(fp40.5, bf6/fp60.75) whereround(DepthU × bpe) × 8can under-estimate (e.g.DU=257,bpe=0.5→ true LBSPP 1280, wrong formula gives 1024).Solution.pyStage 1 (TDMIterateMode: -1)LdsBlockSizePerPad{A,B}when not-1(not always computed from VW) — fixes xccremap false iterate rejection.needsIterateModeForPadinstead oflbspp > 1024alone.Solution.pyStage 2 (VW halving preserved, with guards)Stage 2 halving is kept for pad_interval encoding edge cases, but gated:
_TDMIterateMode)._inputVW*_was_auto) — do not silently rewrite yaml VW.LdsBlockSizePerPad != -1— halving VW does not change explicit LBSPP.TDMIterateModeis explicit (not-1) andneedsIterateModeForPad→ force iterate instead of halving.vw //= 2loop until pad_interval encoding fits.Coexists with develop's TLU rejection (#10935): iterate mode requires K-major global layout (
TLU{A,B}must be false).Characterization yaml
streamk_tdm_prefetchgl2.yaml:DepthU: 256 → 128so the mxf8 SGPR-budget test stays non-iterate atVWB=8,LBSPPB=1024.DU=256would force iterate and overflow the SGPR budget for that test harness.Test plan
cd projects/hipblaslt/tensilelite python -m pytest Tensile/Tests/unit/test_TensorDataMoverPad.py -q python -m pytest Tensile/Tests/unit/characterization/_codegen/ -qtest_TensorDataMoverPad.py(9 cases):DU=128direct and halved-VW paths,DU=64safe cases,LBSPP=0bpe=0.5): non-integerDepthU×bpe(DU=257), cap boundary (DU=256)bpe=0.75): halved-VW path (DU=171)