Skip to content

Commit ebd5db8

Browse files
alex391aalex391a
authored
feat(tensilelite): add gfx1250 Gate residual support (#11887)
JIRA ID : AIHPSPAR-274 ## Motivation Add gfx1250 support for Gate residual cache hints and correct Gate address updates for compact loop store (CLS). ## Technical Details - Add gfx1250 packed FMA mappings in stinkytofu. - Add TemporalHintGate and NonVolatileGate. - Fix Gate row-pointer and SRD updates for CLS. - Add FP8 and FP16 sparse GEMM test configurations. ## Test Plan Add `Tensile/Tests/common/sparse/gfx1250/spmm_gate.yaml` covering FP8 and FP16 Gate residual cases. ## Test Result <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. [AIHPSPAR-274]: https://amd-hub.atlassian.net/browse/AIHPSPAR-274?atlOrigin=eyJpIjoiNWRkNTljNzYxNjVmNDY3MDlhMDU5Y2ZhYzA5YTRkZjUiLCJwIjoiZ2l0aHViLWNvbS1KU1cifQ
1 parent 4e41ad0 commit ebd5db8

63 files changed

Lines changed: 469 additions & 268 deletions

File tree

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

‎projects/hipblaslt/tensilelite/Tensile/AsmAddressCalculation.py‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -195,7 +195,7 @@ def emitRowPtrAdvance(self, kernel, ss, tmpS01, rowInc, lookahead=False):
195195
sgpr(strideW1), rowInc, tmpS01, bcomment))
196196
if kw.vgprs.coutRowPtrGate != -1:
197197
module.add(self.addScaled(vgpr(kw.vgprs.coutRowPtrGate), vgpr(kw.vgprs.coutRowPtrGate), \
198-
sgpr("GateStride+0"), self.rowInc, tmpS01, "Move coutRowPtrGate to next row"))
198+
sgpr("GateStride+0"), rowInc, tmpS01, "Move coutRowPtrGate to next row"))
199199
elif len(kernel["PackedC1IndicesX"]) > 1:
200200
module.add(kw.extractPackedCoord1ToRowStart(kernel, kernel["PackedC1IndicesX"] , self.coord1Vgpr, 'D'))
201201
return module

‎projects/hipblaslt/tensilelite/Tensile/Common/GlobalParameters.py‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -557,6 +557,7 @@
557557
{"NonTemporal": [-1]},
558558
{"TemporalHint": [-1]},
559559
{"TemporalHintE": [0]},
560+
{"TemporalHintGate": [0]},
560561
{"TemporalHintD": [0]},
561562
{"TemporalHintC": [0]},
562563
{"TemporalHintA": [0]},
@@ -567,6 +568,7 @@
567568
{"TemporalHintMetadata": [0]},
568569
{"NonVolatile": [-1]},
569570
{"NonVolatileE": [0]},
571+
{"NonVolatileGate": [0]},
570572
{"NonVolatileD": [0]},
571573
{"NonVolatileC": [0]},
572574
{"NonVolatileA": [0]},

‎projects/hipblaslt/tensilelite/Tensile/Common/RequiredParameters.py‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -89,6 +89,7 @@ def getRequiredParametersMin() -> set:
8989
'NonTemporalC',
9090
'NonTemporalD',
9191
'NonTemporalE',
92+
'NonTemporalGate',
9293
'NonTemporalMetadata',
9394
'NonTemporalMXSA',
9495
'NonTemporalMXSB',
@@ -99,6 +100,7 @@ def getRequiredParametersMin() -> set:
99100
'NonVolatileC',
100101
'NonVolatileD',
101102
'NonVolatileE',
103+
'NonVolatileGate',
102104
'NonVolatileMetadata',
103105
'NonVolatileMXSA',
104106
'NonVolatileMXSB',
@@ -145,6 +147,7 @@ def getRequiredParametersMin() -> set:
145147
'TemporalHintC',
146148
'TemporalHintD',
147149
'TemporalHintE',
150+
'TemporalHintGate',
148151
'TemporalHintMetadata',
149152
'TemporalHintMXSA',
150153
'TemporalHintMXSB',

