[boringmorning]

AIHPBLAS-1332

Motivation

In gfx1250, we cannot generate mxf8 kernel with MT=256x256x256 with PLR=1 due to vgpr limitation. By enabling HalfPLR, we can use less vgprs to achieve local data prefetch, with the price of additional scheduling constraint.

Technical Details

Concept

With PrefetchLocalRead = 1, each operand normally needs 2 VGPR buffers (current + prefetch). HalfPLR splits each K-iteration's local reads and MFMAs into two halves and uses a rotating 3-group buffer of half the per-iter size (i.e. 1.5× buffer instead of 2×), trading scheduling complexity for ~25% lower VGPR pressure on the selected operand.

New parameter — HalfPLR

Bitmask in {0, 1, 2} (Common/ValidParameters.py, GlobalParameters.py, RequiredParameters.py). Solution.py decodes it into:

• HalfPLRA = HalfPLR & 0x1
• HalfPLRB = HalfPLR & 0x2

VGPR layout — KernelWriter.py / KernelWriterAssembly.py

• valuBlocksA/B = 1.5 (was numVgprBufferA/B) when HalfPLRA/B is on.
• Allocates 3 group aliases vgprValu{A,B}_G0/G1/G2 over the same base VGPR range, each numVgprValuPerBlock // 2 wide.

Group rotation

def getHalfPLRGroups(self, kernel, lc, u):
    g0 = ((lc * kernel["LoopIters"] + u) * 2) % 3
    g1 = (g0 + 1) % 3
    return [g0, g1]

Each sub-iter uses two of the three groups; getHalfPLRValuStr(groups, offset, tc) picks G[gIdx] from offset // (numVgprValuPerBlock//2). All Valu{A,B}_X*_I* destination strings in Components/LocalRead.py and all source strings in generateSrcStrForMFMA / generateSrcStrForMFMAshiftK (now plumbed with lc=unrollLoopIdx) are rewritten through this helper when HalfPLR is active.

Sub-iter schedule — _makeSubIterSchedule (SIA = 0)

New localReadCodeSecondHalf arg. New ordering when HalfPLR:

[globalRead (if PGR=1)] -> waitLW -> sync -> localRead(1st half) -> localWrite
   -> waitCode
   -> first half of MFMAs
   -> localRead(2nd half) -> pointerLW -> pointerLR
   -> [globalRead (if PGR=2)]
   -> second half of MFMAs

localReadDo output for the HalfPLR'd operand is split by counting countLocalRead and assigning the first N/2 to localReads, the rest to localReadsSecondHalf. MFMAs are split symmetrically via countInstruction(macIterCode) / 2. The reason is that the 2nd half of localRead use the same vgpr group as the first half of MFMAs.

Loop copies

loopCopies = 3 if kernel["HalfPLR"] else 2 if expand else 1 — three unroll copies are required because the 3-group rotation has period 3 over lc. NGLL & NLL will not be generated to reduce code size, so endCounter for unrolled loop is 0 .

Tail-loop / unrolled loop GR handling — KernelWriterAssembly.py

• New graIncrementMask(kernel, tPA, tPB): in globalReadIncrement at prefetchIndex == 0, if loopCounter <= 1, conditionally zeroes tdmABIncs / tdmMXSAMXSBIncs (multi-wavegroup) or per-tensor GlobalReadIncs{A,B,MXSA,MXSB} (single wavegroup) so the last two PGR iters don't advance addresses.
• globalReadDo header (mode=1, tPA, SuppressNoLoadLoop): if HalfPLR, when loopCounter <= PrefetchGlobalRead, conditionally clears the TDM enable bits (tdmAGroup0, tdmBGroup0, tdmMXSAGroup0, tdmMXSBGroup0) instead of the original SrdA/B+2 = 0 path.
• Before the tail loop (still under SuppressNoLoadLoop): re-enables TDM and calls tdmSwapLdsOffset on A/B (and MXS where applicable) to realign LDS, terminated by s_waitcnt dscnt(0) + s_barrier ("prevent dummy ds_load from polluting LDS buffer").

