Skip to content
Closed
Show file tree
Hide file tree
Changes from all commits
Commits
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
Original file line number Diff line number Diff line change
Expand Up @@ -94,14 +94,11 @@ def __call__(self, writer, kernel, bufferIdx, iui, epsi, tP):
# tP["bpe"], \
# paramList[-1]))
# paramTuple = tuple(paramList)
num = paramList[0] //65536
num = paramList[0] // 65536
paramList[0] = paramList[0] - num * 65536
srcVgpr=vgpr("LocalReadAddr%s+%d"%(tc,num))
srcAddr = vgpr("LocalReadAddr%s+%d"%(tc, num))

if numOffsets == 1:
addrIdx = paramList[0] // 65536
srcAddr=vgpr("LocalReadAddr%s+%u"%(tc, addrIdx))
paramList[0] -= addrIdx * 65536
ds = DSModifiers(na=1, offset=paramList[0])
if numOffsets == 2:
ds = DSModifiers(na=2, offset0=paramList[0], offset1=paramList[1])
Expand Down
8 changes: 7 additions & 1 deletion projects/hipblaslt/tensilelite/Tensile/Components/SIA.py
Original file line number Diff line number Diff line change
Expand Up @@ -535,7 +535,13 @@ def noSchedGlobalRead(writer, kernel, globalReadIncACode, globalReadIncBCode):
tdmLoadIter = min(localWriteEndIter + 1, kernel["LoopIters"] - 1)

if kernel["PrefetchGlobalRead"] == 2:
imod = writer.codes.perIterGlobalRead[0].add(Module())
# SIA0 does not schedule GR/LW instruction-by-instruction. If global reads
# are emitted at iter 0, they clobber vgprG2L* before the later local-write
# iteration stores the previous prefetch to LDS. Place non-TDM reads in the
# local-write iteration; SIA0 sub-iteration order emits localWriteCode before
# globalReadCode, preserving the G2L value until it is consumed.
grIter = localWriteEndIter if kernel["_ScheduleIterAlg"] == 0 and not kernel["NoLdsWriteCode"] and not tdmDeferLoad else 0
imod = writer.codes.perIterGlobalRead[grIter].add(Module())
imod.addComment1("Global Read IncA")
imod.add(globalReadIncACode)
imod.addComment1("Global Read IncB")
Expand Down
22 changes: 19 additions & 3 deletions projects/hipblaslt/tensilelite/Tensile/KernelWriter.py
Original file line number Diff line number Diff line change
Expand Up @@ -2527,6 +2527,7 @@ def calculateRangeAndUpdateCounter(itemCounter, writeCounters, length):
localWrites += skipPreIterLW
dscnt += localWrites
else:
sawLocalRead = False
for item in list(iterCode.items()):
localReads = countWeightedLocalRead(item)
localWrites = countWeightedLocalWrite(item)
Expand All @@ -2545,7 +2546,13 @@ def calculateRangeAndUpdateCounter(itemCounter, writeCounters, length):
# we need to wait for all preceding reads before the macs
# so only opportunity for optimization is if the writes are at the end
if localReads:
sawLocalRead = True
dscnt = 0 # reset to wait for all reads
elif sawLocalRead:
# SIA0 + PGR>=2 emits localReadCode and localWriteCode as separate
# modules. Once local reads require a full drain, a following
# write-only module must not downgrade dscnt to localWrites-only.
pass
else:
dscnt = localWrites # this only survives if writes are at the end

Expand Down Expand Up @@ -5781,6 +5788,14 @@ def _restoreNtabState():
if not kernel["enableTDMB"]:
tempLWCodeModB = self.localWriteDo(kernel, tensorParametersB)
module.add(tempLWCodeModB)
# recalcLocalReadAddressesAB() below resets numReadsIterCoalesced{A,B} to 1,
# so capture the wider-local-read state first.
tdm = kernel["enableTDMA"] and kernel["enableTDMB"]
tdmTailWasWiderLR = (self.states.numReadsIterCoalescedA > 1 or
self.states.numReadsIterCoalescedB > 1)
# TDM tail may keep using whichever LDS buffer the swap parity left it in
# (no forced buffer 0), unless wider local read needs the offset recomputed.
needResetLROffsets = not kernel["1LDSBuffer"] and (not tdm or tdmTailWasWiderLR)
# change local read policy from wider local read to one unit of K at a time
# DirectToVgpr case, use original wider local read instead of recalculating local read address
if not (kernel["DirectToVgprA"] or kernel["DirectToVgprB"]):
Expand All @@ -5804,9 +5819,10 @@ def _restoreNtabState():