‎projects/hipblaslt/tensilelite/Tensile/Common/ValidParameters.py‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1054,6 +1054,7 @@ def makeValidMatrixInstructions():
10541054
# gfx1250-only temporal-hint modifier.
10551055
"TemporalHint": list(range(-1, 8)),
10561056
"TemporalHintE": list(range(0, 8)),
1057+
"TemporalHintGate": list(range(0, 8)),
10571058
"TemporalHintD": list(range(0, 8)),
10581059
"TemporalHintC": list(range(0, 8)),
10591060
"TemporalHintA": list(range(0, 8)),
@@ -1065,6 +1066,7 @@ def makeValidMatrixInstructions():
10651066
# gfx1250-only non-volatile memory modifier.
10661067
"NonVolatile": [-1, 0, 1],
10671068
"NonVolatileE": [0, 1],
1069+
"NonVolatileGate": [0, 1],
10681070
"NonVolatileD": [0, 1],
10691071
"NonVolatileC": [0, 1],
10701072
"NonVolatileA": [0, 1],

‎projects/hipblaslt/tensilelite/Tensile/Components/GlobalWriteBatch.py‎

Lines changed: 8 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -856,6 +856,8 @@ def _emitElt0LdsPreambleBeforeBanner(self, module: Module, bufferOOB,
856856
# Primer SGPRs for delayed incrementToNextRow (NonEdge / optSrdIncForRow).
857857
module.add(SMovB32(dst=sgpr(self.tmpS01), src=0, comment="Init sgpr offset"))
858858
module.add(SMovB32(dst=sgpr(self.tmpS01+1), src=0, comment="Init sgpr offset"))
859+
if self.parentWriter.states.useGateResidual:
860+
module.add(SMovB32(dst=sgpr("CLSGateRowInc"), src=0, comment="Init Gate CLS row offset"))
859861
if self.kernel["StoreRemapVectorWidth"] and self.kernel["CompactLoopStore"]:
860862
# Batch 0 checks out; later batches reuse parentWriter.compactLoopStoreVgpr.
861863
self.CompactLoopStoreVgpr = self.parentWriter.vgprPool.checkOut(1, tag="CompactLoopStoreVgpr_tmpVgpr")
@@ -1189,7 +1191,8 @@ def _prolog(self, module: Module):
11891191
gateLoadMod = self.parentWriter.readInput(
11901192
self.kernel, self.ss, 'Gate',
11911193
_prologLoadDtype,
1192-
addrCalc, vc0, dataGate, self.gwvw, addrGateVgpr, self.tmpS01)
1194+
addrCalc, vc0, dataGate, self.gwvw, addrGateVgpr, self.tmpS01, elementIdx, self.batchIdx,
1195+
overrideAfterPrimerRows=_emitOverrideRows)
11931196
_glTgt.add(gateLoadMod)
11941197
else:
11951198
# no-opt (edge) multi-dtype: per-dtype dispatcher per element (gate
@@ -1237,7 +1240,8 @@ def _prolog(self, module: Module):
12371240
(elementIdx == 0), self.tmpVgpr, tmpInrSgpr, addrGateVgpr, self.addrD, 0))
12381241
module.add(self.parentWriter.readInput(
12391242
self.kernel, self.ss, 'Gate', gDtype,
1240-
addrCalc, vc0, dataGate, self.gwvw, addrGateVgpr, self.tmpS01))
1243+
addrCalc, vc0, dataGate, self.gwvw, addrGateVgpr, self.tmpS01, elementIdx, self.batchIdx,
1244+
overrideAfterPrimerRows=_emitOverrideRows))
12411245
# Restore bpe/offset for the next branch in this elem.
12421246
self.parentWriter.states.bpeGate = _savedBpeGate
12431247
addrCalc.globalOffsetGate = _savedGlobalOffsetGate
@@ -1616,7 +1620,8 @@ def _emitLoadsFor(gDtype):
16161620
bufferOOB, (ei == 0), self.tmpVgpr, self.tmpSgpr, addrGateVgpr, self.addrD, 0))
16171621
module.add(self.parentWriter.readInput(
16181622
self.kernel, self.ss, 'Gate', gDtype, addrCalc, element[3], dataGate,
1619-
self.gwvw, addrGateVgpr, self.tmpS01))
1623+
self.gwvw, addrGateVgpr, self.tmpS01, ei, self.batchIdx,
1624+
overrideAfterPrimerRows=self._lookaheadRowInc(ei)))
16201625

16211626
if not multi:
16221627
# single-dtype: no GateType dispatch needed, just the loads.

