From 74e21396303af890d2fa15593e3b545be30dfe17 Mon Sep 17 00:00:00 2001 From: "Joao P. L. de Carvalho" Date: Thu, 1 Oct 2026 18:31:36 +0000 Subject: [PATCH 1/2] fix(tensilelite): generate aux output E in persistent kernels A persistent TileProcessingStrategy (StreamK or DataParallel) sets GlobalSplitU to 0. The epilogue already emitted E stores for persistent kernels, but the E address registers, row offset and row advance only checked GlobalSplitU == 1, so generating any persistent kernel with E failed with a TypeError. Set up the E address and row offset for the persistent final store with isPersistent, the same way bias and ScaleAlphaVec already do. In the StreamK tree-reduction fixup, check out the 3 SGPRs the flag wait uses instead of 2; with E, the third one was the low word of the E buffer address and the kernel faulted. Reject E with StoreRemapVectorWidth for persistent kernels, as for GlobalSplitU == 1. Add gfx950 common tests that validate D and E for StreamK with each WorkAssignment, with and without tree reduction, and for DataParallel, each with a gradient case that reads E. --- .../Tensile/AsmAddressCalculation.py | 6 +- .../tensilelite/Tensile/AsmStoreState.py | 12 +- .../Tensile/Components/ComputeStoreVgprs.py | 5 +- .../tensilelite/Tensile/Components/LSU.py | 3 +- .../tensilelite/Tensile/Components/StreamK.py | 2 +- .../Tensile/SolutionStructs/Solution.py | 4 +- .../data_parallel_static_hgemm_use_e.yaml | 109 +++++++++++++++++ .../common/streamk/gfx950/sk_hgemm_use_e.yaml | 110 ++++++++++++++++++ 8 files changed, 236 insertions(+), 15 deletions(-) create mode 100644 projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx950/data_parallel_static_hgemm_use_e.yaml create mode 100644 projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx950/sk_hgemm_use_e.yaml diff --git a/projects/hipblaslt/tensilelite/Tensile/AsmAddressCalculation.py b/projects/hipblaslt/tensilelite/Tensile/AsmAddressCalculation.py index 059429f886cb..0cf42ea2793b 100644 --- a/projects/hipblaslt/tensilelite/Tensile/AsmAddressCalculation.py +++ b/projects/hipblaslt/tensilelite/Tensile/AsmAddressCalculation.py @@ -184,7 +184,7 @@ def emitRowPtrAdvance(self, kernel, ss, tmpS01, rowInc, lookahead=False): sgpr("StrideC%s"%strideChar), rowInc, tmpS01, "ROWINC- Move cinRowPtr to next row")) module.add(self.addScaled(vgpr(kw.vgprs.coutRowPtrD), vgpr(kw.vgprs.coutRowPtrD), \ sgpr("StrideD%s"%strideChar), rowInc, tmpS01, "Move coutRowPtrD to next row")) - if kernel["ProblemType"]["UseE"] and (kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1): + if kernel["ProblemType"]["UseE"] and ((kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1) or isPersistent(kernel)): ecomment = "coutRowPtrE.la: += StrideE%s*rowInc"%strideChar if lookahead else "Move coutRowPtrE to next row" module.add(self.addScaled(vgpr(kw.vgprs.coutRowPtrE), vgpr(kw.vgprs.coutRowPtrE), \ sgpr("StrideE%s"%strideChar), rowInc, tmpS01, ecomment)) @@ -765,7 +765,7 @@ def emitAddressSetupCode(self, kernel, tPB, ss, tmpVgpr, tmpS01, edge, beta, ato comment="new rowStart address += shift column * StridesD")) module.add(VCndMaskB32(dst=vgpr(kw.vgprs.coutRowPtrD), src0=vgpr(kw.vgprs.coutRowPtrD), src1=vgpr(vTmp1), src2=sgpr(sTmp1,sgprCnt), \ comment="set new rowStart if meet conditions" )) - if kernel["ProblemType"]["UseE"] and (kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1): + if kernel["ProblemType"]["UseE"] and ((kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1) or isPersistent(kernel)): module.add(VMadI32I24(dst=vgpr(vTmp1), src0=sgpr(strideE1), src1=vgpr(vTmp2), src2=vgpr(kw.vgprs.coutRowPtrE), \ comment="new rowStart address += shift column * StridesE")) module.add(VCndMaskB32(dst=vgpr(kw.vgprs.coutRowPtrE), src0=vgpr(kw.vgprs.coutRowPtrE), src1=vgpr(vTmp1), src2=sgpr(sTmp1,sgprCnt), \ @@ -784,7 +784,7 @@ def emitAddressSetupCode(self, kernel, tPB, ss, tmpVgpr, tmpS01, edge, beta, ato comment="new lds write address += shift column * Lds byte Stride")) module.add(VCndMaskB32(dst=vgpr(kw.vgprs.storeRemapLW), src0=vgpr(kw.vgprs.storeRemapLW), src1=vgpr(vTmp1), \ src2=sgpr(sTmp1,sgprCnt), comment="set new rowStart if meet conditions" )) - if kernel["ProblemType"]["UseE"] and (kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1): + if kernel["ProblemType"]["UseE"] and ((kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1) or isPersistent(kernel)): printExit("Output E does not support StoreRemapVectorWidth") if kw.vgprs.coutRowPtrBias != -1: printExit("Bias reduction does not support StoreRemapVectorWidth") diff --git a/projects/hipblaslt/tensilelite/Tensile/AsmStoreState.py b/projects/hipblaslt/tensilelite/Tensile/AsmStoreState.py index 4fcfd35320a0..3494f7cde684 100644 --- a/projects/hipblaslt/tensilelite/Tensile/AsmStoreState.py +++ b/projects/hipblaslt/tensilelite/Tensile/AsmStoreState.py @@ -254,7 +254,7 @@ def __init__(self, kernelWriter, kernel, gwvw, edge, beta, atomic, elements, vec self.sharedColBiasVgprs = kernelWriter.vgprPool.checkOut(self.numAddrVgpr, "sharedColBiasVgprs for packed elements") else: self.sharedColBiasVgprs = None - if kernel["ProblemType"]["UseE"] and (kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1): + if kernel["ProblemType"]["UseE"] and ((kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1) or (isPersistent(kernel) and not isWorkspace)): self.sharedColEVgprs = kernelWriter.vgprPool.checkOut(self.numAddrVgpr, "sharedColEVgprs for packed elements") else: self.sharedColEVgprs = None @@ -308,7 +308,7 @@ def __init__(self, kernelWriter, kernel, gwvw, edge, beta, atomic, elements, vec self.sharedColBiasVgprs = kernelWriter.vgprPool.checkOut(1, "sharedColBiasVgprs for packed elements") else: self.sharedColBiasVgprs = None - if kernel["ProblemType"]["UseE"] and (kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1): + if kernel["ProblemType"]["UseE"] and ((kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1) or (isPersistent(kernel) and not isWorkspace)): self.sharedColEVgprs = kernelWriter.vgprPool.checkOut(1, "sharedColEVgprs for packed elements") else: self.sharedColEVgprs = None @@ -371,7 +371,7 @@ def __init__(self, kernelWriter, kernel, gwvw, edge, beta, atomic, elements, vec # TODO STREAM-K Update numVgprsPerElement for workspace if kernel["GroupLoadStore"] and kernel["ProblemType"]["UseBeta"]: self.numVgprsPerElement += self.cfg.numVgprsPerAddr - if kernel["ProblemType"]["UseE"] and (kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1): + if kernel["ProblemType"]["UseE"] and ((kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1) or (isPersistent(kernel) and not isWorkspace)): self.numVgprsPerElement += self.cfg.numVgprsPerAddr # E address # Only needed in gradient activation if (kernel["ProblemType"]["Gradient"] and kernel["ProblemType"]["ActivationType"] != 'none'): @@ -946,7 +946,7 @@ def setupStoreElementsForBatch(self, kernel, gwvw, batchElements, batchElementSg dataBias = 0 self.elementDataBias.append(dataBias) # Only needed in gradient activation - if (kernel["ProblemType"]["Gradient"] and kernel["ProblemType"]["ActivationType"] != 'none' and kernel["ProblemType"]["UseE"]) and (kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1): + if (kernel["ProblemType"]["Gradient"] and kernel["ProblemType"]["ActivationType"] != 'none' and kernel["ProblemType"]["UseE"]) and ((kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1) or (isPersistent(kernel) and not isWorkspace)): numVgprs = int(ceil(kernel["ProblemType"]["ComputeDataType"].numRegisters())) dataE = kw.vgprPool.checkOutAligned(int(numVgprs*self.cfg.gwvw), \ int(ceil(numVgprs*self.cfg.gwvw)), "e data for ei=%u"%elementIdx, preventOverflow=False) @@ -1032,7 +1032,7 @@ def setupStoreElementsForBatch(self, kernel, gwvw, batchElements, batchElementSg addrBiasVgpr = self.sharedColBiasVgprs+elementCol else: addrBiasVgpr = None - if kernel["ProblemType"]["UseE"] and (kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1): + if kernel["ProblemType"]["UseE"] and ((kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1) or (isPersistent(kernel) and not isWorkspace)): addrEVgpr = self.sharedColEVgprs+elementCol else: addrEVgpr = None @@ -1083,7 +1083,7 @@ def setupStoreElementsForBatch(self, kernel, gwvw, batchElements, batchElementSg else: addrBiasVgpr = None - if kernel["ProblemType"]["UseE"] and (kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1): + if kernel["ProblemType"]["UseE"] and ((kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1) or (isPersistent(kernel) and not isWorkspace)): addrEVgpr = kw.vgprPool.checkOutAligned(self.cfg.numVgprsPerAddr, \ int(ceil(self.cfg.numVgprsPerAddr)), "loadEBatch-addr for ei=%u"%(elementIdx), preventOverflow=not isOptNLL) else: diff --git a/projects/hipblaslt/tensilelite/Tensile/Components/ComputeStoreVgprs.py b/projects/hipblaslt/tensilelite/Tensile/Components/ComputeStoreVgprs.py index 91ddadaaf478..9161dbf9c28f 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Components/ComputeStoreVgprs.py +++ b/projects/hipblaslt/tensilelite/Tensile/Components/ComputeStoreVgprs.py @@ -29,6 +29,7 @@ from ..Component import ComputeStoreVgprs from ..Common import DataDirection, log2 +from ..ExecutionPolicy import isPersistent class ComputeStoreVgprsVALU(ComputeStoreVgprs): kernel = {"EnableMatrixInstruction": False, @@ -201,7 +202,7 @@ def __call__(self, writer, kernel): strideD1 = "StrideD%s" % (writer.states.indexChars[packedC1[0]]) module.add(VMulLOU32(dst=vgpr(writer.vgprs.cinRowPtr), src0=vgpr(lsuTid1), src1=sgpr(strideC1), comment=" offset 1")) module.add(VMulLOU32(dst=vgpr(writer.vgprs.coutRowPtrD), src0=vgpr(lsuTid1), src1=sgpr(strideD1), comment=" offset 1")) - if kernel["ProblemType"]["UseE"] and (kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1): + if kernel["ProblemType"]["UseE"] and ((kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1) or isPersistent(kernel)): module.add(VMovB32(dst=vgpr(writer.vgprs.coutRowPtrE), src=vgpr(lsuTid1), comment=" save offset 1 for E")) if writer.vgprs.coutRowPtrGate != -1: module.add(VMovB32(dst=vgpr(writer.vgprs.coutRowPtrGate), src=vgpr(lsuTid1), comment=" save offset 1 for Gate")) @@ -332,7 +333,7 @@ def __call__(self, writer, kernel): strideD1 = "StrideD%s" % (writer.states.indexChars[packedC1[0]]) module.add(VMulLOU32(dst=vgpr(writer.vgprs.cinRowPtr), src0=vgpr(lsuTid1), src1=sgpr(strideC1), comment=" offset 1")) module.add(VMulLOU32(dst=vgpr(writer.vgprs.coutRowPtrD), src0=vgpr(lsuTid1), src1=sgpr(strideD1), comment=" offset 1")) - if kernel["ProblemType"]["UseE"] and (kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1): + if kernel["ProblemType"]["UseE"] and ((kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1) or isPersistent(kernel)): module.add(VMovB32(dst=vgpr(writer.vgprs.coutRowPtrE), src=vgpr(lsuTid1), comment=" save offset 1 for E")) if writer.vgprs.coutRowPtrGate != -1: module.add(VMovB32(dst=vgpr(writer.vgprs.coutRowPtrGate), src=vgpr(lsuTid1), comment=" save offset 1 for Gate")) diff --git a/projects/hipblaslt/tensilelite/Tensile/Components/LSU.py b/projects/hipblaslt/tensilelite/Tensile/Components/LSU.py index 4e8612d6c2c3..015437a2ce58 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Components/LSU.py +++ b/projects/hipblaslt/tensilelite/Tensile/Components/LSU.py @@ -29,6 +29,7 @@ from copy import deepcopy from ..Common import log2, ceilDivide, DataDirection from ..Component import Component +from ..ExecutionPolicy import isPersistent from ..KernelWriterModules import * from ..AsmStoreState import StoreState, VectorDataTypes #import abc @@ -450,7 +451,7 @@ def globalWriteIndices(self, writer, kernel): strideD1 = "StrideD%s" % (writer.states.indexChars[packedC1[0]]) module.add(VMulLOU32(dst=vgpr(writer.vgprs.cinRowPtr), src0=vgpr(writer.vgprs.coord1InMT), src1=sgpr(strideC1), comment=" offset 1")) module.add(VMulLOU32(dst=vgpr(writer.vgprs.coutRowPtrD), src0=vgpr(writer.vgprs.coord1InMT), src1=sgpr(strideD1), comment=" offset 1")) - if kernel["ProblemType"]["UseE"] and (kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1): + if kernel["ProblemType"]["UseE"] and ((kernel["GlobalSplitU"] == 1 or kernel["GlobalSplitU"] == -1) or isPersistent(kernel)): module.add(VMovB32(dst=vgpr(writer.vgprs.coutRowPtrE), src=vgpr(writer.vgprs.coord1InMT), comment=" save offset 1 for E")) if writer.vgprs.coutRowPtrGate != -1: module.add(VMovB32(dst=vgpr(writer.vgprs.coutRowPtrGate), src=vgpr(writer.vgprs.coord1InMT), comment=" save offset 1 for Gate")) diff --git a/projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py b/projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py index 78371d8a3899..2c245172d9bb 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py +++ b/projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py @@ -1295,7 +1295,7 @@ def storeBranchesCommon(self, writer, kernel, skPartialsLabel, vectorWidths, ele writer.sgprPool.checkIn(tmpSgpr) # check flag - tmpSgpr = writer.sgprPool.checkOut(2, "globalWriteElements") + tmpSgpr = writer.sgprPool.checkOut(3, "globalWriteElements") module.add(SLShiftLeftB32(dst=sgpr(tmpSgpr), src=sgpr(sFlagIdx), shiftHex=log2(4), comment="flag offset based on wg index")) module.add(skFixupWaitForFlag) # loop to wait for flag diff --git a/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py b/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py index 55c9d76ae05b..90e06734477d 100644 --- a/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py +++ b/projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py @@ -7100,8 +7100,8 @@ def calcEpilogueTurns(factorDims: List) -> int: reject(state, printRejectionReason, "Use E does not support len(PackedC1IndicesX) > 1.") if not state["BufferStore"]: reject(state, printRejectionReason, "Use E only supports BufferStore due to no suppress no store.") - if state["StoreRemapVectorWidth"] and (state["GlobalSplitU"] == 1 or state["GlobalSplitU"] == -1): - reject(state, printRejectionReason, "Use E does not support StoreRemapVectorWidth if GSU == 1.") + if state["StoreRemapVectorWidth"] and (state["GlobalSplitU"] == 1 or state["GlobalSplitU"] == -1 or isPersistent(state)): + reject(state, printRejectionReason, "Use E does not support StoreRemapVectorWidth if GSU == 1 or with a persistent TileProcessingStrategy.") if state["GroupLoadStore"]: reject(state, printRejectionReason, "Use E does not support GroupLoadStore.") diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx950/data_parallel_static_hgemm_use_e.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx950/data_parallel_static_hgemm_use_e.yaml new file mode 100644 index 000000000000..2a2c17e9ea21 --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx950/data_parallel_static_hgemm_use_e.yaml @@ -0,0 +1,109 @@ +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT + +TestParameters: + marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx942, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201, skip-gfx1250, skip-gfx1250-strict] + +GlobalParameters: + Architecture: gfx950 + NumElementsToValidate: -1 + BoundsCheck: 0 + KernelTime: False + NumWarmups: 0 + NumBenchmarks: 1 + SyncsPerBenchmark: 0 + EnqueuesPerSync: 1 + DataInitTypeAlpha: 1 + DataInitTypeBeta: 1 + DataInitTypeA: 12 + DataInitTypeB: 13 + DataInitTypeC: 12 + DataInitTypeE: 21 + MaxWorkspaceSize: 134217728 + PrintSolutionRejectionReason: True + +BenchmarkProblems: + # Bias + activation, pre-activation values written to E + - + - OperationType: GEMM + DataType: h + DestDataType: h + ComputeDataType: s + HighPrecisionAccumulate: True + TransposeA: False + TransposeB: False + UseBeta: True + Batched: True + UseBias: 1 + Activation: True + ActivationType: hipblaslt_all + UseE: True + DataTypeE: h + - InitialSolutionParameters: + BenchmarkCommonParameters: + - KernelLanguage: [Assembly] + ForkParameters: + - MatrixInstruction: + - [16, 16, 16, 1, 1, 4, 4, 2, 2] + - DepthU: [64] + - 1LDSBuffer: [1] + - GlobalReadVectorWidthA: [8] + - GlobalReadVectorWidthB: [8] + - PrefetchGlobalRead: [2] + - PrefetchLocalRead: [1] + - ScheduleIterAlg: [3] + - SourceSwap: [False, True] + - TileProcessingStrategy: [DataParallel] + - WorkAssignment: [StaticGrid] + - ActivationFuncCall: [False, True] + BenchmarkJoinParameters: + BenchmarkFinalParameters: + - ProblemSizes: + - Exact: [512, 512, 1, 512] + - Exact: [129, 97, 1, 1030] + - Exact: [1031, 1031, 1, 4103] + - BiasTypeArgs: ['s'] + - ActivationArgs: + - [Enum: Relu] + - [Enum: Gelu] + + # Gradient: E is read back for the activation derivative + - + - OperationType: GEMM + DataType: h + DestDataType: h + ComputeDataType: s + HighPrecisionAccumulate: True + TransposeA: False + TransposeB: False + UseBeta: True + Batched: True + UseBias: 0 + Activation: True + ActivationType: hipblaslt_all + Gradient: True + UseE: True + DataTypeE: h + - InitialSolutionParameters: + BenchmarkCommonParameters: + - KernelLanguage: [Assembly] + ForkParameters: + - MatrixInstruction: + - [16, 16, 16, 1, 1, 4, 4, 2, 2] + - DepthU: [64] + - 1LDSBuffer: [1] + - GlobalReadVectorWidthA: [8] + - GlobalReadVectorWidthB: [8] + - PrefetchGlobalRead: [2] + - PrefetchLocalRead: [1] + - ScheduleIterAlg: [3] + - SourceSwap: [False, True] + - TileProcessingStrategy: [DataParallel] + - WorkAssignment: [StaticGrid] + BenchmarkJoinParameters: + BenchmarkFinalParameters: + - ProblemSizes: + - Exact: [512, 512, 1, 512] + - Exact: [1031, 1031, 1, 4103] + - ActivationArgs: + - [Enum: Dgelu] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx950/sk_hgemm_use_e.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx950/sk_hgemm_use_e.yaml new file mode 100644 index 000000000000..8d28650dd723 --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx950/sk_hgemm_use_e.yaml @@ -0,0 +1,110 @@ +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT + +TestParameters: + marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx942, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201, skip-gfx1250, skip-gfx1250-strict] + +GlobalParameters: + Architecture: gfx950 + NumElementsToValidate: -1 + BoundsCheck: 0 + KernelTime: False + NumWarmups: 0 + NumBenchmarks: 1 + SyncsPerBenchmark: 0 + EnqueuesPerSync: 1 + DataInitTypeAlpha: 1 + DataInitTypeBeta: 1 + DataInitTypeA: 12 + DataInitTypeB: 13 + DataInitTypeC: 12 + DataInitTypeE: 21 + MaxWorkspaceSize: 134217728 + PrintSolutionRejectionReason: True + +BenchmarkProblems: + # Bias + activation, pre-activation values written to E + - + - OperationType: GEMM + DataType: h + DestDataType: h + ComputeDataType: s + HighPrecisionAccumulate: True + TransposeA: False + TransposeB: False + UseBeta: True + Batched: True + UseBias: 1 + Activation: True + ActivationType: hipblaslt_all + UseE: True + DataTypeE: h + - InitialSolutionParameters: + BenchmarkCommonParameters: + - KernelLanguage: [Assembly] + ForkParameters: + - MatrixInstruction: + - [16, 16, 16, 1, 1, 4, 4, 2, 2] + - DepthU: [64] + - 1LDSBuffer: [1] + - GlobalReadVectorWidthA: [8] + - GlobalReadVectorWidthB: [8] + - PrefetchGlobalRead: [2] + - PrefetchLocalRead: [1] + - ScheduleIterAlg: [3] + - SourceSwap: [False, True] + - TileProcessingStrategy: [StreamK] + - WorkAssignment: [StaticGrid, DynamicWorkQueue, Hybrid] + - StreamKFixupTreeReduction: [0, 1] + - ActivationFuncCall: [False, True] + BenchmarkJoinParameters: + BenchmarkFinalParameters: + - ProblemSizes: + - Exact: [512, 512, 1, 512] + - Exact: [129, 97, 1, 1030] + - Exact: [1031, 1031, 1, 4103] + - BiasTypeArgs: ['s'] + - ActivationArgs: + - [Enum: Relu] + - [Enum: Gelu] + + # Gradient: E is read back for the activation derivative + - + - OperationType: GEMM + DataType: h + DestDataType: h + ComputeDataType: s + HighPrecisionAccumulate: True + TransposeA: False + TransposeB: False + UseBeta: True + Batched: True + UseBias: 0 + Activation: True + ActivationType: hipblaslt_all + Gradient: True + UseE: True + DataTypeE: h + - InitialSolutionParameters: + BenchmarkCommonParameters: + - KernelLanguage: [Assembly] + ForkParameters: + - MatrixInstruction: + - [16, 16, 16, 1, 1, 4, 4, 2, 2] + - DepthU: [64] + - 1LDSBuffer: [1] + - GlobalReadVectorWidthA: [8] + - GlobalReadVectorWidthB: [8] + - PrefetchGlobalRead: [2] + - PrefetchLocalRead: [1] + - ScheduleIterAlg: [3] + - SourceSwap: [False, True] + - TileProcessingStrategy: [StreamK] + - WorkAssignment: [StaticGrid] + BenchmarkJoinParameters: + BenchmarkFinalParameters: + - ProblemSizes: + - Exact: [512, 512, 1, 512] + - Exact: [1031, 1031, 1, 4103] + - ActivationArgs: + - [Enum: Dgelu] From e49d7000874ffded9023403a324fe02a206b97ac Mon Sep 17 00:00:00 2001 From: "Joao P. L. de Carvalho" Date: Mon, 5 Oct 2026 16:53:51 +0000 Subject: [PATCH 2/2] test(tensilelite): cover Stream-K parallel reduction with aux output E The sk_hgemm_use_e sizes all select tree reduction. Add 512x512x8192: with 16 tiles and 128 iterations per tile, origami selects parallel reduction for StaticGrid and Hybrid, so the GEMM kernel writes split partials to the workspace and the post kernel applies bias and the activation and writes E. --- .../Tensile/Tests/common/streamk/gfx950/sk_hgemm_use_e.yaml | 3 +++ 1 file changed, 3 insertions(+) diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx950/sk_hgemm_use_e.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx950/sk_hgemm_use_e.yaml index 8d28650dd723..976543158f2b 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx950/sk_hgemm_use_e.yaml +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx950/sk_hgemm_use_e.yaml @@ -63,6 +63,8 @@ BenchmarkProblems: - Exact: [512, 512, 1, 512] - Exact: [129, 97, 1, 1030] - Exact: [1031, 1031, 1, 4103] + # Few tiles and a long K: parallel reduction with the post kernel + - Exact: [512, 512, 1, 8192] - BiasTypeArgs: ['s'] - ActivationArgs: - [Enum: Relu] @@ -106,5 +108,6 @@ BenchmarkProblems: - ProblemSizes: - Exact: [512, 512, 1, 512] - Exact: [1031, 1031, 1, 4103] + - Exact: [512, 512, 1, 8192] - ActivationArgs: - [Enum: Dgelu]