[chicyang]

JIRA ID: AIHPBLAS-4308
AIHPBLAS-4308

Motivation

TDM iterate mode (iterate_enable) repeats one fetch and advances the global address
between repeats. The number of repeats was a compile-time constant sized for a full
tile, so a workgroup landing on a partial tile kept walking past the end of the tensor
and issued out-of-bounds reads.

Technical Details

• Derive the iteration count at runtime from the rows the wave actually holds
  (size - MacroTile * workgroupId, capped at what one issueLoad covers) instead of a
  compile-time constant. _emitTdmIterCount folds the -1 the field encodes into the
  divide; the clamp to 1 precedes the subtract so an empty wave still gets its one walk
  without the subtract wrapping.

• Narrow the descriptor's dim1 to the same row count. The wave-separated path already
  did this; the non-wave-separated path left the whole tensor extent there.

• TDMSplit issues two loads from one descriptor with the base advanced between them, so
  the second load gets its own count (_emitTdmSplitIterCount). In wave-separated mode
  A and B share an aliased descriptor and both parities read it, so the parity that is
  not in iterate mode is explicitly given zero.

• _tdmIterDivMagic supplies an exhaustively validated multiply-shift pair for a
  non-power-of-two tile_dim1. The shared helper in rocisa overflows int for
  divisor 3, so it cannot be used here.

• Reject TDMIterateMode on a tensor whose TLU is true. The walk steps along the tile
  dimension, which is what dim1 carries only when the tensor is unroll-major in global
  memory; TransposeLDS: 2 sets UnrollMajorLDS without that being true.

• The walk steps a whole tile_dim1 rows at a time, so a wave left with a row count that
  is not a whole number of steps still overruns on its last step. Raise
  AssertFree{0,1}ElementMultiple to lcm(current, tile_dim1) on iterate kernels, which
  becomes a Free0/1SizeMultiple predicate at runtime, so those sizes select a different
  kernel instead. Reject outright when MacroTile itself is not a multiple of tile_dim1.

Test Plan

tox

Test Result

The failure of segment_interleave_gfx1250.yaml is not caused by the PR. ## Submission Checklist

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

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	⚠️ Warning	Error: Source/code files changed without an accompanying unit test. Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*). Current: code file(s) changed: projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py, projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py; no test file found
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

[therock-pr-bot]

🎉 All checks passed! This PR is ready for review.

[codecov]

Codecov Report

❌ Patch coverage is 14.00000% with 86 lines in your changes missing coverage. Please review.

Files with missing lines	Patch %	Lines
...blaslt/tensilelite/Tensile/KernelWriterAssembly.py	13.75%	66 Missing and 3 partials ⚠️
...lt/tensilelite/Tensile/SolutionStructs/Solution.py	15.00%	15 Missing and 2 partials ⚠️

Additional details and impacted files

@@             Coverage Diff             @@
##           develop   #10935      +/-   ##
===========================================
- Coverage    69.35%   69.34%   -0.01%     
===========================================
  Files         2777     2777              
  Lines       457906   458002      +96     
  Branches     67515    67536      +21     
===========================================
+ Hits        317548   317561      +13     
- Misses      117224   117302      +78     
- Partials     23134    23139       +5     

Flag	Coverage Δ		*Carryforward flag
TensileLite-CPP	38.06% <ø> (-<0.01%)	⬇️
TensileLite-Unit	76.41% <14.00%> (-0.09%)	⬇️
hipBLAS	90.62% <ø> (ø)		Carriedforward from d133f09
hipBLASLt	34.89% <ø> (ø)
hipCUB	82.68% <ø> (ø)		Carriedforward from d133f09
hipDNN	86.78% <ø> (ø)		Carriedforward from d133f09
hipFFT	43.60% <ø> (ø)		Carriedforward from d133f09
hipRAND	76.12% <ø> (ø)		Carriedforward from d133f09
hipSOLVER	69.03% <ø> (ø)		Carriedforward from d133f09
hipSPARSE	86.99% <ø> (ø)		Carriedforward from d133f09
rocBLAS	48.26% <ø> (ø)		Carriedforward from d133f09
rocFFT	47.64% <ø> (ø)		Carriedforward from d133f09
rocRAND	57.05% <ø> (ø)		Carriedforward from d133f09
rocSOLVER	77.12% <ø> (ø)		Carriedforward from d133f09
rocSPARSE	74.60% <ø> (ø)		Carriedforward from d133f09
rocThrust	91.60% <ø> (ø)		Carriedforward from d133f09

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

Files with missing lines	Coverage Δ
...lt/tensilelite/Tensile/SolutionStructs/Solution.py	58.67% <15.00%> (-0.24%)	⬇️
...blaslt/tensilelite/Tensile/KernelWriterAssembly.py	70.35% <13.75%> (-0.32%)	⬇️

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

[jichangjichang]

Please add the test cases

[chicyang]

Please add the test cases

I'll add test in other PR.