Outer/inner MFMA loop order

generateSrcStrForMFMA force outer=0, loopSwap=True when HalfPLRA so the A index is the outer loop (reuses the rotating group selection across the B sweep).

Bug fix folded in

Fix wrong VgprG2L allocation condition for tensor B: should not allocate VgprG2L if TDM is enabled.

SCBranchSCC1(labelName=lastIterEnd, ...) -> SCBranchSCC1(labelName=lastIterEnd.getLabelName(), ...) (was passing a Label object instead of its name string).

Test Plan

Add HalfPLR kernels in tensilelite/Tensile/Tests/common/gemm/gfx12/mxf8_gfx1250.yaml:

• Adds MI [16,16,128,1,1,8,8,2,2] for the targeted MT 256x256x256.
• Adds HalfPLR: [0, 1, 2].
• Expands PrefetchGlobalRead: [1, 2].
• Adds problem sizes 256x256x1x{128,384,640,896} to exercise small-K behavior.

Test Result

Passed

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 2.33333% with 293 lines in your changes missing coverage. Please review.

Files with missing lines	Patch %	Lines
...blaslt/tensilelite/Tensile/KernelWriterAssembly.py	3.97%	121 Missing ⚠️
...ects/hipblaslt/tensilelite/Tensile/KernelWriter.py	1.96%	100 Missing ⚠️
...lt/tensilelite/Tensile/SolutionStructs/Solution.py	0.00%	39 Missing ⚠️
...blaslt/tensilelite/Tensile/Components/LocalRead.py	0.00%	33 Missing ⚠️

❌ 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    #7453      +/-   ##
===========================================
- Coverage    61.76%   61.69%   -0.07%     
===========================================
  Files         2082     2083       +1     
  Lines       356042   356740     +698     
  Branches     53628    53812     +184     
===========================================
+ Hits        219883   220072     +189     
- Misses      117417   117917     +500     
- Partials     18742    18751       +9     

Flag	Coverage Δ		*Carryforward flag
TensileLite	26.13% <2.33%> (+0.01%)	⬆️
hipBLAS	90.65% <ø> (ø)		Carriedforward from d5bf18f
hipBLASLt	41.24% <ø> (ø)
hipCUB	82.21% <ø> (ø)		Carriedforward from d5bf18f
hipDNN	85.61% <ø> (ø)		Carriedforward from d5bf18f
hipFFT	51.12% <ø> (ø)		Carriedforward from d5bf18f
hipRAND	76.12% <ø> (ø)		Carriedforward from d5bf18f
hipSOLVER	69.24% <ø> (ø)		Carriedforward from d5bf18f
hipSPARSE	85.09% <ø> (ø)		Carriedforward from d5bf18f
rocBLAS	48.11% <ø> (ø)		Carriedforward from d5bf18f
rocFFT	51.26% <ø> (ø)		Carriedforward from d5bf18f
rocRAND	57.04% <ø> (ø)		Carriedforward from d5bf18f
rocSOLVER	77.83% <ø> (ø)		Carriedforward from d5bf18f
rocSPARSE	72.68% <ø> (ø)		Carriedforward from d5bf18f

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

Files with missing lines	Coverage Δ
...slt/tensilelite/Tensile/Common/GlobalParameters.py	84.86% <ø> (ø)
...t/tensilelite/Tensile/Common/RequiredParameters.py	100.00% <ø> (ø)
...aslt/tensilelite/Tensile/Common/ValidParameters.py	76.92% <ø> (ø)
...blaslt/tensilelite/Tensile/Components/LocalRead.py	2.89% <0.00%> (-0.09%)	⬇️
...lt/tensilelite/Tensile/SolutionStructs/Solution.py	5.75% <0.00%> (-0.38%)	⬇️
...ects/hipblaslt/tensilelite/Tensile/KernelWriter.py	9.46% <1.96%> (-0.10%)	⬇️
...blaslt/tensilelite/Tensile/KernelWriterAssembly.py	7.46% <3.97%> (-0.04%)	⬇️

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

