diff --git a/tensilelite/Tensile/BenchmarkProblems.py b/tensilelite/Tensile/BenchmarkProblems.py index 82b472c8d1..b758fc5be9 100644 --- a/tensilelite/Tensile/BenchmarkProblems.py +++ b/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: diff --git a/tensilelite/Tensile/CodeObjectName.py b/tensilelite/Tensile/CodeObjectName.py new file mode 100644 index 0000000000..e70370f282 --- /dev/null +++ b/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 \ No newline at end of file diff --git a/tensilelite/Tensile/Common/Parallel.py b/tensilelite/Tensile/Common/Parallel.py index 77adfd26c1..46c6415102 100644 --- a/tensilelite/Tensile/Common/Parallel.py +++ b/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 diff --git a/tensilelite/Tensile/CustomYamlLoader.py b/tensilelite/Tensile/CustomYamlLoader.py index bab8c68750..94bffe40b2 100644 --- a/tensilelite/Tensile/CustomYamlLoader.py +++ b/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): diff --git a/tensilelite/Tensile/KernelWriterActivationEnumHeader.py b/tensilelite/Tensile/KernelWriterActivationEnumHeader.py index b7b4f563f0..cbc43c8dbc 100644 --- a/tensilelite/Tensile/KernelWriterActivationEnumHeader.py +++ b/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): diff --git a/tensilelite/Tensile/KernelWriterActivationFunction.py b/tensilelite/Tensile/KernelWriterActivationFunction.py index f8433d30ac..1e68950a34 100644 --- a/tensilelite/Tensile/KernelWriterActivationFunction.py +++ b/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, \ diff --git a/tensilelite/Tensile/KernelWriterActivationOnly.py b/tensilelite/Tensile/KernelWriterActivationOnly.py index c4dea91e04..f35f2e5217 100644 --- a/tensilelite/Tensile/KernelWriterActivationOnly.py +++ b/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 diff --git a/tensilelite/Tensile/KernelWriterAssembly.py b/tensilelite/Tensile/KernelWriterAssembly.py index 2215dbbc48..16fb0cfbd6 100644 --- a/tensilelite/Tensile/KernelWriterAssembly.py +++ b/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) diff --git a/tensilelite/Tensile/KernelWriterBetaOnly.py b/tensilelite/Tensile/KernelWriterBetaOnly.py index fb0ab1e802..d4d97ab4f1 100644 --- a/tensilelite/Tensile/KernelWriterBetaOnly.py +++ b/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 diff --git a/tensilelite/Tensile/KernelWriterConversion.py b/tensilelite/Tensile/KernelWriterConversion.py index 2542d99f82..0ac473f8f5 100644 --- a/tensilelite/Tensile/KernelWriterConversion.py +++ b/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 diff --git a/tensilelite/Tensile/KernelWriterReduction.py b/tensilelite/Tensile/KernelWriterReduction.py index 44fdf8a7e5..e82bc2c76c 100644 --- a/tensilelite/Tensile/KernelWriterReduction.py +++ b/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 diff --git a/tensilelite/Tensile/LibraryIO.py b/tensilelite/Tensile/LibraryIO.py index 09b7a0f493..5914b834d9 100644 --- a/tensilelite/Tensile/LibraryIO.py +++ b/tensilelite/Tensile/LibraryIO.py @@ -25,6 +25,7 @@ from .CustomKernels import getCustomKernelConfig from . import SolutionLibrary from .CustomYamlLoader import load_yaml_stream +from Tensile.CodeObjectName import codeObjectFileBaseName from Tensile import __version__ from Tensile.Common import printExit, printWarning, print2, \ versionIsCompatible, IsaInfo, DepthUConfig @@ -32,6 +33,7 @@ from Tensile.SolutionStructs import Solution, ProblemSizes from Tensile.SolutionStructs.Problem import ProblemType +from enum import IntEnum from typing import NamedTuple, List, Dict import os import sys @@ -275,6 +277,22 @@ def parseSolutionsData( problemSizes = ProblemSizes(problemType, problemSizesConfig) return (problemSizes, solutions) +class DataIndex(IntEnum): + MINIMUM_REQUIRED_VERSION=0 + SCHEDULE_NAME=1 + DEVICE_PROPERTIES=2 + DEVICE_NAMES=3 + PROBLEM_TYPE=4 + SOLUTIONS=5 + INDEX_ORDER=6 + EXACT_LOGIC=7 + RANGE_LOGIC=8 + PERF_METRIC=10 + LIBRARY_TYPE=11 + CODE_OBJECT_NAME=12 + + def __index__(self): + return self.value class LibraryLogic(NamedTuple): """Return tuple for parseLibraryLogicData()""" @@ -284,6 +302,7 @@ class LibraryLogic(NamedTuple): solutions: list exactLogic: list library: SolutionLibrary.MasterSolutionLibrary + srcFile: str def parseLibraryLogicFile( filename, @@ -367,6 +386,7 @@ def solutionStateToSolution(solutionState, assembler, isaInfoMap) -> Solution: isaInfoMap, srcFile ) + solutionObject["LogicFileName"] = srcFile solutionProblemType = solutionObject["ProblemType"] if problemType != solutionProblemType: # find the mismatched items in ProblemType @@ -380,7 +400,8 @@ def solutionStateToSolution(solutionState, assembler, isaInfoMap) -> Solution: solutions = [solutionStateToSolution(solutionState, assembler, isaInfoMap) for solutionState in data["Solutions"]] - newLibrary, _ = SolutionLibrary.MasterSolutionLibrary.FromOriginalState( + n = codeObjectFileBaseName(data) + newLibrary, codeObjectFile = SolutionLibrary.MasterSolutionLibrary.FromOriginalState( data, solutions, splitGSU, @@ -391,9 +412,13 @@ def solutionStateToSolution(solutionState, assembler, isaInfoMap) -> Solution: isaInfoMap, lazyLibraryLoading ) + assert n == codeObjectFile, f"codeObjectFileBaseName computed different name than FromOriginalState \n {n}\n {codeObjectFile}" + + for solution in solutions: + solution["codeObjectFile"] = codeObjectFile return LibraryLogic(data["ScheduleName"], data["ArchitectureName"], problemType, solutions, \ - data.get("ExactLogic"), newLibrary) + data.get("ExactLogic"), newLibrary, srcFile) def parseLibraryLogicList(data, srcFile="?"): @@ -401,33 +426,32 @@ def parseLibraryLogicList(data, srcFile="?"): if len(data) < 9: printExit("Library logic file {} is missing required fields (len = {} < 9)" \ .format(srcFile, len(data))) - rv = {} - rv["MinimumRequiredVersion"] = data[0]["MinimumRequiredVersion"] - rv["ScheduleName"] = data[1] - rv["DeviceNames"] = data[3] - rv["ProblemType"] = data[4] - rv["Solutions"] = data[5] - - if type(data[2]) is dict: - rv["ArchitectureName"] = data[2]["Architecture"] - rv["CUCount"] = data[2]["CUCount"] + rv["MinimumRequiredVersion"] = data[DataIndex.MINIMUM_REQUIRED_VERSION]["MinimumRequiredVersion"] + rv["ScheduleName"] = data[DataIndex.SCHEDULE_NAME] + rv["DeviceNames"] = data[DataIndex.DEVICE_NAMES] + rv["ProblemType"] = data[DataIndex.PROBLEM_TYPE] + rv["Solutions"] = data[DataIndex.SOLUTIONS] + + if type(data[DataIndex.DEVICE_PROPERTIES]) is dict: + rv["ArchitectureName"] = data[DataIndex.DEVICE_PROPERTIES]["Architecture"] + rv["CUCount"] = data[DataIndex.DEVICE_PROPERTIES]["CUCount"] else: - rv["ArchitectureName"] = data[2] + rv["ArchitectureName"] = data[DataIndex.DEVICE_PROPERTIES] rv["CUCount"] = None # TODOBEN: figure out what to do with these... - rv["ExactLogic"] = data[7] - rv["RangeLogic"] = data[8] + rv["ExactLogic"] = data[DataIndex.EXACT_LOGIC] + rv["RangeLogic"] = data[DataIndex.RANGE_LOGIC] # optional fields - if len(data) > 10 and data[10]: - rv["PerfMetric"] = data[10] + if len(data) > 10 and data[DataIndex.PERF_METRIC]: + rv["PerfMetric"] = data[DataIndex.PERF_METRIC] # library logic fields libraryType = None - if len(data) > 11 and data[11]: - libraryType = data[11] + if len(data) > 11 and data[DataIndex.LIBRARY_TYPE]: + libraryType = data[DataIndex.LIBRARY_TYPE] else: printExit("Library logic file {} is missing required field matching property." \ .format(srcFile)) @@ -435,13 +459,13 @@ def parseLibraryLogicList(data, srcFile="?"): rv["LibraryType"] = "FreeSize" rv["Library"] = {} rv["Library"]["indexOrder"] = None - rv["Library"]["table"] = [0, len(data[5])] + rv["Library"]["table"] = [0, len(data[DataIndex.SOLUTIONS])] rv["Library"]["distance"] = None else: rv["LibraryType"] = "Matching" rv["Library"] = {} - rv["Library"]["indexOrder"] = data[6] - rv["Library"]["table"] = data[7] + rv["Library"]["indexOrder"] = data[DataIndex.INDEX_ORDER] + rv["Library"]["table"] = data[DataIndex.EXACT_LOGIC] rv["Library"]["distance"] = libraryType return rv @@ -449,20 +473,20 @@ def parseLibraryLogicList(data, srcFile="?"): def rawLibraryLogic(data): """Returns a tuple of the data in a library logic file.""" - versionString = data[0] - scheduleName = data[1] - architectureName = data[2] - deviceNames = data[3] - problemTypeState = data[4] - solutionStates = data[5] - indexOrder = data[6] - exactLogic = data[7] - rangeLogic = data[8] + versionString = data[DataIndex.MINIMUM_REQUIRED_VERSION] + scheduleName = data[DataIndex.SCHEDULE_NAME] + architectureName = data[DataIndex.DEVICE_PROPERTIES] + deviceNames = data[DataIndex.DEVICE_NAMES] + problemTypeState = data[DataIndex.PROBLEM_TYPE] + solutionStates = data[DataIndex.SOLUTIONS] + indexOrder = data[DataIndex.INDEX_ORDER] + exactLogic = data[DataIndex.EXACT_LOGIC] + rangeLogic = data[DataIndex.RANGE_LOGIC] otherFields = [] dataLength = len(data) - if dataLength > 9: - for idx in range(9, dataLength): + if dataLength > 10: + for idx in range(10, dataLength): otherFields.append(data[idx]) return (versionString, scheduleName, architectureName, deviceNames,\ diff --git a/tensilelite/Tensile/SolutionLibrary.py b/tensilelite/Tensile/SolutionLibrary.py index ee69c97ac3..e2b866457e 100644 --- a/tensilelite/Tensile/SolutionLibrary.py +++ b/tensilelite/Tensile/SolutionLibrary.py @@ -591,7 +591,7 @@ def merge(self, other, startIndex=0): self.lazyLibraries[name] = lib reIndexMap = {} - for k, s in other.solutions.items(): + for k, s in list(other.solutions.items()): reIndexMap[s.index] = curIndex s.index = curIndex self.solutions[curIndex] = s diff --git a/tensilelite/Tensile/SolutionStructs/Solution.py b/tensilelite/Tensile/SolutionStructs/Solution.py index 250810f738..47f17698b3 100644 --- a/tensilelite/Tensile/SolutionStructs/Solution.py +++ b/tensilelite/Tensile/SolutionStructs/Solution.py @@ -25,7 +25,6 @@ from Tensile.TensileInstructions.Base import fastdeepcopy as deepcopy import collections import math - from enum import Enum from typing import List, Dict @@ -54,6 +53,97 @@ from .Utilities import reject, roundupRatio, pvar + +def conversionKernelObjectsNames(solution): + # need to check that the fields mutated in the the init function aren't used in the naming function + # e.g. usebias got set to zero if gradient which changes behavior of name function. + conversionKernelObjectsNames = [] + load_vector_width = [1, 2] if solution["ProblemType"]["DataType"].isDouble() else [1, 2, 4] + gsuList = [internalParameters["GlobalSplitUPGR"]] + if solution["GlobalSplitUAlgorithm"] == "SingleBuffer": + gsuList = [1] + elif solution["GlobalSplitUAlgorithm"] == "MultipleBufferSingleKernel": + return + for vw in load_vector_width: + for _ in gsuList: + if solution["ProblemType"]["UseBias"]: + typeList = solution["ProblemType"]["BiasDataTypeList"] + if solution["ProblemType"]["Gradient"]: + # # If gradient + bias D, generates a normal GSU kernel for bias D = nullptr case + conversionKernelObjectsNames.append(KernelWriterConversion._getKernelName(solution, vw)) + for btype in typeList: + conversionKernelObjectsNames.append(KernelWriterConversion._getKernelName(solution, vw, btype)) + else: + #printExit("don't use bias") + conversionKernelObjectsNames.append(KernelWriterConversion._getKernelName(solution, vw)) + return conversionKernelObjectsNames if conversionKernelObjectsNames else [] + +def activationEnumHeaderObjectsNames(solution): + activationEnumHeaderObjectsNames = [] + if solution["ProblemType"]["ActivationType"] in ['all', 'hipblaslt_all']: + activationEnumHeaderObjectsNames.append(KernelWriterActivationEnumHeader._getKernelName(solution)) + return activationEnumHeaderObjectsNames + +def activationFunctionObjectsNames(solution): + activationFunctionObjectsNames = [] + if solution["ProblemType"]["ActivationType"] in ['all', 'hipblaslt_all']: + activationFunctionObjectsNames.append(KernelWriterActivationFunction._getKernelName(solution)) + return activationFunctionObjectsNames + +def activationOnlyKernelObjectsNames(solution): + activationOnlyKernelObjectsNames = [] + if (solution["ActivationFused"] == False) and (solution["ProblemType"]["ActivationType"] != 'none'): + activationOnlyKernelObjectsNames.append(KernelWriterActivationOnly._getKernelName(solution)) + return activationOnlyKernelObjectsNames + +def reductionKernelObjectsNames(solution): + reductionKernelObjectsNames = [] + if solution["ProblemType"]["Gradient"] and solution["ProblemType"]["UseBias"]: + for btype in solution["ProblemType"]["BiasDataTypeList"]: + reductionKernelObjectsNames.append(KernelWriterReduction._getKernelName(solution, btype)) + return reductionKernelObjectsNames + +def betaOnlyKernelObjectsNames(solution): + betaOnlyKernelObjectsNames = [] + if solution["GlobalSplitU"] > 1 or (solution["StreamK"] > 0 and solution["StreamKAtomic"] == 1): + if solution["ProblemType"]["UseBias"]: + for btype in solution["ProblemType"]["BiasDataTypeList"]: + betaOnlyKernelObjectsNames.append(KernelWriterBetaOnly._getKernelName(solution, btype)) + else: + betaOnlyKernelObjectsNames.append(KernelWriterBetaOnly._getKernelName(solution)) + return betaOnlyKernelObjectsNames + + +def kernelObjectNames(solution): + result = [] + temp = conversionKernelObjectsNames(solution) + if temp: + result.extend(temp) + temp = activationEnumHeaderObjectsNames(solution) + if temp: + result.extend(temp) + temp = activationFunctionObjectsNames(solution) + if temp: + result.extend(temp) + temp = activationOnlyKernelObjectsNames(solution) + if temp: + result.extend(temp) + temp = reductionKernelObjectsNames(solution) + if temp: + result.extend(temp) + temp = betaOnlyKernelObjectsNames(solution) + if temp: + result.extend(temp) + return result + + +# print a labled variable +def pvar(state, field): + return field + "=" + str(state[field]) + +def roundupRatio(dividend, divisor): + return int(math.ceil(float(dividend) / float(divisor))) + class Fbs(Enum): Free=0 # Expect to be free dimension Batch=1 # Expect to be batch dimension @@ -235,7 +325,7 @@ def __init__( ) self._name = config["CustomKernelName"] if "CustomKernelName" in config and config["CustomKernelName"] else None - self.initHelperKernelObjects(targetIsas) + #self.initHelperKernelObjects(targetIsas) # these keys are copied from ProblemType to internal that may be overridden InternalKeys = ["UseSgprForGRO","VectorStore"] @@ -254,17 +344,38 @@ def getKernels(self): ######################################## # create Helper Kernels def initHelperKernelObjects(self, supportedISA: List[IsaVersion]): - self.initBetaOnlyKernelObjects() - self.initConversionKernelObjects(supportedISA) - self.initActivationEnumHeaderObjects() - self.initActivationFunctionObjects(supportedISA) - self.initActivationOnlyKernelObjects() - self.initReductionKernelObjects() + result = [] + + temp = self.initActivationEnumHeaderObjects() + if temp: + result.extend(temp) + + temp = self.initActivationFunctionObjects(supportedISA) + if temp: + result.extend(temp) + + temp = self.initBetaOnlyKernelObjects() + if temp: + result.extend(temp) + + temp = self.initConversionKernelObjects(supportedISA) + if temp: + result.extend(temp) + + temp = self.initActivationOnlyKernelObjects() + if temp: + result.extend(temp) + + temp = self.initReductionKernelObjects() + if temp: + result.extend(temp) + + return result ######################################## # create BetaOnly Kernels def initBetaOnlyKernelObjects(self): - self.betaOnlyKernelObjects = [] + betaOnlyKernelObjects = [] if self["GlobalSplitU"] > 1 or (self["StreamK"] > 0 and self["StreamKAtomic"] == 1): if self["ProblemType"]["UseBias"]: for btype in self["ProblemType"]["BiasDataTypeList"]: @@ -275,20 +386,21 @@ def initBetaOnlyKernelObjects(self): state["ProblemType"]["BiasDataType"] = deepcopy(btype) state["KernelLanguage"] = "Source" state["_GlobalAccumulation"] = self["_GlobalAccumulation"] - self.betaOnlyKernelObjects.append(KernelWriterBetaOnly(state)) + betaOnlyKernelObjects.append(KernelWriterBetaOnly(state)) else: state = {} state["ProblemType"] = deepcopy(self["ProblemType"]) state["ProblemType"]["GroupedGemm"] = False state["KernelLanguage"] = "Source" state["_GlobalAccumulation"] = self["_GlobalAccumulation"] - self.betaOnlyKernelObjects.append(KernelWriterBetaOnly(state)) + betaOnlyKernelObjects.append(KernelWriterBetaOnly(state)) + return betaOnlyKernelObjects ######################################## # create Conversion Kernels - def initConversionKernelObjects(self, supportedArchs: List[tuple]): - self.conversionKernelObjects = [] + def initConversionKernelObjects(self, supportedArchs): + conversionKernelObjects = [] load_vector_width = [1, 2] if self["ProblemType"]["DataType"].isDouble() else [1, 2, 4] genPGRPostKernels = True gsuList = [internalParameters["GlobalSplitUPGR"]] @@ -314,7 +426,7 @@ def initConversionKernelObjects(self, supportedArchs: List[tuple]): state["UnrollOnly"] = unrollOnly state["_GlobalAccumulation"] = self["_GlobalAccumulation"] state["ActivationFused"] = self["ActivationFused"] - self.conversionKernelObjects.append(KernelWriterConversion(state, vw, supportedArchs, self.isaInfoMap)) + conversionKernelObjects.append(KernelWriterConversion(state, vw, supportedArchs, self.isaInfoMap)) for btype in typeList: state = {} state["ProblemType"] = deepcopy(self["ProblemType"]) @@ -327,7 +439,7 @@ def initConversionKernelObjects(self, supportedArchs: List[tuple]): state["UnrollOnly"] = unrollOnly state["_GlobalAccumulation"] = self["_GlobalAccumulation"] state["ActivationFused"] = self["ActivationFused"] - self.conversionKernelObjects.append(KernelWriterConversion(state, vw, supportedArchs, self.isaInfoMap)) + conversionKernelObjects.append(KernelWriterConversion(state, vw, supportedArchs, self.isaInfoMap)) else: state = {} state["ProblemType"] = deepcopy(self["ProblemType"]) @@ -338,31 +450,34 @@ def initConversionKernelObjects(self, supportedArchs: List[tuple]): state["UnrollOnly"] = unrollOnly state["_GlobalAccumulation"] = self["_GlobalAccumulation"] state["ActivationFused"] = self["ActivationFused"] - self.conversionKernelObjects.append(KernelWriterConversion(state, vw, supportedArchs, self.isaInfoMap)) + conversionKernelObjects.append(KernelWriterConversion(state, vw, supportedArchs, self.isaInfoMap)) + return conversionKernelObjects def initActivationEnumHeaderObjects(self): - self.activationEnumHeaderObjects = [] + activationEnumHeaderObjects = [] if self["ProblemType"]["ActivationType"] in ['all', 'hipblaslt_all']: state = {} state["ProblemType"] = deepcopy(self["ProblemType"]) state["ProblemType"]["GroupedGemm"] = False state["KernelLanguage"] = "Source" - self.activationEnumHeaderObjects.append(KernelWriterActivationEnumHeader(state)) + activationEnumHeaderObjects.append(KernelWriterActivationEnumHeader(state)) + return activationEnumHeaderObjects def initActivationFunctionObjects(self, supportedISA: List[IsaVersion]): - self.activationFunctionObjects = [] + activationFunctionObjects = [] if self["ProblemType"]["ActivationType"] in ['all', 'hipblaslt_all']: state = {} state["ProblemType"] = deepcopy(self["ProblemType"]) state["ProblemType"]["GroupedGemm"] = False state["KernelLanguage"] = "Source" state["Kernel"] = {"WavefrontSize": self["WavefrontSize"], "ISA": tuple(self["ISA"])} - if not isinstance(supportedISA, list): - raise Exception(f"{type(supportedISA)}") - self.activationFunctionObjects.append(KernelWriterActivationFunction(state, str(self.assembler.path), supportedISA)) + activationFunctionObjects.append(KernelWriterActivationFunction(state, + str(self.assembler.path), + supportedISA)) + return activationFunctionObjects def initActivationOnlyKernelObjects(self): - self.activationOnlyKernelObjects = [] + activationOnlyKernelObjects = [] if (self["ActivationFused"] == False) and (self["ProblemType"]["ActivationType"] != 'none') : state = {} state["ProblemType"] = deepcopy(self["ProblemType"]) @@ -372,10 +487,11 @@ def initActivationOnlyKernelObjects(self): state["KernelLanguage"] = "Source" state["_GlobalAccumulation"] = self["_GlobalAccumulation"] state["ActivationFused"] = self["ActivationFused"] - self.activationOnlyKernelObjects.append(KernelWriterActivationOnly(state)) + activationOnlyKernelObjects.append(KernelWriterActivationOnly(state)) + return activationOnlyKernelObjects def initReductionKernelObjects(self): - self.reductionKernelObjects = [] + reductionKernelObjects = [] if self["ProblemType"]["Gradient"] and self["ProblemType"]["UseBias"]: for btype in self["ProblemType"]["BiasDataTypeList"]: state = {} @@ -383,14 +499,37 @@ def initReductionKernelObjects(self): state["ProblemType"]["GroupedGemm"] = False state["ProblemType"]["BiasDataTypeList"] = [] state["ProblemType"]["BiasDataType"] = deepcopy(btype) - self.reductionKernelObjects.append(KernelWriterReduction(state)) + reductionKernelObjects.append(KernelWriterReduction(state)) + return reductionKernelObjects + + def initBetaOnlyKernelObjects(self): + betaOnlyKernelObjects = [] + if self["GlobalSplitU"] > 1 or (self["StreamK"] > 0 and self["StreamKAtomic"] == 1): + if self["ProblemType"]["UseBias"]: + for btype in self["ProblemType"]["BiasDataTypeList"]: + state = {} + state["ProblemType"] = deepcopy(self["ProblemType"]) + state["ProblemType"]["GroupedGemm"] = False + state["ProblemType"]["BiasDataTypeList"] = [] + state["ProblemType"]["BiasDataType"] = deepcopy(btype) + state["KernelLanguage"] = "Source" + state["_GlobalAccumulation"] = self["_GlobalAccumulation"] + betaOnlyKernelObjects.append(KernelWriterBetaOnly(state)) + else: + state = {} + state["ProblemType"] = deepcopy(self["ProblemType"]) + state["ProblemType"]["GroupedGemm"] = False + state["KernelLanguage"] = "Source" + state["_GlobalAccumulation"] = self["_GlobalAccumulation"] + betaOnlyKernelObjects.append(KernelWriterBetaOnly(state)) + return betaOnlyKernelObjects ######################################## # get Helper Kernels - def getHelperKernelObjects(self): - return self.activationEnumHeaderObjects + self.activationFunctionObjects + \ - self.betaOnlyKernelObjects + self.conversionKernelObjects + \ - self.activationOnlyKernelObjects + self.reductionKernelObjects + #def getHelperKernelObjects(self): + # return self.activationEnumHeaderObjects + self.activationFunctionObjects + \ + # self.betaOnlyKernelObjects + self.conversionKernelObjects + \ + # self.activationOnlyKernelObjects + self.reductionKernelObjects ######################################## @@ -3349,3 +3488,5 @@ def __ne__(self, other): if result is NotImplemented: return result return not result + + diff --git a/tensilelite/Tensile/Tensile.py b/tensilelite/Tensile/Tensile.py index 90b8af31a4..bd5902d789 100644 --- a/tensilelite/Tensile/Tensile.py +++ b/tensilelite/Tensile/Tensile.py @@ -200,6 +200,10 @@ def splitExtraParameters(par): action="store", default=ToolchainDefaults.ASSEMBLER, help="select which assembler to use") argParser.add_argument("--offload-bundler", dest="OffloadBundler", \ action="store", default=ToolchainDefaults.OFFLOAD_BUNDLER, help="select which offload bundler to use") + argParser.add_argument("--roc-obj-extract", dest="RocObjExtract", \ + action="store", default=ToolchainDefaults.ROC_OBJ_EXTRACT) + argParser.add_argument("--roc-obj-ls", dest="RocObjLs", \ + action="store", default=ToolchainDefaults.ROC_OBJ_LS) argParser.add_argument("--device-enumerator", dest="DeviceEnumerator", \ action="store", default=ToolchainDefaults.DEVICE_ENUMERATOR, help="select which device enumerator to use") argParser.add_argument("--logic-format", dest="LogicFormat", choices=["yaml", "json"], \ @@ -456,9 +460,13 @@ def Tensile(userArgs): cxxCompiler, \ cCompiler, \ offloadBundler, \ + ls, \ + extract, \ enumerator = validateToolchain(args.CxxCompiler, - args.CCompiler, + args.CCompiler, args.OffloadBundler, + args.RocObjLs, + args.RocObjExtract, ToolchainDefaults.DEVICE_ENUMERATOR) asmToolchain = makeAssemblyToolchain( cxxCompiler, @@ -468,8 +476,13 @@ def Tensile(userArgs): srcToolchain = makeSourceToolchain( cxxCompiler, offloadBundler, + ls, + extract, + cpu_threads=4, ) + + currentIsa = detectGlobalCurrentISA(device_id, enumerator) if currentIsa == IsaVersion(9,5,0): printWarning("HardwareMonitor currently disabled for gfx950") diff --git a/tensilelite/Tensile/TensileCreateLibrary/IO.py b/tensilelite/Tensile/TensileCreateLibrary/IO.py new file mode 100644 index 0000000000..841defd27f --- /dev/null +++ b/tensilelite/Tensile/TensileCreateLibrary/IO.py @@ -0,0 +1,113 @@ +################################################################################ +# +# Copyright (C) 2025 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +# SOFTWARE. +# +################################################################################ + +from Tensile.Common import CHeader, printExit, state, DepthUConfig +from Tensile.LibraryIO import parseLibraryLogicFile, write +from Tensile.SolutionLibrary import MasterSolutionLibrary +from Tensile.Utilities.RequiredParameters import getRequiredParametersMin + +from os import getpid +from pathlib import Path +from typing import Dict, Union + + +def generateSolutionsAndLibraries(assembler, isaInfoMap, lazy, logicFiles): + solutions = [] + libraries = [] + for logicFileGroup in logicFiles: + for logicFile in logicFileGroup[1]: + libraryLogic = parseLibraryLogicFile(logicFile, assembler, False, False, False, DepthUConfig(), isaInfoMap, lazy) + solutions.extend(libraryLogic.solutions) + libraries.append((libraryLogic.architecture, libraryLogic.library)) + + return solutions, libraries + + +def writeAssembly(asmPath: Path, result: tuple): + if result[0]: + printExit(f"Failed to generate kernel {result[3]} because it has error code {result[0]}") + filepath = asmPath / f"{result.name}.s" + isa = result[5] + wfsize = result[6] + with open(filepath, "w", encoding="utf-8") as f: + f.write(result[1]) + del result # result.src is very large so let gc know to clean up asap + + return filepath, isa, wfsize + + +def writeHelper(outputPath, kernelHelperObj) -> str: + name = kernelHelperObj.getKernelName() + KERNEL_HELPER_FILENAME_CPP = name + ".cpp" + KERNEL_HELPER_FILENAME_H = name + ".h" + kernelSourceFilename = str(Path(outputPath) / "Kernels" / KERNEL_HELPER_FILENAME_CPP) + kernelHeaderFilename = str(Path(outputPath) / "Kernels" / KERNEL_HELPER_FILENAME_H) + + with open(kernelHeaderFilename, "w", encoding="utf-8") as kernelHeaderFile, \ + open(kernelSourceFilename, "w", encoding="utf-8") as kernelSourceFile: + kernelSourceFile.write(CHeader) + kernelHeaderFile.write(CHeader) + kernelSourceFile.write("#include \"{}.h\"\n".format(name)) + kernelHeaderFile.write("#pragma once\n") + kernelHeaderFile.write("#include \n") + kernelHeaderFile.write("#include \n\n") + kernelHeaderFile.write("#include \"KernelHeader.h\"\n\n") + + if "Enum" not in name: + kernelHeaderFile.write("#include \"Kernels/TensileActivationEnum_S.h\"\n") + kernelHeaderFile.write("#include \"Kernels/TensileActivationEnum_I.h\"\n") + if hasattr(kernelHelperObj, "actGradientPrefix") and kernelHelperObj.actGradientPrefix == "Gradient": + kernelHeaderFile.write("#include \"Kernels/TensileGradientActivationEnum_S.h\"\n") + if "TensileActivation_S" not in name and "TensileActivation_I" not in name and "Enum" not in name: + kernelHeaderFile.write("#include \"Kernels/TensileActivation_S_Hipblaslt_all.h\"\n") + kernelHeaderFile.write("#include \"Kernels/TensileActivation_I_Hipblaslt_all.h\"\n") + if "TensileGradientActivation_S" not in name and "Enum" not in name: + if hasattr(kernelHelperObj, "actGradientPrefix") and kernelHelperObj.actGradientPrefix == "Gradient": + kernelHeaderFile.write("#include \"Kernels/TensileGradientActivation_S_Hipblaslt_all.h\"\n") + + HeaderText = "" + (err, src) = kernelHelperObj.getSourceFileString() + kernelSourceFile.write(src) + if err: + print("*** warning: invalid kernel#%u" % name) + HeaderText += kernelHelperObj.getHeaderFileString() + kernelHeaderFile.write(HeaderText) + + return kernelSourceFilename + + +def genLazyMasterSolutionLibrary(libraryPath, libraryFormat, masterLib): + # Can we do this asynchronously before the call to writeSolutionsAndKernels? + for name, lib in list(masterLib.lazyLibraries.items()): + catalogPath = libraryPath / name + lib.applyNaming(False) + write(str(catalogPath), state(lib), libraryFormat) + + +def generateParentLibrary(libraryFormat: str, libraryPath: Union[Path, str], masterLibs: Dict[str, MasterSolutionLibrary], lazyLibraryLoading: bool): + base = "TensileLibrary_lazy_" if lazyLibraryLoading else "TensileLibrary_" + for arch, masterLib in masterLibs.items(): + name = base + arch + masterLib.applyNaming(False) + write(str(libraryPath / name), state(masterLib), libraryFormat) diff --git a/tensilelite/Tensile/TensileCreateLibrary/Logic.py b/tensilelite/Tensile/TensileCreateLibrary/Logic.py new file mode 100644 index 0000000000..5a6ce63dd7 --- /dev/null +++ b/tensilelite/Tensile/TensileCreateLibrary/Logic.py @@ -0,0 +1,100 @@ +################################################################################ +# +# Copyright (C) 2025 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +# SOFTWARE. +# +################################################################################ + +from Tensile.Common import print2 +from Tensile.LibraryIO import DataIndex +from Tensile.CustomYamlLoader import load_logic_gfx_arch, load_yaml_sequence_item +from Tensile.CodeObjectName import codeObjectFileBaseName + +from glob import iglob +from os.path import getsize +from pathlib import Path +from typing import List +from subprocess import PIPE, run + +def logicFileList(archs, logicPath: Path, logicFilter: str, experimental: bool): + def archMatch(arch: str, archs: List[str]): + return (arch in archs) or any(a.startswith(arch) for a in archs) + def validLogicFile(p: Path): + return p.suffix == ".yaml" and ("all" in archs or archMatch(load_logic_gfx_arch(p), archs)) + + assert logicPath.exists(), f"LogicPath {str(logicPath)} doesn't exist" + + globPattern = str(logicPath / f"**/{logicFilter}.yaml") + logicFiles = (str(logicPath / file) for file in iglob(globPattern, recursive=True)) + if not experimental: + logicFiles = [file for file in logicFiles if "experimental" not in map(str.lower, Path(file).parts)] + + logicFiles = [file for file in logicFiles if validLogicFile(Path(file))] + + print(f"# LogicFilter: {globPattern}") + print(f"# Experimental: {experimental}") + print2(f"# LibraryLogicFiles: {len(logicFiles)}") + for logicFile in logicFiles: + print2("# %s" % logicFile) + + return logicFiles + + +def distribute(lst, n): + import heapq + list_of_lists = [[] for _ in range(n)] + totals = [(0, i) for i in range(n)] + heapq.heapify(totals) + for value, f in lst: + total, index = heapq.heappop(totals) + list_of_lists[index].append((value, f)) + heapq.heappush(totals, (total + value, index)) + return sorted(list_of_lists, key=lambda x: sum(first for first, _ in x), reverse=True) + + +def getCoFileNames(logicFile): + from yaml import Loader + data = {} + data["ProblemType"] = load_yaml_sequence_item(logicFile, Loader, DataIndex.PROBLEM_TYPE.value) + properties = load_yaml_sequence_item(logicFile, Loader, DataIndex.DEVICE_PROPERTIES.value) + if isinstance(properties, dict): + data["ArchitectureName"] = properties["Architecture"] + data["CUCount"] = properties["CUCount"] + else: + data["ArchitectureName"] = properties + data["CUCount"] = None + data["PerfMetric"] = load_yaml_sequence_item(logicFile, Loader, DataIndex.PERF_METRIC.value) + return codeObjectFileBaseName(data), logicFile + + +def schedule(cofiles: list, numberOfTasks: int, procs: int): + problemMap = {} + for codeObjectFile, logicFile in cofiles: + if codeObjectFile in problemMap: + problemMap[codeObjectFile].append(logicFile) + else: + problemMap[codeObjectFile] = [logicFile] + + result = [] + for codeObjectFile, logicFiles in problemMap.items(): + count = sum(getsize(logicFile) for logicFile in logicFiles) + result.append((count, logicFiles)) + + return list(filter(lambda x: x != [], distribute(result, numberOfTasks))) # need to convert list of list of tuples to list of list of strings diff --git a/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py b/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py index 7065b424a8..abf6845fef 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py +++ b/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py @@ -72,6 +72,8 @@ def parseArguments(input: Optional[List[str]] = None) -> Dict[str, Any]: argParser.add_argument( "--assembler", dest="Assembler", action="store", default=ToolchainDefaults.ASSEMBLER ) + argParser.add_argument("--roc-obj-extract", dest="RocObjExtract", action="store", default=ToolchainDefaults.ROC_OBJ_EXTRACT) + argParser.add_argument("--roc-obj-ls", dest="RocObjLs", action="store", default=ToolchainDefaults.ROC_OBJ_LS) argParser.add_argument( "--code-object-version", dest="CodeObjectVersion", @@ -214,6 +216,8 @@ def parseArguments(input: Optional[List[str]] = None) -> Dict[str, Any]: arguments["CxxCompiler"] = args.CxxCompiler arguments["CCompiler"] = args.CCompiler arguments["OffloadBundler"] = args.OffloadBundler + arguments["RocObjExtract"] = args.RocObjExtract + arguments["RocObjLs"] = args.RocObjLs arguments["Assembler"] = args.Assembler arguments["LogicPath"] = args.LogicPath arguments["LogicFilter"] = args.LogicFilter diff --git a/tensilelite/Tensile/TensileCreateLibrary/Run.py b/tensilelite/Tensile/TensileCreateLibrary/Run.py index 308e24e181..73abdc950c 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Run.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Run.py @@ -25,59 +25,53 @@ import rocisa import functools -import glob -import itertools -import os import shutil from pathlib import Path +from os import path, getpid from timeit import default_timer as timer -from typing import List, NamedTuple, Optional, Union - -from Tensile import SOURCE_PATH, LibraryIO -from Tensile.Common import ( - CHeader, - DebugConfig, - DepthUConfig, - ensurePath, - HR, - IsaVersion, - ParallelMap2, - print1, - print2, - printWarning, - printExit, - printWarning, - state, - tqdm, - setVerbosity, - getVerbosity -) +from typing import Dict, List, NamedTuple, Optional + +from Tensile import SOURCE_PATH +from Tensile.Common import DebugConfig, HR, IsaVersion, ParallelMap2, ParallelMapConfig, setVerbosity from Tensile.Common.Architectures import gfxToIsa, isaToGfx, SUPPORTED_GFX from Tensile.Common.Capabilities import makeIsaInfoMap from Tensile.Common.GlobalParameters import assignGlobalParameters, globalParameters -from Tensile.SolutionStructs.Naming import getKernelFileBase, getKeyNoInternalArgs, getSerialNaming, getKernelNameMin - -from Tensile.CustomYamlLoader import load_logic_gfx_arch +from Tensile.SolutionStructs.Naming import getKernelFileBase, getKernelNameMin from Tensile.KernelWriterAssembly import KernelWriterAssembly -from Tensile.KernelWriterBase import ( - KERNEL_HELPER_FILENAME_CPP, - KERNEL_HELPER_FILENAME_H, -) from Tensile.SolutionLibrary import MasterSolutionLibrary -from Tensile.SolutionStructs import Solution +from Tensile.SolutionStructs import kernelObjectNames from Tensile.Toolchain.Assembly import makeAssemblyToolchain, buildAssemblyCodeObjectFiles -from Tensile.Toolchain.Source import makeSourceToolchain, SourceToolchain, buildSourceCodeObjectFiles -from Tensile.Toolchain.Validators import ( - ToolchainDefaults, - validateToolchain, -) +from Tensile.Toolchain.Source import makeSourceToolchain, SourceToolchain, buildSourceCodeObjectFile +from Tensile.Toolchain.Validators import validateToolchain from Tensile.Toolchain.Component import Assembler from Tensile.Utilities.Decorators.Profile import profile from Tensile.Utilities.Decorators.Timing import timing +from .IO import generateSolutionsAndLibraries, genLazyMasterSolutionLibrary, \ + generateParentLibrary, writeAssembly, writeHelper +from .Logic import logicFileList, schedule, getCoFileNames from .ParseArguments import parseArguments +def copyStaticFiles(outputPath=None): + if outputPath is None: + outputPath = globalParameters["WorkingPath"] + libraryStaticFiles = [ + "TensileTypes.h", + "tensile_bfloat16.h", + "tensile_float8_bfloat8.h", + "tensile_float8_bfloat8_bc.h", + "KernelHeader.h", + "ReductionTemplate.h", + "memory_gfx.h" + ] + + for fileName in libraryStaticFiles: + # copy file + shutil.copy( path.join(SOURCE_PATH, fileName), outputPath) + + return libraryStaticFiles + class KernelCodeGenResult(NamedTuple): err: int src: str @@ -91,7 +85,12 @@ class KernelCodeGenResult(NamedTuple): mathclk: int -def processKernelSource(kernelWriterAssembly, data, useShortNames, splitGSU, kernelSerialNaming, kernel) -> KernelCodeGenResult: +def _processKernelSource(kernelWriterAssembly, + data, + useShortNames, + splitGSU, + kernelSerialNaming, + kernel) -> KernelCodeGenResult: """ Generate source for a single kernel. Returns (error, source, header, kernelName). @@ -102,450 +101,142 @@ def processKernelSource(kernelWriterAssembly, data, useShortNames, splitGSU, ker err, src = kernelWriter.getSourceFileString(kernel, useShortNames) header = kernelWriter.getHeaderFileString(kernel) objFilename = kernel._state.get("codeObjectFile", None) - pgr = int(kernel["PrefetchGlobalRead"]) - return KernelCodeGenResult( - err, src, header, asmFilename, objFilename, tuple(kernel["ISA"]), \ - kernel["WavefrontSize"], kernel["CUOccupancy"], \ - pgr, kernel["MathClocksUnrolledLoop"] - ) - - -def removeInvalidSolutionsAndKernels(results, kernels, solutions, errorTolerant, printLevel: bool, splitGSU: bool): - removeKernels = [] - removeKernelNames = [] - removeSolutions = [] - removeResults = [] - - for kernIdx, r in ( - tqdm(enumerate(results)) if printLevel > 1 else enumerate(results) - ): - if r.err != 0: - if not errorTolerant: - print( - "\nKernel generation failed for kernel: {}".format( - kernels[kernIdx]["SolutionIndex"] - ) - ) - print(kernels[kernIdx]["SolutionNameMin"]) - removeKernels.append(kernels[kernIdx]) - kName = getKeyNoInternalArgs(kernels[kernIdx], splitGSU) - if kName not in removeKernelNames: - removeKernelNames.append(kName) - removeResults.append(results[kernIdx]) - - if len(removeKernels) > 0 and not errorTolerant: - printExit("** kernel generation failure **") - - for kern in removeKernels: - kernels.remove(kern) - - for solution in ( - tqdm(solutions, "Finding invalid solutions") - if printLevel > 1 - else solutions - ): - solutionKernels = solution.getKernels() - for kernel in solutionKernels: - kName = getKeyNoInternalArgs(kernel, splitGSU) - if kName in removeKernelNames: - removeSolutions.append(solution) - break - - for solut in removeSolutions: - solutions.remove(solut) - - for rel in removeResults: - results.remove(rel) - -def passPostKernelInfoToSolution(results, kernels, solutions, splitGSU: bool): - resultDict = {} - for kernIdx, r in enumerate(results): - kName = getKeyNoInternalArgs(kernels[kernIdx], splitGSU) - resultDict["%s"%kName] = r - for solution in solutions: - solutionKernels = solution.getKernels() - for kernel in solutionKernels: - kName = getKeyNoInternalArgs(kernel, splitGSU) - result = resultDict["%s"%kName] - solution._state["CUOccupancy"] = result.cuoccupancy - solution._state["PrefetchGlobalRead"] = result.pgr - solution._state["MathClocksUnrolledLoop"] = result.mathclk - -def writeAssembly(asmPath: Union[Path, str], result: KernelCodeGenResult): - if result.err: - printExit(f"Failed to build kernel {result.name} because it has error code {result.err}") - path = Path(asmPath) / f"{result.name}.s" - isa = result.isa - wfsize = result.wavefrontSize - with open(path, "w", encoding="utf-8") as f: - f.write(result.src) - - # result.src is very large so let garbage collector know to clean up - del result - - return path, isa, wfsize - - -def writeHelpers( - outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H -): - kernelSourceFilename = os.path.join(os.path.normcase(outputPath), KERNEL_HELPER_FILENAME_CPP) - kernelHeaderFilename = os.path.join(os.path.normcase(outputPath), KERNEL_HELPER_FILENAME_H) - - with open(kernelHeaderFilename, "w", encoding="utf-8") as kernelHeaderFile, open( - kernelSourceFilename, "w", encoding="utf-8" - ) as kernelSourceFile: - kernelSourceFile.write(CHeader) - kernelHeaderFile.write(CHeader) - kernelSourceFile.write('#include "Kernels.h"\n') - kernelHeaderFile.write("#pragma once\n") - kernelHeaderFile.write("#include \n") - kernelHeaderFile.write("#include \n\n") - kernelHeaderFile.write('#include "KernelHeader.h"\n\n') - HeaderText = "" - for ko in kernelHelperObjs: - kernelName = ko.getKernelName() - (err, src) = ko.getSourceFileString() - kernelSourceFile.write(src) - if err: - print("*** warning: invalid kernel#%u" % kernelName) - HeaderText += ko.getHeaderFileString() - kernelHeaderFile.write(HeaderText) - - -def writeSolutionsAndKernels( - outputPath, - asmToolchain, - srcToolchain, - solutions, - kernels, - kernelHelperObjs, - kernelWriterAssembly, - splitGSU: bool, - cmdlineArchs: List[str], - kernelSerialNaming, - errorTolerant=False, - generateSourcesAndExit=False, - compress=True, - useShortNames=False, -): - codeObjectFiles = [] - - outputPath = Path(outputPath) - destLibPath = ensurePath( - outputPath / "library" - ) # Destination for code object library files (.co) - buildTmpPath = ensurePath(outputPath / "build_tmp" / outputPath.stem.upper()) # - assemblyTmpPath = ensurePath( - buildTmpPath / "assembly" - ) # Temp path for generated assembly files (.s) - objectTmpPath = ensurePath( - buildTmpPath / "code_object_tmp" - ) # Temp path for HSA code object files (.hsaco) - - asmKernels = [k for k in kernels if k["KernelLanguage"] == "Assembly"] - - visited = set() - duplicates = 0 - for k in asmKernels: - base = getKernelFileBase(useShortNames, splitGSU, kernelSerialNaming, k) - k.duplicate = True if base in visited else False - if not k.duplicate: - k["BaseName"] = base - duplicates += k.duplicate - print2(f"Duplicate: {base}") - visited.add(base) - print1(f"Number of duplicate kernels: {duplicates}") - - numAsmKernels = len(asmKernels) - numKernels = len(asmKernels) - assert numKernels == numAsmKernels, "Only assembly kernels are supported in TensileLite" - asmIter = zip( - itertools.repeat(kernelWriterAssembly), - itertools.repeat(rocisa.rocIsa.getInstance().getData()), - itertools.repeat(useShortNames), - itertools.repeat(splitGSU), - itertools.repeat(kernelSerialNaming), - asmKernels - ) - asmResults = ParallelMap2(processKernelSource, asmIter, "Generating assembly kernels", return_as="list") - removeInvalidSolutionsAndKernels( - asmResults, asmKernels, solutions, errorTolerant, getVerbosity(), splitGSU - ) - passPostKernelInfoToSolution( - asmResults, asmKernels, solutions, splitGSU - ) - - def assemble(ret): - p, isa, wavefrontsize = ret - asmToolchain.assembler(isaToGfx(isa), wavefrontsize, str(p), str(p.with_suffix(".o"))) - - unaryWriteAssembly = functools.partial(writeAssembly, assemblyTmpPath) - compose = lambda *F: functools.reduce(lambda f, g: lambda x: f(g(x)), F) - ret = ParallelMap2( - compose(assemble, unaryWriteAssembly), - asmResults, - "Writing assembly kernels", - return_as="list", - multiArg=False, - ) - - writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H) - srcKernelFile = Path(outputPath) / "Kernels.cpp" - - if not generateSourcesAndExit: - codeObjectFiles += buildAssemblyCodeObjectFiles( - asmToolchain.linker, - asmToolchain.bundler, - globalParameters["ROCmLdPath"], - asmKernels, - destLibPath, - assemblyTmpPath, - compress, - ) - buildSourceCodeObjectFiles( - srcToolchain.compiler, - srcToolchain.bundler, - destLibPath, - objectTmpPath, - outputPath, - srcKernelFile, - cmdlineArchs, - ) - - return codeObjectFiles, numKernels - - -def writeSolutionsAndKernelsTCL( - outputPath, - asmToolchain, - srcToolchain, - kernels, - kernelHelperObjs, - kernelWriterAssembly, - cmdlineArchs: List[str], - kernelSerialNaming, - compress=True, - useShortNames=False, -): - outputPath = Path(outputPath) - destLibPath = ensurePath( - outputPath / "library" - ) # Destination for code object library files (.co) - buildTmpPath = ensurePath(outputPath / "build_tmp" / outputPath.stem.upper()) - assemblyTmpPath = ensurePath( - buildTmpPath / "assembly" - ) # Temp path for generated assembly files (.s) - objectTmpPath = ensurePath( - buildTmpPath / "code_object_tmp" - ) # Temp path for HSA code object files (.hsaco) - - asmKernels = [k for k in kernels if k["KernelLanguage"] == "Assembly"] - - visited = set() - duplicates = 0 - splitGSU = False - for k in asmKernels: - base = getKernelFileBase(useShortNames, splitGSU, kernelSerialNaming, k) - k["BaseName"] = base - k.duplicate = True if base in visited else False - duplicates += k.duplicate - print2(f"Duplicate: {base}") - visited.add(base) - print1(f"Number of duplicate kernels: {duplicates}") - - uniqueAsmKernels = [k for k in asmKernels if not k.duplicate] - - def assemble(ret): - p, isa, wavefrontsize = ret - asmToolchain.assembler(isaToGfx(isa), wavefrontsize, str(p), str(p.with_suffix(".o"))) - - unaryProcessKernelSource = functools.partial( - processKernelSource, - kernelWriterAssembly, - rocisa.rocIsa.getInstance().getData(), - useShortNames, - splitGSU, - kernelSerialNaming - ) - - unaryWriteAssembly = functools.partial(writeAssembly, assemblyTmpPath) - compose = lambda *F: functools.reduce(lambda f, g: lambda x: f(g(x)), F) - ret = ParallelMap2( - compose(assemble, unaryWriteAssembly, unaryProcessKernelSource), - uniqueAsmKernels, - "Generating assembly kernels", - multiArg=False, - return_as="list" - ) - buildAssemblyCodeObjectFiles( - asmToolchain.linker, - asmToolchain.bundler, - globalParameters["ROCmLdPath"], - asmKernels, - destLibPath, - assemblyTmpPath, - compress, - ) - - writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H) - srcKernelFile = Path(outputPath) / "Kernels.cpp" - buildSourceCodeObjectFiles( - srcToolchain.compiler, - srcToolchain.bundler, - destLibPath, - objectTmpPath, - outputPath, - srcKernelFile, - cmdlineArchs, - ) - - return len(uniqueAsmKernels) - + return KernelCodeGenResult(err, + src, + header, + asmFilename, + objFilename, + tuple(kernel["ISA"]), + kernel["WavefrontSize"], + kernel["CUOccupancy"], + int(kernel["PrefetchGlobalRead"]), + kernel["MathClocksUnrolledLoop"]) @timing -def copyStaticFiles(outputPath): - libraryStaticFiles = [ - "TensileTypes.h", - "tensile_bfloat16.h", - "tensile_float8_bfloat8.h", - "tensile_float8_bfloat8_bc.h", - "KernelHeader.h", - "ReductionTemplate.h", - "memory_gfx.h", - ] +def buildAssemblyKernels(asmPath: Path, + assembler: Assembler, + kernelWriterAssembly: KernelWriterAssembly, + data, + removeTemporaries: bool, + solnLibs: List[tuple]): + kernels = [s.getKernels()[0] for s in solnLibs[0]] + + visited = set() + duplicates = 0 + for k in kernels: + base = getKernelNameMin(k, False) + k.duplicate = True if base in visited else False + duplicates += k.duplicate + visited.add(base) + uniqueAsmKernels = [k for k in kernels if not k.duplicate] + #uniqueAsmKernels = [k for k in kernels if "BuildKernels" in k] + + pksResults = [_processKernelSource(kernelWriterAssembly, data, False, False, None, k) for k in uniqueAsmKernels] + asmPidPath = asmPath / str(getpid()) + asmPidPath.mkdir(exist_ok=True) + for p, isa, wavefrontsize in set([writeAssembly(asmPidPath, k) for k in pksResults]): + assembler(isaToGfx(isa), wavefrontsize, str(p), str(p.with_suffix(".o"))) # TODO: arguments + if removeTemporaries: + p.unlink() + return uniqueAsmKernels, solnLibs[1] + + +def generateKernelHelperObjects(solutions, isaInfoMap): + khos = [] + visited = set() + for solution in solutions: + build = False + names = kernelObjectNames(solution) + for name in names: + if name not in visited: + visited.add(name) + build = True + if build: + khos.extend(solution.initHelperKernelObjects(list(isaInfoMap.keys()))) + print(f"khos: {len(khos)}") + return list(dict.fromkeys(khos)) + + +#def generateMatchTable(masterSolutionLibraries): +# def libraryIter(lib: MasterSolutionLibrary): +# if len(lib.solutions): +# for i, s in enumerate(lib.solutions.items()): +# yield (i, *s) +# else: +# for _, lazyLib in lib.lazyLibraries.items(): +# yield from libraryIter(lazyLib) +# +# matchTable = {} +# for library in masterSolutionLibraries: +# # Match yaml file solutions to solution index +# for localIdx, _, s in libraryIter(library): +# matchTable[s.index] = [srcFile, localIdx] +# LibraryIO.write("MatchTable", matchTable) + + +def updateParentMasterLibrary( + gfxName: str, + masterLib: MasterSolutionLibrary, + masterLibraries: Dict[str, MasterSolutionLibrary], + nextIdx: Dict[str, int], +) -> None: + if gfxName in masterLibraries: + nextIdx= masterLibraries[gfxName].merge(masterLib, nextIdx) + else: + masterLibraries[gfxName] = masterLib - for fileName in libraryStaticFiles: - shutil.copy(os.path.join(SOURCE_PATH, fileName), outputPath) - return libraryStaticFiles +def updateMasterLibrary( + currMasterLib: MasterSolutionLibrary, + prevMasterLib: MasterSolutionLibrary, + nextIdx: int, +) -> None: + assert prevMasterLib is not None or currMasterLib is not None + if prevMasterLib is not None: + nextIdx = prevMasterLib.merge(currMasterLib, nextIdx) + else: + prevMasterLib = currMasterLib + nextIdx = 0 + return prevMasterLib, nextIdx -@timing -def generateKernelObjectsFromSolutions(solutions): - kernels = [] - kernelHelperObjs = [] - kernelNames = set() - kernelHelperNames = set() - splitGSU = False - for solution in solutions: - solutionKernels = solution.getKernels() - for kernel in solutionKernels: - kName = getKeyNoInternalArgs(kernel, splitGSU) - if kName not in kernelNames: - kernels.append(kernel) - kernelNames.add(kName) - solutionHelperKernels = solution.getHelperKernelObjects() - kernelHelperObjs += solutionHelperKernels - for ko in solutionHelperKernels: - kernelHelperNames.add(ko.getKernelName()) - - # remove duplicates while preserving order - numKhos = len(kernelHelperObjs) - kernelHelperObjs = list(dict.fromkeys(kernelHelperObjs)) - - print1(f"Number of kernel helper objects: {numKhos}") - print1(f"Number of unique kernel helper objects: {len(kernelHelperObjs)}") - - return (kernels, kernelHelperObjs, kernelHelperNames) +def processMsl(libraryPath, libraryFormat, solnLibTup): + solns, libraries = solnLibTup + if len(libraries) > 0: + masterLib = None + nextSolutionIdx = 0 + for _, lib in libraries: + masterLib, nextSolutionIdx = updateMasterLibrary(lib, masterLib, nextSolutionIdx) + genLazyMasterSolutionLibrary(libraryPath, libraryFormat, masterLib) # can we make this async + return solns, libraries -@timing -def generateLogicDataAndSolutions(logicFiles, args, assembler: Assembler, isaInfoMap): - if ";" in args["Architecture"]: - archs = args["Architecture"].split(";") # user arg list format - else: - archs = args["Architecture"].split("_") # workaround for cmake list in list issue - - solutions = [] - masterLibraries = {} - nextSolIndex = 0 - splitGSU = False - printSolutionRejectionReason = False - printIndexAssignmentInfo = False - - fIter = zip( - logicFiles, - itertools.repeat(assembler), - itertools.repeat(splitGSU), - itertools.repeat(printSolutionRejectionReason), - itertools.repeat(printIndexAssignmentInfo), - itertools.repeat(DepthUConfig()), - itertools.repeat(isaInfoMap), - itertools.repeat(args["LazyLibraryLoading"]), - ) - - def libraryIter(lib: MasterSolutionLibrary): - if len(lib.solutions): - for i, s in enumerate(lib.solutions.items()): - yield (i, *s) - else: - for _, lazyLib in lib.lazyLibraries.items(): - yield from libraryIter(lazyLib) - - for library in ParallelMap2( - LibraryIO.parseLibraryLogicFile, fIter, "Loading Logics...", return_as="generator_unordered" - ): - _, architectureName, _, _, _, newLibrary = library - - if architectureName == "": - continue - - if architectureName in masterLibraries: - nextSolIndex = masterLibraries[architectureName].merge(newLibrary, nextSolIndex) - else: - masterLibraries[architectureName] = newLibrary - masterLibraries[architectureName].version = args["CodeObjectVersion"] - - # Sort masterLibraries to make global soln index values deterministic - solnReIndex = 0 - masterLibraries = dict(sorted(masterLibraries.items())) - for _, masterLibrary in masterLibraries.items(): - for _, sol in masterLibrary.solutions.items(): - sol.index = solnReIndex - solnReIndex += 1 - # Sort masterLibrary to make global soln index values deterministic - masterLibrary.lazyLibraries = dict(sorted(masterLibrary.lazyLibraries.items())) - for name, lib in masterLibrary.lazyLibraries.items(): - # Sort solns by the lib logic file they were generated from - lib.solutions = { - k: lib.solutions[k] - for k in sorted(lib.solutions, key=lambda idx: lib.solutions[idx].srcName) - } - for _, sol in lib.solutions.items(): - sol.index = solnReIndex - solnReIndex += 1 - - if args["GenSolTable"]: - matchTable = {} - # Match yaml file solutions to solution index - for _, masterLibrary in masterLibraries.items(): - for _, _, s in libraryIter(masterLibrary): - matchTable[s.index] = [s.srcName, s.libraryLogicIndex] - LibraryIO.write("MatchTable", matchTable) - - if "fallback" in masterLibraries.keys(): - for key, value in masterLibraries.items(): - if key != "fallback": - value.merge(masterLibraries["fallback"]) - masterLibraries.pop("fallback") - for _, masterLibrary in masterLibraries.items(): - for _, sol in masterLibrary.solutions.items(): - solutions.append(sol.originalSolution) - for name, lib in masterLibrary.lazyLibraries.items(): - for _, sol in lib.solutions.items(): - sol.originalSolution._state["codeObjectFile"] = name - solutions.append(sol.originalSolution) - - # remove duplicates while preserving order - numSoln = len(solutions) - solutions = dict.fromkeys(solutions).keys() - - print1(f"Number of solutions parsed: {numSoln}") - print1(f"Number of unique solutions: {len(solutions)}") - - return solutions, masterLibraries +def makeDirectories(outputPath): + outputPath = Path(outputPath) + outputPath.mkdir(parents=True, exist_ok=True) + buildTmp = outputPath / "build_tmp" / str(outputPath.name).upper() + srcCodeObjectPath = buildTmp / "code_object_tmp" + srcCodeObjectPath.mkdir(parents=True, exist_ok=True) + asmPath = buildTmp / "assembly" + asmPath.mkdir(parents=True, exist_ok=True) + kernelsIncludePath = outputPath / "Kernels" + kernelsIncludePath.mkdir(parents=True, exist_ok=True) + libraryPath = outputPath / "library" + libraryPath.mkdir(exist_ok=True) + return outputPath, buildTmp, asmPath, srcCodeObjectPath, libraryPath + + +def extractBuildResults(result): + masterLibs = {} + nextIdx = 0 + flattenedList = [] + for l, library in result: + if len(library) > 0: + for gfxName, lib in library: + updateParentMasterLibrary(gfxName, lib, masterLibs, nextIdx) # I think we can move this out and do this later + flattenedList.extend(l) + return flattenedList, masterLibs ################################################################################ @@ -553,154 +244,152 @@ def libraryIter(lib: MasterSolutionLibrary): ################################################################################ @profile def run(): - start = timer() - print1("") - print1(HR) - print1("# Tensile Create Library") - print2(HR) - print2("") - - arguments = parseArguments() - setVerbosity(arguments["PrintLevel"]) - outputPath = Path(ensurePath(os.path.abspath(arguments["OutputPath"]))) - cxxCompiler, _, offloadBundler, _, _ = validateToolchain( - arguments["CxxCompiler"], - arguments["CCompiler"], - arguments["OffloadBundler"], - arguments["Assembler"], - ToolchainDefaults.HIP_CONFIG, - ) - - if ";" in arguments["Architecture"]: - archs = arguments["Architecture"].split(";") - else: - archs = arguments["Architecture"].split("_") - archs = SUPPORTED_GFX if "all" in archs else archs - - targetIsas = [gfxToIsa(a) for a in archs] - isaInfoMap = makeIsaInfoMap(targetIsas, cxxCompiler) - assignGlobalParameters(arguments, isaInfoMap) - - asmToolchain = makeAssemblyToolchain( - cxxCompiler, - offloadBundler, - arguments["CodeObjectVersion"], - arguments["BuildIdKind"] - ) - srcToolchain = makeSourceToolchain( - cxxCompiler, - offloadBundler, - arguments["AsanBuild"], - arguments["BuildIdKind"], - save_temps=False - ) - - print1(asmToolchain.assembler) - print1(asmToolchain.bundler) - - if not os.path.exists(arguments["LogicPath"]): - printExit(f"LogicPath {arguments['LogicPath']} doesn't exist") - - logicExtFormat = ".yaml" - if arguments["LogicFormat"] == "yaml": - pass - elif arguments["LogicFormat"] == "json": - logicExtFormat = ".json" + start = timer() + print("") + print(HR) + print("# Tensile Create Library") + + arguments = parseArguments() + arguments["OutputPath"] = Path(arguments["OutputPath"]).resolve() + outputPath, buildTmp, assemblyPath, srcCodeObjectPath, libraryPath = makeDirectories(arguments["OutputPath"]) + + setVerbosity(arguments["PrintLevel"]) + cxxCompiler, offloadBundler, ls, extract = validateToolchain( + arguments["CxxCompiler"], + arguments["OffloadBundler"], + arguments["RocObjLs"], + arguments["RocObjExtract"] + ) + + if ";" in arguments["Architecture"]: + archs = arguments["Architecture"].split(";") + else: + archs = arguments["Architecture"].split("_") + archs = SUPPORTED_GFX if "all" in archs else archs + + targetIsas = [gfxToIsa(a) for a in archs] + isaInfoMap = makeIsaInfoMap(targetIsas, cxxCompiler) + assignGlobalParameters(arguments, isaInfoMap) + + asmToolchain = makeAssemblyToolchain( + cxxCompiler, + offloadBundler, + arguments["CodeObjectVersion"], + arguments["BuildIdKind"] + ) + srcToolchain = makeSourceToolchain( + cxxCompiler, + offloadBundler, + ls, + extract, + arguments["CpuThreads"], + arguments["AsanBuild"], + arguments["BuildIdKind"], + save_temps=False + ) + + print(asmToolchain.assembler) + print(asmToolchain.bundler) + + # List logic files + unsortedLogic = logicFileList(archs, + Path(arguments["LogicPath"]), + arguments["LogicFilter"], + arguments["Experimental"]) + cofiles = ParallelMap2(getCoFileNames, + ParallelMapConfig(message="Scheduling work. ", + procs=arguments["CpuThreads"]), + unsortedLogic) + logicFiles = schedule(cofiles, 2*arguments["CpuThreads"], arguments["CpuThreads"]) + + # Phase1: Build assembly and master solution libraries + writerAsm = KernelWriterAssembly(None, asmToolchain.assembler, DebugConfig()) + unaryGenSolutions = functools.partial(generateSolutionsAndLibraries, + asmToolchain.assembler, + isaInfoMap, + arguments["LazyLibraryLoading"]) + + #result = ParallelMap2(unaryGenSolutions, ParallelMapConfig(message="blah"), logicFiles) + #d = {} + #for ss, _ in result: + # for s in ss: + # name = getKernelNameMin(s, False) + # if name in d: + # d[name].add(s["codeObjectFile"]) + # else: + # d[name] = set() + # d[name].add(s["codeObjectFile"]) + #import pprint + #pprint.pprint(d, indent=2) + #count = 0 + #for k, v in d.items(): + # if len(v) > 1: + # print(k,v) + # count += 1 + #print(count) + #assert False + + unaryProcessMsl = functools.partial(processMsl, libraryPath, arguments["LibraryFormat"]) + unaryBuildAsmKernels = functools.partial(buildAssemblyKernels, + assemblyPath, + asmToolchain.assembler, + writerAsm, + rocisa.rocIsa.getInstance().getData(), + not arguments["KeepBuildTmp"]) + unaryBuildCOFile = functools.partial(buildAssemblyCodeObjectFiles, + asmToolchain.linker, + asmToolchain.bundler, + globalParameters["ROCmLdPath"], + libraryPath, + assemblyPath, + arguments["UseCompression"]) + def buildCoAndHelpers(input): + uniqueAsmKernels, libraries = input + unaryBuildCOFile(uniqueAsmKernels) + return uniqueAsmKernels, libraries + + compose = lambda *F: functools.reduce(lambda f, g: lambda x: f(g(x)), F) + if arguments["LazyLibraryLoading"]: + assembly = compose(buildCoAndHelpers, unaryBuildAsmKernels, unaryProcessMsl, unaryGenSolutions) + parMap = functools.partial(ParallelMap2, assembly, ParallelMapConfig(message="Building Lazy Libraries")) + phase1 = compose(extractBuildResults, parMap) + else: + assembly = compose(unaryBuildAsmKernels, unaryGenSolutions) + parMap = functools.partial(ParallelMap2, assembly, ParallelMapConfig(message="Building Objects")) + phase1 = compose(buildCoAndHelpers, extractBuildResults, parMap) + + kernels, masterLibs = phase1(logicFiles) + + # Phase2: Build Master Solution Library + generateParentLibrary(arguments["LibraryFormat"], libraryPath, masterLibs, arguments["LazyLibraryLoading"]) + del masterLibs + + # Phase3: Write Kernel helpers and build Kernels.co + copyStaticFiles(outputPath) + unaryWriteHelpers = functools.partial(writeHelper, outputPath) + srcFiles = ParallelMap2(unaryWriteHelpers, + ParallelMapConfig(message="Generating Kernels code", return_as="list"), + generateKernelHelperObjects(kernels, isaInfoMap)) + kernelsLib = str(srcCodeObjectPath / "Kernels.so") + srcToolchain.compiler(srcFiles, kernelsLib, str(outputPath), archs) + #buildSourceCodeObjectFile(srcToolchain, libraryPath, kernelsLib) + + if not arguments["KeepBuildTmp"]: + if buildTmp.exists() and buildTmp.is_dir(): + shutil.rmtree(buildTmp) else: - printExit(f"Unrecognized LogicFormat: {arguments['LogicFormat']}") - - def archMatch(arch: str, archs: List[str]): - return (arch in archs) or any(a.startswith(arch) for a in archs) - - def validLogicFile(p: Path): - return p.suffix == logicExtFormat and ( - "all" in archs or archMatch(load_logic_gfx_arch(p), archs) - ) - - globPattern = os.path.join( - arguments["LogicPath"], f"**/{arguments['LogicFilter']}{logicExtFormat}" - ) - print1(f"# LogicFilter: {globPattern}") - logicFiles = ( - os.path.join(arguments["LogicPath"], file) - for file in glob.iglob(globPattern, recursive=True) - ) - logicFiles = [file for file in logicFiles if validLogicFile(Path(file))] - - print1(f"# Experimental: {arguments['Experimental']}") - if not arguments["Experimental"]: - logicFiles = [ - file for file in logicFiles if "experimental" not in map(str.lower, Path(file).parts) - ] - - print2(f"# LibraryLogicFiles: {len(logicFiles)}") - - for logicFile in logicFiles: - print2("# %s" % logicFile) - - solutions, masterLibraries = generateLogicDataAndSolutions( - logicFiles, arguments, asmToolchain.assembler, isaInfoMap - ) - - kernels, kernelHelperObjs, _ = generateKernelObjectsFromSolutions(solutions) - kernelSerialNaming = getSerialNaming(kernels) - kernelWriterAssembly = KernelWriterAssembly( - kernelSerialNaming, - asmToolchain.assembler, - DebugConfig(), - ) - - copyStaticFiles(outputPath) - - numKernels = writeSolutionsAndKernelsTCL( - outputPath, - asmToolchain, - srcToolchain, - kernels, - kernelHelperObjs, - kernelWriterAssembly, - archs, - kernelSerialNaming, - useShortNames=arguments["ShortNames"], - compress=arguments["UseCompression"], - ) - - archs = [ # is this really different than the other archs above? - isaToGfx(arch) - for arch in targetIsas - if isaInfoMap[arch].asmCaps["SupportedISA"] - ] - newLibraryDir = ensurePath(os.path.join(outputPath, "library")) - splitGSU = False - for archName, newMasterLibrary in masterLibraries.items(): - if archName in archs: - if arguments["LazyLibraryLoading"]: - masterFile = os.path.join(newLibraryDir, "TensileLibrary_lazy_" + archName) - else: - masterFile = os.path.join(newLibraryDir, "TensileLibrary_" + archName) - newMasterLibrary.applyNaming(splitGSU) - LibraryIO.write(masterFile, state(newMasterLibrary), arguments["LibraryFormat"]) - for name, lib in newMasterLibrary.lazyLibraries.items(): - filename = os.path.join(newLibraryDir, name) - lib.applyNaming(splitGSU) - LibraryIO.write(filename, state(lib), arguments["LibraryFormat"]) - - if not arguments["KeepBuildTmp"]: - buildTmp = Path(arguments["OutputPath"]).parent / "library" / "build_tmp" - if buildTmp.exists() and buildTmp.is_dir(): - shutil.rmtree(buildTmp) - buildTmp = Path(arguments["OutputPath"]) / "build_tmp" - if buildTmp.exists() and buildTmp.is_dir(): - shutil.rmtree(buildTmp) - else: - printWarning(f"Cannot remove build_tmp") - - print1("# Tensile Library Writer DONE") - print1(HR) - print1("") - - stop = timer() - - print1(f"Total time (s): {(stop-start):3.2f}") - print1(f"Total kernels processed: {numKernels}") - print1(f"Kernels processed per second: {(numKernels/(stop-start)):3.2f}") + print(f"Warning: Cannot remove build_tmp") + + print("# Tensile Library Writer DONE") + print(HR) + print("") + + stop = timer() + print(f"Total time (s): {(stop-start):3.2f}") + numKernels = len(kernels) + print(f"Total kernels: {numKernels}") + #print(f"Total kernels processed: {numUniqueKernels}") + #print(f"Duplicate kernels removed: {numDuplicateKernels}") + print(f"Kernels processed per second: {(numKernels/(stop-start)):3.2f}") + #print(f"Total solutions processed: {numSoln}") + #print(f"Duplicate solutions: {numDuplicateSoln}") diff --git a/tensilelite/Tensile/TensileCreateLibrary/Tuning.py b/tensilelite/Tensile/TensileCreateLibrary/Tuning.py new file mode 100644 index 0000000000..938f97f980 --- /dev/null +++ b/tensilelite/Tensile/TensileCreateLibrary/Tuning.py @@ -0,0 +1,227 @@ +################################################################################ +# +# Copyright (C) 2025 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +# SOFTWARE. +# +################################################################################ +import rocisa + +from Tensile.Common import CHeader, print1, print2, printExit, \ + ParallelMap2, ensurePath, tqdm, ParallelMapConfig, getVerbosity +from Tensile.Common.GlobalParameters import globalParameters +from Tensile.KernelWriterBase import KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H +from Tensile.SolutionStructs import Solution, getKernelFileBase, getKeyNoInternalArgs +from Tensile.Toolchain.Assembly import buildAssemblyCodeObjectFiles +from Tensile.Toolchain.Source import buildSourceCodeObjectFile +from .IO import writeAssembly +from .Run import _processKernelSource + +from functools import partial, reduce +from itertools import repeat +from pathlib import Path +from typing import List + +def _writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H): + kernelSourceFilename = str(Path(outputPath) / KERNEL_HELPER_FILENAME_CPP) + kernelHeaderFilename = str(Path(outputPath) / KERNEL_HELPER_FILENAME_H) + + with open(kernelHeaderFilename, "w", encoding="utf-8") as kernelHeaderFile, \ + open(kernelSourceFilename, "w", encoding="utf-8") as kernelSourceFile: + kernelSourceFile.write(CHeader) + kernelHeaderFile.write(CHeader) + kernelSourceFile.write("#include \"Kernels.h\"\n") + kernelHeaderFile.write("#pragma once\n") + kernelHeaderFile.write("#include \n") + kernelHeaderFile.write("#include \n\n") + kernelHeaderFile.write("#include \"KernelHeader.h\"\n\n") + HeaderText = "" + for ko in kernelHelperObjs: + kernelName = ko.getKernelName() + (err, src) = ko.getSourceFileString() + kernelSourceFile.write(src) + if err: + print("*** warning: invalid kernel#%u" % kernelName) + HeaderText += ko.getHeaderFileString() + kernelHeaderFile.write(HeaderText) + + +def removeInvalidSolutionsAndKernels(results, kernels, solutions, errorTolerant, printLevel: bool, splitGSU: bool): + removeKernels = [] + removeKernelNames = [] + removeSolutions = [] + removeResults = [] + + for kernIdx, r in ( + tqdm(enumerate(results)) if printLevel > 1 else enumerate(results) + ): + if r.err != 0: + if not errorTolerant: + print( + "\nKernel generation failed for kernel: {}".format( + kernels[kernIdx]["SolutionIndex"] + ) + ) + print(kernels[kernIdx]["SolutionNameMin"]) + removeKernels.append(kernels[kernIdx]) + kName = getKeyNoInternalArgs(kernels[kernIdx], splitGSU) + if kName not in removeKernelNames: + removeKernelNames.append(kName) + removeResults.append(results[kernIdx]) + + if len(removeKernels) > 0 and not errorTolerant: + printExit("** kernel generation failure **") + + for kern in removeKernels: + kernels.remove(kern) + + for solution in ( + tqdm(solutions, "Finding invalid solutions") + if printLevel > 1 + else solutions + ): + solutionKernels = solution.getKernels() + for kernel in solutionKernels: + kName = getKeyNoInternalArgs(kernel, splitGSU) + if kName in removeKernelNames: + removeSolutions.append(solution) + break + + for solut in removeSolutions: + solutions.remove(solut) + + for rel in removeResults: + results.remove(rel) + + +def passPostKernelInfoToSolution(results, kernels, solutions, splitGSU: bool): + resultDict = {} + for kernIdx, r in enumerate(results): + kName = getKeyNoInternalArgs(kernels[kernIdx], splitGSU) + resultDict["%s"%kName] = r + for solution in solutions: + solutionKernels = solution.getKernels() + for kernel in solutionKernels: + kName = getKeyNoInternalArgs(kernel, splitGSU) + result = resultDict["%s"%kName] + solution._state["CUOccupancy"] = result.cuoccupancy + solution._state["PrefetchGlobalRead"] = result.pgr + solution._state["MathClocksUnrolledLoop"] = result.mathclk + + +def writeSolutionsAndKernels( + outputPath, + asmToolchain, + srcToolchain, + solutions, + kernels, + kernelHelperObjs, + kernelWriterAssembly, + splitGSU: bool, + cmdlineArchs: List[str], + kernelSerialNaming, + errorTolerant=False, + generateSourcesAndExit=False, + compress=True, + useShortNames=False, +): + codeObjectFiles = [] + + outputPath = Path(outputPath) + destLibPath = ensurePath( + outputPath / "library" + ) # Destination for code object library files (.co) + buildTmpPath = ensurePath(outputPath / "build_tmp" / outputPath.stem.upper()) # + assemblyTmpPath = ensurePath( + buildTmpPath / "assembly" + ) # Temp path for generated assembly files (.s) + objectTmpPath = ensurePath( + buildTmpPath / "code_object_tmp" + ) # Temp path for HSA code object files (.hsaco) + + asmKernels = [k for k in kernels if k["KernelLanguage"] == "Assembly"] + + visited = set() + duplicates = 0 + for k in asmKernels: + base = getKernelFileBase(useShortNames, splitGSU, kernelSerialNaming, k) + print1(base) + k.duplicate = True if base in visited else False + if not k.duplicate: + k["BaseName"] = base + duplicates += k.duplicate + print2(f"Duplicate: {base}") + visited.add(base) + print1(f"Number of duplicate kernels: {duplicates}") + + numAsmKernels = len(asmKernels) + numKernels = len(asmKernels) + assert numKernels == numAsmKernels, "Only assembly kernels are supported in TensileLite" + asmIter = zip( + repeat(kernelWriterAssembly), + repeat(rocisa.rocIsa.getInstance().getData()), + repeat(useShortNames), + repeat(splitGSU), + repeat(kernelSerialNaming), + asmKernels + ) + config = ParallelMapConfig(message="Generating assembly kernels", return_as="list", multiArg=True) + asmResults = ParallelMap2(_processKernelSource, config, asmIter) + removeInvalidSolutionsAndKernels( + asmResults, asmKernels, solutions, errorTolerant, getVerbosity(), splitGSU + ) + passPostKernelInfoToSolution( + asmResults, asmKernels, solutions, splitGSU + ) + print1(f"After removal: {len(asmKernels)}") + def assemble(ret): + p, isa, wavefrontsize = ret + asmToolchain.assembler(rocisa.isaToGfx(isa), wavefrontsize, str(p), str(p.with_suffix(".o"))) + + unaryWriteAssembly = partial(writeAssembly, assemblyTmpPath) + compose = lambda *F: reduce(lambda f, g: lambda x: f(g(x)), F) + config = ParallelMapConfig(message="Writing assembly kernels", return_as="list", multiArg=False) + ret = ParallelMap2( + compose(assemble, unaryWriteAssembly), + config, + asmResults, + ) + + _writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H) + + if not generateSourcesAndExit: + codeObjectFiles += buildAssemblyCodeObjectFiles( + asmToolchain.linker, + asmToolchain.bundler, + globalParameters["ROCmLdPath"], + destLibPath, + assemblyTmpPath, + compress, + asmKernels, + ) + kernelsLib = str(objectTmpPath / "Kernels.so") + kernelsSrc = [str(outputPath / "Kernels.cpp")] + srcToolchain.compiler(kernelsSrc, str(kernelsLib), str(outputPath), cmdlineArchs) + buildSourceCodeObjectFile( + srcToolchain, + destLibPath, + kernelsLib, + ) + + return codeObjectFiles, numKernels diff --git a/tensilelite/Tensile/TensileCreateLibrary/__init__.py b/tensilelite/Tensile/TensileCreateLibrary/__init__.py index eb1674916c..5d830030d7 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/__init__.py +++ b/tensilelite/Tensile/TensileCreateLibrary/__init__.py @@ -1 +1,3 @@ -from .Run import copyStaticFiles, run, writeSolutionsAndKernels +from .Run import run +from .Run import copyStaticFiles +from .Tuning import writeSolutionsAndKernels diff --git a/tensilelite/Tensile/Toolchain/Assembly.py b/tensilelite/Tensile/Toolchain/Assembly.py index 96d8ee1ee0..d7b16ac0f2 100644 --- a/tensilelite/Tensile/Toolchain/Assembly.py +++ b/tensilelite/Tensile/Toolchain/Assembly.py @@ -26,13 +26,16 @@ import math import shutil import subprocess +from functools import partial +from os import getpid from pathlib import Path from typing import List, Union, NamedTuple from Tensile.Common import print2 from Tensile.Common.Architectures import isaToGfx -from ..SolutionStructs import Solution +from Tensile.SolutionStructs import Solution, getKernelFileBase +from Tensile.Utilities.RequiredParameters import getRequiredParametersMin from .Component import Assembler, Linker, Bundler @@ -77,10 +80,10 @@ def buildAssemblyCodeObjectFiles( linker: Linker, bundler: Bundler, ldPath: str, - kernels: List[Solution], destDir: Union[Path, str], asmDir: Union[Path, str], - compress: bool=True, + compress: bool, + kernels: List[Solution], ): """Builds code object files from assembly files @@ -107,15 +110,17 @@ def buildAssemblyCodeObjectFiles( continue gfx = isaToGfx(arch) - - objectFiles = [str(asmDir / (k["BaseName"] + extObj)) for k in archKernels if 'codeObjectFile' not in k] + pid = str(getpid()) + baseName = partial(getKernelFileBase, False, False, None) + objectFiles = [str(asmDir / pid / (baseName(k) + extObj)) for k in archKernels if 'codeObjectFile' not in k or k['codeObjectFile'] == "TensileLibrary"] coFileMap = collections.defaultdict(set) if len(objectFiles): coFileMap[asmDir / ("TensileLibrary_"+ gfx + extCoRaw)] = objectFiles - for kernel in archKernels: - coName = kernel.get("codeObjectFile", None) - if coName: - coFileMap[asmDir / (coName + extCoRaw)].add(str(asmDir / (kernel["BaseName"] + extObj))) + else: + for kernel in archKernels: + coName = kernel.get("codeObjectFile", None) + if coName: + coFileMap[asmDir / (coName + extCoRaw)].add(str(asmDir / pid / (baseName(kernel) + extObj))) for coFileRaw, objFiles in coFileMap.items(): objFiles = _batchObjectFiles(ldPath, objFiles, coFileRaw) diff --git a/tensilelite/Tensile/Toolchain/Component.py b/tensilelite/Tensile/Toolchain/Component.py index 5faf4c1ec3..04a2d5ee8a 100644 --- a/tensilelite/Tensile/Toolchain/Component.py +++ b/tensilelite/Tensile/Toolchain/Component.py @@ -6,10 +6,10 @@ from subprocess import check_output, STDOUT, CalledProcessError, PIPE, run from typing import List -from Tensile.Common import SemanticVersion, print1 +from Tensile.Common import SemanticVersion, printExit from .Validators import ToolchainDefaults, validateToolchain -def _invoke(args: List[str], desc: str=""): +def _invoke(args: List[str], desc: str="", working_dir=None): """Invokes a command with the provided arguments in a subprocess. Args: args: A list of arguments to pass to the subprocess. @@ -21,10 +21,10 @@ def _invoke(args: List[str], desc: str=""): """ #print1(f"{desc}: {' '.join(args)}") try: - out = check_output(args, stderr=STDOUT) - except CalledProcessError as err: + out = check_output(args, stderr=STDOUT, cwd=working_dir) + except: raise RuntimeError( - f"Error with {desc}: {err.output}\n" + #f"Error with {desc}: {err.output}\n" f"Failed command: {' '.join(args)}" ) return out @@ -50,6 +50,7 @@ def _getVersion(executable: str, versionFlag: str, regex: str) -> str: if match: result = match.group(1) return SemanticVersion(*[int(c.split("-")[0]) for c in result.split(".")[:3]]) + return None raise Exception(f"No version from {output} matches regex {regex}") except Exception as e: raise RuntimeError(f"Failed to get version when calling {args}: {e}") @@ -63,8 +64,10 @@ def get_rocm_version() -> str: Return: ROCm version string """ - return _getVersion(ToolchainDefaults.HIP_CONFIG, "--version", r'(.+)') - + try: + return _getVersion(ToolchainDefaults.HIP_CONFIG, "--version", r'(.+)') + except: + return None class Component: """A class used to represent a ROCm toolchain component such as clang++""" @@ -180,18 +183,22 @@ class Compiler(Component): Invokes the compiler on the provided arguments """ - def __init__(self, compiler_path: Path, build_id_kind: str, asan_build: bool=False, save_temps: bool=False): + def __init__(self, compiler_path: Path, build_id_kind: str, cpu_threads: int, asan_build: bool=False, save_temps: bool=False): """Constructs and instance of a Compiler.""" super(Compiler, self).__init__(compiler_path) self.default_args = [ *split(environ.get("Tensile_CXX_COMPILER_LAUNCHER", "")), compiler_path, + "-shared", + "-fPIC", + "-fgpu-rdc", + "-Xoffload-linker", + f"--build-id={build_id_kind}", "-D__HIP_HCC_COMPAT_MODE__=1", - "--offload-device-only", "-x", "hip", "-O3", - "-Xoffload-linker", f"--build-id={build_id_kind}", "-std=c++17", + f"-parallel-jobs={cpu_threads}" ] if asan_build: @@ -201,8 +208,7 @@ def __init__(self, compiler_path: Path, build_id_kind: str, asan_build: bool=Fal if os_name == "nt": # should we use fPIIC on all arches? self.default_args.extend(["-fms-extensions", "-fms-compatibility", "-fPIC", "-Wno-deprecated-declarations"]) - - def __call__(self, include_path: str, target_list: List[str], srcPath: str, destPath: str): + def __call__(self, src_paths: List[str], dest_path: str, include_path: str, gfxs: List[str]): """Compiles a source file into an object file. Args: @@ -213,10 +219,12 @@ def __call__(self, include_path: str, target_list: List[str], srcPath: str, dest Raises: RuntimeError: If the compilation command fails. """ - archFlags = [f"--offload-arch={gfx}" for gfx in target_list] - args = [ - *(self.default_args), "-I", include_path, *archFlags, srcPath, "-c", "-o", destPath - ] + + args = list(self.default_args) + args += [f"--offload-arch={gfx}" for gfx in gfxs] + args += ["-I", include_path] + args += src_paths + args +=["-o", dest_path] return _invoke(args, f"Compiling HIP source kernels into objects (.cpp -> .o)") @@ -343,36 +351,72 @@ def __call__(self, srcPaths: List[str], destPath: str): return _invoke(args, "Linking assembly object files into code object (*.o -> .co)") -# class DeviceEnumerator(Component): -# """ -# ROCm amdgpu-arch or rocm_agent_enumerator class used to inspect system for current architecture. -# ... +class RocObjLs(Component): + """ + ROCm roc-obj-ls class used to list code objects. + + ... + + Attributes + ---------- + version : str + the version of the component + rocm_version : str + the ROCm version + Methods + ------- + __call__(self, srcPaths: List[str], destPath: str): + Invokes roc-obj-ls on the provided arguments + """ -# Attributes -# ---------- -# version : str -# the version of the component -# rocm_version : str -# the ROCm version -# path : str -# path to component + def __init__(self, ls_path: Path): + """Constructs and instance of roc-obj-ls.""" + super(RocObjLs, self).__init__(ls_path) -# Methods -# ------- -# __call__(self) -# Invokes component. -# """ -# def __init__(self, component_path: Path): -# """Constructs instance of SystemInterogator. + def __call__(self, sharedObjFile): + """Lists the code objects in shared object. -# Args: -# component_path: The path to amdgpu-arch or rocm_agent_enumeraor. -# """ -# super(Assembler, self).__init__(component_path) -# self._default_args = [str(component_path)] + Args: + sharedObjFile: Name of object file to list. -# def __call__(self): -# """Run component without args to detect system settings.""" + Returns: + List of objects embedded in shared object file. + """ + args = [self._component_path, sharedObjFile] + return [e.strip().split()[1:] for e in _invoke(args, f"Listing code objects in shared object", Path(sharedObjFile).parent).decode().split("\n") if e.strip().split()[1:]] -# return _invoke(self._default_args, "running system inspection") + +class RocObjExtract(Component): + """ + ROCm roc-obj-extract class used to get code objects bundle from shared library. + + ... + + Attributes + ---------- + version : str + the version of the component + rocm_version : str + the ROCm version + Methods + ------- + __call__(self, srcPaths: List[str], destPath: str): + Invokes roc-obj-extract on the provided arguments + """ + + def __init__(self, extract_path: Path): + """Constructs and instance of roc-obj-extract.""" + super(RocObjExtract, self).__init__(extract_path) + + def __call__(self, filename: str): + """Extracts code objects from a shared object. + + Args: + objFile: Name pf object file to extract. + + Returns: + Code object file. + """ + args = [self._component_path, filename] + return _invoke(args, f"Extracting code object.", Path(filename.replace("file:","")).parent) \ No newline at end of file diff --git a/tensilelite/Tensile/Toolchain/Source.py b/tensilelite/Tensile/Toolchain/Source.py index 3cd54c5b35..164ec0e0f0 100644 --- a/tensilelite/Tensile/Toolchain/Source.py +++ b/tensilelite/Tensile/Toolchain/Source.py @@ -32,17 +32,21 @@ from ..Common import print1, ensurePath -from .Component import Compiler, Bundler +from .Component import Compiler, Bundler, RocObjLs, RocObjExtract class SourceToolchain(NamedTuple): compiler: Compiler bundler: Bundler + rocObjLs: RocObjLs + rocObjExtract: RocObjExtract -def makeSourceToolchain(compiler_path, bundler_path, asan_build=False, build_id_kind="sha1", save_temps=False): - compiler = Compiler(compiler_path, build_id_kind, asan_build, save_temps) +def makeSourceToolchain(compiler_path, bundler_path, ls_path, extract_path, cpu_threads, asan_build=False, build_id_kind="sha1", save_temps=False): + compiler = Compiler(compiler_path, build_id_kind, cpu_threads, asan_build, save_temps) bundler = Bundler(bundler_path) - return SourceToolchain(compiler, bundler) + ls = RocObjLs(ls_path) + extract = RocObjExtract(extract_path) + return SourceToolchain(compiler, bundler, ls, extract) def _computeSourceCodeObjectFilename(target: str, base: str, buildPath: Union[Path, str], arch: str) -> Union[Path, None]: @@ -103,16 +107,16 @@ def buildSourceCodeObjectFiles( coPaths= [] objPath = str(tmpObjDir / objFilename) - compiler(str(includeDir), cmdlineArchs, str(kernelPath), objPath) - + compiler([str(kernelPath)], objPath, str(includeDir), cmdlineArchs) for target in bundler.targets(objPath): match = re.search("gfx.*$", target) if match: arch = re.sub(":", "-", match.group()) coPathRaw = _computeSourceCodeObjectFilename(target, kernelPath.stem, tmpObjDir, arch) if not coPathRaw: continue - bundler(target, objPath, str(coPathRaw)) - + buildSourceCodeObjectFile() + #bundler(target, objPath, str(coPathRaw)) + assert False coPath = str(destDir / coPathRaw.stem) coPathsRaw.append(coPathRaw) coPaths.append(coPath) @@ -124,3 +128,30 @@ def buildSourceCodeObjectFiles( print1(f"buildSourceCodeObjectFile time (s): {(stop-start):3.2f}") return coPaths + + +def buildSourceCodeObjectFile(toolchain: SourceToolchain, + destPath: Union[Path, str], + sharedObjPath: Union[Path, str], ) -> List[str]: + """Compiles a HIP source code file into a code object file. + + Args: + toolchain: The source toolchain. + destDir: The destination directory where HSA code object files are placed. + tmpObjDir: The directory where HIP source object files are created. + includeDir: The include directory path. + kernelPath: The path to the kernel source file. + + Returns: + List of paths to the created code objects. + """ + + for target, filename in toolchain.rocObjLs(sharedObjPath): + match = re.search("gfx.*$", target) + if match: + print(f"Generating Kernels.co for {target.split('-')[-1]}") + arch = re.sub(":", "-", match.group()) + toolchain.rocObjExtract(filename) + src = str(Path(sharedObjPath).parent / (str(Path(filename).name).replace("#","-").replace("=","").replace("&","-") + ".co")) + dst = str(destPath / f"Kernels.so-000-{arch}.hsaco") + shutil.move(src, dst) \ No newline at end of file diff --git a/tensilelite/Tensile/Toolchain/Validators.py b/tensilelite/Tensile/Toolchain/Validators.py index 316a0fc29b..16bd16026f 100644 --- a/tensilelite/Tensile/Toolchain/Validators.py +++ b/tensilelite/Tensile/Toolchain/Validators.py @@ -108,9 +108,11 @@ def _posixSearchPaths() -> List[Path]: class ToolchainDefaults(NamedTuple): - CXX_COMPILER = osSelect(linux="amdclang++", windows="clang++.exe") - C_COMPILER = osSelect(linux="amdclang", windows="clang.exe") - OFFLOAD_BUNDLER = osSelect(linux="clang-offload-bundler", windows="clang-offload-bundler.exe") + CXX_COMPILER= osSelect(linux="amdclang++", windows="clang++.exe") + C_COMPILER= osSelect(linux="amdclang", windows="clang.exe") + OFFLOAD_BUNDLER= osSelect(linux="clang-offload-bundler", windows="clang-offload-bundler.exe") + ROC_OBJ_EXTRACT= osSelect(linux="roc-obj-extract", windows="roc-obj-extract.exe") + ROC_OBJ_LS= osSelect(linux="roc-obj-ls", windows="roc-obj-ls.exe") DEVICE_ENUMERATOR = osSelect(linux="rocm_agent_enumerator" if isRhel8() else "amdgpu-arch", windows="hipinfo") ASSEMBLER = osSelect(linux="amdclang++", windows="clang++.exe") HIP_CONFIG = osSelect(linux="hipconfig", windows="hipconfig") @@ -123,6 +125,18 @@ def _supportedComponent(component: str, targets: List[str]) -> bool: return isSupported +def supportedBundler(bundler: str) -> bool: + """Determine if an offload bundler is supported by Tensile. + + Args: + bundler: The name of an roc-obj-extract to test for support. + + Return: + If supported True; otherwise, False. + """ + return _supportedComponent(bundler, ["clang-offload-bundler"]) + + def supportedCCompiler(compiler: str) -> bool: """ Determine if a C compiler/assembler is supported by Tensile. @@ -149,12 +163,36 @@ def supportedCxxCompiler(compiler: str) -> bool: return _supportedComponent(compiler, ["amdclang++", "clang++"]) +def supportedExtract(extract: str) -> bool: + """Determine if an object extracter is supported by Tensile. + + Args: + extract: The name of an roc-obj-extract to test for support. + + Return: + If supported True; otherwise, False. + """ + return _supportedComponent(extract, [ToolchainDefaults.ROC_OBJ_EXTRACT]) + + +def supportedLs(ls: str) -> bool: + """Determine if an object lister is supported by Tensile. + + Args: + ls: The name of an roc-obj-ls to test for support. + + Return: + If supported True; otherwise, False. + """ + return _supportedComponent(ls, [ToolchainDefaults.ROC_OBJ_LS]) + + def supportedOffloadBundler(bundler: str) -> bool: """ Determine if an offload bundler is supported by Tensile. Args: - bundler: The name of an offload bundler to test for support. + bundler: The name of an roc-obj-extract to test for support. Return: If supported True; otherwise, False. @@ -215,11 +253,13 @@ def _validateExecutable(file: str, searchPaths: List[Path]) -> str: The validated executable with an absolute path. """ if not any(( - supportedCxxCompiler(file), - supportedCCompiler(file), - supportedOffloadBundler(file), - supportedHip(file), - supportedDeviceEnumerator(file) + supportedCxxCompiler(file), + supportedCCompiler(file), + supportedBundler(file), + supportedDeviceEnumerator(file), + supportedExtract(file), + supportedLs(file), + supportedHip(file) )): raise ValueError(f"`{file}` is not a supported toolchain component on {'Windows' if os.name == 'nt' else 'Linux'}") diff --git a/tensilelite/Tensile/Utilities/RequiredParameters.py b/tensilelite/Tensile/Utilities/RequiredParameters.py new file mode 100644 index 0000000000..f411bd7e72 --- /dev/null +++ b/tensilelite/Tensile/Utilities/RequiredParameters.py @@ -0,0 +1,118 @@ +_requiredParametersFull = { + "InnerUnroll": True, + "KernelLanguage": True, + "LdsPadA": True, + "LdsPadB": True, + "LdsPadMetadata": True, + "LdsBlockSizePerPadA": True, + "LdsBlockSizePerPadB": True, + "LdsBlockSizePerPadMetadata": True, + "TransposeLDS": True, + "MaxOccupancy": True, + "VectorWidthA": True, + "VectorWidthB": True, + "VectorStore": True, + "StoreVectorWidth": True, + "GlobalReadVectorWidthA": True, + "GlobalReadVectorWidthB": True, + "LocalReadVectorWidth": True, + "WaveSeparateGlobalReadA": True, + "WaveSeparateGlobalReadB": True, + "UnrollLoopSwapGlobalReadOrder": True, + "PrefetchGlobalRead": True, + "PrefetchLocalRead": True, + "ClusterLocalRead": True, + "SuppressNoLoadLoop": True, + "ExpandPointerSwap": True, + "ScheduleGlobalRead": True, + "ScheduleLocalWrite": True, + "ScheduleIterAlg": True, + "GlobalReadPerMfma": True, + "LocalWritePerMfma": True, + "InterleaveAlpha": True, + "OptNoLoadLoop": True, + "BufferLoad": True, + "BufferStore": True, + "DirectToVgprA": True, + "DirectToVgprB": True, + "DirectToVgprSparseMetadata": True, + "DirectToLds": True, + "UseSgprForGRO": True, + "UseInstOffsetForGRO": True, + "AssertSummationElementMultiple": True, + "AssertFree0ElementMultiple": True, + "AssertFree1ElementMultiple": True, + "AssertAIGreaterThanEqual": True, + "AssertAILessThanEqual": True, + "StaggerU": True, + "StaggerUStride": True, + "StaggerUMapping": True, + "MagicDivAlg": True, + "GlobalSplitU": True, + "GlobalSplitUAlgorithm": True, + "GlobalSplitUCoalesced": True, + "GlobalSplitUWorkGroupMappingRoundRobin": True, + "Use64bShadowLimit": True, + "NumLoadsCoalescedA": True, + "NumLoadsCoalescedB": True, + "WorkGroup": True, + "WorkGroupMapping": True, + "WorkGroupMappingXCC": True, + "WorkGroupMappingXCCGroup": True, + "ThreadTile": True, + "WavefrontSize": True, + "MatrixInstruction": True, + "1LDSBuffer": True, + "DepthU": True, + "NonTemporalE": True, + "NonTemporalD": True, + "NonTemporalC": True, + "NonTemporalA": True, + "NonTemporalB": True, + "NonTemporalWS": True, + "NonTemporalMetadata": True, + "NonTemporal": True, + "PreloadKernArgs": True, + "CustomKernelName": True, + "NoReject": True, + "MinVgprNumber": True, + "MaxVgprNumber": True, + "TotalVgprNumber": True, + "StoreRemapVectorWidth": True, + "SourceSwap": True, + "StorePriorityOpt": True, + "NumElementsPerBatchStore": True, + "StoreSyncOpt": True, + "GroupLoadStore": True, + "MIArchVgpr": True, + "StreamK": True, + "StreamKAtomic": True, + "StreamKXCCMapping": True, + "DebugStreamK": True, + "ActivationFused": True, + "ActivationFuncCall": True, + "ActivationAlt": True, + "WorkGroupReduction": True, + "ConvertAfterDS": True, + "ForceDisableShadowInit": True, + "ISA": True, + "MIWaveTile": True, +} + + +_requiredParametersMin = {'ProblemType': False, 'InternalSupportParams': False, 'InnerUnroll': True, 'KernelLanguage': False, 'LdsPadA': True, 'LdsPadB': True, 'LdsPadMetadata': True, 'LdsBlockSizePerPadA': True, 'LdsBlockSizePerPadB': True, 'LdsBlockSizePerPadMetadata': True, 'TransposeLDS': True, 'MaxOccupancy': True, 'VectorWidthA': True, 'VectorWidthB': True, 'VectorStore': True, 'StoreVectorWidth': True, 'GlobalReadVectorWidthA': True, 'GlobalReadVectorWidthB': True, 'LocalReadVectorWidth': True, 'WaveSeparateGlobalReadA': True, 'WaveSeparateGlobalReadB': True, 'WaveSeparateGlobalReadMetadata': False, 'UnrollLoopSwapGlobalReadOrder': True, 'PrefetchGlobalRead': True, 'PrefetchLocalRead': True, 'ClusterLocalRead': True, 'SuppressNoLoadLoop': False, 'ExpandPointerSwap': True, 'ScheduleGlobalRead': False, 'ScheduleLocalWrite': False, 'ScheduleIterAlg': True, 'GlobalReadPerMfma': True, 'LocalWritePerMfma': True, 'InterleaveAlpha': False, 'OptNoLoadLoop': True, 'BufferLoad': False, 'BufferStore': False, 'DirectToVgprA': True, 'DirectToVgprB': False, 'DirectToVgprSparseMetadata': False, 'DirectToLds': False, 'UseSgprForGRO': True, 'UseInstOffsetForGRO': True, 'AssertSummationElementMultiple': True, 'AssertFree0ElementMultiple': True, 'AssertFree1ElementMultiple': True, 'AssertAIGreaterThanEqual': False, 'AssertAILessThanEqual': False, 'StaggerU': True, 'StaggerUStride': True, 'StaggerUMapping': True, 'MagicDivAlg': False, 'GlobalSplitU': True, 'GlobalSplitUAlgorithm': True, 'GlobalSplitUCoalesced': True, 'GlobalSplitUWorkGroupMappingRoundRobin': True, 'Use64bShadowLimit': True, 'NumLoadsCoalescedA': True, 'NumLoadsCoalescedB': True, 'WorkGroup': True, 'WorkGroupMapping': True, 'WorkGroupMappingXCC': True, 'WorkGroupMappingXCCGroup': True, 'ThreadTile': False, 'WavefrontSize': True, 'MatrixInstruction': False, '1LDSBuffer': True, 'DepthU': False, 'NonTemporalE': False, 'NonTemporalD': True, 'NonTemporalC': True, 'NonTemporalA': True, 'NonTemporalB': True, 'NonTemporalWS': False, 'NonTemporalMetadata': True, 'NonTemporal': True, 'PreloadKernArgs': True, 'CustomKernelName': False, 'NoReject': False, 'MinVgprNumber': False, 'MaxVgprNumber': False, 'TotalVgprNumber': False, 'StoreRemapVectorWidth': True, 'SourceSwap': True, 'StorePriorityOpt': True, 'NumElementsPerBatchStore': True, 'StoreSyncOpt': True, 'GroupLoadStore': False, 'MIArchVgpr': True, 'StreamK': False, 'StreamKAtomic': False, 'StreamKXCCMapping': False, 'DebugStreamK': False, 'ActivationFused': False, 'ActivationFuncCall': True, 'ActivationAlt': False, 'WorkGroupReduction': False, 'ConvertAfterDS': True, 'ForceDisableShadowInit': False, 'ISA': True, 'CodeObjectVersion': False, 'AssignedDerivedParameters': False, 'AssignedProblemIndependentDerivedParameters': False, 'CUCount': False, 'DepthULdsDivisor': False, 'DirectToLdsA': False, 'DirectToLdsB': False, 'EdgeType': False, 'EnableF32XdlMathOp': False, 'EnableMatrixInstruction': False, 'GlobalWriteVectorWidth': False, 'GuaranteeNoPartialA': False, 'GuaranteeNoPartialB': False, 'GuaranteeNoPartialMetadata': False, 'LSCA': False, 'LSCB': False, 'LSPA': False, 'LSPB': False, 'LVCA': False, 'LVCB': False, 'LVPA': False, 'LVPB': False, 'LdsBlockSizePerPad': False, 'LdsInitCVgprs': False, 'LdsNumBytes': False, 'LdsNumElements': False, 'LdsNumElementsAlignedA': False, 'LdsNumElementsAlignedB': False, 'LdsNumElementsAlignedMetadata': False, 'LdsOffsetA': False, 'LdsOffsetA_Blk': False, 'LdsOffsetB': False, 'LdsOffsetB_Blk': False, 'LdsOffsetBias': False, 'LdsOffsetBiasGSU': False, 'LdsOffsetBiasNonGSU': False, 'LdsOffsetMetadata': False, 'LdsOffsetMetadata_Blk': False, 'LocalSplitU': False, 'LocalWriteUseSgprA': False, 'LocalWriteUseSgprB': False, 'LoopIters': False, 'LoopUnroll': False, 'MFMA_BF16_1K': False, 'MIBlock': False, 'MIInputPerThread': False, 'MIInputPerThreadA': False, 'MIInputPerThreadB': False, 'MIOutputVectorWidth': False, 'MIRegPerOut': False, 'MIWaveGroup': False, 'MIWaveTile': True, 'MIWaveTileA': False, 'MIWaveTileB': False, 'MIWaveTileMetadata': False, 'MacroTile0': False, 'MacroTile1': False, 'MacroTileA': False, 'MacroTileB': False, 'MatrixInstB': False, 'MatrixInstBM': False, 'MatrixInstBN': False, 'MatrixInstK': False, 'MatrixInstM': False, 'MatrixInstN': False, 'NoLdsWriteCode': False, 'NoTailLoop': False, 'NumElementsPerThread': False, 'NumGlobalWriteVectorsPerThread': False, 'NumLoadsA': False, 'NumLoadsB': False, 'NumLoadsPerpendicularA': False, 'NumLoadsPerpendicularB': False, 'NumThreads': False, 'PackedC0IdxChars': False, 'PackedC0IndicesX': False, 'PackedC1IdxChars': False, 'PackedC1IndicesX': False, 'SolutionIndex': False, 'SolutionNameMin': False, 'SubGroup0': False, 'SubGroup1': False, 'SubGroupA': False, 'SubGroupB': False, 'ThreadTile0': False, 'ThreadTile1': False, 'ThreadTileA': False, 'ThreadTileB': False, 'TransposeLDSMetadata': False, 'UnrollMajorLDSA': False, 'UnrollMajorLDSB': False, 'UnrollMajorLDSMetadata': False, 'Valid': False, 'VectorWidth': False, 'WorkspaceCheck': False, '_DepthU': False, '_DepthUA': False, '_DepthUB': False, '_DepthULds': False, '_DepthUMetadata': False, '_GlobalAccumulation': False, '_UseSgprForGRO': False, '_VectorStore': False, '_WorkspaceSizePerElemBias': False, '_WorkspaceSizePerElemC': False, '_staggerStrideShift': False, 'allowLRVWforTLUandMI': False, 'LdsBytesNoAmax': False, 'ULSGRODoubleG2L': False, 'codeObjectFile': False, 'Kernel': True} + + +def getRequiredParametersFull(): + return _requiredParametersFull + +def getRequiredParametersMin(): + return _requiredParametersMin + + +# Begin comparison tests +if __name__ == "__main__": + from deepdiff import DeepDiff + reqParamA = {'ProblemType': False, 'InternalSupportParams': False, 'InnerUnroll': True, 'KernelLanguage': False, 'LdsPadA': True, 'LdsPadB': True, 'LdsPadMetadata': True, 'LdsBlockSizePerPadA': True, 'LdsBlockSizePerPadB': True, 'LdsBlockSizePerPadMetadata': True, 'TransposeLDS': True, 'MaxOccupancy': True, 'VectorWidthA': True, 'VectorWidthB': True, 'VectorStore': True, 'StoreVectorWidth': True, 'GlobalReadVectorWidthA': True, 'GlobalReadVectorWidthB': True, 'LocalReadVectorWidth': True, 'WaveSeparateGlobalReadA': True, 'WaveSeparateGlobalReadB': True, 'WaveSeparateGlobalReadMetadata': False, 'UnrollLoopSwapGlobalReadOrder': True, 'PrefetchGlobalRead': True, 'PrefetchLocalRead': True, 'ClusterLocalRead': True, 'SuppressNoLoadLoop': False, 'ExpandPointerSwap': True, 'ScheduleGlobalRead': False, 'ScheduleLocalWrite': False, 'ScheduleIterAlg': True, 'GlobalReadPerMfma': True, 'LocalWritePerMfma': True, 'InterleaveAlpha': False, 'OptNoLoadLoop': True, 'BufferLoad': False, 'BufferStore': False, 'DirectToVgprA': True, 'DirectToVgprB': False, 'DirectToVgprSparseMetadata': False, 'DirectToLds': False, 'UseSgprForGRO': True, 'UseInstOffsetForGRO': True, 'AssertSummationElementMultiple': True, 'AssertFree0ElementMultiple': True, 'AssertFree1ElementMultiple': True, 'AssertAIGreaterThanEqual': False, 'AssertAILessThanEqual': False, 'StaggerU': True, 'StaggerUStride': True, 'StaggerUMapping': True, 'MagicDivAlg': False, 'GlobalSplitU': True, 'GlobalSplitUAlgorithm': True, 'GlobalSplitUCoalesced': True, 'GlobalSplitUWorkGroupMappingRoundRobin': True, 'Use64bShadowLimit': True, 'NumLoadsCoalescedA': True, 'NumLoadsCoalescedB': True, 'WorkGroup': True, 'WorkGroupMapping': True, 'WorkGroupMappingXCC': True, 'WorkGroupMappingXCCGroup': True, 'ThreadTile': False, 'WavefrontSize': True, 'MatrixInstruction': False, '1LDSBuffer': True, 'DepthU': False, 'NonTemporalE': False, 'NonTemporalD': True, 'NonTemporalC': True, 'NonTemporalA': True, 'NonTemporalB': True, 'NonTemporalWS': False, 'NonTemporalMetadata': True, 'NonTemporal': True, 'PreloadKernArgs': True, 'CustomKernelName': False, 'NoReject': False, 'MinVgprNumber': False, 'MaxVgprNumber': False, 'TotalVgprNumber': False, 'StoreRemapVectorWidth': True, 'SourceSwap': True, 'StorePriorityOpt': True, 'NumElementsPerBatchStore': True, 'StoreSyncOpt': True, 'GroupLoadStore': False, 'MIArchVgpr': True, 'StreamK': False, 'StreamKAtomic': False, 'StreamKXCCMapping': False, 'DebugStreamK': False, 'ActivationFused': False, 'ActivationFuncCall': True, 'ActivationAlt': False, 'WorkGroupReduction': False, 'ConvertAfterDS': True, 'ForceDisableShadowInit': False, 'ISA': True, 'CodeObjectVersion': False, 'AssignedDerivedParameters': False, 'AssignedProblemIndependentDerivedParameters': False, 'CUCount': False, 'DepthULdsDivisor': False, 'DirectToLdsA': False, 'DirectToLdsB': False, 'EdgeType': False, 'EnableF32XdlMathOp': False, 'EnableMatrixInstruction': False, 'GlobalWriteVectorWidth': False, 'GuaranteeNoPartialA': False, 'GuaranteeNoPartialB': False, 'GuaranteeNoPartialMetadata': False, 'LSCA': False, 'LSCB': False, 'LSPA': False, 'LSPB': False, 'LVCA': False, 'LVCB': False, 'LVPA': False, 'LVPB': False, 'LdsBlockSizePerPad': False, 'LdsInitCVgprs': False, 'LdsNumBytes': False, 'LdsNumElements': False, 'LdsNumElementsAlignedA': False, 'LdsNumElementsAlignedB': False, 'LdsNumElementsAlignedMetadata': False, 'LdsOffsetA': False, 'LdsOffsetA_Blk': False, 'LdsOffsetB': False, 'LdsOffsetB_Blk': False, 'LdsOffsetBias': False, 'LdsOffsetBiasGSU': False, 'LdsOffsetBiasNonGSU': False, 'LdsOffsetMetadata': False, 'LdsOffsetMetadata_Blk': False, 'LocalSplitU': False, 'LocalWriteUseSgprA': False, 'LocalWriteUseSgprB': False, 'LoopIters': False, 'LoopUnroll': False, 'MFMA_BF16_1K': False, 'MIBlock': False, 'MIInputPerThread': False, 'MIInputPerThreadA': False, 'MIInputPerThreadB': False, 'MIOutputVectorWidth': False, 'MIRegPerOut': False, 'MIWaveGroup': False, 'MIWaveTile': True, 'MIWaveTileA': False, 'MIWaveTileB': False, 'MIWaveTileMetadata': False, 'MacroTile0': False, 'MacroTile1': False, 'MacroTileA': False, 'MacroTileB': False, 'MatrixInstB': False, 'MatrixInstBM': False, 'MatrixInstBN': False, 'MatrixInstK': False, 'MatrixInstM': False, 'MatrixInstN': False, 'NoLdsWriteCode': False, 'NoTailLoop': False, 'NumElementsPerThread': False, 'NumGlobalWriteVectorsPerThread': False, 'NumLoadsA': False, 'NumLoadsB': False, 'NumLoadsPerpendicularA': False, 'NumLoadsPerpendicularB': False, 'NumThreads': False, 'PackedC0IdxChars': False, 'PackedC0IndicesX': False, 'PackedC1IdxChars': False, 'PackedC1IndicesX': False, 'SolutionIndex': False, 'SolutionNameMin': False, 'SubGroup0': False, 'SubGroup1': False, 'SubGroupA': False, 'SubGroupB': False, 'ThreadTile0': False, 'ThreadTile1': False, 'ThreadTileA': False, 'ThreadTileB': False, 'TransposeLDSMetadata': False, 'UnrollMajorLDSA': False, 'UnrollMajorLDSB': False, 'UnrollMajorLDSMetadata': False, 'Valid': False, 'VectorWidth': False, 'WorkspaceCheck': False, '_DepthU': False, '_DepthUA': False, '_DepthUB': False, '_DepthULds': False, '_DepthUMetadata': False, '_GlobalAccumulation': False, '_UseSgprForGRO': False, '_VectorStore': False, '_WorkspaceSizePerElemBias': False, '_WorkspaceSizePerElemC': False, '_staggerStrideShift': False, 'allowLRVWforTLUandMI': False, 'LdsBytesNoAmax': False, 'ULSGRODoubleG2L': False, 'codeObjectFile': False, 'Kernel': True} + + print(DeepDiff(reqParamA, _requiredParametersMin, ignore_order=True).pretty()) \ No newline at end of file