Add TLU support for DirectToVgpr via global_load_tr

projects/hipblaslt/tensilelite/Tensile/Components/LocalRead.py

@@ -772,8 +772,12 @@ def __call__(self, writer, kernel, bufferIdx, iui, epsi, tP):
 
                                 elif kernel["ProblemType"]["DataType"].isSingle():
                                     localReadCode.add(writer.assert_eq( dbgVgpr, 1.0) )
-        # DTV case, do not return local read code. Return pack code only.
+        # DTV case, do not return local read code.
         if (tP["isA"] or tP["isB"]) and kernel["DirectToVgpr%s"%tc]:
-          imod = Module("LocalReadDo%s_I%s (Empty)" % (tP["tensorChar"],iui))
+            imod = Module("LocalReadDo%s_I%s (Empty)" % (tP["tensorChar"],iui))
+
+        # DTV and Tr Load case, do not return pack code
+        if (tP["isA"] or tP["isB"]) and kernel["enableGLTr%s"%tc]:
+            pack = Module("Pack%s_I%s (Empty)" % (tP["tensorChar"],iui))
 
         return imod, pack

projects/hipblaslt/tensilelite/Tensile/Components/SIA.py

@@ -369,7 +369,7 @@ def getNumLocalWritePerMfma(writer, kernel, lwStartMfmaIndex):
     #     mfma--24-mfma--24-mfma--24-mfma--24-mfma--24-mfma--24-mfma--24-mfma--24-mfma
     # we need LWPM to get precise LWPM
     # so we iterate formula 10 times to get LWPM
-    while oldValue != newValue and loop < 10:
+    while oldValue != newValue and newValue > 0 and loop < 10:
         loop += 1
         oldValue = newValue
         newValue = roundUp((writesToSched+1 + (oldValue - (writesToSched+1) % oldValue) + oldValue%PRECISION) / numMfmaCanSched)

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

