From 06e5763644fad4b51ff6645416b4916b979e6e2f Mon Sep 17 00:00:00 2001 From: mengzcai Date: Sun, 14 Jun 2026 15:20:51 +0000 Subject: [PATCH 1/5] [tox] Reduce test size/kernels for mxf8/mxf4 yamls --- .../Tests/common/gemm/gfx12/mxf4_gfx1250.yaml | 14 +- .../Tests/common/gemm/gfx12/mxf8_gfx1250.yaml | 126 ++++++++++++++++-- 2 files changed, 119 insertions(+), 21 deletions(-) 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..610895250f0a 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 @@ -62,9 +62,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 +133,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 +203,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 +272,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 +341,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 +410,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/mxf8_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/mxf8_gfx1250.yaml index 6b6781018d50..d9c331825b39 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 @@ -55,7 +55,7 @@ BenchmarkProblems: DataTypeMXSA: e8 DataTypeMXSB: e8 - - # BenchmarkProblemSizeGroup - Standard + - # BenchmarkProblemSizeGroup - 1a: InitialSolutionParameters: BenchmarkCommonParameters: - KernelLanguage: ["Assembly"] @@ -82,7 +82,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 +166,7 @@ BenchmarkProblems: - Exact: [256, 257, 1, 2048] #edge-n - ActivationArgs: - [Enum: none] + ####################################### # TN - standard, MXF8SS Block32, E8 ####################################### @@ -132,7 +189,7 @@ BenchmarkProblems: DataTypeMXSA: e8 DataTypeMXSB: e8 - - # BenchmarkProblemSizeGroup - Standard + - # BenchmarkProblemSizeGroup - 2a: MI x TDMIterateMode x TDMSplit (TDM load paths) InitialSolutionParameters: BenchmarkCommonParameters: - KernelLanguage: ["Assembly"] @@ -142,7 +199,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 +214,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 +294,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 From 2a7e693e296890a5523058f1419f58daa1f929bc Mon Sep 17 00:00:00 2001 From: mengzcai Date: Mon, 15 Jun 2026 08:00:19 +0000 Subject: [PATCH 2/5] [Test yamls] set SIA and 1LDSBuffer to 0, remove deviceLDS and maxLDS see default LDS size in hardware_caps.hpp --- .../common/gemm/gfx12/1024_vgpr_gfx1250.yaml | 24 +++++++++---------- .../common/gemm/gfx12/b6f4ss_gfx1250.yaml | 16 ++++++------- .../common/gemm/gfx12/b8b8s_gfx1250.yaml | 16 ++++++------- .../Tests/common/gemm/gfx12/bf16_gfx1250.yaml | 16 ++++++------- .../Tests/common/gemm/gfx12/bf6_gfx1250.yaml | 14 +++++------ .../Tests/common/gemm/gfx12/f32_gfx1250.yaml | 16 ++++++------- .../Tests/common/gemm/gfx12/f4_gfx1250.yaml | 14 +++++------ .../common/gemm/gfx12/f4b8ss_gfx1250.yaml | 16 ++++++------- .../Tests/common/gemm/gfx12/f64_gfx1250.yaml | 16 ++++++------- .../Tests/common/gemm/gfx12/f6_gfx1250.yaml | 14 +++++------ .../common/gemm/gfx12/f6b6ss_gfx1250.yaml | 2 +- .../common/gemm/gfx12/f6b8ss_gfx1250.yaml | 16 ++++++------- .../common/gemm/gfx12/f8b6ss_gfx1250.yaml | 16 ++++++------- .../common/gemm/gfx12/f8b8ss_gfx1250.yaml | 16 ++++++------- .../common/gemm/gfx12/f8f4ss_gfx1250.yaml | 16 ++++++------- .../common/gemm/gfx12/f8f8s_gfx1250.yaml | 16 ++++++------- .../common/gemm/gfx12/f8f8s_sr_gfx1250.yaml | 16 ++++++------- .../Tests/common/gemm/gfx12/fp16_gfx1250.yaml | 16 ++++++------- .../common/gemm/gfx12/fp16_use_e_gfx1250.yaml | 4 ++-- .../Tests/common/gemm/gfx12/fp8_gfx1250.yaml | 16 ++++++------- .../Tests/common/gemm/gfx12/i8ii_gfx1250.yaml | 16 ++++++------- .../common/gemm/gfx12/largeLds_gfx1250.yaml | 4 ++-- .../Tests/common/gemm/gfx12/mxf4_gfx1250.yaml | 2 -- .../common/gemm/gfx12/mxf6_tdm_gfx1250.yaml | 2 -- .../gemm/gfx12/mxf6f4ss_tdm_gfx1250.yaml | 2 -- .../Tests/common/gemm/gfx12/mxf8_gfx1250.yaml | 2 -- .../common/gemm/gfx12/nt_th_nv_gfx1250.yaml | 4 ++-- .../gemm/gfx12/subtile_bf16_gfx1250.yaml | 2 +- .../common/gemm/gfx12/xfp32_gfx1250.yaml | 4 ++-- .../gradient/gfx1250/bbs_bgrada_gfx1250.yaml | 16 ++++++------- .../gradient/gfx1250/bbs_bgradb_gfx1250.yaml | 16 ++++++------- .../gradient/gfx1250/bbs_bgradd_gfx1250.yaml | 16 ++++++------- .../gradient/gfx1250/bbs_dgelu_gfx1250.yaml | 16 ++++++------- .../gradient/gfx1250/hhs_bgrada_gfx1250.yaml | 16 ++++++------- .../gradient/gfx1250/hhs_bgradb_gfx1250.yaml | 16 ++++++------- .../gradient/gfx1250/hhs_bgradd_gfx1250.yaml | 16 ++++++------- .../gradient/gfx1250/hhs_dgelu_gfx1250.yaml | 16 ++++++------- 37 files changed, 226 insertions(+), 236 deletions(-) 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 610895250f0a..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: 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 d9c331825b39..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: 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] From 22d50c8be92f5d982804b4db9d69bd59e524dd3c Mon Sep 17 00:00:00 2001 From: Chuang-Yu Cheng Date: Mon, 15 Jun 2026 08:22:22 +0000 Subject: [PATCH 3/5] [tensile] Fix SIA0 PGR2 global-read placement Defer SIA0 non-TDM PGR2 global reads until after the local-write iteration consumes the previous vgprG2L values, and keep tail local-read offsets reset before tail MACs. This prevents next-prefetch buffer loads from clobbering data before it is stored to LDS. --- .../tensilelite/Tensile/Components/SIA.py | 8 ++++- .../tensilelite/Tensile/KernelWriter.py | 33 +++++++++++-------- 2 files changed, 27 insertions(+), 14 deletions(-) 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..7bef8dd364a8 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 @@ -5804,19 +5811,19 @@ 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"]: - module.addComment1("Tail: local read reset offsets a") - module.add(self.localReadResetOffsets(kernel, tensorParametersA)) - if kernel["ProblemType"]["MXBlockA"]: - module.addComment1("Tail: local read reset offsets mxsa") - module.add(self.localReadResetOffsets(kernel, tensorParametersA["MX"])) - if kernel["ProblemType"]["MXBlockB"]: - module.addComment1("Tail: local read reset offsets mxsb") - module.add(self.localReadResetOffsets(kernel, tensorParametersB["MX"])) - module.addComment1("Tail: local read reset offsets b") - module.add(self.localReadResetOffsets(kernel, tensorParametersB)) + # 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. + module.addComment1("Tail: local read reset offsets a") + module.add(self.localReadResetOffsets(kernel, tensorParametersA)) + if kernel["ProblemType"]["MXBlockA"]: + module.addComment1("Tail: local read reset offsets mxsa") + module.add(self.localReadResetOffsets(kernel, tensorParametersA["MX"])) + if kernel["ProblemType"]["MXBlockB"]: + module.addComment1("Tail: local read reset offsets mxsb") + module.add(self.localReadResetOffsets(kernel, tensorParametersB["MX"])) + module.addComment1("Tail: local read reset offsets b") + module.add(self.localReadResetOffsets(kernel, tensorParametersB)) module.addComment1("Tail: local read init pointers a") module.add(self.localReadInitPointers(kernel, tensorParametersA, tensorParametersA)) From f435fc7b36787bb3a6e839db087f8cb3c732d027 Mon Sep 17 00:00:00 2001 From: mengzcai Date: Mon, 15 Jun 2026 11:17:53 +0000 Subject: [PATCH 4/5] Fix TDM tail local-read using wrong LDS buffer Before: TDM + double-buffer tail forced the tail LR offsets back to buffer0, but the TDM write pointer may stay at buffer1. (after an odd number of main-loop iters) After: No longer forces buffer 0 in the tail, unless wider local read needs the offset recomputed (numReadsIterCoalesced>1). Repro/affected tests: TDM kernels with a tail e.g. bf6_tdm_gfx1250, tdm_gfx1250, ... --- .../tensilelite/Tensile/KernelWriter.py | 29 ++++++++++++------- 1 file changed, 19 insertions(+), 10 deletions(-) diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py index 7bef8dd364a8..6fc21f3783fa 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py @@ -5788,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"]): @@ -5814,16 +5822,17 @@ def _restoreNtabState(): # 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. - module.addComment1("Tail: local read reset offsets a") - module.add(self.localReadResetOffsets(kernel, tensorParametersA)) - if kernel["ProblemType"]["MXBlockA"]: - module.addComment1("Tail: local read reset offsets mxsa") - module.add(self.localReadResetOffsets(kernel, tensorParametersA["MX"])) - if kernel["ProblemType"]["MXBlockB"]: - module.addComment1("Tail: local read reset offsets mxsb") - module.add(self.localReadResetOffsets(kernel, tensorParametersB["MX"])) - module.addComment1("Tail: local read reset offsets b") - module.add(self.localReadResetOffsets(kernel, tensorParametersB)) + if needResetLROffsets: + module.addComment1("Tail: local read reset offsets a") + module.add(self.localReadResetOffsets(kernel, tensorParametersA)) + if kernel["ProblemType"]["MXBlockA"]: + module.addComment1("Tail: local read reset offsets mxsa") + module.add(self.localReadResetOffsets(kernel, tensorParametersA["MX"])) + if kernel["ProblemType"]["MXBlockB"]: + module.addComment1("Tail: local read reset offsets mxsb") + module.add(self.localReadResetOffsets(kernel, tensorParametersB["MX"])) + module.addComment1("Tail: local read reset offsets b") + module.add(self.localReadResetOffsets(kernel, tensorParametersB)) module.addComment1("Tail: local read init pointers a") module.add(self.localReadInitPointers(kernel, tensorParametersA, tensorParametersA)) From afcc1889e771633b3483fdf513afe3754b6ab6f2 Mon Sep 17 00:00:00 2001 From: mengzcai Date: Mon, 15 Jun 2026 18:35:07 +0000 Subject: [PATCH 5/5] Fix LocalReadVALU overwriting correct LDS>=64KB read address 1. Set LDS>=64KB and 1LDSB and SIA0 to reproduce the issue 2. Fix test yaml: f64_gfx1250 (using LocalReadVALU) --- .../hipblaslt/tensilelite/Tensile/Components/LocalRead.py | 7 ++----- 1 file changed, 2 insertions(+), 5 deletions(-) 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])