Skip to content

Enable plr min #1208

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
b-shi merged 2 commits into from
Oct 21, 2025
Merged

Enable plr min #1208

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
9fd396c
Enable PLR when DepthU=MI_K
adityalj Aug 12, 2025
e668710
Fix CI failures, add more tests, enable for ldsTR
b-shi Oct 13, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions projects/hipblaslt/tensilelite/Tensile/Common/RequiredParameters.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -60,6 +60,7 @@ def getRequiredParametersMin() -> set:
'LdsPadA',
'LdsPadB',
'LdsPadMetadata',
'LDSTrInst',
'LocalReadVectorWidth',
'LocalWritePerMfma',
'MIArchVgpr',
Expand Down
40 changes: 22 additions & 18 deletions projects/hipblaslt/tensilelite/Tensile/Components/CustomSchedule.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -427,22 +427,24 @@ def hasCustomSchedule(kernel):
optSchedule = dict()
syncCode = []

plr = 3 if kernel["ForceUnrollSubIter"] else 1

if isTN and TLDS == 1:
optSchedule = {
'SYNC' : [[6,7, 20,21, 46,47, 61]],
'GRIncA' : [[0,1,2,3,4,4,4,4,4]],
'GRIncB' : [[5,5,5,5,5,6,6,6,6]],
'LRA0' : [[0,0, 1,1, 2,2, 3,3]],
'GRA' : [[8,8,9,9,10,10,11,11,12,12, 23,23,24,24,25,25]],
'LRB0' : [[13,13,14,14,15,15,16,16]],
'LRA1' : [[48,48,49,49,50,50,51,51]],
'LRB1' : [[52,52,54,54,55,55,56,56]],
'GRB' : [[26,26,27,27, 39,39,40,40,41,41,42,42,43,43, 53,53]],
'LCC' : [[60, 60]],
'LRSA' : [[17]],
'LRSB' : [[17]],
'LWSA' : [[57]],
'LWSB' : [[57]],
'SYNC' : [[6,7, 20,21, 46,47, 61]],
'GRIncA' : [[0,1,2,3,4,4,4,4,4]],
'GRIncB' : [[5,5,5,5,5,6,6,6,6]],
'LRA0' : [[0,0, 1,1, 2,2, 3,3]],
'GRA' : [[8,8,9,9,10,10,11,11,12,12, 23,23,24,24,25,25]],
'LRB0' : [[13,13,14,14,15,15,16,16]],
'LRA%u'%plr : [[48,48,49,49,50,50,51,51]],
'LRB%u'%plr : [[52,52,54,54,55,55,56,56]],
'GRB' : [[26,26,27,27, 39,39,40,40,41,41,42,42,43,43, 53,53]],
'LCC' : [[60, 60]],
'LRSA' : [[17]],
'LRSB' : [[17]],
'LWSA' : [[57]],
'LWSB' : [[57]],
}
syncCode = [SWaitCnt(dscnt=0, vlcnt=-1, vscnt=-1, comment="Wait for LRA0/LRB0 to complete"),
SBarrier(comment=""),
Expand All @@ -456,10 +458,12 @@ def hasCustomSchedule(kernel):

numMfma = 64
# B0A0, B0A1, B1A0, B1A1
mfmaReorder = [0,1,2,3, 8,9,10,11, 16,17,18,19, 24,25,26,27,
4,5,6,7, 12,13,14,15, 20,21,22,23, 28,29,30,31,
32,33,34,35, 40,41,42,43, 48,49,50,51, 56,57,58,59,
36,37,38,39, 44,45,46,47, 52,53,54,55, 60,61,62,63]
mfmaReorder = []
if not kernel["ForceUnrollSubIter"]:
mfmaReorder = [0,1,2,3, 8,9,10,11, 16,17,18,19, 24,25,26,27,
4,5,6,7, 12,13,14,15, 20,21,22,23, 28,29,30,31,
32,33,34,35, 40,41,42,43, 48,49,50,51, 56,57,58,59,
36,37,38,39, 44,45,46,47, 52,53,54,55, 60,61,62,63]
opt1 = ScheduleInfo(1, numMfma, optSchedule, syncCode, mfmaReorder)
return True, opt1
elif is192x256x64DTL and is16bit and not isMixed and ([GRVWA, GRVWB, LRVW] == [8, 8, 8]) and MI == [16,16,32,1] and MIWG == [2,2]:
Expand Down
246 changes: 138 additions & 108 deletions projects/hipblaslt/tensilelite/Tensile/Components/LocalRead.py
View file
Open in desktop

Large diffs are not rendered by default.

6 changes: 5 additions & 1 deletion projects/hipblaslt/tensilelite/Tensile/Components/SIA.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,7 +33,6 @@

from copy import deepcopy
from typing import Tuple

PRECISION = 100
class SIA3(SIA):
kernel = {"ScheduleIterAlg": 3}
Expand Down Expand Up @@ -258,6 +257,11 @@ def calculateLatencyLeft(numReads, localReadBlockWidth, localReadLatency):
# final index definition
writer.states.numMfmaForNextLoopLR = min(writer.states.numMfmaForNextLoopLR,numMfmaPerIter-1)
writer.states.syncPlrMfmaIndex = numMfmaPerIter*(kernel["LoopIters"]-writer.states.numItersPLR+1) - writer.states.numMfmaForNextLoopLR - 1 if writer.states.numItersPLR else 0

if kernel["ForceUnrollSubIter"]:
if ( kernel["ProblemType"]["DataType"].isComplex()):
writer.states.syncPlrMfmaIndex = writer.states.syncPlrMfmaIndex *4 # Complex

numMfmaBetweenLWandBarrier = 2 if kernel["MatrixInstM"] == 32 else 3
writer.states.lwEndMfmaIndex = max(writer.states.syncPlrMfmaIndex - numMfmaBetweenLWandBarrier,0) if writer.states.numItersPLR else numMfmaPerIter*kernel["LoopIters"] - 1
if kernel["DirectToLds"] and kernel["PrefetchGlobalRead"] == 2:
Expand Down
129 changes: 102 additions & 27 deletions projects/hipblaslt/tensilelite/Tensile/KernelWriter.py
View file
Open in desktop

Large diffs are not rendered by default.

27 changes: 25 additions & 2 deletions projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6790,8 +6790,31 @@ def findSparseOffset(isA:bool):
outer = 0
loopSwap = True
inner = 1 - outer # inner is the opposite of outer
for idxOuter in range(0, kernel["MIWaveTile"][outer]):
for idxInner in range(0, kernel["MIWaveTile"][inner]):

idxOuter_start = 0
idxInner_start = 0
idxOuter_stop = kernel["MIWaveTile"][outer]
idxInner_stop = kernel["MIWaveTile"][inner]
numSubTiles = kernel["numSubTiles"]
if numSubTiles > 1 and not self.states.inTailLoop:
# iter (idxOuter_start, idxOuter_stop) (idxInner_start, idxInner_stop) MFMA
# 0 (0,4) (0,4) MFMA(A0,B0)
# 1 (0,4) (4,8) MFMA(A1,B0)
# 2 (4,8) (0,4) MFMA(A0,B1)
# 3 (4,8) (4,8) MFMA(A1,B1)
outerBy2=(kernel["MIWaveTile"][outer]//numSubTiles)
innerBy2=(kernel["MIWaveTile"][inner]//numSubTiles)
outerMod2=(kernel["MIWaveTile"][outer]%numSubTiles)
innerMod2=(kernel["MIWaveTile"][inner]%numSubTiles)
idxHalfO = u//numSubTiles
idxHalfI = u % numSubTiles
idxOuter_start = (outerBy2 + outerMod2)*idxHalfO
idxInner_start = (innerBy2 + innerMod2)*idxHalfI
idxOuter_stop = kernel["MIWaveTile"][outer] - (1-idxHalfO)* outerBy2
idxInner_stop = kernel["MIWaveTile"][inner] - (1-idxHalfI)* innerBy2

for idxOuter in range(idxOuter_start, idxOuter_stop):
Comment thread
b-shi marked this conversation as resolved.
for idxInner in range(idxInner_start, idxInner_stop):
idx0 = idxInner
idx1 = idxOuter
if loopSwap:
Expand Down
36 changes: 31 additions & 5 deletions projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1426,6 +1426,37 @@ def assignDerivedParameters(
if "ValidDepthU" in state:
del state["ValidDepthU"]

#################################################################
# ForceUnrollSubIter requirements
# - Needs PGR > 0, double buffer
# - MIWaveTile must be even and larger than 2
# - TLU{A,B} cases only supported if using LdsTR or if VPerm not needed (size{A,B} >= 4)
#
# - Not supported for mixed precision cases currently
sizeDataTypeA = state["ProblemType"]["DataTypeA"].numBytes()
sizeDataTypeB = state["ProblemType"]["DataTypeB"].numBytes()
sizeDataType = state["ProblemType"]["DataType"].numBytes()
TLUA = state["ProblemType"]["TLUA"]
TLUB = state["ProblemType"]["TLUB"]
if (
Comment thread
b-shi marked this conversation as resolved.
state["EnableMatrixInstruction"] and not state["ExpandPointerSwap"] and
state["DepthU"] == state["MatrixInstK"] and state["PrefetchGlobalRead"] and not state["1LDSBuffer"]
and (state["MIWaveTile"][0] > 2 and state["MIWaveTile"][1] > 2)
and (state["MIWaveTile"][0] % 2 == 0 and state["MIWaveTile"][1] % 2 == 0)
and (sizeDataTypeA == sizeDataType) and (sizeDataTypeB == sizeDataType)
and ((TLUA == False or state["enableLDSTrA"] or sizeDataTypeA >= 4) and (TLUB == False or state["enableLDSTrB"] or sizeDataTypeB >= 4) )
):
state["ForceUnrollSubIter"] = True
state["numSubTiles"] = 2
state["PrefetchLocalRead"] = 0 if state["ClusterLocalRead"] == 0 else state["PrefetchLocalRead"]
else:
state["ForceUnrollSubIter"] = False
state["numSubTiles"] = 1

# Check if CMS is available for this solution
hasCMS,_ = hasCustomSchedule(state)
state["UseCustomMainLoopSchedule"] = hasCMS

# 0: Normal mode. Hardware applies all of the normal data dependency checks
# 1: Full expert mode (not suppoeted yet). Disable hardware checks against: VA_VDST, VA_SDST, VA_SSRC, VA_VCC, VM_VSRC and SA_SDST.
# 2: Disable only VA_VDST and VM_VSRC checks.
Expand Down Expand Up @@ -3307,11 +3338,6 @@ def calcEpilogueTurns(factorDims: List) -> int:
#print("Force to Disable PreloadKernArgs since this hipcc version doesn't support",)
state["PreloadKernArgs"] = 0

hasCMS,_ = hasCustomSchedule(state)
state["UseCustomMainLoopSchedule"] = hasCMS

state["InternalSupportParams"]["UseSFC"] = (len(state["SpaceFillingAlgo"]) > 0)

########################################
@ staticmethod
def getParametersIndented(state, indent):
Expand Down
Loading