[Serge45]

Motivation

JIRA ID - AIHPBLAS-3993
Fixes correctness and codegen bugs in TDMSplit (split tensor_load_to_lds into two half-tile loads) for gfx1250 TDM kernels, and adds dedicated BF16/MXF8 test coverage.

Technical Details

• TDMSplit edge-M correctness fix: Compute per-half dim1 boundary correctly for wave-separated loads. The component term now accounts for dim1Divisor, and the second split half (dim1_h1 = dim1_h0 - mt//dim1Divisor) is clamped to 0 on underflow. Previously a shared/incorrect dim1 at the M mod MT0 == MT0/2 edge caused one half to load nothing and another to over-read OOB, corrupting ~half the output. dim1 is stored as a raw 32-bit row count and packed viasetTensorDim1.
• Single-wave SGPR allocation fix: Add a separate NumWaves == 1 allocation block for per-tensor split-increment SGPRs
  (tdm{A,B}GlobalSplitIncs/tdm{A,B}LdsSplitIncs), replacing dead else code that left them undefined and caused an assembler "expected absolute expression" error.
• WAR barrier preservation (SIA4): Wrap the cross-wave TDM swap sync in _PreservedBarrierModule so the SIA4 scheduling pass does not strip the current-LR WAR barrier for TDMIterateMode.
• TDM tensor-wait fixes: Emit s_wait_tensorcnt at ordinary NLL entry for PLR0 GSU kernels; defer MX-scale tensor loads in SIA0 PGR2 to fix a self-wave LDS WAR hazard.
• Parity-indexed split memtokens: Assign parity-indexed split memtokens to TDM half-writes and ds_reads; fix B-tile read half-token for single-vector tiles.
• Tests: Add tdm_split_gfx1250.yaml covering BF16 and MXF8 (Block16/Block32) with TDMSplit=[False,True], SIA=[0,4], and edge-M problem sizes; remove redundant TDMSplit combos from mxf8_gfx1250.yaml.

Test Plan

• tdm_split_gfx1250.yaml passes on gfx1250 (BF16 + MXF8 Block16/Block32), including non-aligned edge-M sizes (M=129, M=257)
• M-sweep (128/192/256/384/512) all PASS with no incorrect values

Test Result

Above tests passed on GPU.

Submission Checklist

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

[codecov-commenter]

Codecov Report

❌ Patch coverage is 31.00000% with 69 lines in your changes missing coverage. Please review.

Files with missing lines	Patch %	Lines
...blaslt/tensilelite/Tensile/KernelWriterAssembly.py	12.90%	51 Missing and 3 partials ⚠️
...rojects/hipblaslt/tensilelite/Tensile/Component.py	38.46%	7 Missing and 1 partial ⚠️
...ects/hipblaslt/tensilelite/Tensile/KernelWriter.py	66.67%	3 Missing and 1 partial ⚠️
...ts/hipblaslt/tensilelite/Tensile/Components/SIA.py	70.00%	2 Missing and 1 partial ⚠️

❌ Your project status has failed because the head coverage (76.92%) 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    #8658      +/-   ##
===========================================
- Coverage    71.33%   71.32%   -0.01%     
===========================================
  Files         2628     2628              
  Lines       413043   413121      +78     
  Branches     61875    61886      +11     
===========================================
+ Hits        294613   294635      +22     
- Misses       96656    96706      +50     
- Partials     21774    21780       +6     

Flag	Coverage Δ		*Carryforward flag
TensileLite	76.59% <31.00%> (-0.06%)	⬇️
hipBLAS	90.81% <ø> (ø)		Carriedforward from 28275cc
hipBLASLt	41.35% <ø> (ø)
hipCUB	82.68% <ø> (ø)		Carriedforward from 28275cc
hipDNN	85.91% <ø> (ø)		Carriedforward from 28275cc
hipFFT	50.17% <ø> (ø)		Carriedforward from 28275cc
hipRAND	76.12% <ø> (ø)		Carriedforward from 28275cc
hipSOLVER	69.18% <ø> (ø)		Carriedforward from 28275cc
hipSPARSE	86.55% <ø> (ø)		Carriedforward from 28275cc
rocBLAS	48.06% <ø> (ø)		Carriedforward from 28275cc
rocFFT	46.30% <ø> (ø)		Carriedforward from 28275cc
rocRAND	57.07% <ø> (ø)		Carriedforward from 28275cc
rocSOLVER	76.92% <ø> (ø)		Carriedforward from 28275cc
rocSPARSE	72.37% <ø> (ø)		Carriedforward from 28275cc
rocThrust	91.36% <ø> (ø)		Carriedforward from 28275cc

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

Files with missing lines	Coverage Δ
...blaslt/tensilelite/Tensile/Components/LocalRead.py	55.49% <100.00%> (+0.08%)	⬆️
...ts/hipblaslt/tensilelite/Tensile/Components/SIA.py	85.46% <70.00%> (-0.22%)	⬇️
...ects/hipblaslt/tensilelite/Tensile/KernelWriter.py	70.82% <66.67%> (-<0.01%)	⬇️
...rojects/hipblaslt/tensilelite/Tensile/Component.py	92.54% <38.46%> (-5.86%)	⬇️
...blaslt/tensilelite/Tensile/KernelWriterAssembly.py	68.82% <12.90%> (-0.21%)	⬇️

... and 1 file 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.

[hcman2]

Will you add GSU > 1 test case?

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
🌿 Branch Name	✅ Pass	—
📝 PR Title/Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 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.

[therock-pr-bot]

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