Skip to content

fix(tensilelite): generate aux output E in persistent kernels #12869

New issue

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

Sign up for GitHub

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

Already on GitHub? Sign in to your account

Merged
jaopaulolc merged 9 commits into from
Oct 7, 2026
Merged

fix(tensilelite): generate aux output E in persistent kernels #12869

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
9 commits
Select commit Hold shift + click to select a range
74e2139
fix(tensilelite): generate aux output E in persistent kernels
jaopaulolc Oct 1, 2026
5bff6e3
Merge branch 'develop' into users/jolabega/fix-tensilelite-streamk-au…
jaopaulolc Oct 2, 2026
e49d700
test(tensilelite): cover Stream-K parallel reduction with aux output E
jaopaulolc Oct 5, 2026
878ab6b
Merge branch 'develop' into users/jolabega/fix-tensilelite-streamk-au…
jaopaulolc Oct 5, 2026
23cb111
Merge branch 'develop' into users/jolabega/fix-tensilelite-streamk-au…
jaopaulolc Oct 5, 2026
7cdcaab
Merge branch 'develop' into users/jolabega/fix-tensilelite-streamk-au…
jaopaulolc Oct 6, 2026
88df9d7
Merge branch 'develop' into users/jolabega/fix-tensilelite-streamk-au…
jaopaulolc Oct 6, 2026
df06c92
Merge branch 'develop' into users/jolabega/fix-tensilelite-streamk-au…
jaopaulolc Oct 6, 2026
75e73e5
Merge branch 'develop' into users/jolabega/fix-tensilelite-streamk-au…
jaopaulolc Oct 7, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 3 additions & 3 deletions projects/hipblaslt/tensilelite/Tensile/AsmAddressCalculation.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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))
Expand Down Expand Up @@ -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), \
Expand All @@ -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")
Expand Down
12 changes: 6 additions & 6 deletions projects/hipblaslt/tensilelite/Tensile/AsmStoreState.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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'):
Expand Down Expand Up @@ -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)
Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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:
Expand Down
5 changes: 3 additions & 2 deletions projects/hipblaslt/tensilelite/Tensile/Components/ComputeStoreVgprs.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,6 +29,7 @@

from ..Component import ComputeStoreVgprs
from ..Common import DataDirection, log2
from ..ExecutionPolicy import isPersistent

class ComputeStoreVgprsVALU(ComputeStoreVgprs):
kernel = {"EnableMatrixInstruction": False,
Expand Down Expand Up @@ -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"))
Expand Down Expand Up @@ -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"))
Expand Down
3 changes: 2 additions & 1 deletion projects/hipblaslt/tensilelite/Tensile/Components/LSU.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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"))
Expand Down
2 changes: 1 addition & 1 deletion projects/hipblaslt/tensilelite/Tensile/Components/StreamK.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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")
Comment thread
jaopaulolc marked this conversation as resolved.
Resolved
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
Expand Down
4 changes: 2 additions & 2 deletions projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7117,8 +7117,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.")

Expand Down
109 changes: 109 additions & 0 deletions ...slt/tensilelite/Tensile/Tests/common/streamk/gfx950/data_parallel_static_hgemm_use_e.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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]
113 changes: 113 additions & 0 deletions projects/hipblaslt/tensilelite/Tensile/Tests/common/streamk/gfx950/sk_hgemm_use_e.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,113 @@
# 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]
# Few tiles and a long K: parallel reduction with the post kernel
- Exact: [512, 512, 1, 8192]
- 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]
- Exact: [512, 512, 1, 8192]
- ActivationArgs:
- [Enum: Dgelu]
Loading