CustomKernels don't need tuning arguments

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

@@ -423,7 +423,7 @@
     {"NonTemporalMetadata": [0]},
     {"NonTemporal": [-1]},
     {"PreloadKernArgs": [True]},
-    {"CustomKernel": [{"name": "", "args": [], "macrotile": [0,0,0], "threads": [0,0,0], "grid": [0,0,0]}]},
+    # {"CustomKernel": [{"name": "", "args": [], "macrotile": [0,0,0], "threads": [0,0,0], "grid": [0,0,0]}]},
     {"NoReject": [False]},
     {"StoreRemapVectorWidth": [0]},
     {"SourceSwap": [False]},

projects/hipblaslt/tensilelite/Tensile/Contractions.py

@@ -589,11 +589,6 @@ def FromOriginalState(cls, d):
         return cls(name=d['name'], args=d['args'], macrotile=d['macrotile'], threads=d['threads'], grid=d['grid'])
 
     def __init__(self, **kwargs):
-        self.name = ""
-        self.args = []
-        self.macrotile = [0,0,0]
-        self.threads = [0,0,0]
-        self.grid = [0,0,0]
         for (key, value) in list(kwargs.items()):
             setattr(self, key, value)
 
@@ -623,7 +618,6 @@ class SizeMapping:
                  'workspaceSizePerElemC',
                  'workspaceSizePerElemBias',
                  'activationFused',
-                 'customKernel',
                  'workGroupMappingXCC',
                  'workGroupMappingXCCGroup',
                  'globalSplitUCoalesced',
@@ -686,11 +680,6 @@ def convertSFCWGMListToHex(value):
         dtlA = bool(d['DirectToLdsA'])
         dtlB = bool(d['DirectToLdsB'])
 
-        if 'CustomKernel' in d:
-            customKernel = CustomKernel.FromOriginalState(d['CustomKernel'])
-        else:
-            customKernel = CustomKernel(name="", args=[], macrotile=[0,0,0], threads=[0,0,0], grid=[0,0,0])
-
         return cls(waveNum                  = d['NumThreads'] // d['WavefrontSize'],
                    workGroup                = d['WorkGroup'],
                    macroTile                = cls.ReadOriginalMacroTile(d),
@@ -716,7 +705,6 @@ def convertSFCWGMListToHex(value):
                    workspaceSizePerElemC    = d['_WorkspaceSizePerElemC'],
                    workspaceSizePerElemBias = d['_WorkspaceSizePerElemBias'],
                    activationFused          = d['ActivationFused'],
-                   customKernel             = customKernel,
                    workGroupMappingXCC      = d['WorkGroupMappingXCC'],
                    workGroupMappingXCCGroup = d['WorkGroupMappingXCCGroup'],
                    globalSplitUCoalesced    = d['GlobalSplitUCoalesced'],
@@ -782,17 +770,18 @@ def __init__(self, **kwargs):
 class Solution:
     StateKeys = ['name',
                  'kernelName',
-                'problemType',
-                'hardwarePredicate',
-                'problemPredicate',
-                'taskPredicate',
-                'sizeMapping',
-                'internalArgsSupport',
-                'debugKernel',
-                'libraryLogicIndex',
-                'index',
-                'ideals',
-                'linearModel']
+                 'problemType',
+                 'hardwarePredicate',
+                 'problemPredicate',
+                 'taskPredicate',
+                 'sizeMapping',
+                 'customKernel',
+                 'internalArgsSupport',
+                 'debugKernel',
+                 'libraryLogicIndex',
+                 'index',
+                 'ideals',
+                 'linearModel']
     HiddenKeys = ['originalSolution']
 
     @classmethod
