diff --git a/projects/hipblaslt/tensilelite/Tensile/Component.py b/projects/hipblaslt/tensilelite/Tensile/Component.py index 04b3b47179db..9639eced4bef 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Component.py +++ b/projects/hipblaslt/tensilelite/Tensile/Component.py @@ -218,13 +218,31 @@ class LocalRead(Component): """ Local read block. """ - def _getLdsReadMemToken(self, writer, kernel, tP): + def _getLdsReadMemToken(self, writer, kernel, tP, ldsByteOffset=None, bothHalves=False): from rocisa.container import MemTokenData - tok = writer.states.ldsReadTokenIdx + useSplit = (kernel["TDMSplit"] and not kernel["ProblemType"]["Sparse"] + and ldsByteOffset is not None and not tP.get("isM", False)) + if useSplit: + parity = writer.states.ldsReadTokenIdx + if bothHalves: + # Tile whose per-wave reads do not statically separate the two TDMSplit halves + # (numVectorsPerTile==1, e.g. B with MIWaveTile[N] == VectorWidth): a single read's + # combined region is fed by BOTH half tensor_loads, but the byte-offset classifier + # below can never reach the half1 boundary so it would tag everything half0 and the + # half1 load would go un-waited. Carry BOTH half tokens so StinkyTofu waits on the + # youngest of the two (fine-grained min(count-1), not a full drain). + toks = [writer.states.memTokenLdsSplit[parity][0], + writer.states.memTokenLdsSplit[parity][1]] + return MemTokenData(toks), toks[0] + inBuf = ldsByteOffset - tP["localReadSwapByteOffset"] + half = 1 if inBuf >= writer.tdmSplitLdsBoundary(kernel, tP) else 0 + tok = writer.states.memTokenLdsSplit[parity][half] + else: + tok = writer.states.ldsReadTokenIdx return MemTokenData([tok]), tok - def _emitLdsRead(self, writer, kernel, tP, LocalReadX, dst, src, ds, module, comment=""): - ldsMemToken, ldsMemTokenIdx = self._getLdsReadMemToken(writer, kernel, tP) + def _emitLdsRead(self, writer, kernel, tP, LocalReadX, dst, src, ds, module, ldsByteOffset=None, bothHalves=False, comment=""): + ldsMemToken, ldsMemTokenIdx = self._getLdsReadMemToken(writer, kernel, tP, ldsByteOffset, bothHalves) fullComment = "%s sync LDS%u" % (comment, ldsMemTokenIdx) if comment else "sync LDS%u" % ldsMemTokenIdx inst = LocalReadX(dst=dst, src=src, ds=ds, comment=fullComment) inst.setMemToken(ldsMemToken) diff --git a/projects/hipblaslt/tensilelite/Tensile/Components/LocalRead.py b/projects/hipblaslt/tensilelite/Tensile/Components/LocalRead.py index af2202094f4b..b9f4336e810a 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Components/LocalRead.py +++ b/projects/hipblaslt/tensilelite/Tensile/Components/LocalRead.py @@ -1652,6 +1652,7 @@ def calcGfx1250LdsOffset(): addrIdx = paramList[0] // 65536 srcAddr=vgpr("LocalReadAddr%s+%u"%(tc, addrIdx)) + tdmFullLdsOffset = paramList[0] paramList[0] -= addrIdx * 65536 if numOffsets == 1: @@ -1667,7 +1668,14 @@ def calcGfx1250LdsOffset(): # indexTranpose case, disable index conversion for local read destVgpr = self.getVgprForEmu(writer, kernel, tc, bufferIdx, iui, index, lrvwTile, vgprLen=numVgpr, dst=False, localRead=True) - self._emitLdsRead(writer, kernel, tP, LocalReadX, dst=destVgpr, src=srcAddr, ds=ds, module=localReadCodeT, comment=comment) + # When numVectorsPerTile==1 the per-wave reads never cross the TDMSplit + # half boundary (only vIdx=0 exists), so the byte-offset half classifier + # would tag every read half0 and leave the half1 tensor_load un-waited. + # Such reads' combined region depends on BOTH half loads -> carry both + # half tokens. + tdmBothHalves = (kernel["TDMSplit"] and not kernel["ProblemType"]["Sparse"] + and not tP.get("isM", False) and numVectorsPerTile == 1) + self._emitLdsRead(writer, kernel, tP, LocalReadX, dst=destVgpr, src=srcAddr, ds=ds, module=localReadCodeT, ldsByteOffset=tdmFullLdsOffset, bothHalves=tdmBothHalves, comment=comment) # TODO - handle vector-load with writer.allocTmpSgpr(1, tag="LocalReadVALU_tmpSgprInfo2") as tmpSgprInfo: tmpSgpr = tmpSgprInfo.idx diff --git a/projects/hipblaslt/tensilelite/Tensile/Components/SIA.py b/projects/hipblaslt/tensilelite/Tensile/Components/SIA.py index d0f0c0351a6f..f592bbf68346 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Components/SIA.py +++ b/projects/hipblaslt/tensilelite/Tensile/Components/SIA.py @@ -555,10 +555,17 @@ def noSchedGlobalRead(writer, kernel, globalReadIncACode, globalReadIncBCode): imod.add(nonTdmMod) imod.addComment1("Global Read MXSA") imod.add(writer.codes.dtlsM0UpdateMXSA) - imod.add(writer.codes.globalReadMXSA) + # MXSA/MXSB share A/B's XOR double-buffered MX-scale LDS, so their + # tensor_load_to_lds must defer in lockstep with A/B (after the LDS swap). + # Emitting them eagerly lets the async TDM overwrite the MX-scale LDS while + # this wave's own MX ds_loads are still reading it (self-wave WAR; s_wait_dscnt + # does not gate TDM, which retires on tensorcnt). + nonTdmModMXSA, tdmLoadModMXSA = _splitTdmLoad(writer.codes.globalReadMXSA) + imod.add(nonTdmModMXSA) imod.addComment1("Global Read MXSB") imod.add(writer.codes.dtlsM0UpdateMXSB) - imod.add(writer.codes.globalReadMXSB) + nonTdmModMXSB, tdmLoadModMXSB = _splitTdmLoad(writer.codes.globalReadMXSB) + imod.add(nonTdmModMXSB) imod.addComment1("Global Read B") imod.add(writer.codes.dtlsM0UpdateB) nonTdmModB, tdmLoadModB = _splitTdmLoad(writer.codes.globalReadB) @@ -569,6 +576,12 @@ def noSchedGlobalRead(writer, kernel, globalReadIncACode, globalReadIncBCode): if tdmLoadModA.itemsSize() > 0: deferMod.addComment1("Global Read A (TDM deferred after LDS swap)") deferMod.add(tdmLoadModA) + if tdmLoadModMXSA.itemsSize() > 0: + deferMod.addComment1("Global Read MXSA (TDM deferred after LDS swap)") + deferMod.add(tdmLoadModMXSA) + if tdmLoadModMXSB.itemsSize() > 0: + deferMod.addComment1("Global Read MXSB (TDM deferred after LDS swap)") + deferMod.add(tdmLoadModMXSB) if tdmLoadModB.itemsSize() > 0: # TODO: For the 1-wave case, schedule B's TDM load independently for better tensor load balance. deferMod.addComment1("Global Read B (TDM deferred after LDS swap)") diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py index a63a57da1e7c..33301ed420dc 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriter.py @@ -4078,6 +4078,8 @@ def noLoadLoop( self, kernel, tensorParametersA, tensorParametersB, isOptNLL, is module.add(self._wait(kernel, tensorParametersA, tensorParametersB, vlcntVal, -1, -1, "10wait for global read")) if not kernel["NoLdsWriteCode"]: module.add(self._wait(kernel, tensorParametersA, tensorParametersB, -1, 0, -1, "4wait for local write")) + elif kernel["enableTDMA"] and kernel["enableTDMB"]: + module.add(self._wait(kernel, tensorParametersA, tensorParametersB, 0, -1, -1, "wait for TDM tensor loads")) module.add(self._syncThreads(kernel, "wait for local write done, sync")) papPriorSync = True elif kernel["enableTDMA"] and kernel["enableTDMB"]: @@ -6238,7 +6240,7 @@ def _restoreNtabState(): module.addComment1("Recalc local read offsets") module.add(self.recalcLocalReadAddressesAB(kernel, tensorParametersA, tensorParametersB)) module.add(self._wait(kernel, tensorParametersA, tensorParametersB, -1, 0, -1, "5wait for local write")) - module.add(self._syncThreads(kernel, "Tail loop LW->LR, sync LDS0", memoryToken=[self.states.memTokenLdsBuffer0])) + module.add(self._syncThreads(kernel, "Tail loop LW->LR, sync LDS0", memoryToken=self._tailLoopBarrierTokens(kernel))) #module.add(self.dumpLds(kernel, 0, 8)) # tail: free G2L Vgpr @@ -6466,7 +6468,7 @@ def _restoreNtabState(): # generate wait code for early exit module.add(self._wait(kernel, tensorParametersA, tensorParametersB, 0, -1, -1, " tailloopInNll: wait for global read")) module.add(self._wait(kernel, tensorParametersA, tensorParametersB, -1, 0, -1, " tailloopInNll: wait for local read")) - module.add(self._syncThreads(kernel, "tailloopInNll: wait until LR done, sync LDS0", memoryToken=[self.states.memTokenLdsBuffer0])) + module.add(self._syncThreads(kernel, "tailloopInNll: wait until LR done, sync LDS0", memoryToken=self._tailLoopBarrierTokens(kernel))) if self.states.lastValuMXSAB: self.vgprPool.add(0 , self.states.lastValuMXSAB, "ValuMXSAB") @@ -7183,9 +7185,12 @@ def initSubTileInfo(tc): if kernel["1LDSBuffer"]: self.states.memTokenLdsBuffer0 = 0 self.states.memTokenLdsBuffer1 = 0 + self.states.memTokenLdsSplit = [[1, 2], [1, 2]] else: self.states.memTokenLdsBuffer0 = 0 self.states.memTokenLdsBuffer1 = 1 + # [bufferParity][half]; half-0 reuses buffer token so MXSA/MXSB share it + self.states.memTokenLdsSplit = [[self.states.memTokenLdsBuffer0, 2], [self.states.memTokenLdsBuffer1, 6]] self.states.ldsReadTokenIdx = self.states.memTokenLdsBuffer0 self.states.ldsTensorTokenIdx = self.states.memTokenLdsBuffer0 self.states.ldsDirectToLDSTokenIdx = self.states.memTokenLdsBuffer0 @@ -10400,6 +10405,14 @@ def _syncThreads(self, kernel, comment="", skipForceWaitcnt0=False, memoryToken= return syncThreads(kernel, self.states.archCaps, self.states.asmCaps, comment, skipForceWaitcnt0=skipForceWaitcnt0, memoryToken=memoryToken) return Module("SyncThreads (Empty)") + def _tailLoopBarrierTokens(self, kernel): + # Explicit tail barriers must list every token; the auto-barrier pass doesn't cover them. + tokens = [self.states.memTokenLdsBuffer0] + if kernel["TDMSplit"] and not kernel["ProblemType"]["Sparse"]: + for row in self.states.memTokenLdsSplit: + tokens.extend(row) + return sorted(set(tokens)) + ############################################################################## # PostMainLoopBarrierCheckAndReset # Rebuild SBarrier placement by memory-token read/write phase transitions. diff --git a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py index eb25ac4b68e4..20cebee14d3e 100644 --- a/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py +++ b/projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py @@ -976,17 +976,21 @@ def defineTdmSgprs(self, kernel): module.add(self.defineSgpr("tdmABIncs", 1)) if kernel["TDMSplit"] and not kernel["ProblemType"]["Sparse"]: - if kernel["NumWaves"] > 1: - module.add(self.defineSgpr("tdmABGlobalSplitIncs", 1)) - module.add(self.defineSgpr("tdmABLdsSplitIncs", 1)) - else: - module.add(self.defineSgpr("tdmAGlobalSplitIncs", 1)) - module.add(self.defineSgpr("tdmALdsSplitIncs", 1)) - module.add(self.defineSgpr("tdmBGlobalSplitIncs", 1)) - module.add(self.defineSgpr("tdmBLdsSplitIncs", 1)) + module.add(self.defineSgpr("tdmABGlobalSplitIncs", 1)) + module.add(self.defineSgpr("tdmABLdsSplitIncs", 1)) + # dim1 half boundaries (H0/H1) are recomputed in globalReadDo, not persisted. if kernel["ProblemType"]["MXBlockA"] and kernel["ProblemType"]["MXBlockB"]: module.add(self.defineSgpr("tdmMXSAMXSBIncs", 1)) + + # Single-wave (NumWaves == 1): descriptors are independent, so TDMSplit + # uses per-tensor increment SGPRs instead of the aliased AB pair. + if (kernel["enableTDMA"] and kernel["enableTDMB"] and kernel["NumWaves"] == 1 + and kernel["TDMSplit"] and not kernel["ProblemType"]["Sparse"]): + module.add(self.defineSgpr("tdmAGlobalSplitIncs", 1)) + module.add(self.defineSgpr("tdmALdsSplitIncs", 1)) + module.add(self.defineSgpr("tdmBGlobalSplitIncs", 1)) + module.add(self.defineSgpr("tdmBLdsSplitIncs", 1)) if kernel["enableTDMMetadata"]: module.add(self.defineSgpr("tdmMetadataGroup0", 4, 4)) @@ -11116,7 +11120,12 @@ def globalReadDo(self, kernel, mode, tP, unrollLoopIdx=-1, g2lBufIdx=0, \ if tc == "A" and kernel["enableTDMA"]: comp: TensorDataMoverLoad = TensorDataMoverLoad.find(self) - comp.setMemToken([self.states.ldsTensorTokenIdx]) + useSplitTokens = bool(kernel["TDMSplit"]) and not kernel["ProblemType"]["Sparse"] + tdmParity = self.states.ldsTensorTokenIdx + if useSplitTokens: + comp.setMemToken([self.states.memTokenLdsSplit[tdmParity][0]]) + else: + comp.setMemToken([self.states.ldsTensorTokenIdx]) if self.states.inTailLoop and not kernel["1LDSBuffer"] and kernel["StreamK"]: ldsAddrSgprName = comp.getLdsAddrSgprName("tdmAGroup0") if self.isPrefetchAcrossPersistentEnabled(kernel): @@ -11148,7 +11157,39 @@ def globalReadDo(self, kernel, mode, tP, unrollLoopIdx=-1, g2lBufIdx=0, \ imod.middle.add(SAddU32(sgpr(f"tdm{tc}Group0+1"), sgpr(f"tdm{tc}Group0+1"), sgpr(ldsIncSgprName))) imod.middle.add(SAddU32(sgpr(f"tdm{tc}Group0+2"), sgpr(f"tdm{tc}Group0+2"), sgpr(globalIncSgprName))) imod.middle.add(SAddCU32(sgpr(f"tdm{tc}Group0+3"), sgpr(f"tdm{tc}Group0+3"), 0, f"tdm{tc} split carry")) - imod.middle.add(comp.issueLoad("tdmAGroup0", "tdmAGroup1", tdmAGroup2, tdmAGroup3)) + if numWaves > 1: + # Recompute the two per-half dim1 boundaries from the live descriptor + # instead of persisting them (saves 2 SGPRs). halfRows is per-wave: even + # waves load A (MacroTile0//2), odd load B (MacroTile1//2). + group1 = f"tdm{tc}Group1" + halfRowsA = kernel["MacroTile0"] // 2 + halfRowsB = kernel["MacroTile1"] // 2 + wavelen = kernel["WavefrontSize"] + with self.allocTmpSgpr(3, tag="tdmSplitDim1Recompute") as tmpSgprRes: + h0 = tmpSgprRes.idx + h1 = tmpSgprRes.idx + 1 + hr = tmpSgprRes.idx + 2 + if halfRowsA == halfRowsB: + imod.middle.add(SMovB32(sgpr(hr), halfRowsA, "halfRows")) + else: + # WaveIdx is not live in the loop; recompute parity from Serial. + with self.allocTmpSgpr(1, tag="tdmSplitDim1Parity") as waveIdTmp: + imod.middle.add(VReadfirstlaneB32(dst=sgpr(waveIdTmp.idx), src=vgpr("Serial"), comment="get tId")) + imod.middle.add(SLShiftRightB32(dst=sgpr(waveIdTmp.idx), shiftHex=ceil(log2(wavelen)), src=sgpr(waveIdTmp.idx), comment="waveId")) + imod.middle.add(SBitcmp1B32(src0=sgpr(waveIdTmp.idx), src1=0, comment="wave parity")) + imod.middle.add(SCSelectB32(sgpr(hr), halfRowsB, halfRowsA, "halfRows = parity ? B : A")) + imod.middle.add(SLShiftRightB32(sgpr(h0), hex(16), sgpr(f"{group1}+2"), "H0 = dim1 lo")) + imod.middle.add(SLShiftLeftB32(sgpr(h1), hex(16), sgpr(f"{group1}+3"), "H0 hi << 16")) + imod.middle.add(SOrB32(sgpr(h0), sgpr(h0), sgpr(h1), "H0 = full dim1")) + imod.middle.add(SSubU32(sgpr(h1), sgpr(h0), sgpr(hr), "H1 = H0 - halfRows")) + imod.middle.add(SCSelectB32(sgpr(h1), 0, sgpr(h1), "clamp H1 to 0")) + imod.middle.add(comp.setTensorDim1(group1, h1, self)) + comp.setMemToken([self.states.memTokenLdsSplit[tdmParity][1]]) + imod.middle.add(comp.issueLoad("tdmAGroup0", "tdmAGroup1", tdmAGroup2, tdmAGroup3)) + imod.middle.add(comp.setTensorDim1(group1, h0, self)) + else: + comp.setMemToken([self.states.memTokenLdsSplit[tdmParity][1]]) + imod.middle.add(comp.issueLoad("tdmAGroup0", "tdmAGroup1", tdmAGroup2, tdmAGroup3)) return imod if tc == "MXSA" and kernel["enableTDMA"]: @@ -11179,7 +11220,12 @@ def globalReadDo(self, kernel, mode, tP, unrollLoopIdx=-1, g2lBufIdx=0, \ #TODO: TDM refactor, wave separated TDM only issues 1 tensor load if numWaves == 1: comp: TensorDataMoverLoad = TensorDataMoverLoad.find(self) - comp.setMemToken([self.states.ldsTensorTokenIdx]) + useSplitTokens = bool(kernel["TDMSplit"]) and not kernel["ProblemType"]["Sparse"] + tdmParity = self.states.ldsTensorTokenIdx + if useSplitTokens: + comp.setMemToken([self.states.memTokenLdsSplit[tdmParity][0]]) + else: + comp.setMemToken([self.states.ldsTensorTokenIdx]) if self.states.inTailLoop and not kernel["1LDSBuffer"] and kernel["StreamK"]: ldsAddrSgprName = comp.getLdsAddrSgprName("tdmBGroup0") if self.isPrefetchAcrossPersistentEnabled(kernel): @@ -11207,6 +11253,7 @@ def globalReadDo(self, kernel, mode, tP, unrollLoopIdx=-1, g2lBufIdx=0, \ imod.middle.add(SAddU32(sgpr(f"tdm{tc}Group0+1"), sgpr(f"tdm{tc}Group0+1"), sgpr(ldsIncSgprName))) imod.middle.add(SAddU32(sgpr(f"tdm{tc}Group0+2"), sgpr(f"tdm{tc}Group0+2"), sgpr(globalIncSgprName))) imod.middle.add(SAddCU32(sgpr(f"tdm{tc}Group0+3"), sgpr(f"tdm{tc}Group0+3"), 0, f"tdm{tc} split carry")) + comp.setMemToken([self.states.memTokenLdsSplit[tdmParity][1]]) imod.middle.add(comp.issueLoad("tdmBGroup0", "tdmBGroup1", tdmBGroup2, tdmBGroup3)) return imod @@ -18682,6 +18729,19 @@ def _emitTdmIterateInit(self, mod, kernel, tc, dtype, du, mt, perIssueLoadRowDiv f"TDM iter_count = rows_per_il({rows_per_il}) / tile_dim1({tile_dim1}) - 1")) mod.add(comp.setIterations(descSgprName(2), sIter)) + def tdmSplitLdsBoundary(self, kernel: Mapping, tP: Mapping) -> int: + tc: str = tP['tensorChar'] + ti: int = tP["idx"] + mt: int = kernel[f"MacroTile{ti}"] + du: int = kernel["DepthU"] + bpe: float = tP["bpeGR"] if not tP["isM"] else 1 + dim1Divisor = 2 + ldsBlockSizePerPad: int = kernel[f"LdsBlockSizePerPad{tc}"] + ldsPadSize: int = int(kernel[f"LdsPad{tc}"] * bpe) + half = round(mt * du * bpe // dim1Divisor) + extraPadSize = half // ldsBlockSizePerPad * ldsPadSize if ldsBlockSizePerPad != 0 and ldsPadSize != 0 else 0 + return half + extraPadSize + def initTDMDescriptor(self, kernel: Mapping, tP: Mapping) -> Module: comp: TensorDataMoverLoad = TensorDataMoverLoad.find(self) tc: str = tP['tensorChar'] @@ -18822,8 +18882,8 @@ def maskSgprName(tc: str) -> str: mod.add(comp.setTensorStride0(descSgprName(1), strideRefName(), sizeShifter)) if (kernel["TDMSplit"] and not ("MXS" in tc) and not kernel["ProblemType"]["Sparse"]): - extraPadSize: int = round(mt * du * bpe // dim1Divisor) // ldsBlockSizePerPad * ldsPadSize if ldsBlockSizePerPad != 0 and ldsPadSize != 0 else 0 - mod.add(SMovB32(sgpr(f"tdm{tc}LdsSplitIncs"), round(mt * du * bpe // dim1Divisor) + extraPadSize, comment=f"tdm{tc} Lds Split Incs({mt * du * bpe // dim1Divisor})")) + splitBoundary: int = self.tdmSplitLdsBoundary(kernel, tP) + mod.add(SMovB32(sgpr(f"tdm{tc}LdsSplitIncs"), splitBoundary, comment=f"tdm{tc} Lds Split Incs({round(mt * du * bpe // dim1Divisor)})")) mod.add(SMulI32(sgpr(f"tdm{tc}GlobalSplitIncs"), strideRefName(), round(mt * bpe) // dim1Divisor, comment=f"tdm{tc} Global Split Incs(stride * {mt * bpe // dim1Divisor})")) if isTdmIter: @@ -18971,10 +19031,12 @@ def maskSgprName(tc: str) -> str: if unrolledMajor: mod.add(VReadfirstlaneB32(sgpr(tmpSgprWaveOffset), vgpr("Serial"), "first tId")) mod.add(SLShiftRightB32(sgpr(tmpSgprWaveOffset), ceil(log2(wavelen)) + 1, sgpr(tmpSgprWaveOffset), "wId=fTid // wavelen // 2")) - mod.add(SMulI32(sgpr(tmpSgprWaveOffset), sgpr(tmpSgprWaveOffset), round(mt // numComp), "woffset = wId * (mt // numComp)")) + mod.add(SMulI32(sgpr(tmpSgprWaveOffset), sgpr(tmpSgprWaveOffset), round(mt // numComp // dim1Divisor), "woffset = wId * (mt // numComp // dim1Divisor)")) mod.add(SSubU32(sgpr(dim1), sgpr(dim1), sgpr(tmpSgprWaveOffset), "consider multiple waves")) mod.add(SCMovB32(sgpr(dim1), 0, "set to 0 for waves that no enough data to load")) mod.add(comp.setTensorDim1(descSgprName(1), dim1, self, 0, False, isSparseTrack if not unrolledMajor else False, isMetadata if not unrolledMajor else False)) + # The descriptor now holds the full per-wave dim1; TDMSplit recomputes + # the half boundaries from it in globalReadDo. if tc.startswith("MX"): #reset to 0 since scale of sizeTile0 and stride for MX is not required @@ -19013,12 +19075,11 @@ def maskSgprName(tc: str) -> str: mod.add(comp.setTensorStride0(descSgprName(1), strideRefName(), sizeShifter)) if (kernel["TDMSplit"] and not ("MXS" in tc) and not kernel["ProblemType"]["Sparse"]): - extraPadSize: int = round(mt * du * bpe // dim1Divisor) // ldsBlockSizePerPad * ldsPadSize if ldsBlockSizePerPad != 0 and ldsPadSize != 0 else 0 - mod.add(SMovB32(sgpr("tdmABLdsSplitIncs"), round(mt * du * bpe // dim1Divisor) + extraPadSize, comment=f"tdm{tc} Lds Split Incs({mt * du * bpe // dim1Divisor})")) + splitBoundary: int = self.tdmSplitLdsBoundary(kernel, tP) + mod.add(SMovB32(sgpr("tdmABLdsSplitIncs"), splitBoundary, comment=f"tdm{tc} Lds Split Incs({round(mt * du * bpe // dim1Divisor)})")) strRefSplit = strideRefName() strArgSplit = strRefSplit if isinstance(strRefSplit, RegisterContainer) else sgpr(strRefSplit) mod.add(SMulI32(sgpr("tdmABGlobalSplitIncs"), strArgSplit, round(mt * bpe) // dim1Divisor, comment=f"tdm{tc} Global Split Incs(stride * {mt * bpe // dim1Divisor})")) - if isTdmIter: self._emitTdmIterateInit(mod, kernel, tc, dtype, du, mt, perIssueLoadRowDivisor=numComp * dim1Divisor, 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 61a090a178e8..71885689a4f5 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 @@ -141,7 +141,6 @@ BenchmarkProblems: - DepthU: [256] - TDMInst: [3] - TDMIterateMode: [-1,0,1,2,3] - - TDMSplit: [False, True] - TransposeLDS: [-1] - LdsPadA: [-1] - LdsPadB: [-1] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/tdm_split_gfx1250.yaml b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/tdm_split_gfx1250.yaml new file mode 100644 index 000000000000..36c3aae62720 --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/tdm_split_gfx1250.yaml @@ -0,0 +1,254 @@ +TestParameters: + 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] # not supported by arch +GlobalParameters: + SyncsPerBenchmark: 0 + NumElementsToValidate: -1 + KernelTime: False + BoundsCheck: 0 + PrintSolutionRejectionReason: True + ValidationMaxToPrint: 4 + ValidationPrintValids: False + PruneSparseMode: 0 + DataInitTypeA: 3 + DataInitTypeB: 3 + DataInitTypeD: 2 + DataInitTypeMXSA: 3 + DataInitTypeMXSB: 3 + DataInitTypeAlpha: 1 + DataInitTypeBeta: 0 + NumWarmups: 0 + EnqueuesPerSync: 1 + PrintWinnersOnly: True + Architecture: gfx1250 + KeepBuildTmp: True + PrintLevel: 2 + + +BenchmarkProblems: + ####################################### + # TN - BF16, TDMSplit + ####################################### + - + - # ProblemType, TN + OperationType: GEMM + DataType: b + DestDataType: b + ComputeDataType: s + HighPrecisionAccumulate: True + TransposeA: True + TransposeB: False + UseBeta: True + Batched: True + Activation: False + SupportUserArgs: True + UseBias: 0 + + - # BenchmarkProblemSizeGroup - BF16 TDMSplit + InitialSolutionParameters: + BenchmarkCommonParameters: + - KernelLanguage: ["Assembly"] + ForkParameters: + - UseSgprForGRO: [0] + - ForceDisableShadowInit: [True] + - MatrixInstruction: + - [16, 16, 32, 1, 1, 1, 1, 1, 1] + - [16, 16, 32, 1, 1, 1, 1, 2, 2] + - [16, 16, 32, 1, 1, 8, 4, 2, 2] + - DepthU: [64, 128] + - TDMInst: [3] + - TDMIterateMode: [0] + - TDMSplit: [False, True] + - TransposeLDS: [-1] + - LdsPadA: [-1] + - LdsPadB: [-1] + - LdsBlockSizePerPadA: [-1] + - LdsBlockSizePerPadB: [-1] + - LdsPadMetadata: [0] + - LDSTrInst: [False] + - StaggerU: [0] + - ScheduleIterAlg: [0, 4] + - PrefetchLocalRead: [1] + - PrefetchGlobalRead: [2] + - StoreRemapVectorWidth: [0] + - GlobalSplitU: [1] + - GlobalSplitUAlgorithm: [MultipleBuffer] + - GlobalReadVectorWidthA: [-1] + - GlobalReadVectorWidthB: [-1] + - VectorWidthA: [-1] + - VectorWidthB: [-1] + - LocalReadVectorWidth: [8] + - 1LDSBuffer: [0] + - SourceSwap: [False] + - DirectToVgprSparseMetadata: [False] + - WorkGroupMapping: [1] + - StoreVectorWidth: [-1] + - DirectToVgprA: [False] + - DirectToVgprB: [False] + - WavefrontSize: [32] + - ExpandPointerSwap: [False] + BenchmarkForkParameters: + JoinParameters: + BenchmarkJoinParameters: + BenchmarkFinalParameters: + - ProblemSizes: + - Exact: [128, 128, 1, 128] + - Exact: [128, 128, 1, 192] #tail + - Exact: [129, 128, 1, 128] #edge-m + - Exact: [128, 129, 1, 128] #edge-n + - ActivationArgs: + - [Enum: none] + ####################################### + # TN - MXF8SS Block16, E8, TDMSplit + ####################################### + - + - # ProblemType, TN + OperationType: GEMM + DataType: F8 + DestDataType: s + ComputeDataType: s + HighPrecisionAccumulate: True + TransposeA: True + TransposeB: False + UseBeta: True + Batched: True + Activation: False + SupportUserArgs: True + UseBias: 0 + MXBlockA: 16 + MXBlockB: 16 + DataTypeMXSA: e8 + DataTypeMXSB: e8 + + - # BenchmarkProblemSizeGroup - MXF8 Block16 TDMSplit + InitialSolutionParameters: + BenchmarkCommonParameters: + - KernelLanguage: ["Assembly"] + ForkParameters: + - UseSgprForGRO: [0] + - ForceDisableShadowInit: [True] + - 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] + - TDMSplit: [False] + - TransposeLDS: [-1] + - LdsPadA: [-1] + - LdsPadB: [-1] + - LdsBlockSizePerPadA: [-1] + - LdsBlockSizePerPadB: [-1] + - LdsPadMetadata: [0] + - LDSTrInst: [False] + - StaggerU: [0] + - ScheduleIterAlg: [0, 4] + - PrefetchLocalRead: [1] + - PrefetchGlobalRead: [2] + - StoreRemapVectorWidth: [0] + - GlobalSplitU: [1] + - GlobalSplitUAlgorithm: [MultipleBuffer] + - GlobalReadVectorWidthA: [-1] + - GlobalReadVectorWidthB: [-1] + - VectorWidthA: [-1] + - VectorWidthB: [-1] + - LocalReadVectorWidth: [-1] + - 1LDSBuffer: [0] + - SourceSwap: [False] + - DirectToVgprSparseMetadata: [False] + - WorkGroupMapping: [1] + - StoreVectorWidth: [-1] + - DirectToVgprA: [False] + - DirectToVgprB: [False] + - WavefrontSize: [32] + - ExpandPointerSwap: [False] + - AssertSummationElementMultiple: [16] + BenchmarkForkParameters: + JoinParameters: + BenchmarkJoinParameters: + BenchmarkFinalParameters: + - ProblemSizes: + - Exact: [256, 256, 1, 2048] + - Exact: [256, 256, 1, 2064] #tail + - Exact: [257, 256, 1, 2048] #edge-m + - Exact: [256, 257, 1, 2048] #edge-n + - ActivationArgs: + - [Enum: none] + ####################################### + # TN - MXF8SS Block32, E8, TDMSplit + ####################################### + - + - # ProblemType, TN + OperationType: GEMM + DataType: F8 + DestDataType: s + ComputeDataType: s + HighPrecisionAccumulate: True + TransposeA: True + TransposeB: False + UseBeta: True + Batched: True + Activation: False + SupportUserArgs: True + UseBias: 0 + MXBlockA: 32 + MXBlockB: 32 + DataTypeMXSA: e8 + DataTypeMXSB: e8 + + - # BenchmarkProblemSizeGroup - MXF8 Block32 TDMSplit + InitialSolutionParameters: + BenchmarkCommonParameters: + - KernelLanguage: ["Assembly"] + ForkParameters: + - UseSgprForGRO: [0] + - ForceDisableShadowInit: [True] + - 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, True] + - TransposeLDS: [-1] + - LdsPadA: [-1] + - LdsPadB: [-1] + - LdsBlockSizePerPadA: [-1] + - LdsBlockSizePerPadB: [-1] + - LdsPadMetadata: [0] + - LDSTrInst: [False] + - StaggerU: [0] + - ScheduleIterAlg: [0, 4] + - PrefetchLocalRead: [1] + - PrefetchGlobalRead: [1, 2] + - StoreRemapVectorWidth: [0] + - GlobalSplitU: [1] + - GlobalSplitUAlgorithm: [MultipleBuffer] + - GlobalReadVectorWidthA: [-1] + - GlobalReadVectorWidthB: [-1] + - VectorWidthA: [-1] + - VectorWidthB: [-1] + - LocalReadVectorWidth: [-1] + - 1LDSBuffer: [0] + - SourceSwap: [False] + - DirectToVgprSparseMetadata: [False] + - WorkGroupMapping: [1] + - StoreVectorWidth: [-1] + - DirectToVgprA: [False] + - DirectToVgprB: [False] + - WavefrontSize: [32] + - ExpandPointerSwap: [False] + - AssertSummationElementMultiple: [32] + BenchmarkForkParameters: + JoinParameters: + BenchmarkJoinParameters: + BenchmarkFinalParameters: + - ProblemSizes: + - Exact: [256, 256, 1, 2048] + - Exact: [256, 256, 1, 2080] #tail + - Exact: [257, 256, 1, 2048] #edge-m + - Exact: [256, 257, 1, 2048] #edge-n + - ActivationArgs: + - [Enum: none] diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/Component/test_component_char.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/Component/test_component_char.py index af5de9b57fe5..70b95649b935 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/Component/test_component_char.py +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/characterization/Component/test_component_char.py @@ -165,7 +165,7 @@ def test_findall_recurses_into_abstract(snapshot): def test_localread_get_lds_read_mem_token(): writer = types.SimpleNamespace(states=types.SimpleNamespace(ldsReadTokenIdx=3)) - token, idx = LocalRead._getLdsReadMemToken(types.SimpleNamespace(), writer, None, None) + token, idx = LocalRead._getLdsReadMemToken(types.SimpleNamespace(), writer, {"TDMSplit": False}, None) assert idx == 3 @@ -182,7 +182,7 @@ def setMemToken(self, t): module = types.SimpleNamespace(add=added.append) fake_self = types.SimpleNamespace( - _getLdsReadMemToken=lambda w, k, t: ("TOKEN", 5) + _getLdsReadMemToken=lambda w, k, t, u, v: ("TOKEN", 5) ) # comment set -> the "%s sync LDS%u" branch. LocalRead._emitLdsRead(fake_self, None, None, None, _Inst, "dst", "src", "ds",