diff --git a/projects/hipblaslt/tensilelite/Tensile/Components/LocalRead.py b/projects/hipblaslt/tensilelite/Tensile/Components/LocalRead.py index 8f9e52e8c01f..e9d150b44917 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Components/LocalRead.py +++ b/projects/hipblaslt/tensilelite/Tensile/Components/LocalRead.py @@ -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]) diff --git a/projects/hipblaslt/tensilelite/Tensile/Components/SIA.py b/projects/hipblaslt/tensilelite/Tensile/Components/SIA.py index c43ed5b663f0..b5c1671253cf 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Components/SIA.py +++ b/projects/hipblaslt/tensilelite/Tensile/Components/SIA.py @@ -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") diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py index 236446991687..6fc21f3783fa 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py @@ -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) @@ -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 @@ -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"]): @@ -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"]: diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/1024_vgpr_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/1024_vgpr_gfx1250.yaml index 2d354220f8e5..b8ffd6c8e9b2 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/1024_vgpr_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/1024_vgpr_gfx1250.yaml @@ -5,7 +5,6 @@ GlobalParameters: SyncsPerBenchmark: 0 NumElementsToValidate: -1 Architecture: gfx1250 - DeviceLDS: 327680 PrintSolutionRejectionReason: True BenchmarkProblems: @@ -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] @@ -53,7 +51,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] - DirectToVgprA: [False] @@ -102,7 +100,7 @@ BenchmarkProblems: - LdsPadA: [-1] - LdsPadB: [-1] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - ClusterLocalRead: [1] @@ -111,7 +109,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] - DirectToVgprA: [False] @@ -160,7 +158,7 @@ BenchmarkProblems: - LdsPadA: [-1] - LdsPadB: [-1] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - ClusterLocalRead: [1] @@ -169,7 +167,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] - DirectToVgprA: [False] @@ -218,7 +216,7 @@ BenchmarkProblems: - LdsPadA: [-1] - LdsPadB: [-1] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - ClusterLocalRead: [1] @@ -227,7 +225,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] - DirectToVgprA: [False] @@ -271,7 +269,7 @@ BenchmarkProblems: - DepthU: [128,512] - UseSgprForGRO: [0] - LocalReadVectorWidth: [32] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [2] @@ -332,7 +330,7 @@ BenchmarkProblems: - LdsPadB: [-1] - LdsPadMetadata: [0] - StaggerU: [0] - - ScheduleIterAlg: [1] ########SIA2 and SIA3 only support MatrixInstruction + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] # - PrefetchGlobalRead: [2] # - StoreRemapVectorWidth: [0] @@ -340,7 +338,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [1] # - GlobalReadVectorWidthB: [1] # - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/b6f4ss_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/b6f4ss_gfx1250.yaml index 3b2d1c0ee31f..998cd4e3e42f 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/b6f4ss_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/b6f4ss_gfx1250.yaml @@ -59,7 +59,7 @@ BenchmarkProblems: - DepthU: [128, 256] # - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1] @@ -75,7 +75,7 @@ BenchmarkProblems: - LdsPadB: [0] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] @@ -121,7 +121,7 @@ BenchmarkProblems: - DepthU: [128, 256] # - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1] @@ -137,7 +137,7 @@ BenchmarkProblems: - LdsPadB: [0] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] @@ -183,7 +183,7 @@ BenchmarkProblems: - DepthU: [128, 256] # - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1] @@ -199,7 +199,7 @@ BenchmarkProblems: - LdsPadB: [0] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] @@ -245,7 +245,7 @@ BenchmarkProblems: - DepthU: [128, 256] # - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1] @@ -261,7 +261,7 @@ BenchmarkProblems: - LdsPadB: [0] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/b8b8s_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/b8b8s_gfx1250.yaml index 6ebf0f854035..ce27edf8c555 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/b8b8s_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/b8b8s_gfx1250.yaml @@ -61,7 +61,7 @@ BenchmarkProblems: - WorkGroupMapping: [1] - WorkGroupMappingXCC: [1] - LocalReadVectorWidth: [16] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [1] - LdsBlockSizePerPadA: [-1] @@ -72,7 +72,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [1] - ActivationFuncCall: [true] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - StaggerU: [32] - StaggerUMapping: [0] - StaggerUStride: [256] @@ -128,7 +128,7 @@ BenchmarkProblems: - WorkGroupMapping: [1] - WorkGroupMappingXCC: [1] - LocalReadVectorWidth: [16] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [1] - LdsBlockSizePerPadA: [-1] @@ -139,7 +139,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [1] - ActivationFuncCall: [true] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - StaggerU: [32] - StaggerUMapping: [0] - StaggerUStride: [256] @@ -194,7 +194,7 @@ BenchmarkProblems: - WorkGroupMapping: [1] - WorkGroupMappingXCC: [1] - LocalReadVectorWidth: [16] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] @@ -203,7 +203,7 @@ BenchmarkProblems: - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - ActivationFuncCall: [true] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - StaggerU: [32] - StaggerUMapping: [0] - StaggerUStride: [256] @@ -259,7 +259,7 @@ BenchmarkProblems: - WorkGroupMapping: [1] - WorkGroupMappingXCC: [1] - LocalReadVectorWidth: [16] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [1] - LdsBlockSizePerPadA: [-1] @@ -270,7 +270,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [1] - ActivationFuncCall: [true] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - StaggerU: [32] - StaggerUMapping: [0] - StaggerUStride: [256] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/bf16_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/bf16_gfx1250.yaml index 3e83f87d4e90..1a3184c71db1 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/bf16_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/bf16_gfx1250.yaml @@ -76,7 +76,7 @@ BenchmarkProblems: - LdsPadB: [-1] - LdsPadMetadata: [0] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [0] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -84,7 +84,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - - 1LDSBuffer: [-1, 1] + - 1LDSBuffer: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] @@ -150,7 +150,7 @@ BenchmarkProblems: - LdsPadB: [-1] - LdsPadMetadata: [0] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [0] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -158,7 +158,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - - 1LDSBuffer: [-1, 1] + - 1LDSBuffer: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] @@ -224,7 +224,7 @@ BenchmarkProblems: - LdsPadB: [-1] - LdsPadMetadata: [0] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [0] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -232,7 +232,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - - 1LDSBuffer: [-1, 1] + - 1LDSBuffer: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] @@ -298,7 +298,7 @@ BenchmarkProblems: - LdsPadB: [-1] - LdsPadMetadata: [0] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [0] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -306,7 +306,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - - 1LDSBuffer: [-1, 1] + - 1LDSBuffer: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/bf6_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/bf6_gfx1250.yaml index 373fe5c1cccc..4313ff41afc6 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/bf6_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/bf6_gfx1250.yaml @@ -57,7 +57,7 @@ BenchmarkProblems: - DepthU: [128, 256] # - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [32] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1, 2] @@ -119,7 +119,7 @@ BenchmarkProblems: - LDSTrInst: [True] - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [32] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1] @@ -134,7 +134,7 @@ BenchmarkProblems: - LdsPadB: [0] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] @@ -182,7 +182,7 @@ BenchmarkProblems: - LDSTrInst: [True] - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [32] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1] @@ -197,7 +197,7 @@ BenchmarkProblems: - LdsPadB: [0] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] @@ -245,7 +245,7 @@ BenchmarkProblems: - LDSTrInst: [True] - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [32] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1] @@ -260,7 +260,7 @@ BenchmarkProblems: - LdsPadB: [0] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f32_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f32_gfx1250.yaml index 049d9094b9aa..8336d82d62b0 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f32_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f32_gfx1250.yaml @@ -65,7 +65,7 @@ BenchmarkProblems: - LdsPadB: [-1] - LdsPadMetadata: [0] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] # - PrefetchGlobalRead: [2] # - StoreRemapVectorWidth: [0] @@ -73,7 +73,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBufferSingleKernel] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] @@ -137,7 +137,7 @@ BenchmarkProblems: - LdsPadB: [-1] - LdsPadMetadata: [0] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] # - PrefetchGlobalRead: [2] # - StoreRemapVectorWidth: [0] @@ -145,7 +145,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] @@ -208,7 +208,7 @@ BenchmarkProblems: - LdsPadB: [-1] - LdsPadMetadata: [0] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] # - PrefetchGlobalRead: [2] # - StoreRemapVectorWidth: [0] @@ -216,7 +216,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBufferSingleKernel] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] @@ -279,7 +279,7 @@ BenchmarkProblems: - LdsPadB: [-1] - LdsPadMetadata: [0] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] # - PrefetchGlobalRead: [2] # - StoreRemapVectorWidth: [0] @@ -287,7 +287,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f4_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f4_gfx1250.yaml index bc0c875d7ac0..62fdcba3d003 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f4_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f4_gfx1250.yaml @@ -58,7 +58,7 @@ BenchmarkProblems: # - [16, 16, 128, 1, 1, 1, 4, 4, 1] - DepthU: [128, 256] - LocalReadVectorWidth: [32] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1, 2] @@ -121,7 +121,7 @@ BenchmarkProblems: - DepthU: [128, 256] - LDSTrInst: [True] - LocalReadVectorWidth: [32] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1, 2] @@ -136,7 +136,7 @@ BenchmarkProblems: - LdsPadB: [-1] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] @@ -184,7 +184,7 @@ BenchmarkProblems: - DepthU: [128, 256] - LDSTrInst: [True] - LocalReadVectorWidth: [32] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1, 2] @@ -199,7 +199,7 @@ BenchmarkProblems: - LdsPadB: [-1] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] @@ -247,7 +247,7 @@ BenchmarkProblems: - DepthU: [128, 256] - LDSTrInst: [True] - LocalReadVectorWidth: [32] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1, 2] @@ -262,7 +262,7 @@ BenchmarkProblems: - LdsPadB: [-1] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f4b8ss_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f4b8ss_gfx1250.yaml index 518ad418b1bb..35e2a3582195 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f4b8ss_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f4b8ss_gfx1250.yaml @@ -66,7 +66,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True, False] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -77,7 +77,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0,1] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] @@ -133,7 +133,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -144,7 +144,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0,1] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] @@ -200,7 +200,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -211,7 +211,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0,1] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] @@ -267,7 +267,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -278,7 +278,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0,1] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f64_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f64_gfx1250.yaml index b62df1ea599b..dd051699c278 100755 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f64_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f64_gfx1250.yaml @@ -63,7 +63,7 @@ BenchmarkProblems: - LdsPadB: [-1] - LdsPadMetadata: [0] - StaggerU: [0] - - ScheduleIterAlg: [1] ########SIA2 and SIA3 only support MatrixInstruction + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] # - PrefetchGlobalRead: [2] # - StoreRemapVectorWidth: [0] @@ -71,7 +71,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [1] # - GlobalReadVectorWidthB: [1] # - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] @@ -128,7 +128,7 @@ BenchmarkProblems: - LdsPadB: [-1] - LdsPadMetadata: [0] - StaggerU: [0] - - ScheduleIterAlg: [1] ########SIA2 and SIA3 only support MatrixInstruction + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] # - PrefetchGlobalRead: [2] # - StoreRemapVectorWidth: [0] @@ -136,7 +136,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [1] # - GlobalReadVectorWidthB: [1] # - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] @@ -193,7 +193,7 @@ BenchmarkProblems: - LdsPadB: [-1] - LdsPadMetadata: [0] - StaggerU: [0] - - ScheduleIterAlg: [1] ########SIA2 and SIA3 only support MatrixInstruction + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] # - PrefetchGlobalRead: [2] # - StoreRemapVectorWidth: [0] @@ -201,7 +201,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [1] # - GlobalReadVectorWidthB: [1] # - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] @@ -258,7 +258,7 @@ BenchmarkProblems: - LdsPadB: [-1] - LdsPadMetadata: [0] - StaggerU: [0] - - ScheduleIterAlg: [1] ########SIA2 and SIA3 only support MatrixInstruction + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] # - PrefetchGlobalRead: [2] # - StoreRemapVectorWidth: [0] @@ -266,7 +266,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [1] # - GlobalReadVectorWidthB: [1] # - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f6_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f6_gfx1250.yaml index 333ff06bd244..63ab7de50349 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f6_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f6_gfx1250.yaml @@ -57,7 +57,7 @@ BenchmarkProblems: - DepthU: [128, 256] - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [32] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1, 2] @@ -119,7 +119,7 @@ BenchmarkProblems: - LDSTrInst: [True] # - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [32] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1] @@ -134,7 +134,7 @@ BenchmarkProblems: - LdsPadB: [0] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] @@ -182,7 +182,7 @@ BenchmarkProblems: - LDSTrInst: [True] - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [32] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1] @@ -197,7 +197,7 @@ BenchmarkProblems: - LdsPadB: [0] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] @@ -245,7 +245,7 @@ BenchmarkProblems: - LDSTrInst: [True] - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [32] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1] @@ -260,7 +260,7 @@ BenchmarkProblems: - LdsPadB: [0] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f6b6ss_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f6b6ss_gfx1250.yaml index 71e50cb69c90..88514eadf2ee 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f6b6ss_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f6b6ss_gfx1250.yaml @@ -59,7 +59,7 @@ BenchmarkProblems: - DepthU: [128, 256] - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [32] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1, 2] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f6b8ss_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f6b8ss_gfx1250.yaml index 693e66e2ff2d..bafa81c1a664 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f6b8ss_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f6b8ss_gfx1250.yaml @@ -59,7 +59,7 @@ BenchmarkProblems: - DepthU: [128, 256] # - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1] @@ -75,7 +75,7 @@ BenchmarkProblems: - LdsPadB: [0] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] @@ -121,7 +121,7 @@ BenchmarkProblems: - DepthU: [128, 256] # - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1] @@ -137,7 +137,7 @@ BenchmarkProblems: - LdsPadB: [0] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] @@ -183,7 +183,7 @@ BenchmarkProblems: - DepthU: [128, 256] # - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1] @@ -199,7 +199,7 @@ BenchmarkProblems: - LdsPadB: [0] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] @@ -245,7 +245,7 @@ BenchmarkProblems: - DepthU: [128, 256] # - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1] @@ -261,7 +261,7 @@ BenchmarkProblems: - LdsPadB: [0] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8b6ss_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8b6ss_gfx1250.yaml index 8f171d2d8d34..740160c84a7f 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8b6ss_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8b6ss_gfx1250.yaml @@ -59,7 +59,7 @@ BenchmarkProblems: - DepthU: [128, 256] # - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1] @@ -75,7 +75,7 @@ BenchmarkProblems: - LdsPadB: [0] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] @@ -121,7 +121,7 @@ BenchmarkProblems: - DepthU: [128, 256] # - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1] @@ -137,7 +137,7 @@ BenchmarkProblems: - LdsPadB: [0] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] @@ -183,7 +183,7 @@ BenchmarkProblems: - DepthU: [128, 256] # - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1] @@ -199,7 +199,7 @@ BenchmarkProblems: - LdsPadB: [0] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] @@ -245,7 +245,7 @@ BenchmarkProblems: - DepthU: [128, 256] # - UseSgprForGRO: [0, 1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchGlobalRead: [2] - PrefetchLocalRead: [1] - VectorWidthA: [1] @@ -261,7 +261,7 @@ BenchmarkProblems: - LdsPadB: [0] - WaveSeparateGlobalReadA: [0] - WaveSeparateGlobalReadB: [0] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - GlobalReadPerMfma: [1] - LocalWritePerMfma: [-1] - StoreVectorWidth: [-1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8b8ss_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8b8ss_gfx1250.yaml index 1796c61a270d..14d41da2a33a 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8b8ss_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8b8ss_gfx1250.yaml @@ -64,7 +64,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True, False] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -75,7 +75,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0,1] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] @@ -130,7 +130,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True, False] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -141,7 +141,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0,1] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] @@ -196,7 +196,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True, False] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -207,7 +207,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0,1] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] @@ -262,7 +262,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True, False] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -273,7 +273,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0,1] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8f4ss_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8f4ss_gfx1250.yaml index 60627997bc79..ac507ec816b5 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8f4ss_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8f4ss_gfx1250.yaml @@ -66,7 +66,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True, False] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -77,7 +77,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0,1] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] @@ -133,7 +133,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -144,7 +144,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0,1] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] @@ -200,7 +200,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -211,7 +211,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0,1] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] @@ -267,7 +267,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -278,7 +278,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0,1] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8f8s_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8f8s_gfx1250.yaml index 3b96feb078bb..3f5797185ffb 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8f8s_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8f8s_gfx1250.yaml @@ -60,7 +60,7 @@ BenchmarkProblems: - WorkGroupMapping: [1] - WorkGroupMappingXCC: [1] - LocalReadVectorWidth: [16] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [1] # - LdsBlockSizePerPadA: [-1] @@ -71,7 +71,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [1] - ActivationFuncCall: [true] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - StaggerU: [32] - StaggerUMapping: [0] - StaggerUStride: [256] @@ -128,7 +128,7 @@ BenchmarkProblems: - WorkGroupMapping: [1] - WorkGroupMappingXCC: [1] - LocalReadVectorWidth: [16] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [1] - LdsBlockSizePerPadA: [-1] @@ -139,7 +139,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [1] - ActivationFuncCall: [true] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - StaggerU: [32] - StaggerUMapping: [0] - StaggerUStride: [256] @@ -194,7 +194,7 @@ BenchmarkProblems: - WorkGroupMapping: [1] - WorkGroupMappingXCC: [1] - LocalReadVectorWidth: [16] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [0] - LdsBlockSizePerPadA: [-1] @@ -205,7 +205,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] - ActivationFuncCall: [true] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - StaggerU: [32] - StaggerUMapping: [0] - StaggerUStride: [-1] @@ -260,7 +260,7 @@ BenchmarkProblems: - WorkGroupMapping: [1] - WorkGroupMappingXCC: [1] - LocalReadVectorWidth: [16] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [1] - LdsBlockSizePerPadA: [-1] @@ -271,7 +271,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [1] - ActivationFuncCall: [true] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - StaggerU: [32] - StaggerUMapping: [0] - StaggerUStride: [-1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8f8s_sr_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8f8s_sr_gfx1250.yaml index 9b7f6903ae73..02fdb991ebcb 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8f8s_sr_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/f8f8s_sr_gfx1250.yaml @@ -61,7 +61,7 @@ BenchmarkProblems: - WorkGroupMapping: [1] - WorkGroupMappingXCC: [1] - LocalReadVectorWidth: [16] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [1] # - LdsBlockSizePerPadA: [-1] @@ -72,7 +72,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [1] - ActivationFuncCall: [true] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - StaggerU: [32] - StaggerUMapping: [0] - StaggerUStride: [256] @@ -129,7 +129,7 @@ BenchmarkProblems: - WorkGroupMapping: [1] - WorkGroupMappingXCC: [1] - LocalReadVectorWidth: [16] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [1] # - LdsBlockSizePerPadA: [-1] @@ -140,7 +140,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [1] - ActivationFuncCall: [true] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - StaggerU: [32] - StaggerUMapping: [0] - StaggerUStride: [256] @@ -197,7 +197,7 @@ BenchmarkProblems: - WorkGroupMapping: [1] - WorkGroupMappingXCC: [1] - LocalReadVectorWidth: [16] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [0] - LdsBlockSizePerPadA: [-1] @@ -208,7 +208,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] - ActivationFuncCall: [true] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - StaggerU: [32] - StaggerUMapping: [0] - WavefrontSize: [32] @@ -264,7 +264,7 @@ BenchmarkProblems: - WorkGroupMapping: [1] - WorkGroupMappingXCC: [1] - LocalReadVectorWidth: [16] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [1] # - LdsBlockSizePerPadA: [-1] @@ -275,7 +275,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [1] - ActivationFuncCall: [true] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - StaggerU: [32] - StaggerUMapping: [0] - StaggerUStride: [256] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/fp16_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/fp16_gfx1250.yaml index 3715292d76b8..223eb4ea11ff 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/fp16_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/fp16_gfx1250.yaml @@ -57,7 +57,7 @@ BenchmarkProblems: - LdsPadB: [-1] - LdsPadMetadata: [0] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [0] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -65,7 +65,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - - 1LDSBuffer: [-1, 1] + - 1LDSBuffer: [0] - VectorWidthA: [1, 2] - SourceSwap: [0, 1] - DirectToVgprSparseMetadata: [0] @@ -127,7 +127,7 @@ BenchmarkProblems: - LdsPadB: [-1] - LdsPadMetadata: [0] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [0] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -135,7 +135,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - - 1LDSBuffer: [-1, 1] + - 1LDSBuffer: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] @@ -195,7 +195,7 @@ BenchmarkProblems: - LdsPadB: [-1] - LdsPadMetadata: [0] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [0] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -203,7 +203,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - - 1LDSBuffer: [-1, 1] + - 1LDSBuffer: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] @@ -264,7 +264,7 @@ BenchmarkProblems: - LdsPadB: [-1] - LdsPadMetadata: [0] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [0] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -272,7 +272,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - - 1LDSBuffer: [-1, 1] + - 1LDSBuffer: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] - StoreVectorWidth: [-1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/fp16_use_e_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/fp16_use_e_gfx1250.yaml index d11c607aa570..758425b1a281 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/fp16_use_e_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/fp16_use_e_gfx1250.yaml @@ -61,7 +61,7 @@ BenchmarkProblems: - WorkGroupMapping: [1] - WorkGroupMappingXCC: [1] - LocalReadVectorWidth: [8] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [1] - LdsBlockSizePerPadA: [-1] @@ -72,7 +72,7 @@ BenchmarkProblems: - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [1] - ActivationFuncCall: [true] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - StaggerU: [0] - StaggerUMapping: [0] - StaggerUStride: [256] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/fp8_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/fp8_gfx1250.yaml index 6e0b9be72927..f4494faedef6 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/fp8_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/fp8_gfx1250.yaml @@ -64,7 +64,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True, False] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -75,7 +75,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] @@ -127,7 +127,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True, False] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -138,7 +138,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] @@ -189,7 +189,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True, False] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -200,7 +200,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] @@ -251,7 +251,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True, False] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -262,7 +262,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/i8ii_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/i8ii_gfx1250.yaml index 7658127618bc..b7ade69c57ff 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/i8ii_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/i8ii_gfx1250.yaml @@ -70,7 +70,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True, False] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -81,7 +81,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] @@ -142,7 +142,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True, False] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -153,7 +153,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] @@ -214,7 +214,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True, False] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -225,7 +225,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] @@ -286,7 +286,7 @@ BenchmarkProblems: - LdsPadMetadata: [0] - LDSTrInst: [True, False] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] @@ -297,7 +297,7 @@ BenchmarkProblems: - VectorWidthA: [-1] - VectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - 1LDSBuffer: [1] + - 1LDSBuffer: [0] - SourceSwap: [0] - DirectToVgprSparseMetadata: [0] - WorkGroupMapping: [1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/largeLds_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/largeLds_gfx1250.yaml index 521adf8d82e7..6c4bf846130d 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/largeLds_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/largeLds_gfx1250.yaml @@ -42,11 +42,11 @@ BenchmarkProblems: - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - StaggerU: [0] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] - PrefetchGlobalRead: [1,2] - ExpandPointerSwap: [1] - - 1LDSBuffer: [0,1] + - 1LDSBuffer: [0] - ClusterLocalRead: [0] - StoreRemapVectorWidth: [0] - GlobalSplitU: [1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf4_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf4_gfx1250.yaml index 0344f2ba1790..79ae6b041697 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf4_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf4_gfx1250.yaml @@ -26,8 +26,6 @@ GlobalParameters: Architecture: gfx1250 # KeepBuildTmp: True PrintLevel: 2 - MaxLDS: 163840 - DeviceLDS: 163840 BenchmarkProblems: @@ -62,9 +60,7 @@ BenchmarkProblems: - ForceDisableShadowInit: [1] - MatrixInstruction: - [16, 16, 128, 1, 1, 1, 1, 2, 2] - - [16, 16, 128, 1, 1, 1, 2, 2, 1] - - [16, 16, 128, 1, 1, 1, 4, 2, 1] - - [16, 16, 128, 1, 1, 1, 8, 4, 1] + - [16, 16, 128, 1, 1, 1, 4, 4, 1] - [16, 16, 128, 1, 1, 2, 2, 2, 2] - [16, 16, 128, 1, 1, 4, 4, 2, 2] - [16, 16, 128, 1, 1, 8, 8, 2, 2] @@ -135,8 +131,6 @@ BenchmarkProblems: - ForceDisableShadowInit: [1] - MatrixInstruction: - [16, 16, 128, 1, 1, 1, 1, 2, 2] - - [16, 16, 128, 1, 1, 1, 2, 2, 1] - - [16, 16, 128, 1, 1, 1, 4, 2, 1] - [16, 16, 128, 1, 1, 1, 4, 4, 1] - [16, 16, 128, 1, 1, 2, 2, 2, 2] - [16, 16, 128, 1, 1, 4, 4, 2, 2] @@ -207,8 +201,6 @@ BenchmarkProblems: - ForceDisableShadowInit: [1] - MatrixInstruction: - [16, 16, 128, 1, 1, 1, 1, 2, 2] - - [16, 16, 128, 1, 1, 1, 2, 2, 1] - - [16, 16, 128, 1, 1, 1, 4, 2, 1] - [16, 16, 128, 1, 1, 1, 4, 4, 1] - [16, 16, 128, 1, 1, 2, 2, 2, 2] - [16, 16, 128, 1, 1, 4, 4, 2, 2] @@ -278,8 +270,6 @@ BenchmarkProblems: - ForceDisableShadowInit: [1] - MatrixInstruction: - [16, 16, 128, 1, 1, 1, 1, 2, 2] - - [16, 16, 128, 1, 1, 1, 2, 2, 1] - - [16, 16, 128, 1, 1, 1, 4, 2, 1] - [16, 16, 128, 1, 1, 1, 4, 4, 1] - [16, 16, 128, 1, 1, 2, 2, 2, 2] - [16, 16, 128, 1, 1, 4, 4, 2, 2] @@ -349,8 +339,6 @@ BenchmarkProblems: - ForceDisableShadowInit: [1] - MatrixInstruction: - [16, 16, 128, 1, 1, 1, 1, 2, 2] - - [16, 16, 128, 1, 1, 1, 2, 2, 1] - - [16, 16, 128, 1, 1, 1, 4, 2, 1] - [16, 16, 128, 1, 1, 1, 4, 4, 1] - [16, 16, 128, 1, 1, 2, 2, 2, 2] - [16, 16, 128, 1, 1, 4, 4, 2, 2] @@ -420,8 +408,6 @@ BenchmarkProblems: - ForceDisableShadowInit: [1] - MatrixInstruction: - [16, 16, 128, 1, 1, 1, 1, 2, 2] - - [16, 16, 128, 1, 1, 1, 2, 2, 1] - - [16, 16, 128, 1, 1, 1, 4, 2, 1] - [16, 16, 128, 1, 1, 1, 4, 4, 1] - [16, 16, 128, 1, 1, 2, 2, 2, 2] - [16, 16, 128, 1, 1, 4, 4, 2, 2] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf6_tdm_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf6_tdm_gfx1250.yaml index 2538356a85a5..afa5fe000e65 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf6_tdm_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf6_tdm_gfx1250.yaml @@ -23,8 +23,6 @@ GlobalParameters: MinKForGSU: 1 Architecture: gfx1250 PrintLevel: 2 - MaxLDS: 163840 - DeviceLDS: 163840 BenchmarkProblems: ####################################### diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf6f4ss_tdm_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf6f4ss_tdm_gfx1250.yaml index 3744c4172132..a952729f911c 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf6f4ss_tdm_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf6f4ss_tdm_gfx1250.yaml @@ -23,8 +23,6 @@ GlobalParameters: MinKForGSU: 1 Architecture: gfx1250 PrintLevel: 2 - MaxLDS: 163840 - DeviceLDS: 163840 BenchmarkProblems: ####################################### diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf8_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf8_gfx1250.yaml index 6b6781018d50..6338b278d836 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf8_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf8_gfx1250.yaml @@ -28,8 +28,6 @@ GlobalParameters: KeepBuildTmp: True PrintLevel: 2 UseGPUTimer: False - MaxLDS: 163840 - DeviceLDS: 163840 BenchmarkProblems: @@ -55,7 +53,7 @@ BenchmarkProblems: DataTypeMXSA: e8 DataTypeMXSB: e8 - - # BenchmarkProblemSizeGroup - Standard + - # BenchmarkProblemSizeGroup - 1a: InitialSolutionParameters: BenchmarkCommonParameters: - KernelLanguage: ["Assembly"] @@ -82,7 +80,63 @@ BenchmarkProblems: - PrefetchLocalRead: [1] - PrefetchGlobalRead: [2] - StoreRemapVectorWidth: [0] - - GlobalSplitU: [1,2,4,8] + - GlobalSplitU: [1] + - GlobalSplitUAlgorithm: [MultipleBuffer] + - GlobalReadVectorWidthA: [-1] + - GlobalReadVectorWidthB: [-1] + - VectorWidthA: [-1] + - VectorWidthB: [-1] + - LocalReadVectorWidth: [-1] + - 1LDSBuffer: [0] + - SourceSwap: [0] + - DirectToVgprSparseMetadata: [0] + - WorkGroupMapping: [1] + - StoreVectorWidth: [-1] + - DirectToVgprA: [False] + - DirectToVgprB: [False] + - WavefrontSize: [32] + - ExpandPointerSwap: [0] + - AssertSummationElementMultiple: [ 16 ] + BenchmarkForkParameters: + JoinParameters: + BenchmarkJoinParameters: + BenchmarkFinalParameters: + - ProblemSizes: + - Exact: [256, 256, 1, 768] + - Exact: [256, 256, 1, 832] #tail + - Exact: [257, 256, 1, 768] #edge-m + - Exact: [256, 257, 1, 768] #edge-n + - ActivationArgs: + - [Enum: none] + + - # BenchmarkProblemSizeGroup - 1b: GSU + InitialSolutionParameters: + BenchmarkCommonParameters: + - KernelLanguage: ["Assembly"] + ForkParameters: + - UseSgprForGRO: [0] + - ForceDisableShadowInit: [1] + - Groups: + - + - MatrixInstruction: [16, 16, 128, 1, 1, 4, 4, 2, 1] + WorkGroup: [16, 16, 2] + - MatrixInstruction: [16, 16, 128, 1, 1, 8, 4, 2, 2] + - DepthU: [256] + - TDMInst: [3] + - TDMIterateMode: [0, 3] + - TransposeLDS: [-1] + - LdsPadA: [-1] + - LdsPadB: [-1] + - LdsBlockSizePerPadA: [-1] + - LdsBlockSizePerPadB: [-1] + - LdsPadMetadata: [0] + - LDSTrInst: [False] + - StaggerU: [0] + - ScheduleIterAlg: [0] + - PrefetchLocalRead: [1] + - PrefetchGlobalRead: [2] + - StoreRemapVectorWidth: [0] + - GlobalSplitU: [2,4,8] - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] @@ -110,6 +164,7 @@ BenchmarkProblems: - Exact: [256, 257, 1, 2048] #edge-n - ActivationArgs: - [Enum: none] + ####################################### # TN - standard, MXF8SS Block32, E8 ####################################### @@ -132,7 +187,7 @@ BenchmarkProblems: DataTypeMXSA: e8 DataTypeMXSB: e8 - - # BenchmarkProblemSizeGroup - Standard + - # BenchmarkProblemSizeGroup - 2a: MI x TDMIterateMode x TDMSplit (TDM load paths) InitialSolutionParameters: BenchmarkCommonParameters: - KernelLanguage: ["Assembly"] @@ -142,7 +197,7 @@ BenchmarkProblems: - MatrixInstruction: - [16, 16, 128, 1, 1, 8, 4, 2, 2] - [16, 16, 128, 1, 1, 8, 8, 2, 2] - - HalfPLR: [0,1,2,3] + - HalfPLR: [0,3] - DepthU: [256] - TDMInst: [3] - TDMIterateMode: [-1,0,1,2,3] @@ -157,9 +212,65 @@ BenchmarkProblems: - StaggerU: [0] - ScheduleIterAlg: [0] - PrefetchLocalRead: [1] + - PrefetchGlobalRead: [2] + - StoreRemapVectorWidth: [0] + - GlobalSplitU: [1] + - GlobalSplitUAlgorithm: [MultipleBuffer] + - GlobalReadVectorWidthA: [-1] + - GlobalReadVectorWidthB: [-1] + - VectorWidthA: [-1] + - VectorWidthB: [-1] + - LocalReadVectorWidth: [-1] + - 1LDSBuffer: [0] + - SourceSwap: [0] + - DirectToVgprSparseMetadata: [0] + - WorkGroupMapping: [1] + - StoreVectorWidth: [-1] + - DirectToVgprA: [False] + - DirectToVgprB: [False] + - WavefrontSize: [32] + - ExpandPointerSwap: [0] + - AssertSummationElementMultiple: [ 32 ] + BenchmarkForkParameters: + JoinParameters: + BenchmarkJoinParameters: + BenchmarkFinalParameters: + - ProblemSizes: + - Exact: [256, 256, 1, 512] + - Exact: [256, 256, 1, 576] #tail + - Exact: [257, 256, 1, 512] #edge-m + - Exact: [256, 257, 1, 512] #edge-n + - ActivationArgs: + - [Enum: none] + + - # BenchmarkProblemSizeGroup - 2b: MI x HalfPLR x PGR x GSU (K-path + local-read) + InitialSolutionParameters: + BenchmarkCommonParameters: + - KernelLanguage: ["Assembly"] + ForkParameters: + - UseSgprForGRO: [0] + - ForceDisableShadowInit: [1] + - MatrixInstruction: + - [16, 16, 128, 1, 1, 8, 4, 2, 2] + - [16, 16, 128, 1, 1, 8, 8, 2, 2] + - HalfPLR: [0,1,2,3] + - DepthU: [256] + - TDMInst: [3] + - TDMIterateMode: [0] + - TDMSplit: [False] + - TransposeLDS: [-1] + - LdsPadA: [-1] + - LdsPadB: [-1] + - LdsBlockSizePerPadA: [-1] + - LdsBlockSizePerPadB: [-1] + - LdsPadMetadata: [0] + - LDSTrInst: [False] + - StaggerU: [0] + - ScheduleIterAlg: [0] + - PrefetchLocalRead: [1] - PrefetchGlobalRead: [1,2] - StoreRemapVectorWidth: [0] - - GlobalSplitU: [1,2,4,8] + - GlobalSplitU: [2,4,8] - GlobalSplitUAlgorithm: [MultipleBuffer] - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] @@ -181,10 +292,7 @@ BenchmarkProblems: BenchmarkJoinParameters: BenchmarkFinalParameters: - ProblemSizes: - - Exact: [256, 256, 1, 128] - - Exact: [256, 256, 1, 384] - Exact: [256, 256, 1, 640] - - Exact: [256, 256, 1, 896] - Exact: [256, 256, 1, 2048] - Exact: [256, 256, 1, 2080] #tail - Exact: [257, 256, 1, 2048] #edge-m diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/nt_th_nv_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/nt_th_nv_gfx1250.yaml index 965d7f0c7cc0..cfb423e52e95 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/nt_th_nv_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/nt_th_nv_gfx1250.yaml @@ -51,7 +51,7 @@ BenchmarkProblems: - DepthU: [32] - PrefetchGlobalRead: [1] - PrefetchLocalRead: [1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - TransposeLDS: [-1] - LdsPadA: [-1] - LdsPadB: [-1] @@ -60,7 +60,7 @@ BenchmarkProblems: - GlobalReadVectorWidthB: [-1] - StoreVectorWidth: [-1] - WorkGroupMapping: [1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - Groups: diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/subtile_bf16_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/subtile_bf16_gfx1250.yaml index 3e43e260b454..19b3dab34c35 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/subtile_bf16_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/subtile_bf16_gfx1250.yaml @@ -48,7 +48,7 @@ BenchmarkProblems: - DepthU: [64, 128] - WavefrontSize: [32] - PrefetchGlobalRead: [2] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - UseSubtileImpl: [1] - StreamK: [3] - TDMInst: [3] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/xfp32_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/xfp32_gfx1250.yaml index 6f70185364ba..e190a3517a71 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/xfp32_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/xfp32_gfx1250.yaml @@ -118,7 +118,7 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [1,2] - GlobalReadVectorWidthB: [1] - 1LDSBuffer: [0] - - ScheduleIterAlg: [0, 1] + - ScheduleIterAlg: [0] - SourceSwap: [1] - DirectToVgprA: [False] - StoreVectorWidth: [-1] @@ -174,7 +174,7 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [1] - GlobalReadVectorWidthB: [1,2] - 1LDSBuffer: [0] - - ScheduleIterAlg: [1] + - ScheduleIterAlg: [0] - SourceSwap: [1] - DirectToVgprA: [False] - StoreVectorWidth: [-1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/bbs_bgrada_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/bbs_bgrada_gfx1250.yaml index 1b6177aaab39..f21144da0926 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/bbs_bgrada_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/bbs_bgrada_gfx1250.yaml @@ -77,14 +77,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -139,14 +139,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -201,14 +201,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -263,14 +263,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/bbs_bgradb_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/bbs_bgradb_gfx1250.yaml index 8d55f5b7c41f..07b5be2be3c4 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/bbs_bgradb_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/bbs_bgradb_gfx1250.yaml @@ -77,14 +77,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -139,14 +139,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -201,14 +201,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -263,14 +263,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/bbs_bgradd_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/bbs_bgradd_gfx1250.yaml index 95c567f44692..92413c01e5d2 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/bbs_bgradd_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/bbs_bgradd_gfx1250.yaml @@ -77,14 +77,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -139,14 +139,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -201,14 +201,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -263,14 +263,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/bbs_dgelu_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/bbs_dgelu_gfx1250.yaml index a0fd802bbda2..9f25c1b8ce41 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/bbs_dgelu_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/bbs_dgelu_gfx1250.yaml @@ -77,14 +77,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -139,14 +139,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -201,14 +201,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -263,14 +263,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/hhs_bgrada_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/hhs_bgrada_gfx1250.yaml index a7ff5a44cc41..8d69d9b32ccc 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/hhs_bgrada_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/hhs_bgrada_gfx1250.yaml @@ -77,14 +77,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -139,14 +139,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -201,14 +201,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -263,14 +263,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/hhs_bgradb_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/hhs_bgradb_gfx1250.yaml index 774d20a3655c..8eecf403e325 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/hhs_bgradb_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/hhs_bgradb_gfx1250.yaml @@ -77,14 +77,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -139,14 +139,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -201,14 +201,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -263,14 +263,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/hhs_bgradd_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/hhs_bgradd_gfx1250.yaml index e476b5fbeca8..ecad467eb4cd 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/hhs_bgradd_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/hhs_bgradd_gfx1250.yaml @@ -77,14 +77,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -139,14 +139,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -201,14 +201,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -263,14 +263,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/hhs_dgelu_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/hhs_dgelu_gfx1250.yaml index f6692324b933..2c0895c03c0a 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/hhs_dgelu_gfx1250.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gradient/gfx1250/hhs_dgelu_gfx1250.yaml @@ -77,14 +77,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -139,14 +139,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -201,14 +201,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0] @@ -263,14 +263,14 @@ BenchmarkProblems: - GlobalReadVectorWidthA: [-1] - GlobalReadVectorWidthB: [-1] - LocalReadVectorWidth: [-1] - - ScheduleIterAlg: [3] + - ScheduleIterAlg: [0] - ExpandPointerSwap: [0] - TransposeLDS: [-1] - LdsBlockSizePerPadA: [-1] - LdsBlockSizePerPadB: [-1] - LdsPadA: [-1] - LdsPadB: [-1] - - 1LDSBuffer: [-1] + - 1LDSBuffer: [0] - GlobalSplitU: [1] - GlobalSplitUAlgorithm: ["MultipleBuffer"] - SourceSwap: [0]