‎projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -16436,6 +16436,8 @@ def globalWriteElementBatch(self, kernel, tPA, tPB, activation, \
1643616436
if kernel["CompactLoopStore"]:
1643716437
edgeModule.add(self.defineSgpr("CLSm0Base", 1))
1643816438
edgeModule.add(self.defineSgpr("CLSLoopCounter", 1))
16439+
if self.states.useGateResidual:
16440+
edgeModule.add(self.defineSgpr("CLSGateRowInc", 1))
1643916441

1644016442
# for storeRemap edge case, non-beta still can enable vector stores
1644116443
gwvw = vectorWidth
@@ -16716,6 +16718,8 @@ def _emit_batch_loop():
1671616718

1671716719
# Free dedicated CLS SGPRs after all batches / activation branches.
1671816720
if kernel["CompactLoopStore"]:
16721+
if self.states.useGateResidual:
16722+
edgeModule.add(self.undefineSgpr("CLSGateRowInc"))
1671916723
edgeModule.add(self.undefineSgpr("CLSLoopCounter"))
1672016724
edgeModule.add(self.undefineSgpr("CLSm0Base"))
1672116725

@@ -17304,6 +17308,9 @@ def readInput(self, kernel, ss, tc: str, dataType, addrCalc, vc0, data, gwvw, ad
1730417308
# CLS: seed the SRD chain on elt0/batch0. C uses tmpS01+1 so D's primer is kept.
1730517309
if (ss.optSrdIncForRow and (addrCalc.rowInc or (kernel["CompactLoopStore"] and elementIdx == 0 and batchIdx == 0))) and not isWorkspace:
1730617310
_stmp = (tmpS01 + 1) if (tc == 'C' and kernel["CompactLoopStore"]) else tmpS01
17311+
# Gate has its own stride and must preserve C/D's delayed increments.
17312+
if tc == 'Gate' and kernel["CompactLoopStore"]:
17313+
_stmp = "CLSGateRowInc"
1730717314
module.add(addrCalc.incrementToNextRow(kernel, tc, ss, _stmp, forceinitrow0=1, bpeType=bpeType,
1730817315
overrideAfterPrimerRows=overrideAfterPrimerRows))
1730917316

‎projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -341,8 +341,8 @@ def _validateStreamKMulticast(state, printRejectionReason, isaInfoMap):
341341
# module (Tensile/Tests/unit/test_validateParameterTypes.py) that imports
342342
# them from Solution.
343343

344-
_cacheHintTensors = ("A", "B", "C", "D", "E", "MXSA", "MXSB", "WS", "Metadata")
345-
_cacheHintLoadTensors = ("A", "B", "C", "E", "MXSA", "MXSB", "WS", "Metadata")
344+
_cacheHintTensors = ("A", "B", "C", "D", "E", "Gate", "MXSA", "MXSB", "WS", "Metadata")
345+
_cacheHintLoadTensors = ("A", "B", "C", "E", "Gate", "MXSA", "MXSB", "WS", "Metadata")
346346

347347
# Module-level collector that accumulates type mismatches across all Solution
348348
# instances during a build. Key is (param_name, actual_type_name,

‎projects/hipblaslt/tensilelite/Tensile/Tests/common/sparse/gfx1250/spmm_gate.yaml‎

Lines changed: 135 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,135 @@
1+
TestParameters:
2+
marks: [skip-gfx900, skip-gfx906, skip-gfx908, skip-gfx90a, skip-gfx940, skip-gfx941, skip-gfx942, skip-gfx950, skip-gfx1010, skip-gfx1011, skip-gfx1012, skip-gfx1030, skip-gfx1100, skip-gfx1101, skip-gfx1102, skip-gfx1200, skip-gfx1201]
3+
GlobalParameters:
4+
SyncsPerBenchmark: 0
5+
NumElementsToValidate: -1
6+
KernelTime: True
7+
BoundsCheck: 0
8+
PrintSolutionRejectionReason: True
9+
PrintWinnersOnly: True
10+
ValidationMaxToPrint: 4
11+
ValidationPrintValids: False
12+
Architecture: gfx1250
13+
14+
BenchmarkProblems:
15+
-
16+
- # ProblemType
17+
OperationType: GEMM
18+
DataType: f8
19+
DestDataType: s
20+
ComputeDataType: s
21+
HighPrecisionAccumulate: True
22+
TransposeA: False
23+
TransposeB: False
24+
UseBeta: True
25+
UseGateResidual: True
26+
GateResidualDataTypeList: ['s', 'f8']
27+
Batched: True
28+
Activation: True
29+
ActivationType: all
30+
UseBias: 3
31+
UseScaleAlphaVec: 3
32+
Sparse: 1
33+
MetadataLayout: 0
34+
35+
- # BenchmarkProblemSizeGroup - Standard
36+
InitialSolutionParameters:
37+
BenchmarkCommonParameters:
38+
- KernelLanguage: ["Assembly"]
39+
ForkParameters:
40+
- MatrixInstruction:
41+
- [16, 16, 128, 1, 1, 2, 2, 2, 2]
42+
- WavefrontSize: [32]
43+
- DepthU: [128]
44+
- PrefetchGlobalRead: [2]
45+
- PrefetchLocalRead: [1]
46+
- ClusterLocalRead: [1]
47+
- ScheduleIterAlg: [3]
48+
- TransposeLDS: [0]
49+
- SourceSwap: [false]
50+
- LdsPadA: [-1]
51+
- LdsPadB: [-1]
52+
- LdsPadMetadata: [-1]
53+
- GlobalSplitU: [1]
54+
- GlobalSplitUAlgorithm: [MultipleBuffer]
55+
- 1LDSBuffer: [-1]
56+
- DirectToVgprSparseMetadata: [false]
57+
- WorkGroupMapping: [1]
58+
- GlobalReadVectorWidthA: [-1]
59+
- GlobalReadVectorWidthB: [-1]
60+
- CompactLoopStore: [False, True]
61+
- NonTemporalGate: [0, 1]
62+
- TemporalHintGate: [0, 1]
63+
- NonVolatileGate: [0, 1]
64+
BenchmarkForkParameters:
65+
JoinParameters:
66+
BenchmarkJoinParameters:
67+
BenchmarkFinalParameters:
68+
- ProblemSizes:
69+
- Exact: [256, 256, 1, 256]
70+
- Exact: [264, 264, 1, 392]
71+
- Exact: [1024, 1024, 1, 1024]
72+
- BiasTypeArgs: ['s']
73+
- GateTypeArgs: ['s', 'f8']
74+
- FactorDimArgs: [0]
75+
- ActivationArgs:
76+
- [Enum: none]
77+
78+
-
79+
- # ProblemType
80+
OperationType: GEMM
81+
DataType: h
82+
DestDataType: h
83+
ComputeDataType: s
84+
HighPrecisionAccumulate: True
85+
TransposeA: False
86+
TransposeB: False
87+
UseBeta: True
88+
UseGateResidual: True
89+
GateResidualDataTypeList: ['h']
90+
Batched: True
91+
Activation: True
92+
ActivationType: all
93+
UseBias: 3
94+
UseScaleAlphaVec: 3
95+
Sparse: 1
96+
97+
- # BenchmarkProblemSizeGroup - Standard
98+
InitialSolutionParameters:
99+
BenchmarkCommonParameters:
100+
- KernelLanguage: ["Assembly"]
101+
ForkParameters:
102+
- MatrixInstruction:
103+
- [16, 16, 64, 1, 1, 2, 2, 1, 1]
104+
- WavefrontSize: [32]
105+
- DepthU: [128]
106+
- PrefetchGlobalRead: [2]
107+
- PrefetchLocalRead: [1]
108+
- ScheduleIterAlg: [3]
109+
- TransposeLDS: [-1]
110+
- SourceSwap: [false]
111+
- LdsPadA: [-1]
112+
- LdsPadB: [-1]
113+
- LdsPadMetadata: [-1]
114+
- GlobalSplitU: [1]
115+
- GlobalSplitUAlgorithm: [MultipleBuffer]
116+
- 1LDSBuffer: [-1]
117+
- GlobalReadVectorWidthA: [-1]
118+
- GlobalReadVectorWidthB: [-1]
119+
- CompactLoopStore: [False, True]
120+
- NonTemporalGate: [0, 1]
121+
- TemporalHintGate: [0, 1]
122+
- NonVolatileGate: [0, 1]
123+
BenchmarkForkParameters:
124+
JoinParameters:
125+
BenchmarkJoinParameters:
126+
BenchmarkFinalParameters:
127+
- ProblemSizes:
128+
- Exact: [256, 256, 1, 256]
129+
- Exact: [264, 264, 1, 392]
130+
- Exact: [1024, 1024, 1, 1024]
131+
- BiasTypeArgs: ['s']
132+
- GateTypeArgs: ['h']
133+
- FactorDimArgs: [0]
134+
- ActivationArgs:
135+
- [Enum: none]

0 commit comments

Comments
 (0)