@@ -851,6 +840,10 @@ def FromOriginalState(
         rv.libraryLogicIndex = int(info.get("SolutionIndex", -1))
 
         rv.sizeMapping = SizeMapping.FromOriginalState(d)
+        if 'CustomKernel' in d:
+            rv.customKernel = CustomKernel.FromOriginalState(d['CustomKernel'])
+        else:
+            rv.customKernel = {}
 
         rv.internalArgsSupport = InternalArgsSupport.FromOriginalState(d)
 
@@ -896,6 +889,7 @@ def __init__(self, **kwargs):
         self.problemPredicate = ProblemPredicate('TruePred')
         self.taskPredicate = TaskPredicate('TruePred')
         self.sizeMapping = None
+        self.customKernel = None
         self.debugKernel = False
         self.libraryLogicIndex = {}
         self.index = None

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

@@ -215,14 +215,17 @@ def __init__(
       self["AssignedProblemIndependentDerivedParameters"] = False
     if "AssignedDerivedParameters" not in self._state:
       self["AssignedDerivedParameters"] = False
-    Solution.assignDerivedParameters(
-      self._state,
-      splitGSU,
-      printSolutionRejectionReason,
-      printIndexAssignmentInfo,
-      isaInfoMap,
-      assembler.rocm_version
-    )
+    if "CustomKernel" not in self._state:
+      Solution.assignDerivedParameters(
+        self._state,
+        splitGSU,
+        printSolutionRejectionReason,
+        printIndexAssignmentInfo,
+        isaInfoMap,
+        assembler.rocm_version
+      )
+    else:
+      Solution.assignCustomKernelParameters(self._state)
     self._name = config["CustomKernel"]["name"] if "CustomKernel" in config and config["CustomKernel"]["name"] else None
 
   # these keys are copied from ProblemType to internal that may be overridden
@@ -935,6 +938,52 @@ def getDivisorName(state, tC):
         divisorName = "LVP{}".format(tC)
     return divisorName
 
+  @staticmethod
+  def assignCustomKernelParameters(state):
+    # Store values in state so they can be saved in SizeMapping
+    # Not all SizeMapping values will be needed (or applicable) for custom kernels
+
+    state["MacroTile0"] = state["CustomKernel"]["macrotile"][0]
+    state["MacroTile1"] = state["CustomKernel"]["macrotile"][1]
+    state["DepthU"] = state["CustomKernel"]["macrotile"][2]
+
+    # TODO Temp values to allow serializing SizeMapping
+    state["_GlobalAccumulation"]    = None
+    state["CUOccupancy"]            = -1
+    state["MathClocksUnrolledLoop"] = 0
+    state["PackedC0IndicesX"] = []
+    state["ThreadTile0"] = 0
+    state["ThreadTile1"] = 0
+    state["NumThreads"] = state["CustomKernel"]["threads"][0] * state["CustomKernel"]["threads"][1] * state["CustomKernel"]["threads"][2]
+
+    # TODO LSU
+    # numElementsPerWorkGroup = state["MacroTile0"]*state["MacroTile1"]*state["NumWaveSplitK"]
+    numElementsPerWorkGroup = state["MacroTile0"] * state["MacroTile1"]
+    state["NumElementsPerThread"] = numElementsPerWorkGroup // state["NumThreads"]
+
+    # TODO Check what DTL affects in SizeMapping
+    state["DirectToLdsA"] = state["DirectToLds"] == 1 or state["DirectToLds"] == 2
+    state["DirectToLdsB"] = state["DirectToLds"] == 1 or state["DirectToLds"] == 3
+
+    # TODO Override workspace info for custom kernel
+    computeBytes = state["ProblemType"]["ComputeDataType"].numBytes()
+    state["_WorkspaceSizePerElemC"] = computeBytes
+    state["_WorkspaceSizePerElemBias"] = 0
+    if state["ProblemType"]["UseBias"] and state["ProblemType"]["Gradient"]:
+      state["_WorkspaceSizePerElemBias"] = computeBytes
+    # state["WorkspaceCheck"] = [state["_WorkspaceSizePerElemC"], state["_WorkspaceSizePerElemBias"], state["GlobalSplitU"] if (state["GlobalSplitUAlgorithm"] == 'MultipleBuffer' or state["_GlobalAccumulation"] == 'MultipleBufferSingleKernel') else 1]
+
+    # state["MIWaveGroup"] = [state["SubGroup0"] // state["WavefrontSize"],  state["SubGroup1"]]
+    state["MIWaveGroup"] = [0, 0]
+
+    # state["LocalSplitU"] = state["WorkGroup"][2]
+    state["LocalSplitU"] = 1
+
+    state["GlobalReadVectorWidthA"] = 1
+    state["GlobalReadVectorWidthB"] = 1
+    state["StoreVectorWidth"] = 1
+
+
   ########################################
   # assign all derived parameters
   @staticmethod

projects/hipblaslt/tensilelite/Tensile/Tests/custom/custom_ck.yaml

@@ -66,44 +66,6 @@ BenchmarkProblems:
             macrotile: [256, 256, 64]
             threads: [256, 1, 1]
             grid: [TilesXYBatch, One, One]
-        - MatrixInstruction:
-          - [16,16,16, 1, 1, 8,8, 2,2] # 256 x 128
-        - PreloadKernArgs: [False]
-        # - PrefetchGlobalRead: [2]
-        # - PrefetchLocalRead: [1]
-        - DepthU: [64]
-        - VectorWidthA: [-1]
-        - VectorWidthB: [-1]
-        - GlobalReadVectorWidthA: [-1]
-        - GlobalReadVectorWidthB: [-1]
-        - LocalReadVectorWidth: [-1]
-        - TransposeLDS: [-1]
-        - LdsBlockSizePerPadA: [-1]
-        - LdsBlockSizePerPadB: [-1]
-        - LdsPadA: [-1]
-        - LdsPadB: [-1]
-        - StaggerU: [16]
-        - StaggerUStride: [-1]
-        - WorkGroupMapping: [1]
-        - StaggerUMapping: [0]
-        - 1LDSBuffer: [-1]
-        - WorkGroupMappingXCC: [8]
-        - WorkGroupMappingXCCGroup: [-1]
-        - GlobalSplitU: [1]
-        - GlobalSplitUAlgorithm: ["MultipleBuffer"]
-        - GlobalReadPerMfma: [1]
-        - LocalWritePerMfma: [-1]
-        - StoreRemapVectorWidth: [0]
-        - StoreVectorWidth: [-1]
-        - SourceSwap: [1]
-        - NumElementsPerBatchStore: [16]
-        - ClusterLocalRead: [1]
-        - NonTemporalA: [0]
-        - NonTemporalB: [0]
-        - DirectToVgprA: [1]
-        - DirectToVgprB: [0]
-        # - WorkGroup: [[32,4,4]]
-
       BenchmarkJoinParameters:
       BenchmarkFinalParameters:
         - ProblemSizes:

projects/hipblaslt/tensilelite/include/Tensile/ContractionSolution.hpp

@@ -216,8 +216,6 @@ namespace TensileLite
 
         bool activationFused = true;
 
-        CustomKernel customKernel;
-
         int  workGroupMappingXCC                    = 0;
         int  workGroupMappingXCCGroup               = 0;
         bool globalSplitUCoalesced                  = false;
@@ -677,6 +675,7 @@ namespace TensileLite
             = std::make_shared<Predicates::True<Hardware>>();
 
         SizeMapping sizeMapping;
+        CustomKernel customKernel;
 
         InternalArgsSupport internalArgsSupport;
 

projects/hipblaslt/tensilelite/include/Tensile/Serialization/ContractionSolution.hpp

@@ -66,6 +66,7 @@ namespace TensileLite
                 iot::mapOptional(io, "linearModel", s.linearModel);
 
                 iot::mapRequired(io, "sizeMapping", s.sizeMapping);
+                iot::mapOptional(io, "customKernel", s.customKernel);
                 iot::mapRequired(io, "internalArgsSupport", s.internalArgsSupport);
                 iot::mapRequired(io, "problemType", s.problemType);
             }
@@ -141,8 +142,6 @@ namespace TensileLite
 
                 iot::mapOptional(io, "activationFused", s.activationFused);
 
-                iot::mapOptional(io, "customKernel", s.customKernel);
-
                 iot::mapRequired(io, "workGroupMappingXCC", s.workGroupMappingXCC);
                 iot::mapRequired(io, "workGroupMappingXCCGroup", s.workGroupMappingXCCGroup);
 

projects/hipblaslt/tensilelite/src/ContractionSolution.cpp

@@ -1562,8 +1562,8 @@ namespace TensileLite
         if(problem.transposeC01())
             std::swap(tiles.x, tiles.y);
 
-        tiles.x = CeilDivide(tiles.x, sizeMapping.customKernel.macrotile.x);
-        tiles.y = CeilDivide(tiles.y, sizeMapping.customKernel.macrotile.y);
+        tiles.x = CeilDivide(tiles.x, customKernel.macrotile.x);
+        tiles.y = CeilDivide(tiles.y, customKernel.macrotile.y);
     }
 
     void ContractionSolution::calculateGrid(dim3&                               workGroupSize,
@@ -1621,19 +1621,19 @@ namespace TensileLite
         assert(pAMDGPU);
 
         int wavefrontSize = pAMDGPU->wavefrontSize;
-        // rv.workGroupSize.x = sizeMapping.customKernel.waves.x * wavefrontSize;
-        // rv.workGroupSize.y = sizeMapping.customKernel.waves.y * wavefrontSize;
-        // rv.workGroupSize.z = sizeMapping.customKernel.waves.z * wavefrontSize;
-        rv.workGroupSize.x = sizeMapping.customKernel.threads.x;
-        rv.workGroupSize.y = sizeMapping.customKernel.threads.y;
-        rv.workGroupSize.z = sizeMapping.customKernel.threads.z;
+        // rv.workGroupSize.x = customKernel.waves.x * wavefrontSize;
+        // rv.workGroupSize.y = customKernel.waves.y * wavefrontSize;
+        // rv.workGroupSize.z = customKernel.waves.z * wavefrontSize;
+        rv.workGroupSize.x = customKernel.threads.x;
+        rv.workGroupSize.y = customKernel.threads.y;
+        rv.workGroupSize.z = customKernel.threads.z;
 
         if(T_Debug)
         {
             std::cout << "Wavefront size: " << wavefrontSize << std::endl;
-            std::cout << "Threads: " << sizeMapping.customKernel.threads.x << ", " << sizeMapping.customKernel.threads.y << ", " << sizeMapping.customKernel.threads.z << std::endl;
+            std::cout << "Threads: " << customKernel.threads.x << ", " << customKernel.threads.y << ", " << customKernel.threads.z << std::endl;
             std::cout << "Work group size: " << rv.workGroupSize.x << ", " << rv.workGroupSize.y << ", " << rv.workGroupSize.z << std::endl;
-            std::cout << "Macrotile: " << sizeMapping.customKernel.macrotile.x << ", " << sizeMapping.customKernel.macrotile.y << ", " << sizeMapping.customKernel.macrotile.z << std::endl;
+            std::cout << "Macrotile: " << customKernel.macrotile.x << ", " << customKernel.macrotile.y << ", " << customKernel.macrotile.z << std::endl;
         }
 
         dim3 tiles;
@@ -1677,9 +1677,9 @@ namespace TensileLite
             }
         };
 