@@ -456,13 +456,16 @@ def setGlobalLoadTileDimClassic(state, tc, numLoads, totalVectorsCoalesced, tota
     if state["WaveSeparateGlobalRead%s"%tc]:
       totalElementsPerp = roundupRatio(totalElementsPerp, state["NumThreads"] // state["WavefrontSize"])
 
+    if (tc == "A" or tc == "B") and state["enableGLTr%s"%tc]:
+      totalVectorsCoalesced, totalElementsPerp = totalElementsPerp, totalVectorsCoalesced
+
     # nlc = 1
-    if state["NumLoadsCoalesced%s"%tc] == 1 :
+    if state["NumLoadsCoalesced%s"%tc] == 1:
       foundValid = False
       nlcStart = 1
       if (tc == "A" or tc == "B") and state["DirectToVgpr%s"%tc]:
         # adjust nlc for DirectToVgprA/B
-        if state["ProblemType"]["TLU%s"%tc]:
+        if state["ProblemType"]["TLU%s"%tc] and not state["enableGLTr%s"%tc]:
           nlcStart = roundupRatio(state["MIWaveTile%s"%tc], state["GlobalReadVectorWidth%s"%tc])
         else:
           nlcStart = roundupRatio(depthU, state["MatrixInstK"] * state["GlobalReadVectorWidth%s"%tc] * state["LocalSplitU"] // state["MIInputPerThread"])
@@ -522,6 +525,9 @@ def setGlobalLoadTileDimClassic(state, tc, numLoads, totalVectorsCoalesced, tota
               % (tc, totalElementsPerp, state["NumLoadsPerpendicular%s"%tc]))
         return False
 
+    if (tc == "A" or tc == "B") and state["enableGLTr%s"%tc]:
+      state["NumLoadsCoalesced%s"%tc], state["NumLoadsPerpendicular%s"%tc] = state["NumLoadsPerpendicular%s"%tc], state["NumLoadsCoalesced%s"%tc]
+
     if state["ProblemType"]["TLU%s"%tc]:
       state["LSC%s"%tc] = state["MacroTile%s"%tc] // state["NumLoadsCoalesced%s"%tc]
       state["LSP%s"%tc] = int(math.ceil(float(depthU) / state["NumLoadsPerpendicular%s"%tc]))
@@ -607,8 +613,10 @@ def isDirectToVgprDoable(state, tc, printRejectionReason: bool, isaInfoMap: Dict
       state["PrefetchLocalRead"] = 0
       # So far, DTVA + DTVB does not perform well (waitcnt is not ideal).
       # Disable it for now (TODO: improve waitcnt and re-enable)
-      reject(state, printRejectionReason, "DirectToVgprA + DirectToVgprB disabled")
-      return False
+      # Exception: allow if enableGLTrA or enableGLTrB are set; this still provides better performance than enabling only one side
+      if not (state["enableGLTrA"] or state["enableGLTrB"]):
+        reject(state, printRejectionReason, "DirectToVgprA + DirectToVgprB disabled")
+        return False
 
     # DTV + input type conversion
     if state["ProblemType"]["DataType%s"%tc] != state["ProblemType"]["DataType"]:
@@ -622,7 +630,7 @@ def isDirectToVgprDoable(state, tc, printRejectionReason: bool, isaInfoMap: Dict
       return False
 
     # Does not work with TLU = False and PrefetchLocalRead = 0
-    if (not state["ProblemType"]["TLU%c"%tc]) and state["PrefetchLocalRead"] == 0:
+    if (not state["ProblemType"]["TLU%c"%tc]) and state["PrefetchLocalRead"] == 0 and not (state["DirectToVgprA"] and state["DirectToVgprB"]):
       reject(state, printRejectionReason, "DirectToVgpr%c does not supports TLU%c = False and PrefetchLocalRead = 0"%(tc, tc))
       return False
 
@@ -639,7 +647,7 @@ def isDirectToVgprDoable(state, tc, printRejectionReason: bool, isaInfoMap: Dict
 
     if numBytes < 4:
       # numBytes < 4 case
-      if state["ProblemType"]["TLU%c"%tc]:
+      if state["ProblemType"]["TLU%c"%tc] and not state["enableGLTr%c"%tc]:
         # use pack logic (with v_perm) same as local read (only if VgprForLocalReadPacking is doable)
         if not Solution.isVgprForLocalReadPackingDoable(state, isaInfoMap):
           reject(state, printRejectionReason, "Does not meet the requirement for DirectToVgpr%c + TLU%c + numByte < 4"%(tc, tc))
@@ -670,7 +678,7 @@ def isDirectToVgprDoable(state, tc, printRejectionReason: bool, isaInfoMap: Dict
       return False
 
     # Does not work with TLU + VectorWidth != GlobalReadVectorWidth (VW = 2 + GRVW = 1 or VW = 1 + GRVW = 2 does not work)
-    if state["ProblemType"]["TLU%c"%tc] and state["VectorWidth%s"%tc] != state["GlobalReadVectorWidth%c"%tc]:
+    if state["ProblemType"]["TLU%c"%tc] and (state["VectorWidth%s"%tc] != state["GlobalReadVectorWidth%c"%tc]) and (not state["enableGLTr%s"%tc]):
       reject(state, printRejectionReason, "DirectToVgpr%c does not supports TLU + VectorWidth%s(=%u) != GlobalReadVectorWidth%c(%u)"%(tc, tc, state["VectorWidth%s"%tc], tc, state["GlobalReadVectorWidth%c"%tc]))
       return False
 
@@ -680,6 +688,11 @@ def isDirectToVgprDoable(state, tc, printRejectionReason: bool, isaInfoMap: Dict
       reject(state, printRejectionReason, "DirectToVgpr%c does not supports TLU=False and NumLoadsCoalesced%c != DepthU//(MatrixInstK*GlobalReadVectorWidth*LocalSplitU//MIInputPerThread(=%u))"%(tc, tc, state["MIInputPerThread"]))
       return False
 
+    # Does not work with enableGLTr and GRVW != 8
+    if state["enableGLTr%c"%tc] and state["GlobalReadVectorWidth%c"%tc] != 8:
+      reject(state, printRejectionReason, "DirectToVgpr%c does not supports enableGLTr%c and GlobalReadVectorWidth != 8"%(tc, tc))
+      return False
+
     # TLU=False case, need GlobalReadVectorWidth == LocalReadVectorWidth
     if (not state["ProblemType"]["TLU%c"%tc]) and \
        state["GlobalReadVectorWidth%c"%tc] != state["LocalReadVectorWidth"]:
@@ -732,7 +745,7 @@ def isDirectToVgprDoable(state, tc, printRejectionReason: bool, isaInfoMap: Dict
         state["AssertSummationElementMultiple"] = max(state["AssertSummationElementMultiple"], state["DepthU"])
 
     # for DTVB, does not work with NN and Tail-loop
-    if  tc == 'B' and (not state["ProblemType"]["TransposeA"] and not state["ProblemType"]["TransposeB"]):
+    if tc == 'B' and (not state["ProblemType"]["TransposeA"] and not state["ProblemType"]["TransposeB"]):
         # Use AssertSummationElementMultiple (BoundSizeMultiple in predicates) to exclude failed tail-loop cases
         state["AssertSummationElementMultiple"] = max(state["AssertSummationElementMultiple"], state["DepthU"])
 
@@ -1306,10 +1319,19 @@ def assignDerivedParameters(
     state["enableLDSTrB"] = state["LDSTrInst"] and isaInfoMap[isa].asmCaps["HasLDSTr"] and numBytes == 2 \
             and not state["UnrollMajorLDSB"] and not state["DirectToVgprB"]
 
-    if state["enableLDSTrA"]:
+    state["enableGLTrA"] = state["DirectToVgprA"] and state["ProblemType"]["TLUA"] \
+      and ((numBytes == 1 and isaInfoMap[isa].asmCaps["HasGLTr8B64"]) \
+        or (numBytes == 2 and isaInfoMap[isa].asmCaps["HasGLTr16B128"]) \
+      )
+    state["enableGLTrB"] = state["DirectToVgprB"] and state["ProblemType"]["TLUB"] \
+      and ((numBytes == 1 and isaInfoMap[isa].asmCaps["HasGLTr8B64"]) \
+        or (numBytes == 2 and isaInfoMap[isa].asmCaps["HasGLTr16B128"]) \
+      )
+
+    if state["enableLDSTrA"] or state["enableGLTrA"]:
       state["VectorWidthA"] = 1
 
-    if state["enableLDSTrB"]:
+    if state["enableLDSTrB"] or state["enableGLTrB"]:
       state["VectorWidthB"] = 1
 
     # if state["EnableMatrixInstruction"] and not state["SourceSwap"] and (state["VectorWidthA"] > 1 or state["VectorWidthB"] > 1):
@@ -1403,7 +1425,6 @@ def evaluateExpertSchedulingMode() -> Literal[0, 1, 2]:
         "ScheduleLocalWrite": state["ScheduleLocalWrite"] == 1,
         "GlobalReadPerMfma": state["GlobalReadPerMfma"] == 1,
         "InterleaveAlpha": not state["InterleaveAlpha"],
-        "DirectToVgprA": not state["DirectToVgprA"],
         "DirectToLds": not state["DirectToLds"],
         "UseSgprForGRO": state["UseSgprForGRO"] == -1,
         "UseInstOffsetForGRO": not state["UseInstOffsetForGRO"],

projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/fp16_tn_gfx12.yaml → projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/fp16_gfx12.yaml

@@ -14,9 +14,6 @@ GlobalParameters:
   Device: 0
 
 BenchmarkProblems:
-  ########################################
-  # NN - standard
-  ########################################
   -
     - # ProblemType
       OperationType: GEMM
@@ -190,3 +187,104 @@ BenchmarkProblems:
         - BiasTypeArgs: ['h']
         - ActivationArgs:
           - [Enum: relu]
+  -
+    - # ProblemType
+      OperationType: GEMM
+      DataType: h
+      DestDataType: h
+      ComputeDataType: s
+      HighPrecisionAccumulate: True
+      TransposeA: 0
+      TransposeB: 1
+      UseBeta: True
+      UseBias: 1
+      Batched: True
+      Activation: True
+      ActivationType: all
+    - # BenchmarkProblemSizeGroup - Standard - All problem
+      InitialSolutionParameters:
+      BenchmarkCommonParameters:
+        - KernelLanguage: ["Assembly"]
+      ForkParameters:
+        - MatrixInstruction:
+          - [16, 16,16, 1,  1,   1, 1,  4,1 ]
+          - [16, 16,16, 1,  1,   2, 1,  4,1 ]
+          - [16, 16,16, 1,  1,   3, 1,  4,1 ]
+          - [16, 16,16, 1,  1,   4, 1,  4,1 ]
+          - [16, 16,16, 1,  1,   1, 1,  2,2 ]
+          - [16, 16,16, 1,  1,   2, 1,  2,2 ]
+          - [16, 16,16, 1,  1,   3, 1,  2,2 ]
+          - [16, 16,16, 1,  1,   5, 1,  2,2 ]
+          - [16, 16,16, 1,  1,   6, 1,  2,2 ]
+          - [16, 16,16, 1,  1,   7, 1,  2,2 ]
+          - [16, 16,16, 1,  1,   1, 1,  1,4 ]
+          - [16, 16,16, 1,  1,   2, 1,  1,4 ]
+          - [16, 16,16, 1,  1,   3, 1,  1,4 ]
+          - [16, 16,16, 1,  1,   5, 1,  1,4 ]
+          - [16, 16,16, 1,  1,   6, 1,  1,4 ]
+          - [16, 16,16, 1,  1,   7, 1,  1,4 ]
+          - [16, 16,16, 1,  1,   9, 1,  1,4 ]
+          - [16, 16,16, 1,  1,  10, 1,  1,4 ]
+          - [16, 16,16, 1,  1,  11, 1,  1,4 ]
+          - [16, 16,16, 1,  1,  13, 1,  1,4 ]
+          - [16, 16,16, 1,  1,  14, 1,  1,4 ]
+          - [16, 16,16, 1,  1,  15, 1,  1,4 ]
+
+        - PrefetchGlobalRead: [2]
+        - PrefetchLocalRead: [1]
+        - ClusterLocalRead: [1]
+        - DepthU: [32,64]
+        - WavefrontSize: [32]
+        - VectorWidthA: [1]
+        - VectorWidthB: [1]
+        - LocalReadVectorWidth: [8]
+        - GlobalReadVectorWidthA: [8]
+        - GlobalReadVectorWidthB: [8]
+        - ScheduleIterAlg: [3]
+        - ExpandPointerSwap: [0]
+        - TransposeLDS: [0,1,2]
+        - LdsBlockSizePerPadA: [-1]
+        - LdsBlockSizePerPadB: [-1]
+        - LdsPadA: [-1]
+        - LdsPadB: [-1]
+        - 1LDSBuffer: [-1]
+        - GlobalSplitU: [1] #,2] #known bug
+        - GlobalSplitUAlgorithm: ["MultipleBuffer"]
+        - SourceSwap: [0,1]
+        - DirectToVgprA: [0,1]
+        - DirectToVgprB: [0,1]
+      BenchmarkJoinParameters:
+      BenchmarkFinalParameters:
+        - ProblemSizes:
+          - Exact: [255, 255, 1, 255]
+          - Exact: [255, 255, 1, 256]
+          - Exact: [255, 255, 1, 257]
+          - Exact: [255, 256, 1, 255]
+          - Exact: [255, 256, 1, 256]
+          - Exact: [255, 256, 1, 257]
+          - Exact: [255, 257, 1, 255]
+          - Exact: [255, 257, 1, 256]
+          - Exact: [255, 257, 1, 257]
+
+          - Exact: [256, 255, 1, 255]
+          - Exact: [256, 255, 1, 256]
+          - Exact: [256, 255, 1, 257]
+          - Exact: [256, 256, 1, 255]
+          - Exact: [256, 256, 1, 256]
+          - Exact: [256, 256, 1, 257]
+          - Exact: [256, 257, 1, 255]
+          - Exact: [256, 257, 1, 256]
+          - Exact: [256, 257, 1, 257]
+
+          - Exact: [257, 255, 1, 255]
+          - Exact: [257, 255, 1, 256]
+          - Exact: [257, 255, 1, 257]
+          - Exact: [257, 256, 1, 255]
+          - Exact: [257, 256, 1, 256]
+          - Exact: [257, 256, 1, 257]
+          - Exact: [257, 257, 1, 255]
+          - Exact: [257, 257, 1, 256]
+          - Exact: [257, 257, 1, 257]
+        - BiasTypeArgs: ['h']
+        - ActivationArgs:
+          - [Enum: relu]

projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/fp8_gfx12.yaml

@@ -47,6 +47,7 @@ BenchmarkProblems:
       ForkParameters:
         - MatrixInstruction:
           - [16,16,16, 1, 1, 1,1, 1,1]  # 16x16
+          - [16,16,16, 1, 1, 2,4, 2,1]
         - DepthU: [ 64 ]
           #- AssertFree0ElementMultiple: [1]
         - PrefetchGlobalRead: [2]
@@ -56,7 +57,7 @@ BenchmarkProblems:
         - LocalReadVectorWidth: [8]
         - VectorWidthA: [1]
         - VectorWidthB: [1]
-        - GlobalReadVectorWidthA: [-1] #[1, 2, 4, 8, 16]
+        - GlobalReadVectorWidthA: [-1,8] #[1, 2, 4, 8, 16]
         - GlobalReadVectorWidthB: [-1] #[1, 2, 4, 8, 16]
         - ScheduleIterAlg: [3]
         - InnerUnroll: [1]
@@ -74,6 +75,7 @@ BenchmarkProblems:
         - LocalWritePerMfma: [-1]
         - StoreVectorWidth: [-1]
         - SourceSwap: [1]
+        - DirectToVgprA: [0,1]
         - DirectToVgprB: [0,1]
         - NumElementsPerBatchStore: [0]
         - StorePriorityOpt: [0]
@@ -228,6 +230,115 @@ BenchmarkProblems:
         - ActivationArgs:
           - [Enum: none]
 
+  ########################################
+  # FP8FP8S-NT
+  ########################################
+  -
+    - # ProblemType
+      OperationType: GEMM
+      DataType: F8
+      ComputeDataType: S
+      HighPrecisionAccumulate: True
+      TransposeA: 0
+      TransposeB: 1
+      UseBeta: True
+      Batched: True
+      Activation:    True
+      UseScaleAB: "Scalar"
+      UseScaleCD: True
+      UseScaleAlphaVec: 1
+      UseBias: 1
+      BiasDataTypeList: [s]
+    - # BenchmarkProblemSizeGroup - Standard
+      InitialSolutionParameters:
+      BenchmarkCommonParameters:
+        - KernelLanguage: ["Assembly"]
+      ForkParameters:
+        - MatrixInstruction:
+          - [16, 16,16, 1,  1,   1, 1,  4,1 ]
+          - [16, 16,16, 1,  1,   2, 1,  4,1 ]
+          - [16, 16,16, 1,  1,   3, 1,  4,1 ]
+          - [16, 16,16, 1,  1,   4, 1,  4,1 ]
+          - [16, 16,16, 1,  1,   1, 1,  2,2 ]
+          - [16, 16,16, 1,  1,   2, 1,  2,2 ]
+          - [16, 16,16, 1,  1,   3, 1,  2,2 ]
+          - [16, 16,16, 1,  1,   5, 1,  2,2 ]
+          - [16, 16,16, 1,  1,   6, 1,  2,2 ]
+          - [16, 16,16, 1,  1,   7, 1,  2,2 ]
+          - [16, 16,16, 1,  1,   1, 1,  1,4 ]
+          - [16, 16,16, 1,  1,   2, 1,  1,4 ]
+          - [16, 16,16, 1,  1,   3, 1,  1,4 ]
+          - [16, 16,16, 1,  1,   5, 1,  1,4 ]
+          - [16, 16,16, 1,  1,   6, 1,  1,4 ]
+          - [16, 16,16, 1,  1,   7, 1,  1,4 ]
+          - [16, 16,16, 1,  1,   9, 1,  1,4 ]
+          - [16, 16,16, 1,  1,  10, 1,  1,4 ]
+          - [16, 16,16, 1,  1,  11, 1,  1,4 ]
+          - [16, 16,16, 1,  1,  13, 1,  1,4 ]
+          - [16, 16,16, 1,  1,  14, 1,  1,4 ]
+          - [16, 16,16, 1,  1,  15, 1,  1,4 ]
+        - DepthU: [ 32,64 ]
+        - PrefetchGlobalRead: [2]
+        - PrefetchLocalRead: [1]
+        - ClusterLocalRead: [1]
+        - VectorWidthA: [1]
+        - VectorWidthB: [1]
+        - WavefrontSize: [32]
+        - LocalReadVectorWidth: [8]
+        - GlobalReadVectorWidthA: [-1,8]
+        - GlobalReadVectorWidthB: [-1,8]
+        - ScheduleIterAlg: [3]
+        - InnerUnroll: [1]
+        - ExpandPointerSwap: [1]
+        - TransposeLDS: [0,1,2]
+        - WaveSeparateGlobalReadB: [0,1]
+        - 1LDSBuffer: [-1]
+        - GlobalSplitU: [1]
+        - GlobalSplitUAlgorithm: ["MultipleBuffer"]
+        - GlobalReadPerMfma: [1]
+        - LocalWritePerMfma: [-1]
+        - StoreVectorWidth: [-1]
+        - SourceSwap: [1]
+        - DirectToVgprA: [0,1]
+        - DirectToVgprB: [0,1]
+        - NumElementsPerBatchStore: [0]
+        - StorePriorityOpt: [0]
+      BenchmarkJoinParameters:
+      BenchmarkFinalParameters:
+        - ProblemSizes:
+          - Exact: [255, 255, 1, 255]
+          - Exact: [255, 255, 1, 256]
+          - Exact: [255, 255, 1, 257]
+          - Exact: [255, 256, 1, 255]
+          - Exact: [255, 256, 1, 256]
+          - Exact: [255, 256, 1, 257]
+          - Exact: [255, 257, 1, 255]
+          - Exact: [255, 257, 1, 256]
+          - Exact: [255, 257, 1, 257]
+
+          - Exact: [256, 255, 1, 255]
+          - Exact: [256, 255, 1, 256]
+          - Exact: [256, 255, 1, 257]
+          - Exact: [256, 256, 1, 255]
+          - Exact: [256, 256, 1, 256]
+          - Exact: [256, 256, 1, 257]
+          - Exact: [256, 257, 1, 255]
+          - Exact: [256, 257, 1, 256]
+          - Exact: [256, 257, 1, 257]
+
+          - Exact: [257, 255, 1, 255]
+          - Exact: [257, 255, 1, 256]
+          - Exact: [257, 255, 1, 257]
+          - Exact: [257, 256, 1, 255]
+          - Exact: [257, 256, 1, 256]
+          - Exact: [257, 256, 1, 257]
+          - Exact: [257, 257, 1, 255]
+          - Exact: [257, 257, 1, 256]
+          - Exact: [257, 257, 1, 257]
+        - BiasTypeArgs: ['s']
+        - ActivationArgs:
+          - [Enum: none]
+
   ########################################
   # F8BS-TN
   ########################################