[Copilot]

Pull request overview

This PR introduces a new HalfPLR option for TensileLite (hipBLASLt) kernels on gfx1250 MXF8, aiming to reduce VGPR pressure when PrefetchLocalRead=1 by using a 3-group “half-sized” rotating VGPR buffer scheme and a split sub-iteration schedule.

Changes:

• Add the HalfPLR tuning parameter (and derived HalfPLRA/HalfPLRB flags) with validity constraints and defaults.
• Update kernel generation (VGPR allocation, macro aliases, MFMA/local-read operand selection, and tail/no-load-loop handling) to support HalfPLR scheduling and address masking.
• Extend gfx12 MXF8 test YAML to exercise HalfPLR and additional small-K sizes / PGR settings.

Reviewed changes

Copilot reviewed 8 out of 8 changed files in this pull request and generated 4 comments.

Show a summary per file

File	Description
projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf8_gfx1250.yaml	Adds HalfPLR and additional MI/PGR/problem sizes to test gfx1250 MXF8 kernels.
projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py	Decodes HalfPLR into A/B flags and rejects unsupported configurations.
projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py	Adds HalfPLR VGPR alias macros, HalfPLR MFMA source selection, and tail-loop / GR increment masking logic.
projects/hipblaslt/tensilelite/Tensile/KernelWriter.py	Implements HalfPLR-aware scheduling, local-read splitting, VGPR allocation adjustments, and loop copy changes.
projects/hipblaslt/tensilelite/Tensile/Components/LocalRead.py	Routes LDS transpose local-read destinations through HalfPLR VGPR group selection when enabled.
projects/hipblaslt/tensilelite/Tensile/Common/ValidParameters.py	Registers HalfPLR as a valid tuning parameter.
projects/hipblaslt/tensilelite/Tensile/Common/RequiredParameters.py	Includes HalfPLR in the minimal required parameter set.
projects/hipblaslt/tensilelite/Tensile/Common/GlobalParameters.py	Adds the default HalfPLR: 0 setting to solution defaults.

Comments suppressed due to low confidence (3)

projects/hipblaslt/tensilelite/Tensile/KernelWriter.py:4090

• halfCnt is computed as countLocalRead(localReadCodeB) / 2, yielding a float. If the local-read count is odd, halfCnt == 0 will never be reached and the split between first/second-half reads will be wrong. Use integer division and validate the read count is even before splitting.

          localReadCodeB, packCodeB, packPreB = self.localReadDo(kernel, bufferIdx, iui*self.states.numReadsIterCoalescedB, 0, tensorParametersB)
          if kernel["HalfPLRB"]:
            halfCnt = countLocalRead(localReadCodeB) / 2
            readItems = localReadCodeB.flatitems()
            while readItems:
              item = readItems.pop(0)
              localReads.add(item)
              readsThisItem = countLocalRead(item)
              if readsThisItem:
                assert readsThisItem==1, "Scheduler assumes 1 read per item"
                halfCnt -= 1
                if halfCnt == 0:
                  break

projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py:2182

• Grammar in rejection reason: "does not supports" should be "does not support" (user-facing message).

      if (state["HalfPLRA"] and not (state["UnrollMajorLDSA"] or state["enableLDSTrA"])) or \
        (state["HalfPLRB"] and not (state["UnrollMajorLDSB"] or state["enableLDSTrB"])):
        reject(state, printRejectionReason, "Currently HalfPLR does not supports packing (need UnrollMajorLDS or use LDSTrInst)")
        return

projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py:2188

• Grammar in rejection reason: "does not supports" should be "does not support" (user-facing message).

        return
      if state["ConvertAfterDS"]:
        reject(state, printRejectionReason, "Currently HalfPLR does not supports ConvertAfterDS")
        return

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.