-        assignGridSize(rv.numWorkGroups.x, sizeMapping.customKernel.grid.x);
-        assignGridSize(rv.numWorkGroups.y, sizeMapping.customKernel.grid.y);
-        assignGridSize(rv.numWorkGroups.z, sizeMapping.customKernel.grid.z);
+        assignGridSize(rv.numWorkGroups.x, customKernel.grid.x);
+        assignGridSize(rv.numWorkGroups.y, customKernel.grid.y);
+        assignGridSize(rv.numWorkGroups.z, customKernel.grid.z);
 
         if(T_Debug)
         {
@@ -1707,7 +1707,7 @@ namespace TensileLite
         if(T_Debug)
             std::cout << "Custom call arguments:" << std::endl;
 
-        for(auto arg : sizeMapping.customKernel.args)
+        for(auto arg : customKernel.args)
         {
             if(T_Debug)
                 std::cout << "Type: " << toString(arg.type) << " Semantic: " << toString(arg.semantic) << std::endl;
@@ -3158,7 +3158,7 @@ namespace TensileLite
             }
         }
 
-        if(sizeMapping.customKernel.name.empty())
+        if(customKernel.name.empty())
         {
             // Regular generated kernel
             if(debug)
@@ -3695,7 +3695,7 @@ namespace TensileLite
         AMDGPU const* pAMDGPU = dynamic_cast<AMDGPU const*>(&hardware);
         assert(pAMDGPU != nullptr && pAMDGPU->computeUnitCount != 0);
 
-        if(!sizeMapping.customKernel.name.empty())
+        if(!customKernel.name.empty())
         {
             // Custom kernel currently only supports single-kernel reduction
             reductionStrat = origami::reduction_t::tree;