# tail: re-init local read addresses
if kernel["PrefetchGlobalRead"]:
# StreamK tail workgroups may start from a partial tile slice, so SIA0
# also needs the LRO reset before tail MACs.
if kernel["_ScheduleIterAlg"] != 0 or kernel["StreamK"]:
# Main-loop pointer swaps can leave LROs in the Blk half. Reset before
# tail MACs for every scheduler; localReadResetOffsets is empty for
# 1LDSBuffer / DirectToVgpr cases.
if needResetLROffsets:
module.addComment1("Tail: local read reset offsets a")
module.add(self.localReadResetOffsets(kernel, tensorParametersA))
if kernel["ProblemType"]["MXBlockA"]:
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -5,7 +5,6 @@ GlobalParameters:
SyncsPerBenchmark: 0
NumElementsToValidate: -1
Architecture: gfx1250
DeviceLDS: 327680
PrintSolutionRejectionReason: True

BenchmarkProblems:
Expand Down Expand Up @@ -39,12 +38,11 @@ BenchmarkProblems:
- [16, 16, 32, 1, 1, 2, 2, 2, 2 ] #64x64x512
- [16, 16, 32, 1, 1, 8, 8, 2, 2 ] #256x256x64
- DepthU: [64,512]
- MaxLDS: [327680]
- TransposeLDS: [-1]
- LdsPadA: [-1]
- LdsPadB: [-1]
- StaggerU: [0]
- ScheduleIterAlg: [3]
- ScheduleIterAlg: [0]
- PrefetchLocalRead: [1]
- PrefetchGlobalRead: [2]
- ClusterLocalRead: [1]
Expand All @@ -53,7 +51,7 @@ BenchmarkProblems:
- GlobalSplitUAlgorithm: [MultipleBuffer]
- GlobalReadVectorWidthA: [-1]
- GlobalReadVectorWidthB: [-1]
- 1LDSBuffer: [-1]
- 1LDSBuffer: [0]
- WorkGroupMapping: [1]
- StoreVectorWidth: [-1]
- DirectToVgprA: [False]
Expand Down Expand Up @@ -102,7 +100,7 @@ BenchmarkProblems:
- LdsPadA: [-1]
- LdsPadB: [-1]
- StaggerU: [0]
- ScheduleIterAlg: [3]
- ScheduleIterAlg: [0]
- PrefetchLocalRead: [1]
- PrefetchGlobalRead: [2]
- ClusterLocalRead: [1]
Expand All @@ -111,7 +109,7 @@ BenchmarkProblems:
- GlobalSplitUAlgorithm: [MultipleBuffer]
- GlobalReadVectorWidthA: [-1]
- GlobalReadVectorWidthB: [-1]
- 1LDSBuffer: [-1]
- 1LDSBuffer: [0]
- WorkGroupMapping: [1]
- StoreVectorWidth: [-1]
- DirectToVgprA: [False]
Expand Down Expand Up @@ -160,7 +158,7 @@ BenchmarkProblems:
- LdsPadA: [-1]
- LdsPadB: [-1]
- StaggerU: [0]
- ScheduleIterAlg: [3]
- ScheduleIterAlg: [0]
- PrefetchLocalRead: [1]
- PrefetchGlobalRead: [2]
- ClusterLocalRead: [1]
Expand All @@ -169,7 +167,7 @@ BenchmarkProblems:
- GlobalSplitUAlgorithm: [MultipleBuffer]
- GlobalReadVectorWidthA: [-1]
- GlobalReadVectorWidthB: [-1]
- 1LDSBuffer: [-1]
- 1LDSBuffer: [0]
- WorkGroupMapping: [1]
- StoreVectorWidth: [-1]
- DirectToVgprA: [False]
Expand Down Expand Up @@ -218,7 +216,7 @@ BenchmarkProblems:
- LdsPadA: [-1]
- LdsPadB: [-1]
- StaggerU: [0]
- ScheduleIterAlg: [3]
- ScheduleIterAlg: [0]
- PrefetchLocalRead: [1]
- PrefetchGlobalRead: [2]
- ClusterLocalRead: [1]
Expand All @@ -227,7 +225,7 @@ BenchmarkProblems:
- GlobalSplitUAlgorithm: [MultipleBuffer]
- GlobalReadVectorWidthA: [-1]
- GlobalReadVectorWidthB: [-1]
- 1LDSBuffer: [-1]
- 1LDSBuffer: [0]
- WorkGroupMapping: [1]
- StoreVectorWidth: [-1]
- DirectToVgprA: [False]
Expand Down Expand Up @@ -271,7 +269,7 @@ BenchmarkProblems:
- DepthU: [128,512]
- UseSgprForGRO: [0]
- LocalReadVectorWidth: [32]
- ScheduleIterAlg: [3]
- ScheduleIterAlg: [0]
- PrefetchGlobalRead: [2]
- PrefetchLocalRead: [1]
- VectorWidthA: [2]
Expand Down Expand Up @@ -332,15 +330,15 @@ BenchmarkProblems:
- LdsPadB: [-1]
- LdsPadMetadata: [0]
- StaggerU: [0]
- ScheduleIterAlg: [1] ########SIA2 and SIA3 only support MatrixInstruction
- ScheduleIterAlg: [0]
- PrefetchLocalRead: [1] #
- PrefetchGlobalRead: [2] #
- StoreRemapVectorWidth: [0]
- GlobalSplitU: [1]
- GlobalSplitUAlgorithm: [MultipleBuffer]
- GlobalReadVectorWidthA: [1] #
- GlobalReadVectorWidthB: [1] #
- 1LDSBuffer: [-1]
- 1LDSBuffer: [0]
- DirectToVgprSparseMetadata: [0]
- WorkGroupMapping: [1]
- StoreVectorWidth: [-1]
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -59,7 +59,7 @@ BenchmarkProblems:
- DepthU: [128, 256]
# - UseSgprForGRO: [0, 1]
- LocalReadVectorWidth: [-1]
- ScheduleIterAlg: [3]
- ScheduleIterAlg: [0]
- PrefetchGlobalRead: [2]
- PrefetchLocalRead: [1]
- VectorWidthA: [1]
Expand All @@ -75,7 +75,7 @@ BenchmarkProblems:
- LdsPadB: [0]
- WaveSeparateGlobalReadA: [0]
- WaveSeparateGlobalReadB: [0]
- 1LDSBuffer: [1]
- 1LDSBuffer: [0]
- GlobalReadPerMfma: [1]
- LocalWritePerMfma: [-1]
- StoreVectorWidth: [-1]
Expand Down Expand Up @@ -121,7 +121,7 @@ BenchmarkProblems:
- DepthU: [128, 256]
# - UseSgprForGRO: [0, 1]
- LocalReadVectorWidth: [-1]
- ScheduleIterAlg: [3]
- ScheduleIterAlg: [0]
- PrefetchGlobalRead: [2]
- PrefetchLocalRead: [1]
- VectorWidthA: [1]
Expand All @@ -137,7 +137,7 @@ BenchmarkProblems:
- LdsPadB: [0]
- WaveSeparateGlobalReadA: [0]
- WaveSeparateGlobalReadB: [0]
- 1LDSBuffer: [1]
- 1LDSBuffer: [0]
- GlobalReadPerMfma: [1]
- LocalWritePerMfma: [-1]
- StoreVectorWidth: [-1]
Expand Down Expand Up @@ -183,7 +183,7 @@ BenchmarkProblems:
- DepthU: [128, 256]
# - UseSgprForGRO: [0, 1]
- LocalReadVectorWidth: [-1]
- ScheduleIterAlg: [3]
- ScheduleIterAlg: [0]
- PrefetchGlobalRead: [2]
- PrefetchLocalRead: [1]
- VectorWidthA: [1]
Expand All @@ -199,7 +199,7 @@ BenchmarkProblems:
- LdsPadB: [0]
- WaveSeparateGlobalReadA: [0]
- WaveSeparateGlobalReadB: [0]
- 1LDSBuffer: [1]
- 1LDSBuffer: [0]
- GlobalReadPerMfma: [1]
- LocalWritePerMfma: [-1]
- StoreVectorWidth: [-1]
Expand Down Expand Up @@ -245,7 +245,7 @@ BenchmarkProblems:
- DepthU: [128, 256]
# - UseSgprForGRO: [0, 1]
- LocalReadVectorWidth: [-1]
- ScheduleIterAlg: [3]
- ScheduleIterAlg: [0]
- PrefetchGlobalRead: [2]
- PrefetchLocalRead: [1]
- VectorWidthA: [1]
Expand All @@ -261,7 +261,7 @@ BenchmarkProblems:
- LdsPadB: [0]
- WaveSeparateGlobalReadA: [0]
- WaveSeparateGlobalReadB: [0]
- 1LDSBuffer: [1]
- 1LDSBuffer: [0]
- GlobalReadPerMfma: [1]
- LocalWritePerMfma: [-1]
- StoreVectorWidth: [-1]
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -61,7 +61,7 @@ BenchmarkProblems:
- WorkGroupMapping: [1]
- WorkGroupMappingXCC: [1]
- LocalReadVectorWidth: [16]
- ScheduleIterAlg: [3]
- ScheduleIterAlg: [0]
- ExpandPointerSwap: [0]
- TransposeLDS: [1]
- LdsBlockSizePerPadA: [-1]
Expand All @@ -72,7 +72,7 @@ BenchmarkProblems:
- GlobalSplitUAlgorithm: ["MultipleBuffer"]
- SourceSwap: [1]
- ActivationFuncCall: [true]
- 1LDSBuffer: [1]
- 1LDSBuffer: [0]
- StaggerU: [32]
- StaggerUMapping: [0]
- StaggerUStride: [256]
Expand Down Expand Up @@ -128,7 +128,7 @@ BenchmarkProblems:
- WorkGroupMapping: [1]
- WorkGroupMappingXCC: [1]
- LocalReadVectorWidth: [16]
- ScheduleIterAlg: [3]
- ScheduleIterAlg: [0]
- ExpandPointerSwap: [0]
- TransposeLDS: [1]
- LdsBlockSizePerPadA: [-1]
Expand All @@ -139,7 +139,7 @@ BenchmarkProblems:
- GlobalSplitUAlgorithm: ["MultipleBuffer"]
- SourceSwap: [1]
- ActivationFuncCall: [true]
- 1LDSBuffer: [1]
- 1LDSBuffer: [0]
- StaggerU: [32]
- StaggerUMapping: [0]
- StaggerUStride: [256]
Expand Down Expand Up @@ -194,7 +194,7 @@ BenchmarkProblems:
- WorkGroupMapping: [1]
- WorkGroupMappingXCC: [1]
- LocalReadVectorWidth: [16]
- ScheduleIterAlg: [3]
- ScheduleIterAlg: [0]
- ExpandPointerSwap: [0]
- LdsBlockSizePerPadA: [-1]
- LdsBlockSizePerPadB: [-1]
Expand All @@ -203,7 +203,7 @@ BenchmarkProblems:
- GlobalSplitU: [1]
- GlobalSplitUAlgorithm: ["MultipleBuffer"]
- ActivationFuncCall: [true]
- 1LDSBuffer: [1]
- 1LDSBuffer: [0]
- StaggerU: [32]
- StaggerUMapping: [0]
- StaggerUStride: [256]
Expand Down Expand Up @@ -259,7 +259,7 @@ BenchmarkProblems:
- WorkGroupMapping: [1]
- WorkGroupMappingXCC: [1]
- LocalReadVectorWidth: [16]
- ScheduleIterAlg: [3]
- ScheduleIterAlg: [0]
- ExpandPointerSwap: [0]
- TransposeLDS: [1]
- LdsBlockSizePerPadA: [-1]
Expand All @@ -270,7 +270,7 @@ BenchmarkProblems:
- GlobalSplitUAlgorithm: ["MultipleBuffer"]
- SourceSwap: [1]
- ActivationFuncCall: [true]
- 1LDSBuffer: [1]
- 1LDSBuffer: [0]
- StaggerU: [32]
- StaggerUMapping: [0]
- StaggerUStride: [256]
Expand Down
Loading
Loading