New build algo

tensilelite/Tensile/BenchmarkProblems.py

@@ -46,7 +46,9 @@
 from .Contractions import ProblemType as ContractionsProblemType
 from .ClientWriter import runClient, writeClientConfig, writeClientConfigIni
 from .KernelWriterAssembly import KernelWriterAssembly
-from .TensileCreateLibrary import copyStaticFiles, writeSolutionsAndKernels
+from .SolutionStructs import Solution, ProblemType, ProblemSizes
+from Tensile.TensileCreateLibrary import copyStaticFiles 
+from Tensile.TensileCreateLibrary.Tuning import writeSolutionsAndKernels
 from .CustomKernels import getCustomKernelConfig
 from .Toolchain.Assembly import AssemblyToolchain
 from .Toolchain.Source import SourceToolchain
@@ -229,7 +231,7 @@ def writeBenchmarkFiles(
                 kernels.append(kernel)
                 kernelNames.add(kName)
 
-        solutionHelperKernels = solution.getHelperKernelObjects()
+        solutionHelperKernels = solution.initHelperKernelObjects(list(isaInfoMap.keys()))
         for ko in solutionHelperKernels:
             kname = ko.getKernelName()
             if kname not in kernelHelperNames:

tensilelite/Tensile/CodeObjectName.py

@@ -0,0 +1,73 @@
+
+
+from Tensile.Properties import Predicate
+from Tensile.Contractions import ProblemType
+
+
+def codeObjectFileBaseName(d: dict, useLazyLibraryLoading: bool=True):
+    """Compute the code object library name from the solution metadata.
+
+    Args:
+        d: Solution metadata used to create the code object name.
+        lazyLibary: Use naming conventions necessary for lazy library loading.
+
+    Returns:
+        A string with the.
+    """
+    placeholderName = "TensileLibrary"
+    problemType = ProblemType.FromOriginalState(d["ProblemType"])
+    if useLazyLibraryLoading:
+        placeholderName += '_' + str(problemType.aType) + str(problemType.bType)
+        placeholderName += '_' + str(problemType.cType) + str(problemType.computeInputType)
+        if problemType.activationType != 'none':
+            if str(problemType.activationType).upper() == 'ALL':
+                placeholderName += "_A"
+            elif str(problemType.activationType).upper() == 'HIPBLASLT_ALL':
+                placeholderName += "_HA"
+            else:
+                placeholderName += "_%s"%str(problemType.activationType).upper()
+        if problemType.swizzleTensorA:
+            placeholderName += '_STA'
+        if problemType.swizzleTensorB:
+            placeholderName += '_STB'
+        if problemType.useBias:
+            placeholderName += '_Bias'
+        if problemType.useE:
+            placeholderName += '_Grad' if problemType.useGradient else '_Aux'
+        if problemType.groupedGemm:
+            placeholderName += "_GG"
+        else:
+            placeholderName += "" if problemType.stridedBatched else "_GB" # legacy
+        if problemType.useScaleAB == "Scalar":
+            placeholderName += '_SAB'
+        elif problemType.useScaleAB == "Vector":
+            placeholderName += '_SABV'
+        if problemType.useScaleCD:
+            placeholderName += '_SCD'
+        if problemType.useScaleAlphaVec:
+            placeholderName += '_SAV'
+        if problemType.sparse:
+            placeholderName += '_SPB' if problemType.sparse == 2 else '_SPA'
+        if not problemType.f32XdlMathOp.isSingle() and problemType.computeInputType.isSingle():
+            placeholderName += '_M' + str(problemType.f32XdlMathOp)
+        if problemType.supportDeviceUserArguments:
+            placeholderName += '_UA'
+
+        placeholderName += problemType.placeholderStr(includeBatch=True, includeType=True)
+
+        if d.get("PerfMetric", "DeviceEfficiency") != "DeviceEfficiency":
+            predicate = Predicate(tag=d["PerfMetric"])
+        else:
+            predicate = Predicate(tag="TruePred")
+        if predicate.tag != "TruePred":
+            placeholderName += "_" + predicate.tag
+
+        operationID = problemType.operationIdentifier
+        placeholderName += "_" + operationID
+
+        devicePart = d["ArchitectureName"]
+        cuCount = d["CUCount"]
+        if cuCount: placeholderName += "_CU" + str(cuCount)
+        placeholderName += "_" + str(devicePart)
+
+    return placeholderName

tensilelite/Tensile/Common/Parallel.py

@@ -29,8 +29,9 @@
 import time
 
 from joblib import Parallel, delayed
+from typing import NamedTuple
 
-from .Utilities import tqdm
+from .Utilities import tqdm, printExit
 
 
 def joblibParallelSupportsGenerator():
@@ -59,13 +60,11 @@ def CPUThreadCount(enable=True):
         return min(cpu_count, cpuThreads)
 
 
-def pcallWithGlobalParamsMultiArg(f, args, newGlobalParameters):
-    OverwriteGlobalParameters(newGlobalParameters)
+def pcallWithGlobalParamsMultiArg(f, args):
     return f(*args)
 
 
-def pcallWithGlobalParamsSingleArg(f, arg, newGlobalParameters):
-    OverwriteGlobalParameters(newGlobalParameters)
+def pcallWithGlobalParamsSingleArg(f, arg):
     return f(arg)
 
 
@@ -195,8 +194,16 @@ def ParallelMapReturnAsGenerator(function, objects, message="", enable=True, mul
             yield result.result()
 
 
+class ParallelMapConfig(NamedTuple):
+  message: str = ""
+  enable: bool = True
+  multiArg: bool = False
+  return_as: str = "generator_unordered"
+  procs: int = None
+
+
 def ParallelMap2(
-    function, objects, message="", enable=True, multiArg=True, return_as="list", procs=None
+    function, config: ParallelMapConfig, objects
 ):
     """
     Generally equivalent to list(map(function, objects)), possibly executing in parallel.
@@ -206,18 +213,17 @@ def ParallelMap2(
       multiArg: True if objects represent multiple arguments
                   (differentiates multi args vs single collection arg)
     """
-    if return_as in ("generator", "generator_unordered") and not joblibParallelSupportsGenerator():
-        return ParallelMapReturnAsGenerator(function, objects, message, enable, multiArg)
+    message = config.message
+    if config.return_as in ("generator", "generator_unordered") and not joblibParallelSupportsGenerator():
+        return ParallelMapReturnAsGenerator(function, objects, message, config.enable, config.multiArg)
 
     from .GlobalParameters import globalParameters
 
-    threadCount = procs if procs else CPUThreadCount(enable)
-
-    threadCount = CPUThreadCount(enable)
-
+    threadCount = config.procs if config.procs else CPUThreadCount(config.enable)
+
     if threadCount <= 1 and globalParameters["ShowProgressBar"]:
         # Provide a progress bar for single-threaded operation.
-        return [function(*args) if multiArg else function(args) for args in tqdm(objects, message)]
+        return [function(*args) if config.multiArg else function(args) for args in tqdm(objects, message)]
 
     countMessage = ""
     try:
@@ -231,16 +237,15 @@ def ParallelMap2(
     sys.stdout.flush()
     currentTime = time.time()
 
-    pcall = pcallWithGlobalParamsMultiArg if multiArg else pcallWithGlobalParamsSingleArg
-    pargs = zip(objects, itertools.repeat(globalParameters))
+    pcall = pcallWithGlobalParamsMultiArg if config.multiArg else pcallWithGlobalParamsSingleArg
 
     if joblibParallelSupportsGenerator():
-        rv = Parallel(n_jobs=threadCount, timeout=99999, return_as=return_as)(
-            delayed(pcall)(function, a, params) for a, params in pargs
+        rv = Parallel(n_jobs=threadCount, timeout=99999, return_as=config.return_as)(
+            delayed(pcall)(function, a) for a in objects
         )
     else:
         rv = Parallel(n_jobs=threadCount, timeout=99999)(
-            delayed(pcall)(function, a, params) for a, params in pargs
+            delayed(pcall)(function, a) for a in objects
         )
 
     totalTime = time.time() - currentTime

tensilelite/Tensile/CustomYamlLoader.py

@@ -139,7 +139,7 @@ def load_yaml_dict_item(yaml_path: Path, loader_type: yaml.Loader, key: str):
 
 def load_logic_gfx_arch(yaml_path: Path, loader_type: yaml.Loader = DEFAULT_YAML_LOADER):
     try:
-        GFX_ARCH_IDX = 2
+        GFX_ARCH_IDX = 2 # DataIndex.DEVICE_PROPERTIES
         arch = load_yaml_sequence_item(yaml_path, loader_type, GFX_ARCH_IDX)
 
         if isinstance(arch, dict):

tensilelite/Tensile/KernelWriterActivationEnumHeader.py

@@ -39,13 +39,19 @@ def __init__(self, state):
 
     # derive parameter
     self.language = "HIP"
-    self.kernelName = self.getKernelName()
+    self.kernelName = KernelWriterActivationEnumHeader.getKernelName(self)
 
   def keys(self):
     return self.getKernelName()
 
+  @staticmethod
+  def _getKernelName(solution):
+    s = "Gradient" if solution._state["ProblemType"]["Gradient"] else ""
+    return "Tensile%sActivationEnum_%s"%(s,
+                                         solution._state["ProblemType"]["ActivationComputeDataType"].toChar())
+
   def getKernelName(self):
-    return "Tensile%sActivationEnum_%s"%(self.actGradientPrefix, \
+    return "Tensile%sActivationEnum_%s"%(self.actGradientPrefix,
                                          self.state["ProblemType"]["ActivationComputeDataType"].toChar())
 
   def getSourceFileString(self):

tensilelite/Tensile/KernelWriterActivationFunction.py

@@ -56,6 +56,17 @@ def __init__(self, state, cxxCompiler: str, supportedISA: List[IsaVersion]):
   def keys(self):
     return self.getKernelName()
 
+
+  @staticmethod
+  def _getKernelName(solution):
+    actGradientPrefix = "Gradient" if solution._state["ProblemType"]["Gradient"] else ""
+    gaurdStr = "NG" if solution._state["ProblemType"]["ActivationNoGuard"] else ""
+    return "Tensile%sActivation%s_%s_%s"%(actGradientPrefix, \
+                                          gaurdStr, \
+                                          solution._state["ProblemType"]["ActivationComputeDataType"].toChar(), \
+                                          solution._state["ProblemType"]["ActivationType"])
+
+
   def getKernelName(self):
     return "Tensile%sActivation%s_%s_%s"%(self.actGradientPrefix, \
                                           self.gaurdStr, \

tensilelite/Tensile/KernelWriterActivationOnly.py

@@ -207,6 +207,25 @@ def kernelBody(self):
 
     return kStr
 
+  @classmethod
+  def _getKernelName(cls, solution):
+    indexChars = INDEX_CHARS
+    # C dimensions
+    name = "D"
+    for i in range(0, solution.state["ProblemType"]["NumIndicesC"]):
+      name += indexChars[i].lower()
+    name += "_"
+    name += solution.state["ProblemType"]["DestDataType"].toChar()
+    if solution.state["ProblemType"]["ActivationType"] != 'none':
+      if solution.state["ProblemType"]["ActivationType"] in ['all', 'hipblaslt_all']:
+        name += "_%s"%"A"
+      else:
+        name += "_%s"%str(solution.state["ProblemType"]["ActivationType"]).upper()
+      name += solution.state["ProblemType"]["ActivationComputeDataType"].toChar()
+    name += ("ng" if solution.state["ProblemType"]["ActivationNoGuard"] else "")
+
+    return name
+
 
   def getKernelName(self):
     indexChars = INDEX_CHARS

tensilelite/Tensile/KernelWriterAssembly.py

@@ -126,10 +126,6 @@ def getSourceFileString(self,
                           kernel,
                           useShortNames: bool=False) -> Tuple[int, str]:
     assert kernel["KernelLanguage"] == "Assembly"
-    # Skip if .o files will have already been built for this file
-    if kernel.duplicate:
-      self.language = "ASM"
-      return (0, "") # should this be an non zero number
 
     try:
       code = self._getCustomKernelSource(useShortNames, kernel, CUSTOM_KERNEL_PATH) if isCustomKernelConfig(kernel) else self._getKernelSource(kernel)

tensilelite/Tensile/KernelWriterBetaOnly.py

@@ -292,6 +292,26 @@ def kernelBodyBetaOnly(self):
     return kStr
 
 
+  @staticmethod
+  def _getKernelName(solution, btype=None):
+    indexChars = INDEX_CHARS
+    # C dimensions
+    name = "C"
+    for i in range(0, solution._state["ProblemType"]["NumIndicesC"]):
+      name += indexChars[i].lower()
+    name += "_"
+    name += solution._state["ProblemType"]["DestDataType"].toChar()
+    if solution._state["ProblemType"]["GroupedGemm"]:
+      name += "_GG"
+    else:
+      name += "" if solution._state["ProblemType"]["StridedBatched"] else "_GB"
+    if solution._state["ProblemType"]["BetaOnlyUseBias"]:
+      name += "_Bias%s"%btype.toChar()
+    name += "_GA" if solution._state["_GlobalAccumulation"] else ""
+
+    return name
+
+
   def getKernelName(self):
     indexChars = INDEX_CHARS
     # C dimensions

tensilelite/Tensile/KernelWriterConversion.py

@@ -785,6 +785,64 @@ def kernelBody(self):
     return kStr
 
 
+  @staticmethod
+  def _getKernelName(solution, num_elements_load, btype=None):
+    indexChars = INDEX_CHARS
+    # C dimensions
+    name = "C"
+    for i in range(0, solution._state["ProblemType"]["NumIndicesC"]):
+      name += indexChars[i].lower()
+    name += "_"
+
+    # add input datatype into kernel name (the datatype of workspace)
+    inputTypeStr = DataType("I").toChar() if solution._state["ProblemType"]["DataType"].isInt8() or solution._state["ProblemType"]["DataType"].isInt32() else \
+                                  (DataType("D").toChar() if solution._state["ProblemType"]["DataType"].isDouble() else DataType("S").toChar())
+
+    name += (inputTypeStr + solution._state["ProblemType"]["DestDataType"].toChar())
+
+    if solution._state["ProblemType"]["GroupedGemm"]:
+      name += "_GG"
+    else:
+      name += "" if solution._state["ProblemType"]["StridedBatched"] else "_GB"
+    if btype:
+      if solution._state["ProblemType"]["Gradient"]:
+        name += "_DBias%s"%(btype.toChar())
+        name += "_BiasSrc%s"%(solution._state["ProblemType"]["BiasSrc"])
+      else:
+        name += "_Bias%s"%btype.toChar()
+
+    factorDim =  0 if solution._state["ProblemType"]["Gradient"] else solution._state["ProblemType"]["UseBias"]
+    factorDim =  max(factorDim, solution._state["ProblemType"]["UseScaleAlphaVec"])
+    if factorDim > 1:
+        name += "_FD%s"%("N" if factorDim == 2 else "MN")
+
+    if solution._state["ProblemType"]["UseE"]:
+      if solution._state["ProblemType"]["Gradient"]:
+        name += "_Grad%s"%solution._state["ProblemType"]["DataTypeE"].toChar()
+      else:
+        name += "_Aux%s"%solution._state["ProblemType"]["DataTypeE"].toChar()
+
+    if ((solution._state["ProblemType"]["ActivationType"] != 'none') and solution._state["ActivationFused"]):
+      if solution._state["ProblemType"]["ActivationType"] == 'all':
+        name += "_A"
+      elif solution._state["ProblemType"]["ActivationType"] == 'hipblaslt_all':
+        name += "_HA"
+      else:
+        name += "_%s"%str(solution._state["ProblemType"]["ActivationType"]).upper()
+      name += solution._state["ProblemType"]["ActivationComputeDataType"].toChar()
+      name += ("ng" if solution._state["ProblemType"]["ActivationNoGuard"] else "")
+    if solution._state["ProblemType"]["UseScaleAB"] == "Scalar":
+      name += "_ScaleAB"
+    elif solution._state["ProblemType"]["UseScaleAB"] == "Vector":
+      name += "_ScaleABVec"
+    name += "_ScaleCD" if solution._state["ProblemType"]["UseScaleCD"] else ""
+    name += "_ScaleAlphaVec" if solution._state["ProblemType"]["UseScaleAlphaVec"] else ""
+    name += "_PostGSU" + str(solution._state["GlobalSplitU"])
+    if num_elements_load != None:
+      name += "_VW" + str(num_elements_load)
+    return name
+
+
   def getKernelName(self):
     indexChars = INDEX_CHARS
     # C dimensions

tensilelite/Tensile/KernelWriterReduction.py

@@ -54,6 +54,21 @@ def kernelBody(self):
         return kStr
 
 
+    @staticmethod
+    def _getKernelName(solution, btype):
+        # C dimensions
+        indexChars = INDEX_CHARS
+        indicesStr = ""
+        for i in range(0, solution._state["ProblemType"]["NumIndicesC"]):
+            c = indexChars[i].lower()
+            indicesStr += indexChars[i].lower()
+        name = "D"
+        name += indicesStr
+        name += "_%s%s"%(btype.toChar(), solution._state["ProblemType"]["ComputeDataType"].toChar())
+        name += "_Reduction"
+        return name
+
+
     def getKernelName(self):
         name = "D"
         name += self.indicesStr