From 7506f877fbbd1b97819f89b6f9ecc06514b00006 Mon Sep 17 00:00:00 2001 From: David Dixon Date: Sun, 5 Jan 2025 18:10:42 -0600 Subject: [PATCH 01/25] Make TensileCreateLibrary a directory and add RCL specific ParseArguments module --- .../TensileCreateLibrary/ParseArguments.py | 135 +++++++++ .../Tensile/TensileCreateLibrary/__init__.py | 1 + .../main.py} | 265 ++++-------------- tensilelite/Tensile/bin/TensileCreateLibrary | 7 +- 4 files changed, 194 insertions(+), 214 deletions(-) create mode 100644 tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py create mode 100644 tensilelite/Tensile/TensileCreateLibrary/__init__.py rename tensilelite/Tensile/{TensileCreateLibrary.py => TensileCreateLibrary/main.py} (60%) diff --git a/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py b/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py new file mode 100644 index 0000000000..f8b8a272f4 --- /dev/null +++ b/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py @@ -0,0 +1,135 @@ +################################################################################ +# +# Copyright (C) 2022-2024 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.Utilities.Toolchain import ToolchainDefaults + +import os +from argparse import ArgumentParser +from typing import Any, Dict, List, Optional + +from Tensile.Common import architectureMap + +def parseArguments(input: Optional[List[str]] = None) -> Dict[str, Any]: + """Parse command line arguments for TensileCreateLibrary. + + Args: + input: List of strings representing command line arguments used when + calling parseArguments prgrammatically e.g. in testing. + + Returns: + A dictionary containing the keys representing options and their values. + """ + + argParser = ArgumentParser( + description="TensileCreateLibrary generates libraries and code object files " + "for a set of supplied logic files.", + ) + + argParser.add_argument("LogicPath", help="Path to LibraryLogic.yaml files.") + argParser.add_argument("OutputPath", help="Where to write library files?") + argParser.add_argument("RuntimeLanguage", help="Which runtime language?", choices=["OCL", "HIP", "HSA"]) + argParser.add_argument("--cxx-compiler", dest="CxxCompiler", action="store", default=ToolchainDefaults.CXX_COMPILER, + help=f"Default: {ToolchainDefaults.CXX_COMPILER}") + argParser.add_argument("--c-compiler", dest="CCompiler", action="store", default=ToolchainDefaults.C_COMPILER) + argParser.add_argument("--cmake-cxx-compiler", dest="CmakeCxxCompiler", action="store") + argParser.add_argument("--offload-bundler", dest="OffloadBundler", action="store", default=ToolchainDefaults.OFFLOAD_BUNDLER) + argParser.add_argument("--assembler", dest="Assembler", action="store", default=ToolchainDefaults.ASSEMBLER) + argParser.add_argument("--code-object-version", dest="CodeObjectVersion", choices=["default", "V4", "V5"], action="store") + argParser.add_argument("--architecture", dest="Architecture", type=str, action="store", default="all", help="Supported archs: " + " ".join(architectureMap.keys())) + argParser.add_argument("--short-file-names", dest="ShortNames", action="store_true") + argParser.add_argument("--no-short-file-names", dest="ShortNames", action="store_false") + argParser.add_argument("--library-print-debug", dest="LibraryPrintDebug", action="store_true") + argParser.add_argument("--no-library-print-debug", dest="LibraryPrintDebug", action="store_false") + argParser.add_argument("--no-compress", dest="NoCompress", action="store_true", help="Don't compress assembly code objects.") + argParser.add_argument("--experimental", dest="Experimental", action="store_true", + help="Include logic files in directories named 'Experimental'.") + argParser.add_argument("--no-enumerate", action="store_true", help="Do not run rocm_agent_enumerator.") + argParser.add_argument("--version", help="Version string to embed into library file.") + argParser.add_argument("--logic-format", dest="LogicFormat", choices=["yaml", "json"], \ + action="store", default="yaml", help="select which logic format to use") + argParser.add_argument("--library-format", dest="LibraryFormat", choices=["yaml", "msgpack"], + action="store", default="msgpack", help="select which library format to use") + argParser.add_argument("--generate-sources-and-exit", dest="GenerateSourcesAndExit", action="store_true", + default=False, help="Output source files only and exit.") + argParser.add_argument("--jobs", "-j", dest="CpuThreads", type=int, + default=-1, help="Number of parallel jobs to launch.") + argParser.add_argument("--verbose", "-v", dest="PrintLevel", type=int, + default=1, help="Set printout verbosity level.") + argParser.add_argument("--print-timing", dest="PrintTiming", + default=False, action="store_true", help="Print duration of each stage.") + argParser.add_argument("--separate-architectures", dest="SeparateArchitectures", action="store_true", + default=False, help="Separates TensileLibrary file by architecture") + argParser.add_argument("--lazy-library-loading", dest="LazyLibraryLoading", action="store_true", + default=False, help="Loads Tensile libraries when needed instead of upfront.") + argParser.add_argument("--enable-marker", dest="EnableMarker", action="store_true", + default=False, help="Enable marker in Tensile.") + argParser.add_argument("--no-generate-solution-table", dest="GenSolTable", action="store_false", default=True, + help="Skip generating solution-yaml matching table") + argParser.add_argument("--asm-debug", dest="AsmDebug", action="store_true", default=False, + help="Keep debug information for built code objects") + argParser.add_argument("--build-id", dest="BuildIdKind", action="store", default="sha1") + argParser.add_argument("--address-sanitizer", dest="AsanBuild", action="store_true", + default=False, help="Enable ASAN build.") + argParser.add_argument("--keep-build-tmp", dest="KeepBuildTmp", action="store_true", + default=False, help="Do not remove the temporary build directory (may required hundreds of GBs of space)"), + argParser.add_argument("--validate-library", dest="ValidateLibrary", action="store_true", default=False) + argParser.add_argument("--logic-filter", dest="LogicFilter", action="store", default="*", type=str, + help="Cutomsized logic filter, default is *, i.e. all logics." + " Example: gfx942/Equality/* for building equality of gfx942 only") + + args = argParser.parse_args() + + arguments = {} + arguments["RuntimeLanguage"] = args.RuntimeLanguage + arguments["CodeObjectVersion"] = args.CodeObjectVersion + arguments["Architecture"] = args.Architecture + arguments["SeparateArchitectures"] = args.SeparateArchitectures + arguments["LazyLibraryLoading"] = args.LazyLibraryLoading + arguments["EnableMarker"] = args.EnableMarker + if args.CmakeCxxCompiler: + os.environ["CMAKE_CXX_COMPILER"] = args.CmakeCxxCompiler + arguments["ShortNames"] = args.ShortNames + arguments["LibraryPrintDebug"] = args.LibraryPrintDebug + arguments["CodeFromFiles"] = False + arguments["LogicFormat"] = args.LogicFormat + arguments["LibraryFormat"] = args.LibraryFormat + if args.no_enumerate: + arguments["AMDGPUArchPath"] = False + arguments["CpuThreads"] = args.CpuThreads + arguments["PrintLevel"] = args.PrintLevel + arguments["PrintTiming"] = args.PrintTiming + arguments["AsmDebug"] = args.AsmDebug + arguments["BuildIdKind"] = args.BuildIdKind + arguments["KeepBuildTmp"] = args.KeepBuildTmp + arguments["AsanBuild"] = args.AsanBuild + arguments["ValidateLibrary"] = args.ValidateLibrary + arguments["UseCompression"] = not args.NoCompress + arguments["CxxCompiler"] = args.CxxCompiler + arguments["CCompiler"] = args.CCompiler + arguments["OffloadBundler"] = args.OffloadBundler + arguments["Assembler"] = args.Assembler + arguments["LogicPath"] = args.LogicPath + arguments["OutputPath"] = args.OutputPath + + return arguments diff --git a/tensilelite/Tensile/TensileCreateLibrary/__init__.py b/tensilelite/Tensile/TensileCreateLibrary/__init__.py new file mode 100644 index 0000000000..68f8e21c64 --- /dev/null +++ b/tensilelite/Tensile/TensileCreateLibrary/__init__.py @@ -0,0 +1 @@ +from .main import run \ No newline at end of file diff --git a/tensilelite/Tensile/TensileCreateLibrary.py b/tensilelite/Tensile/TensileCreateLibrary/main.py similarity index 60% rename from tensilelite/Tensile/TensileCreateLibrary.py rename to tensilelite/Tensile/TensileCreateLibrary/main.py index 41777a4cba..4dd940cc8a 100644 --- a/tensilelite/Tensile/TensileCreateLibrary.py +++ b/tensilelite/Tensile/TensileCreateLibrary/main.py @@ -22,38 +22,27 @@ # ################################################################################ -# This script only gets called by CMake - -if __name__ == "__main__": - print("This file can no longer be run as a script. Run 'Tensile/bin/TensileCreateLibrary' instead.") - exit(1) - -from . import Common -from . import ClientExecutable -from . import EmbeddedData -from . import LibraryIO -from . import Utils -from .TensileInstructions import getGfxName, TensileInstructions -from .Common import globalParameters, HR, print1, print2, printExit, ensurePath, \ - CHeader, CMakeHeader, assignGlobalParameters, \ - architectureMap, printWarning, \ - splitArchs -from .KernelWriterAssembly import KernelWriterAssembly -from .SolutionLibrary import MasterSolutionLibrary -from .SolutionStructs import Solution -from .CustomYamlLoader import load_logic_gfx_arch -from .Utilities.Profile import profile -from .Utilities.Toolchain import getVersion, validateToolchain, ToolchainDefaults -from .BuildCommands import SourceCommands, AssemblyCommands - -import argparse +from .ParseArguments import parseArguments + +from Tensile import Common +from Tensile import LibraryIO +from Tensile import Utils +from Tensile.TensileInstructions import getGfxName, TensileInstructions +from Tensile.Common import globalParameters, HR, print1, print2, printExit, ensurePath, \ + CHeader, assignGlobalParameters, architectureMap +from Tensile.KernelWriterAssembly import KernelWriterAssembly +from Tensile.SolutionLibrary import MasterSolutionLibrary +from Tensile.SolutionStructs import Solution +from Tensile.CustomYamlLoader import load_logic_gfx_arch +from Tensile.Utilities.Profile import profile +from Tensile.Utilities.Toolchain import getVersion, validateToolchain, ToolchainDefaults +from Tensile.BuildCommands import SourceCommands, AssemblyCommands + import collections import glob import itertools import os -import re import shutil -import sys from timeit import default_timer as timer from pathlib import Path from typing import Sequence, List, Union @@ -69,7 +58,8 @@ def wrapper(*args, **kwargs): return res return wrapper -################################################################################ + + def processKernelSource(kernel, kernelWriterAssembly, ti): """ Generate source for a single kernel. @@ -91,8 +81,6 @@ def processKernelSource(kernel, kernelWriterAssembly, ti): return (err, src, header, kernelName, filename) - -################################################################################ def buildKernelSourceAndHeaderFiles(results, outputPath, kernelsWithBuildErrs): """ Logs errors and writes appropriate info to kernelSourceFile and kernelHeaderFile. @@ -108,23 +96,19 @@ def buildKernelSourceAndHeaderFiles(results, outputPath, kernelsWithBuildErrs): sourceFilenames: Array containing source kernel filenames """ - # Find kernels to write kernelsToWrite = [] filesToWrite = collections.defaultdict(list) validKernelCount = 0 for (err,src,header,kernelName, filename) in results: - # Keep track of kernels with errors if err: kernelsWithBuildErrs[kernelName] = err - # Don't create a file for empty kernels if len(src.strip()) == 0: continue kernelsToWrite.append((err, src, header, kernelName)) - # Create list of files if filename: filesToWrite[os.path.join(os.path.normcase(outputPath),filename)].append((err, src, header, kernelName)) else: @@ -139,9 +123,6 @@ def buildKernelSourceAndHeaderFiles(results, outputPath, kernelsWithBuildErrs): kernelSuffix = "" filesToWrite[os.path.join(os.path.normcase(outputPath), "Kernels"+kernelSuffix)] = [] - - # Write kernel data to files - #Parse list of files and write kernels for filename, kernelList in filesToWrite.items(): with open(filename+".h", "w", encoding="utf-8") as kernelHeaderFile, \ open(filename+".cpp", "w", encoding="utf-8") as kernelSourceFile: @@ -163,9 +144,7 @@ def buildKernelSourceAndHeaderFiles(results, outputPath, kernelsWithBuildErrs): return sourceFilenames -################################################################################ -# Write Solutions and Kernels for BenchmarkClient or LibraryClient -################################################################################ + @timing def writeSolutionsAndKernels(outputPath, cxxCompiler, assembler, offloadBundler, solutions, kernels, kernelHelperObjs, \ kernelWriterAssembly, errorTolerant=False, compress=True): @@ -184,9 +163,6 @@ def writeSolutionsAndKernels(outputPath, cxxCompiler, assembler, offloadBundler, kernelSourceFile = None kernelHeaderFile = None - ############################################################################## - # Write Kernels - ############################################################################## kernelsWithBuildErrs = {} # Kernels may be intended for different co files, but generate the same .o file @@ -256,7 +232,6 @@ def success(kernel): kernelHeaderFile = open(kernelHeaderFilename, "a", encoding="utf-8") HeaderText = "" - # handle helper kernel function for ko in kernelHelperObjs: kernelName = ko.getKernelName() @@ -267,7 +242,6 @@ def success(kernel): HeaderText += ko.getHeaderFileString() - # write kernel.h in one shot kernelHeaderFile.write(HeaderText) if kernelSourceFile: @@ -285,24 +259,17 @@ def success(kernel): return codeObjectFiles, numKernels -############################################################################## -# Min Naming / Solution and Kernel Writers -############################################################################## @timing def getSolutionAndKernelWriters(solutions, kernels, assembler): - # if any kernels are assembly, append every ISA supported kernelSerialNaming = Solution.getSerialNaming(kernels) - solutionMinNaming = Solution.getMinNaming(solutions) kernelMinNaming = Solution.getMinNaming(kernels) kernelWriterAssembly = KernelWriterAssembly(kernelMinNaming, kernelSerialNaming, assembler) return (kernelWriterAssembly, kernelMinNaming, solutionMinNaming) -################################################################################ -# copy static cpp files and headers -################################################################################ + @timing def copyStaticFiles(outputPath=None): if outputPath is None: @@ -324,9 +291,6 @@ def copyStaticFiles(outputPath=None): return libraryStaticFiles -################################################################################ -# Generate Kernel Objects From Solutions -################################################################################ @timing def generateKernelObjectsFromSolutions(solutions): # create solution writer and kernel writer @@ -351,17 +315,14 @@ def generateKernelObjectsFromSolutions(solutions): kernelHelperObjs = list(dict.fromkeys(kernelHelperObjs)) return (kernels, kernelHelperObjs, kernelHelperNames) -################################################################################ -# Generate Logic Data and Solutions -################################################################################ + @timing def generateLogicDataAndSolutions(logicFiles, args, cxxCompiler): - # skip the logic which architectureName is not in the build target. - if ";" in args.Architecture: - archs = args.Architecture.split(";") # user arg list format + if ";" in args["Architecture"]: + archs = args["Architecture"].split(";") # user arg list format else: - archs = args.Architecture.split("_") # workaround for cmake list in list issue + archs = args["Architecture"].split("_") # workaround for cmake list in list issue solutions = [] masterLibraries = {} @@ -389,15 +350,15 @@ def libraryIter(lib: MasterSolutionLibrary): nextSolIndex = masterLibraries[architectureName].merge(newLibrary, nextSolIndex) else: masterLibraries[architectureName] = newLibrary - masterLibraries[architectureName].version = args.version + masterLibraries[architectureName].version = args["version"] else: if fullMasterLibrary is None: fullMasterLibrary = newLibrary - fullMasterLibrary.version = args.version + fullMasterLibrary.version = args["version"] else: fullMasterLibrary.merge(newLibrary) - if args.GenSolTable: + if args["GenSolTable"]: # Match yaml file solutions to solution index for localIdx, _, s in libraryIter(newLibrary): matchTable[s.index] = [srcFile, localIdx] @@ -423,7 +384,7 @@ def libraryIter(lib: MasterSolutionLibrary): # remove duplicates while preserving order solutions = dict.fromkeys(solutions).keys() - if args.GenSolTable: + if args["GenSolTable"]: LibraryIO.write("MatchTable", matchTable) return solutions, masterLibraries, fullMasterLibrary @@ -446,11 +407,12 @@ def validateLibrary(masterLibraries: MasterSolutionLibrary, assert ok and "Inconsistent kernel sizes detected!" + ################################################################################ # Tensile Create Library ################################################################################ @profile -def TensileCreateLibrary(): +def run(): start = timer() print1("") print1(HR) @@ -458,122 +420,20 @@ def TensileCreateLibrary(): print2(HR) print2("") - ############################################################################## - # Parse Command Line Arguments - ############################################################################## - def splitExtraParameters(par): - """ - Allows the --global-parameters option to specify any parameters from the command line. - """ + arguments = parseArguments() + logicPath = arguments["LogicPath"] + outputPath = arguments["OutputPath"] + cxxCompiler = arguments["CxxCompiler"] + offloadBundler = arguments["OffloadBundler"] + assembler = arguments["Assembler"] + libraryFormat = arguments["LibraryFormat"] + useCompression = arguments["UseCompression"] - (key, value) = par.split("=") - value = eval(value) - return (key, value) - - print2("Arguments: %s" % sys.argv) - argParser = argparse.ArgumentParser() - argParser.add_argument("LogicPath", help="Path to LibraryLogic.yaml files.") - argParser.add_argument("OutputPath", help="Where to write library files?") - argParser.add_argument("RuntimeLanguage", help="Which runtime language?", choices=["OCL", "HIP", "HSA"]) - argParser.add_argument("--cxx-compiler", dest="CxxCompiler", action="store", default=ToolchainDefaults.CXX_COMPILER, - help=f"Default: {ToolchainDefaults.CXX_COMPILER}") - argParser.add_argument("--c-compiler", dest="CCompiler", action="store", default=ToolchainDefaults.C_COMPILER) - argParser.add_argument("--cmake-cxx-compiler", dest="CmakeCxxCompiler", action="store") - argParser.add_argument("--offload-bundler", dest="OffloadBundler", action="store", default=ToolchainDefaults.OFFLOAD_BUNDLER) - argParser.add_argument("--assembler", dest="Assembler", action="store", default=ToolchainDefaults.ASSEMBLER) - argParser.add_argument("--code-object-version", dest="CodeObjectVersion", choices=["default", "V4", "V5"], action="store") - argParser.add_argument("--architecture", dest="Architecture", type=str, action="store", default="all", help="Supported archs: " + " ".join(architectureMap.keys())) - argParser.add_argument("--short-file-names", dest="ShortNames", action="store_true") - argParser.add_argument("--no-short-file-names", dest="ShortNames", action="store_false") - argParser.add_argument("--library-print-debug", dest="LibraryPrintDebug", action="store_true") - argParser.add_argument("--no-library-print-debug", dest="LibraryPrintDebug", action="store_false") - argParser.add_argument("--no-compress", dest="NoCompress", action="store_true", help="Don't compress assembly code objects.") - argParser.add_argument("--experimental", dest="Experimental", action="store_true", - help="Include logic files in directories named 'Experimental'.") - argParser.add_argument("--no-enumerate", action="store_true", help="Do not run rocm_agent_enumerator.") - argParser.add_argument("--version", help="Version string to embed into library file.") - argParser.add_argument("--logic-format", dest="LogicFormat", choices=["yaml", "json"], \ - action="store", default="yaml", help="select which logic format to use") - argParser.add_argument("--library-format", dest="LibraryFormat", choices=["yaml", "msgpack"], - action="store", default="msgpack", help="select which library format to use") - argParser.add_argument("--generate-sources-and-exit", dest="GenerateSourcesAndExit", action="store_true", - default=False, help="Output source files only and exit.") - argParser.add_argument("--jobs", "-j", dest="CpuThreads", type=int, - default=-1, help="Number of parallel jobs to launch.") - argParser.add_argument("--verbose", "-v", dest="PrintLevel", type=int, - default=1, help="Set printout verbosity level.") - argParser.add_argument("--print-timing", dest="PrintTiming", - default=False, action="store_true", help="Print duration of each stage.") - argParser.add_argument("--separate-architectures", dest="SeparateArchitectures", action="store_true", - default=False, help="Separates TensileLibrary file by architecture") - argParser.add_argument("--lazy-library-loading", dest="LazyLibraryLoading", action="store_true", - default=False, help="Loads Tensile libraries when needed instead of upfront.") - argParser.add_argument("--enable-marker", dest="EnableMarker", action="store_true", - default=False, help="Enable marker in Tensile.") - argParser.add_argument("--global-parameters", nargs="+", type=splitExtraParameters, default=[]) - argParser.add_argument("--no-generate-solution-table", dest="GenSolTable", action="store_false", default=True, - help="Skip generating solution-yaml matching table") - argParser.add_argument("--asm-debug", dest="AsmDebug", action="store_true", default=False, - help="Keep debug information for built code objects") - argParser.add_argument("--build-id", dest="BuildIdKind", action="store", default="sha1") - argParser.add_argument("--address-sanitizer", dest="AsanBuild", action="store_true", - default=False, help="Enable ASAN build.") - argParser.add_argument("--keep-build-tmp", dest="KeepBuildTmp", action="store_true", - default=False, help="Do not remove the temporary build directory (may required hundreds of GBs of space)"), - argParser.add_argument("--validate-library", dest="ValidateLibrary", action="store_true", default=False) - argParser.add_argument("--logic-filter", dest="LogicFilter", action="store", default="*", type=str, - help="Cutomsized logic filter, default is *, i.e. all logics." - " Example: gfx942/Equality/* for building equality of gfx942 only") - - args = argParser.parse_args() - - logicPath = args.LogicPath - outputPath = args.OutputPath - cxxCompiler = args.CxxCompiler - offloadBundler = args.OffloadBundler - assembler = args.Assembler - libraryFormat = args.LibraryFormat - useCompression = not args.NoCompress - print2("OutputPath: %s" % outputPath) - ensurePath(outputPath) outputPath = os.path.abspath(outputPath) - arguments = {} - arguments["RuntimeLanguage"] = args.RuntimeLanguage - arguments["CodeObjectVersion"] = args.CodeObjectVersion - arguments["Architecture"] = args.Architecture - arguments["SeparateArchitectures"] = args.SeparateArchitectures - arguments["LazyLibraryLoading"] = args.LazyLibraryLoading - arguments["EnableMarker"] = args.EnableMarker - if args.CmakeCxxCompiler: - os.environ["CMAKE_CXX_COMPILER"] = args.CmakeCxxCompiler - arguments["ShortNames"] = args.ShortNames - arguments["LibraryPrintDebug"] = args.LibraryPrintDebug - arguments["CodeFromFiles"] = False - arguments["LogicFormat"] = args.LogicFormat - arguments["LibraryFormat"] = args.LibraryFormat - if args.no_enumerate: - arguments["AMDGPUArchPath"] = False - - arguments["GenerateSourcesAndExit"] = args.GenerateSourcesAndExit - if arguments["GenerateSourcesAndExit"]: - # Generated sources are preserved and go into output dir - arguments["WorkingPath"] = outputPath - - arguments["CpuThreads"] = args.CpuThreads - arguments["PrintLevel"] = args.PrintLevel - arguments["PrintTiming"] = args.PrintTiming - arguments["AsmDebug"] = args.AsmDebug - arguments["BuildIdKind"] = args.BuildIdKind - arguments["KeepBuildTmp"] = args.KeepBuildTmp - arguments["AsanBuild"] = args.AsanBuild - arguments["ValidateLibrary"] = args.ValidateLibrary - - for key, value in args.global_parameters: - arguments[key] = value - cxxCompiler, cCompiler, offloadBundler, assembler, hipconfig = validateToolchain( - args.CxxCompiler, args.CCompiler, args.OffloadBundler, args.Assembler, ToolchainDefaults.HIP_CONFIG + arguments["CxxCompiler"], arguments["CCompiler"], arguments["OffloadBundler"], arguments["Assembler"], ToolchainDefaults.HIP_CONFIG ) + print1(f"# HIP Version: {getVersion(hipconfig, regex=r'(.+)')}") print1(f"# Cxx Compiler: {cxxCompiler} (version {getVersion(cxxCompiler)})") print1(f"# C Compiler: {cCompiler} (version {getVersion(cCompiler)})") @@ -589,9 +449,9 @@ def splitExtraParameters(par): printExit("LogicPath %s doesn't exist" % logicPath) if ";" in arguments["Architecture"]: - archs = arguments["Architecture"].split(";") # user arg list format + archs = arguments["Architecture"].split(";") else: - archs = arguments["Architecture"].split("_") # workaround for cmake list in list issue + archs = arguments["Architecture"].split("_") logicArchs = set() for arch in archs: if arch in architectureMap: @@ -599,14 +459,13 @@ def splitExtraParameters(par): else: printExit("Architecture %s not supported" % arch) - # Recursive directory search logicExtFormat = ".yaml" - if args.LogicFormat == "yaml": + if arguments["LogicFormat"] == "yaml": pass - elif args.LogicFormat == "json": + elif arguments["LogicFormat"] == "json": logicExtFormat = ".json" else: - printExit("Unrecognized LogicFormat", args.LogicFormat) + printExit("Unrecognized LogicFormat", arguments["LogicFormat"]) def archMatch(arch: str, archs: List[str]): return (arch in archs) or any(a.startswith(arch) for a in archs) @@ -614,30 +473,21 @@ def archMatch(arch: str, archs: List[str]): def validLogicFile(p: Path): return p.suffix == logicExtFormat and ("all" in archs or archMatch(load_logic_gfx_arch(p), archs)) - globPattern = os.path.join(logicPath, f"**/{args.LogicFilter}{logicExtFormat}") + globPattern = os.path.join(logicPath, f"**/{arguments['LogicFilter']}{logicExtFormat}") print1(f"# LogicFilter: {globPattern}") logicFiles = (os.path.join(logicPath, file) for file in glob.iglob(globPattern, recursive=True)) logicFiles = [file for file in logicFiles if validLogicFile(Path(file))] - print1(f"# Experimental: {args.Experimental}") - if not args.Experimental: + 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) - - ############################################################################## - # Parse config files - ############################################################################## - - # Parse logicData, solutions, and masterLibraries from logic files - solutions, masterLibraries, fullMasterLibrary = generateLogicDataAndSolutions(logicFiles, args, cxxCompiler) - + solutions, masterLibraries, fullMasterLibrary = generateLogicDataAndSolutions(logicFiles, arguments, cxxCompiler) kernels, kernelHelperObjs, _ = generateKernelObjectsFromSolutions(solutions) - - # if any kernels are assembly, append every ISA supported kernelWriterAssembly, kernelMinNaming, _ = getSolutionAndKernelWriters(solutions, kernels, assembler) if globalParameters["ValidateLibrary"]: @@ -645,13 +495,11 @@ def validLogicFile(p: Path): staticFiles = copyStaticFiles(outputPath) - # Make sure to copy the library static files. for fileName in staticFiles: shutil.copy( os.path.join(globalParameters["SourcePath"], fileName), \ outputPath ) - # write solutions and kernels - codeObjectFiles, numKernels = writeSolutionsAndKernels(outputPath, cxxCompiler, assembler, offloadBundler, solutions, + _, numKernels = writeSolutionsAndKernels(outputPath, cxxCompiler, assembler, offloadBundler, solutions, kernels, kernelHelperObjs, kernelWriterAssembly, compress=useCompression) archs = [getGfxName(arch) for arch in globalParameters['SupportedISA'] \ @@ -666,23 +514,18 @@ def validLogicFile(p: Path): else: masterFile = os.path.join(newLibraryDir, "TensileLibrary_"+archName) newMasterLibrary.applyNaming(kernelMinNaming) - LibraryIO.write(masterFile, Utils.state(newMasterLibrary), args.LibraryFormat) + LibraryIO.write(masterFile, Utils.state(newMasterLibrary), libraryFormat) - #Write placeholder libraries for name, lib in newMasterLibrary.lazyLibraries.items(): filename = os.path.join(newLibraryDir, name) - lib.applyNaming(kernelMinNaming) #@TODO Check to see if kernelMinNaming is correct - LibraryIO.write(filename, Utils.state(lib), args.LibraryFormat) + lib.applyNaming(kernelMinNaming) + LibraryIO.write(filename, Utils.state(lib), libraryFormat) else: masterFile = os.path.join(newLibraryDir, "TensileLibrary") fullMasterLibrary.applyNaming = timing(fullMasterLibrary.applyNaming) fullMasterLibrary.applyNaming(kernelMinNaming) - LibraryIO.write(masterFile, Utils.state(fullMasterLibrary), args.LibraryFormat) - - theMasterLibrary = fullMasterLibrary - if globalParameters["SeparateArchitectures"] and len(masterLibraries) > 0: - theMasterLibrary = list(masterLibraries.values())[0] + LibraryIO.write(masterFile, Utils.state(fullMasterLibrary), libraryFormat) print1("# Tensile Library Writer DONE") print1(HR) diff --git a/tensilelite/Tensile/bin/TensileCreateLibrary b/tensilelite/Tensile/bin/TensileCreateLibrary index f2891c73d4..c8bd2e10b0 100755 --- a/tensilelite/Tensile/bin/TensileCreateLibrary +++ b/tensilelite/Tensile/bin/TensileCreateLibrary @@ -27,17 +27,18 @@ # This script only gets called by CMake try: - from Tensile.TensileCreateLibrary import TensileCreateLibrary + from Tensile import TensileCreateLibrary except ImportError: import os.path import sys parentdir = os.path.normpath(os.path.join(os.path.dirname(os.path.realpath(__file__)), "..", "..")) + print(parentdir) sys.path.append(parentdir) - from Tensile.TensileCreateLibrary import TensileCreateLibrary + from Tensile import TensileCreateLibrary ################################################################################ # Main ################################################################################ if __name__ == "__main__": - TensileCreateLibrary() + TensileCreateLibrary.run() From a5d47577556e9f5432f66b32e9f9cdb876d20774 Mon Sep 17 00:00:00 2001 From: root Date: Mon, 6 Jan 2025 00:43:35 +0000 Subject: [PATCH 02/25] Initialize missing variables --- tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py | 5 ++++- tensilelite/Tensile/TensileCreateLibrary/main.py | 6 ++++-- 2 files changed, 8 insertions(+), 3 deletions(-) diff --git a/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py b/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py index f8b8a272f4..2f8d252cbe 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py +++ b/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py @@ -130,6 +130,9 @@ def parseArguments(input: Optional[List[str]] = None) -> Dict[str, Any]: arguments["OffloadBundler"] = args.OffloadBundler arguments["Assembler"] = args.Assembler arguments["LogicPath"] = args.LogicPath - arguments["OutputPath"] = args.OutputPath + arguments["LogicFilter"] = args.LogicFilter + arguments["OutputPath"] = args.OutputPath + arguments["Experimental"] = args.Experimental + arguments["GenSolTable"] = args.GenSolTable return arguments diff --git a/tensilelite/Tensile/TensileCreateLibrary/main.py b/tensilelite/Tensile/TensileCreateLibrary/main.py index 4dd940cc8a..38799038fc 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/main.py +++ b/tensilelite/Tensile/TensileCreateLibrary/main.py @@ -350,11 +350,11 @@ def libraryIter(lib: MasterSolutionLibrary): nextSolIndex = masterLibraries[architectureName].merge(newLibrary, nextSolIndex) else: masterLibraries[architectureName] = newLibrary - masterLibraries[architectureName].version = args["version"] + masterLibraries[architectureName].version = args["CodeObjectVersion"] else: if fullMasterLibrary is None: fullMasterLibrary = newLibrary - fullMasterLibrary.version = args["version"] + fullMasterLibrary.version = args["CodeObjectVersion"] else: fullMasterLibrary.merge(newLibrary) @@ -429,7 +429,9 @@ def run(): libraryFormat = arguments["LibraryFormat"] useCompression = arguments["UseCompression"] + ensurePath(outputPath) outputPath = os.path.abspath(outputPath) + cxxCompiler, cCompiler, offloadBundler, assembler, hipconfig = validateToolchain( arguments["CxxCompiler"], arguments["CCompiler"], arguments["OffloadBundler"], arguments["Assembler"], ToolchainDefaults.HIP_CONFIG ) From 2a48b5ce8954f20546cdbabf1d121d0945d05d09 Mon Sep 17 00:00:00 2001 From: David Dixon Date: Sat, 11 Jan 2025 15:33:50 +0000 Subject: [PATCH 03/25] Add exported functions to module --- tensilelite/Tensile/TensileCreateLibrary/__init__.py | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/tensilelite/Tensile/TensileCreateLibrary/__init__.py b/tensilelite/Tensile/TensileCreateLibrary/__init__.py index 68f8e21c64..8f22e1ee66 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/__init__.py +++ b/tensilelite/Tensile/TensileCreateLibrary/__init__.py @@ -1 +1,3 @@ -from .main import run \ No newline at end of file +from .main import run +from .main import copyStaticFiles +from .main import writeSolutionsAndKernels \ No newline at end of file From 0c055162aeed9a06775700dedaa891742850a8bb Mon Sep 17 00:00:00 2001 From: David Dixon Date: Sat, 11 Jan 2025 17:03:09 +0000 Subject: [PATCH 04/25] Add reasonable defaults and clean up --- tensilelite/Tensile/ClientWriter.py | 14 +- .../TensileCreateLibrary/ParseArguments.py | 15 +- .../Tensile/TensileCreateLibrary/main.py | 168 +++++------------- 3 files changed, 54 insertions(+), 143 deletions(-) diff --git a/tensilelite/Tensile/ClientWriter.py b/tensilelite/Tensile/ClientWriter.py index 4b9f7c1ccc..be16b5f6b3 100644 --- a/tensilelite/Tensile/ClientWriter.py +++ b/tensilelite/Tensile/ClientWriter.py @@ -215,21 +215,11 @@ def getBuildClientLibraryScript(buildPath, libraryLogicPath, cxxCompiler): callCreateLibraryCmd = globalParameters["ScriptPath"] + "/bin/TensileCreateLibrary" - if globalParameters["SeparateArchitectures"]: - callCreateLibraryCmd += " --separate-architectures" - - if globalParameters["LazyLibraryLoading"]: - callCreateLibraryCmd += " --lazy-library-loading" + if not globalParameters["LazyLibraryLoading"]: + callCreateLibraryCmd += " --no-lazy-library-loading" if globalParameters["ShortNames"]: callCreateLibraryCmd += " --short-file-names" - else: - callCreateLibraryCmd += " --no-short-file-names" - - if globalParameters["LibraryPrintDebug"]: - callCreateLibraryCmd += " --library-print-debug" - else: - callCreateLibraryCmd += " --no-library-print-debug" if globalParameters.get("AsmDebug", False): callCreateLibraryCmd += " --asm-debug" diff --git a/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py b/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py index a3e42a62d6..a44df16a35 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py +++ b/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py @@ -57,10 +57,7 @@ def parseArguments(input: Optional[List[str]] = None) -> Dict[str, Any]: argParser.add_argument("--assembler", dest="Assembler", action="store", default=ToolchainDefaults.ASSEMBLER) argParser.add_argument("--code-object-version", dest="CodeObjectVersion", choices=["4", "5"], default="4", action="store") argParser.add_argument("--architecture", dest="Architecture", type=str, action="store", default="all", help="Supported archs: " + " ".join(architectureMap.keys())) - argParser.add_argument("--short-file-names", dest="ShortNames", action="store_true") - argParser.add_argument("--no-short-file-names", dest="ShortNames", action="store_false") - argParser.add_argument("--library-print-debug", dest="LibraryPrintDebug", action="store_true") - argParser.add_argument("--no-library-print-debug", dest="LibraryPrintDebug", action="store_false") + argParser.add_argument("--short-file-names", dest="ShortNames", action="store_true", default=False) argParser.add_argument("--no-compress", dest="NoCompress", action="store_true", help="Don't compress assembly code objects.") argParser.add_argument("--experimental", dest="Experimental", action="store_true", help="Include logic files in directories named 'Experimental'.") @@ -78,10 +75,8 @@ def parseArguments(input: Optional[List[str]] = None) -> Dict[str, Any]: default=1, help="Set printout verbosity level.") argParser.add_argument("--print-timing", dest="PrintTiming", default=False, action="store_true", help="Print duration of each stage.") - argParser.add_argument("--separate-architectures", dest="SeparateArchitectures", action="store_true", - default=False, help="Separates TensileLibrary file by architecture") - argParser.add_argument("--lazy-library-loading", dest="LazyLibraryLoading", action="store_true", - default=False, help="Loads Tensile libraries when needed instead of upfront.") + argParser.add_argument("--no-lazy-library-loading", dest="LazyLibraryLoading", action="store_false", + default=True, help="Loads Tensile libraries when needed instead of upfront.") argParser.add_argument("--enable-marker", dest="EnableMarker", action="store_true", default=False, help="Enable marker in Tensile.") argParser.add_argument("--no-generate-solution-table", dest="GenSolTable", action="store_false", default=True, @@ -93,7 +88,6 @@ def parseArguments(input: Optional[List[str]] = None) -> Dict[str, Any]: default=False, help="Enable ASAN build.") argParser.add_argument("--keep-build-tmp", dest="KeepBuildTmp", action="store_true", default=False, help="Do not remove the temporary build directory (may required hundreds of GBs of space)"), - argParser.add_argument("--validate-library", dest="ValidateLibrary", action="store_true", default=False) argParser.add_argument("--logic-filter", dest="LogicFilter", action="store", default="*", type=str, help="Cutomsized logic filter, default is *, i.e. all logics." " Example: gfx942/Equality/* for building equality of gfx942 only") @@ -104,13 +98,11 @@ def parseArguments(input: Optional[List[str]] = None) -> Dict[str, Any]: arguments["RuntimeLanguage"] = args.RuntimeLanguage arguments["CodeObjectVersion"] = args.CodeObjectVersion arguments["Architecture"] = args.Architecture - arguments["SeparateArchitectures"] = args.SeparateArchitectures arguments["LazyLibraryLoading"] = args.LazyLibraryLoading arguments["EnableMarker"] = args.EnableMarker if args.CmakeCxxCompiler: os.environ["CMAKE_CXX_COMPILER"] = args.CmakeCxxCompiler arguments["ShortNames"] = args.ShortNames - arguments["LibraryPrintDebug"] = args.LibraryPrintDebug arguments["CodeFromFiles"] = False arguments["LogicFormat"] = args.LogicFormat arguments["LibraryFormat"] = args.LibraryFormat @@ -123,7 +115,6 @@ def parseArguments(input: Optional[List[str]] = None) -> Dict[str, Any]: arguments["BuildIdKind"] = args.BuildIdKind arguments["KeepBuildTmp"] = args.KeepBuildTmp arguments["AsanBuild"] = args.AsanBuild - arguments["ValidateLibrary"] = args.ValidateLibrary arguments["UseCompression"] = not args.NoCompress arguments["CxxCompiler"] = args.CxxCompiler arguments["CCompiler"] = args.CCompiler diff --git a/tensilelite/Tensile/TensileCreateLibrary/main.py b/tensilelite/Tensile/TensileCreateLibrary/main.py index 3953f407c3..f9a81a5828 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/main.py +++ b/tensilelite/Tensile/TensileCreateLibrary/main.py @@ -33,15 +33,13 @@ from typing import NamedTuple, List, Optional, Sequence, Union from Tensile import Utils -from Tensile import Common from Tensile.Toolchain.Assembly import AssemblyToolchain, buildAssemblyCodeObjectFiles from Tensile.Toolchain.Source import SourceToolchain, buildSourceCodeObjectFile from Tensile.Toolchain.Validators import validateToolchain, getVersion, ToolchainDefaults from Tensile.TensileInstructions import getGfxName, TensileInstructions from Tensile.Common import globalParameters, HR, print1, print2, printExit, ensurePath, \ - CHeader, assignGlobalParameters, \ - architectureMap, printWarning, \ - IsaVersion + CHeader, assignGlobalParameters, architectureMap, IsaVersion, pushWorkingPath, \ + popWorkingPath, ParallelMap2 from Tensile.KernelWriterAssembly import KernelWriterAssembly from Tensile.KernelWriterBase import KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H from Tensile import LibraryIO @@ -81,12 +79,10 @@ def processKernelSource(kernelWriterAssembly, ti, kernel) -> KernelCodeGenResult Returns (error, source, header, kernelName). """ kernelWriter = kernelWriterAssembly - # get kernel name kernelWriter.setTensileInstructions(ti) asmFilename = kernelWriter.getKernelFileBase(kernel) err, src = kernelWriter.getSourceFileString(kernel) header = kernelWriter.getHeaderFileString(kernel) - # will be put in Kernels.h/cpp if None objFilename = kernel._state.get("codeObjectFile", None) return KernelCodeGenResult(err, src, header, asmFilename, objFilename, tuple(kernel["ISA"]), kernel["WavefrontSize"]) @@ -129,6 +125,7 @@ def removeInvalidSolutionsAndKernels(results, kernels, solutions, errorTolerant, for rel in removeResults: results.remove(rel) + 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}") @@ -156,18 +153,14 @@ def writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNE 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) @@ -175,20 +168,12 @@ def writeSolutionsAndKernels(outputPath, asmToolchain, srcToolchain, solutions, kernelWriterAssembly, errorTolerant=False, compress=True): codeObjectFiles = [] - # Push working path into build_tmp folder because there may be more than - # one process running this script. This is to avoid build directory clashing. - # NOTE: file paths must not contain the lower case word 'kernel' or the - # /opt/rocm/bin/extractkernel will fail. - # See buildSourceCodeObjectFile:167 for the call to this binary. - Common.pushWorkingPath('build_tmp') - Common.pushWorkingPath(os.path.basename(outputPath).upper()) + pushWorkingPath('build_tmp') + pushWorkingPath(os.path.basename(outputPath).upper()) asmPath = ensurePath(os.path.join(globalParameters["WorkingPath"], "assembly")) asmKernels = [k for k in kernels if k['KernelLanguage'] == 'Assembly'] - # Kernels may be intended for different co files, but generate the same .o file - # Mark duplicate kernels to avoid race condition - # @TODO improve organization so this problem doesn't appear visited = set() duplicates = 0 for k in asmKernels: @@ -197,32 +182,28 @@ def writeSolutionsAndKernels(outputPath, asmToolchain, srcToolchain, solutions, duplicates += k.duplicate print2(f"Duplicate: {base}") visited.add(base) - print1(f"Number of duplicates: {duplicates}") numAsmKernels = len(asmKernels) numKernels = len(asmKernels) assert numKernels == numAsmKernels, "Only assembly kernels are supported in TensileLite" asmIter = zip(itertools.repeat(kernelWriterAssembly), itertools.repeat(TensileInstructions()), asmKernels) - asmResults = Common.ParallelMap2(processKernelSource, asmIter, "Generating assembly kernels") + asmResults = ParallelMap2(processKernelSource, asmIter, "Generating assembly kernels") removeInvalidSolutionsAndKernels(asmResults, asmKernels, solutions, errorTolerant, globalParameters) def assemble(ret): p, isa, wavefrontsize = ret asmToolchain.assemble(str(p), str(p.with_suffix(".o")), getGfxName(isa), wavefrontsize) unaryWriteAssembly = functools.partial(writeAssembly, asmPath) compose = lambda *F: functools.reduce(lambda f, g: lambda x: f(g(x)), F) - ret = Common.ParallelMap2(compose(assemble, unaryWriteAssembly), asmResults, "Writing assembly kernels", return_as="list", multiArg=False) + ret = ParallelMap2(compose(assemble, unaryWriteAssembly), asmResults, "Writing assembly kernels", return_as="list", multiArg=False) codeObjectFiles += buildAssemblyCodeObjectFiles(asmToolchain, asmKernels, kernelWriterAssembly, outputPath, compress) - srcKernels = [k for k in kernels if k['KernelLanguage'] != 'Assembly'] - if srcKernels: - raise ValueError(f"Non-helper HIP source kernels are not supported Tensilelite, found {len(srcKernels)}") writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H) srcKernelFile = Path(outputPath) / "Kernels.cpp" buildSourceCodeObjectFile(srcToolchain, outputPath, srcKernelFile) - Common.popWorkingPath() # build_tmp - Common.popWorkingPath() # workingDir + popWorkingPath() # build_tmp + popWorkingPath() # workingDir return codeObjectFiles, numKernels @@ -230,8 +211,8 @@ def assemble(ret): def writeSolutionsAndKernelsTCL(outputPath, asmToolchain, srcToolchain, kernels, kernelHelperObjs, \ kernelWriterAssembly, compress=True): - Common.pushWorkingPath('build_tmp') - Common.pushWorkingPath(os.path.basename(outputPath).upper()) + pushWorkingPath('build_tmp') + pushWorkingPath(os.path.basename(outputPath).upper()) asmPath = ensurePath(os.path.join(globalParameters["WorkingPath"], "assembly")) asmKernels = [k for k in kernels if k['KernelLanguage'] == 'Assembly'] @@ -244,30 +225,27 @@ def writeSolutionsAndKernelsTCL(outputPath, asmToolchain, srcToolchain, kernels, duplicates += k.duplicate print2(f"Duplicate: {base}") visited.add(base) - print1(f"Number of duplicates: {duplicates}") uniqueAsmKernels = [k for k in asmKernels if not k.duplicate] - numAsmKernels = len(asmKernels) numKernels = len(kernels) assert numKernels == numAsmKernels, "Only assembly kernels are supported in TensileLite" - def assemble(ret): p, isa, wavefrontsize = ret asmToolchain.assemble(str(p), str(p.with_suffix(".o")), getGfxName(isa), wavefrontsize) unaryProcessKernelSource = functools.partial(processKernelSource, kernelWriterAssembly, TensileInstructions()) unaryWriteAssembly = functools.partial(writeAssembly, asmPath) compose = lambda *F: functools.reduce(lambda f, g: lambda x: f(g(x)), F) - ret = Common.ParallelMap2(compose(assemble, unaryWriteAssembly, unaryProcessKernelSource), uniqueAsmKernels, "Generating assembly kernels", multiArg=False) + ret = ParallelMap2(compose(assemble, unaryWriteAssembly, unaryProcessKernelSource), uniqueAsmKernels, "Generating assembly kernels", multiArg=False) buildAssemblyCodeObjectFiles(asmToolchain, asmKernels, kernelWriterAssembly, outputPath, compress) writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H) srcKernelFile = Path(outputPath) / "Kernels.cpp" buildSourceCodeObjectFile(srcToolchain, outputPath, srcKernelFile) - Common.popWorkingPath() # build_tmp - Common.popWorkingPath() # workingDir + popWorkingPath() # build_tmp + popWorkingPath() # workingDir return numKernels @@ -306,7 +284,6 @@ def copyStaticFiles(outputPath=None): @timing def generateKernelObjectsFromSolutions(solutions): - # create solution writer and kernel writer kernels = [] kernelHelperObjs = [] kernelNames = set() @@ -339,7 +316,6 @@ def generateLogicDataAndSolutions(logicFiles, args, cxxCompiler): solutions = [] masterLibraries = {} - fullMasterLibrary = None nextSolIndex = 0 matchTable = {} fIter = zip(logicFiles, itertools.repeat(cxxCompiler), itertools.repeat(archs)) @@ -352,47 +328,35 @@ def libraryIter(lib: MasterSolutionLibrary): for _, lazyLib in lib.lazyLibraries.items(): yield from libraryIter(lazyLib) - for library in Common.ParallelMap2(LibraryIO.parseLibraryLogicFile, fIter, "Loading Logics...", return_as="generator_unordered"): + for library in ParallelMap2(LibraryIO.parseLibraryLogicFile, fIter, "Loading Logics...", return_as="generator_unordered"): _, architectureName, _, _, _, newLibrary, srcFile = library if architectureName == "": continue - if globalParameters["SeparateArchitectures"] or globalParameters["LazyLibraryLoading"]: - if architectureName in masterLibraries: - nextSolIndex = masterLibraries[architectureName].merge(newLibrary, nextSolIndex) - else: - masterLibraries[architectureName] = newLibrary - masterLibraries[architectureName].version = args["CodeObjectVersion"] + if architectureName in masterLibraries: + nextSolIndex = masterLibraries[architectureName].merge(newLibrary, nextSolIndex) else: - if fullMasterLibrary is None: - fullMasterLibrary = newLibrary - fullMasterLibrary.version = args["CodeObjectVersion"] - else: - fullMasterLibrary.merge(newLibrary) - + masterLibraries[architectureName] = newLibrary + masterLibraries[architectureName].version = args["CodeObjectVersion"] + if args["GenSolTable"]: # Match yaml file solutions to solution index for localIdx, _, s in libraryIter(newLibrary): matchTable[s.index] = [srcFile, localIdx] - if globalParameters["SeparateArchitectures"] or globalParameters["LazyLibraryLoading"]: - 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(): + 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) - for name, lib in masterLibrary.lazyLibraries.items(): - for _, sol in lib.solutions.items(): - sol.originalSolution._state["codeObjectFile"] = name - solutions.append(sol.originalSolution) - else: - solutions = [sol.originalSolution for _, sol in fullMasterLibrary.solutions.items()] # remove duplicates while preserving order solutions = dict.fromkeys(solutions).keys() @@ -400,25 +364,7 @@ def libraryIter(lib: MasterSolutionLibrary): if args["GenSolTable"]: LibraryIO.write("MatchTable", matchTable) - return solutions, masterLibraries, fullMasterLibrary - - -def validateLibrary(masterLibraries: MasterSolutionLibrary, - kernels: Sequence[Solution], - kernelWriterAssembly: KernelWriterAssembly): - kernelsByCodeObjectFiles = {k: list(g) for k, g in itertools.groupby(kernels, lambda k: k["codeObjectFile"])} - - ok: bool = True - - for _, lib in masterLibraries.items(): - for name, llib in lib.lazyLibraries.items(): - uniqueKernelsInLib = {kernelWriterAssembly.getKernelName(s.originalSolution) for s in llib.solutions.values()} - - if len(uniqueKernelsInLib) != len(kernelsByCodeObjectFiles[name]): - ok = False - print(f"{name} library and co has inconsistent kernel size {len(uniqueKernelsInLib)} vs {len(kernelsByCodeObjectFiles[name])}") - - assert ok and "Inconsistent kernel sizes detected!" + return solutions, masterLibraries ################################################################################ @@ -434,18 +380,9 @@ def run(): print2("") arguments = parseArguments() - logicPath = arguments["LogicPath"] - outputPath = arguments["OutputPath"] - cxxCompiler = arguments["CxxCompiler"] - offloadBundler = arguments["OffloadBundler"] - assembler = arguments["Assembler"] - libraryFormat = arguments["LibraryFormat"] - useCompression = arguments["UseCompression"] - coVersion = arguments["CodeObjectVersion"] - - ensurePath(outputPath) - outputPath = os.path.abspath(outputPath) - + ensurePath(arguments["OutputPath"]) + arguments["OutputPath"] = os.path.abspath(arguments["OutputPath"]) + cxxCompiler, cCompiler, offloadBundler, assembler, hipconfig = validateToolchain( arguments["CxxCompiler"], arguments["CCompiler"], arguments["OffloadBundler"], arguments["Assembler"], ToolchainDefaults.HIP_CONFIG ) @@ -455,17 +392,17 @@ def run(): print1(f"# C Compiler: {cCompiler} (version {getVersion(cCompiler)})") print1(f"# Assembler: {assembler} (version {getVersion(assembler)})") print1(f"# Offload Bundler: {offloadBundler} (version {getVersion(offloadBundler)})") - print1(f"# Code Object Version: {coVersion}") + print1(f"# Code Object Version: {arguments['CodeObjectVersion']}") print1(f"# Architecture(s): {arguments['Architecture']}") - print1(f"# Library Format: {libraryFormat}") + print1(f"# Library Format: {arguments['LibraryFormat']}") assignGlobalParameters(arguments, cxxCompiler) - asmToolchain = AssemblyToolchain(assembler, offloadBundler, globalParameters["BuildIdKind"], coVersion) + asmToolchain = AssemblyToolchain(assembler, offloadBundler, globalParameters["BuildIdKind"], arguments["CodeObjectVersion"]) srcToolchain = SourceToolchain(cxxCompiler, offloadBundler, globalParameters["BuildIdKind"], globalParameters["AsanBuild"], globalParameters["SaveTemps"]) - if not os.path.exists(logicPath): - printExit("LogicPath %s doesn't exist" % logicPath) + if not os.path.exists(arguments["LogicPath"]): + printExit(f"LogicPath {arguments['LogicPath']} doesn't exist") if ";" in arguments["Architecture"]: archs = arguments["Architecture"].split(";") @@ -492,9 +429,9 @@ def archMatch(arch: str, archs: List[str]): def validLogicFile(p: Path): return p.suffix == logicExtFormat and ("all" in archs or archMatch(load_logic_gfx_arch(p), archs)) - globPattern = os.path.join(logicPath, f"**/{arguments['LogicFilter']}{logicExtFormat}") + globPattern = os.path.join(arguments["LogicPath"], f"**/{arguments['LogicFilter']}{logicExtFormat}") print1(f"# LogicFilter: {globPattern}") - logicFiles = (os.path.join(logicPath, file) for file in glob.iglob(globPattern, recursive=True)) + 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']}") @@ -505,25 +442,18 @@ def validLogicFile(p: Path): for logicFile in logicFiles: print2("# %s" % logicFile) - solutions, masterLibraries, fullMasterLibrary = generateLogicDataAndSolutions(logicFiles, arguments, cxxCompiler) + solutions, masterLibraries = generateLogicDataAndSolutions(logicFiles, arguments, cxxCompiler) kernels, kernelHelperObjs, _ = generateKernelObjectsFromSolutions(solutions) kernelWriterAssembly, kernelMinNaming, _ = getSolutionAndKernelWriters(solutions, kernels, assembler) - if globalParameters["ValidateLibrary"]: - validateLibrary(masterLibraries, kernels, kernelWriterAssembly) - - staticFiles = copyStaticFiles(outputPath) - - for fileName in staticFiles: - shutil.copy( os.path.join(globalParameters["SourcePath"], fileName), \ - outputPath ) + copyStaticFiles(arguments["OutputPath"]) - numKernels = writeSolutionsAndKernelsTCL(outputPath, asmToolchain, srcToolchain, kernels, - kernelHelperObjs, kernelWriterAssembly, compress=useCompression) + numKernels = writeSolutionsAndKernelsTCL(arguments["OutputPath"], asmToolchain, srcToolchain, kernels, + kernelHelperObjs, kernelWriterAssembly, compress=arguments["UseCompression"]) archs = [getGfxName(arch) for arch in globalParameters['SupportedISA'] \ if globalParameters["AsmCaps"][arch]["SupportedISA"]] - newLibraryDir = ensurePath(os.path.join(outputPath, 'library')) + newLibraryDir = ensurePath(os.path.join(arguments["OutputPath"], 'library')) for archName, newMasterLibrary in masterLibraries.items(): if archName in archs: @@ -532,11 +462,11 @@ def validLogicFile(p: Path): else: masterFile = os.path.join(newLibraryDir, "TensileLibrary_"+archName) newMasterLibrary.applyNaming(kernelMinNaming) - LibraryIO.write(masterFile, Utils.state(newMasterLibrary), libraryFormat) + LibraryIO.write(masterFile, Utils.state(newMasterLibrary), arguments["LibraryFormat"]) for name, lib in newMasterLibrary.lazyLibraries.items(): filename = os.path.join(newLibraryDir, name) lib.applyNaming(kernelMinNaming) - LibraryIO.write(filename, Utils.state(lib), libraryFormat) + LibraryIO.write(filename, Utils.state(lib), arguments["LibraryFormat"]) print1("# Tensile Library Writer DONE") print1(HR) From dfb684170b1365e4398d3520945a0470daf04065 Mon Sep 17 00:00:00 2001 From: David Dixon Date: Sat, 11 Jan 2025 17:38:39 +0000 Subject: [PATCH 05/25] Can't completely remove generate sources and exit --- tensilelite/Tensile/BenchmarkProblems.py | 3 ++- .../Tensile/TensileCreateLibrary/ParseArguments.py | 2 -- .../Tensile/TensileCreateLibrary/{main.py => Run.py} | 8 +++++--- tensilelite/Tensile/TensileCreateLibrary/__init__.py | 6 +++--- 4 files changed, 10 insertions(+), 9 deletions(-) rename tensilelite/Tensile/TensileCreateLibrary/{main.py => Run.py} (98%) diff --git a/tensilelite/Tensile/BenchmarkProblems.py b/tensilelite/Tensile/BenchmarkProblems.py index 24c89e23d0..1cc2139057 100644 --- a/tensilelite/Tensile/BenchmarkProblems.py +++ b/tensilelite/Tensile/BenchmarkProblems.py @@ -148,7 +148,8 @@ def writeBenchmarkFiles(stepBaseDir, solutions, problemSizes, \ codeObjectFiles, _= writeSolutionsAndKernels( \ globalParameters["WorkingPath"], asmToolchain, srcToolchain, \ solutions, kernels, kernelHelperOjbs, \ - kernelWriterAssembly, errorTolerant=True ) + kernelWriterAssembly, errorTolerant=True, \ + generateSourcesAndExit=globalParameters["GenerateSourcesAndExit"] ) # ^ this is where solutions is mutated newLibraryDir = ensurePath(os.path.join(globalParameters["WorkingPath"], 'library')) diff --git a/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py b/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py index a44df16a35..5377d5b209 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py +++ b/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py @@ -67,8 +67,6 @@ def parseArguments(input: Optional[List[str]] = None) -> Dict[str, Any]: action="store", default="yaml", help="select which logic format to use") argParser.add_argument("--library-format", dest="LibraryFormat", choices=["yaml", "msgpack"], action="store", default="msgpack", help="select which library format to use") - argParser.add_argument("--generate-sources-and-exit", dest="GenerateSourcesAndExit", action="store_true", - default=False, help="Output source files only and exit.") argParser.add_argument("--jobs", "-j", dest="CpuThreads", type=int, default=-1, help="Number of parallel jobs to launch.") argParser.add_argument("--verbose", "-v", dest="PrintLevel", type=int, diff --git a/tensilelite/Tensile/TensileCreateLibrary/main.py b/tensilelite/Tensile/TensileCreateLibrary/Run.py similarity index 98% rename from tensilelite/Tensile/TensileCreateLibrary/main.py rename to tensilelite/Tensile/TensileCreateLibrary/Run.py index f9a81a5828..eb43e66768 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/main.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Run.py @@ -165,7 +165,7 @@ def writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNE def writeSolutionsAndKernels(outputPath, asmToolchain, srcToolchain, solutions, kernels, kernelHelperObjs, \ - kernelWriterAssembly, errorTolerant=False, compress=True): + kernelWriterAssembly, errorTolerant=False, generateSourcesAndExit=False, compress=True): codeObjectFiles = [] pushWorkingPath('build_tmp') @@ -196,11 +196,13 @@ def assemble(ret): unaryWriteAssembly = functools.partial(writeAssembly, asmPath) 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) - codeObjectFiles += buildAssemblyCodeObjectFiles(asmToolchain, asmKernels, kernelWriterAssembly, outputPath, compress) writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H) srcKernelFile = Path(outputPath) / "Kernels.cpp" - buildSourceCodeObjectFile(srcToolchain, outputPath, srcKernelFile) + + if not generateSourcesAndExit: + codeObjectFiles += buildAssemblyCodeObjectFiles(asmToolchain, asmKernels, kernelWriterAssembly, outputPath, compress) + buildSourceCodeObjectFile(srcToolchain, outputPath, srcKernelFile) popWorkingPath() # build_tmp popWorkingPath() # workingDir diff --git a/tensilelite/Tensile/TensileCreateLibrary/__init__.py b/tensilelite/Tensile/TensileCreateLibrary/__init__.py index 8f22e1ee66..16a8e3dc98 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/__init__.py +++ b/tensilelite/Tensile/TensileCreateLibrary/__init__.py @@ -1,3 +1,3 @@ -from .main import run -from .main import copyStaticFiles -from .main import writeSolutionsAndKernels \ No newline at end of file +from .Run import run +from .Run import copyStaticFiles +from .Run import writeSolutionsAndKernels \ No newline at end of file From 209ef14a0224373f2d69ed39850ecd8aee9b8c42 Mon Sep 17 00:00:00 2001 From: David Dixon Date: Sat, 11 Jan 2025 21:24:40 +0000 Subject: [PATCH 06/25] Update cmake support --- install.sh | 24 ++++--------------- .../amd_detail/rocblaslt/src/CMakeLists.txt | 7 ++---- tensilelite/Tensile/cmake/TensileConfig.cmake | 18 +++----------- 3 files changed, 10 insertions(+), 39 deletions(-) diff --git a/install.sh b/install.sh index d542c5a112..3d8eec9cb5 100755 --- a/install.sh +++ b/install.sh @@ -396,8 +396,7 @@ tensile_logic= tensile_cov="4" tensile_threads=$(nproc) tensile_fork= -tensile_lazy_library_loading=true -tensile_separate_architectures=true +tensile_no_lazy_library_loading=false tensile_tag= tensile_test_local_path= tensile_version= @@ -425,7 +424,7 @@ fi # check if we have a modern version of getopt that can handle whitespace and long parameters getopt -T if [[ $? -eq 4 ]]; then - GETOPT_PARSE=$(getopt --name "${0}" --longoptions help,install,clients,dependencies,debug,hip-clang,static,relocatable,codecoverage,relwithdebinfo,address-sanitizer,separate-architectures,lazy-library-loading,no-separate-architectures,no-lazy-library-loading,no_tensile,no-tensile,msgpack,no-msgpack,logic:,cov:,fork:,branch:,test_local_path:,cpu_ref_lib:,build_dir:,use-custom-version:,architecture:,gprof,keep-build-tmp,no-compress,experimental,legacy_hipblas_direct,disable-hipblaslt-marker,enable-tensile-marker,logic-yaml-filter: --options hicdgrka:j:o:l:f:b:nu:t: -- "$@") + GETOPT_PARSE=$(getopt --name "${0}" --longoptions help,install,clients,dependencies,debug,hip-clang,static,relocatable,codecoverage,relwithdebinfo,address-sanitizer,no-lazy-library-loading,no_tensile,no-tensile,msgpack,no-msgpack,logic:,cov:,fork:,branch:,test_local_path:,cpu_ref_lib:,build_dir:,use-custom-version:,architecture:,gprof,keep-build-tmp,no-compress,experimental,legacy_hipblas_direct,disable-hipblaslt-marker,enable-tensile-marker,logic-yaml-filter: --options hicdgrka:j:o:l:f:b:nu:t: -- "$@") else echo "Need a new version of getopt" exit 1 @@ -505,17 +504,8 @@ while true; do -n|--no_tensile|--no-tensile) build_tensile=false shift ;; - --lazy-library-loading) - tensile_lazy_library_loading=true - shift ;; --no-lazy-library-loading) - tensile_lazy_library_loading=false - shift ;; - --separate-architectures) - tensile_separate_architectures=true - shift ;; - --no-separate-architectures) - tensile_separate_architectures=false + tensile_no_lazy_library_loading=false shift ;; -u|--use-custom-version) tensile_version=${2} @@ -773,12 +763,8 @@ pushd . fi fi - if [[ "${tensile_lazy_library_loading}" == false ]]; then - tensile_opt="${tensile_opt} -DTensile_LAZY_LIBRARY_LOADING=OFF" - fi - - if [[ "${tensile_separate_architectures}" == false ]]; then - tensile_opt="${tensile_opt} -DTensile_SEPARATE_ARCHITECTURES=OFF" + if [[ "${tensile_no_lazy_library_loading}" == true ]]; then + tensile_opt="${tensile_opt} -DTensile_NO_LAZY_LIBRARY_LOADING=OFF" fi if [[ "${tensile_msgpack_backend}" == true ]]; then diff --git a/library/src/amd_detail/rocblaslt/src/CMakeLists.txt b/library/src/amd_detail/rocblaslt/src/CMakeLists.txt index 791a7e3a28..e45023bd3f 100644 --- a/library/src/amd_detail/rocblaslt/src/CMakeLists.txt +++ b/library/src/amd_detail/rocblaslt/src/CMakeLists.txt @@ -48,11 +48,8 @@ if( BUILD_WITH_TENSILE ) if(PACKAGE_TENSILE_LIBRARY) set(Tensile_Options ${Tensile_Options} GENERATE_PACKAGE) endif() - if(Tensile_SEPARATE_ARCHITECTURES) - set(Tensile_Options ${Tensile_Options} SEPARATE_ARCHITECTURES) - endif() - if(Tensile_LAZY_LIBRARY_LOADING) - set(Tensile_Options ${Tensile_Options} LAZY_LIBRARY_LOADING) + if(Tensile_NO_LAZY_LIBRARY_LOADING) + set(Tensile_Options ${Tensile_Options} NO_LAZY_LIBRARY_LOADING) endif() if(Tensile_ASAN_BUILD) set(Tensile_Options ${Tensile_Options} ASAN_BUILD) diff --git a/tensilelite/Tensile/cmake/TensileConfig.cmake b/tensilelite/Tensile/cmake/TensileConfig.cmake index 1455e023d0..afddec553f 100644 --- a/tensilelite/Tensile/cmake/TensileConfig.cmake +++ b/tensilelite/Tensile/cmake/TensileConfig.cmake @@ -84,7 +84,7 @@ function(TensileCreateLibraryFiles PRINT_DEBUG GENERATE_PACKAGE SEPARATE_ARCHITECTURES - LAZY_LIBRARY_LOADING + NO_LAZY_LIBRARY_LOADING ASAN_BUILD KEEP_BUILD_TMP NO_COMPRESS @@ -128,12 +128,8 @@ function(TensileCreateLibraryFiles message(STATUS "Tensile script: ${Script}") - if(Tensile_SEPARATE_ARCHITECTURES) - set(Options ${Options} "--separate-architectures") - endif() - - if(Tensile_LAZY_LIBRARY_LOADING) - set(Options ${Options} "--lazy-library-loading") + if(Tensile_NO_LAZY_LIBRARY_LOADING) + set(Options ${Options} "---nolazy-library-loading") endif() if(Tensile_ENABLE_MARKER) @@ -162,14 +158,6 @@ function(TensileCreateLibraryFiles if(Tensile_SHORT_FILE_NAMES) set(Options ${Options} "--short-file-names") - else() - set(Options ${Options} "--no-short-file-names") - endif() - - if(Tensile_PRINT_DEBUG) - set(Options ${Options} "--library-print-debug") - else() - set(Options ${Options} "--no-library-print-debug") endif() if(Tensile_EMBED_LIBRARY) From bdaa17026061faf718c000f298ea7410cc6f1f88 Mon Sep 17 00:00:00 2001 From: David Dixon Date: Tue, 14 Jan 2025 15:45:24 +0000 Subject: [PATCH 07/25] Update docs string --- tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py b/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py index 5377d5b209..3ecc3f4204 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py +++ b/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py @@ -74,7 +74,7 @@ def parseArguments(input: Optional[List[str]] = None) -> Dict[str, Any]: argParser.add_argument("--print-timing", dest="PrintTiming", default=False, action="store_true", help="Print duration of each stage.") argParser.add_argument("--no-lazy-library-loading", dest="LazyLibraryLoading", action="store_false", - default=True, help="Loads Tensile libraries when needed instead of upfront.") + default=True, help="Disable building for lazy library loading.") argParser.add_argument("--enable-marker", dest="EnableMarker", action="store_true", default=False, help="Enable marker in Tensile.") argParser.add_argument("--no-generate-solution-table", dest="GenSolTable", action="store_false", default=True, From 1480cb02d23a964c3a7e23a33dc3576a4f677e7d Mon Sep 17 00:00:00 2001 From: David Dixon Date: Tue, 14 Jan 2025 18:34:31 +0000 Subject: [PATCH 08/25] Refactor to simplify processing logic files --- tensilelite/Tensile/LibraryIO.py | 4 +- .../Tensile/TensileCreateLibrary/Run.py | 135 +++++++----------- 2 files changed, 56 insertions(+), 83 deletions(-) diff --git a/tensilelite/Tensile/LibraryIO.py b/tensilelite/Tensile/LibraryIO.py index fd5b5beafb..3920c037e1 100644 --- a/tensilelite/Tensile/LibraryIO.py +++ b/tensilelite/Tensile/LibraryIO.py @@ -307,7 +307,9 @@ def solutionStateToSolution(solutionState, cxxCompiler) -> Solution: solutions = [solutionStateToSolution(solutionState, cxxCompiler) for solutionState in data["Solutions"]] - newLibrary, _ = SolutionLibrary.MasterSolutionLibrary.FromOriginalState(data, solutions, cxxCompiler) + newLibrary, codeObjectFile = SolutionLibrary.MasterSolutionLibrary.FromOriginalState(data, solutions, cxxCompiler) + for solution in solutions: + solution["codeObjectFile"] = codeObjectFile return LibraryLogic(data["ScheduleName"], data["ArchitectureName"], problemType, solutions, \ data.get("ExactLogic"), newLibrary, srcFile) diff --git a/tensilelite/Tensile/TensileCreateLibrary/Run.py b/tensilelite/Tensile/TensileCreateLibrary/Run.py index eb43e66768..718d1e6680 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Run.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Run.py @@ -210,13 +210,15 @@ def assemble(ret): return codeObjectFiles, numKernels -def writeSolutionsAndKernelsTCL(outputPath, asmToolchain, srcToolchain, kernels, kernelHelperObjs, \ - kernelWriterAssembly, compress=True): +def writeSolutionsAndKernelsTCL(outputPath, asmToolchain, srcToolchain, solutions, assembler, compress=True): pushWorkingPath('build_tmp') pushWorkingPath(os.path.basename(outputPath).upper()) asmPath = ensurePath(os.path.join(globalParameters["WorkingPath"], "assembly")) + kernels, kernelHelperObjs, _ = generateKernelObjectsFromSolutions(solutions) + kernelWriterAssembly, _, _ = getSolutionAndKernelWriters(solutions, kernels, assembler) + asmKernels = [k for k in kernels if k['KernelLanguage'] == 'Assembly'] visited = set() @@ -309,64 +311,33 @@ def generateKernelObjectsFromSolutions(solutions): @timing -def generateLogicDataAndSolutions(logicFiles, args, cxxCompiler): - - 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 - matchTable = {} - fIter = zip(logicFiles, itertools.repeat(cxxCompiler), itertools.repeat(archs)) - - 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, srcFile = library - - if architectureName == "": - continue - - if architectureName in masterLibraries: - nextSolIndex = masterLibraries[architectureName].merge(newLibrary, nextSolIndex) - else: - masterLibraries[architectureName] = newLibrary - masterLibraries[architectureName].version = args["CodeObjectVersion"] - - if args["GenSolTable"]: +def generateSolutions(args, cxxCompiler, logicFile): + 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 = LibraryIO.parseLibraryLogicFile(logicFile, cxxCompiler, archs).solutions + numSoln = len(solutions) + solutions = dict.fromkeys(solutions).keys() + print1(f"Number of duplicate solutions: {numSoln - len(solutions)}") + return solutions + + +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(newLibrary): + for localIdx, _, s in libraryIter(library): matchTable[s.index] = [srcFile, localIdx] - - 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 - solutions = dict.fromkeys(solutions).keys() - - if args["GenSolTable"]: - LibraryIO.write("MatchTable", matchTable) - - return solutions, masterLibraries + LibraryIO.write("MatchTable", matchTable) ################################################################################ @@ -443,32 +414,32 @@ def validLogicFile(p: Path): print2(f"# LibraryLogicFiles: {len(logicFiles)}") for logicFile in logicFiles: print2("# %s" % logicFile) - - solutions, masterLibraries = generateLogicDataAndSolutions(logicFiles, arguments, cxxCompiler) - kernels, kernelHelperObjs, _ = generateKernelObjectsFromSolutions(solutions) - kernelWriterAssembly, kernelMinNaming, _ = getSolutionAndKernelWriters(solutions, kernels, assembler) - + copyStaticFiles(arguments["OutputPath"]) - numKernels = writeSolutionsAndKernelsTCL(arguments["OutputPath"], asmToolchain, srcToolchain, kernels, - kernelHelperObjs, kernelWriterAssembly, compress=arguments["UseCompression"]) - - archs = [getGfxName(arch) for arch in globalParameters['SupportedISA'] \ - if globalParameters["AsmCaps"][arch]["SupportedISA"]] - newLibraryDir = ensurePath(os.path.join(arguments["OutputPath"], 'library')) - - for archName, newMasterLibrary in masterLibraries.items(): - if archName in archs: - if globalParameters["LazyLibraryLoading"]: - masterFile = os.path.join(newLibraryDir, "TensileLibrary_lazy_"+archName) - else: - masterFile = os.path.join(newLibraryDir, "TensileLibrary_"+archName) - newMasterLibrary.applyNaming(kernelMinNaming) - LibraryIO.write(masterFile, Utils.state(newMasterLibrary), arguments["LibraryFormat"]) - for name, lib in newMasterLibrary.lazyLibraries.items(): - filename = os.path.join(newLibraryDir, name) - lib.applyNaming(kernelMinNaming) - LibraryIO.write(filename, Utils.state(lib), arguments["LibraryFormat"]) + unaryGenerateSolutions = functools.partial(generateSolutions, arguments, cxxCompiler) + solutions = ParallelMap2(unaryGenerateSolutions, logicFiles, "Reading logic files and generating solutions", multiArg=False) + solutions = (s for solutionList in solutions for s in solutionList) + + numKernels = writeSolutionsAndKernelsTCL(arguments["OutputPath"], asmToolchain, srcToolchain, solutions, + assembler, compress=arguments["UseCompression"]) + + # archs = [getGfxName(arch) for arch in globalParameters['SupportedISA'] \ + # if globalParameters["AsmCaps"][arch]["SupportedISA"]] + # newLibraryDir = ensurePath(os.path.join(arguments["OutputPath"], 'library')) + + # for archName, newMasterLibrary in masterLibraries.items(): + # if archName in archs: + # if globalParameters["LazyLibraryLoading"]: + # masterFile = os.path.join(newLibraryDir, "TensileLibrary_lazy_"+archName) + # else: + # masterFile = os.path.join(newLibraryDir, "TensileLibrary_"+archName) + # newMasterLibrary.applyNaming(kernelMinNaming) + # LibraryIO.write(masterFile, Utils.state(newMasterLibrary), arguments["LibraryFormat"]) + # for name, lib in newMasterLibrary.lazyLibraries.items(): + # filename = os.path.join(newLibraryDir, name) + # lib.applyNaming(kernelMinNaming) + # LibraryIO.write(filename, Utils.state(lib), arguments["LibraryFormat"]) print1("# Tensile Library Writer DONE") print1(HR) From 279a17e40424c9fdb2878b664ddf8d08e91cde4c Mon Sep 17 00:00:00 2001 From: David Dixon Date: Fri, 24 Jan 2025 18:35:15 +0000 Subject: [PATCH 09/25] memory consumption down to about 80GB for 64 procs --- tensilelite/Tensile/Common.py | 1 + tensilelite/Tensile/KernelWriterAssembly.py | 4 - tensilelite/Tensile/LibraryIO.py | 1 + tensilelite/Tensile/Parallel.py | 2 +- .../Tensile/TensileCreateLibrary/Run.py | 329 +++++++++--------- tensilelite/Tensile/Toolchain/Assembly.py | 6 +- 6 files changed, 169 insertions(+), 174 deletions(-) diff --git a/tensilelite/Tensile/Common.py b/tensilelite/Tensile/Common.py index 115eeffd8f..467d72c28c 100644 --- a/tensilelite/Tensile/Common.py +++ b/tensilelite/Tensile/Common.py @@ -429,6 +429,7 @@ def getArchitectureName(gfxName): validSMFMA["_format9"].append([MI[0],MI[1],MI[2],MI[3],2**bm,tt0,tt1,2**wave_m, 2**wave_n]) validSparseMatrixInstructions = validSMFMA["H"] + validSMFMA["B"] + validSMFMA["4xi8"] validMatrixInstructions = validMatrixInstructions + validSparseMatrixInstructions + validSMFMA["_format9"] +validMatrixInstructions = validMatrixInstructions + validSparseMatrixInstructions diff --git a/tensilelite/Tensile/KernelWriterAssembly.py b/tensilelite/Tensile/KernelWriterAssembly.py index 4f9188de54..db429f3236 100644 --- a/tensilelite/Tensile/KernelWriterAssembly.py +++ b/tensilelite/Tensile/KernelWriterAssembly.py @@ -73,10 +73,6 @@ def __init__(self, kernelMinNaming, kernelSerialNaming, assembler: str): def getSourceFileString(self, kernel) -> 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(kernel, globalParameters["CustomKernelDirectory"]) if isCustomKernelConfig(kernel) else self._getKernelSource(kernel) diff --git a/tensilelite/Tensile/LibraryIO.py b/tensilelite/Tensile/LibraryIO.py index 3920c037e1..57b0ef8b7c 100644 --- a/tensilelite/Tensile/LibraryIO.py +++ b/tensilelite/Tensile/LibraryIO.py @@ -294,6 +294,7 @@ def solutionStateToSolution(solutionState, cxxCompiler) -> Solution: # Therefore, we override the customKernel setting with the ActivationType value from ProblemType to avoid false alarms during subsequent problemType checks. solutionState["ProblemType"]["ActivationType"] = problemType["ActivationType"] solutionObject = Solution(solutionState, cxxCompiler) + solutionObject["LogicFileName"] = srcFile solutionProblemType = solutionObject["ProblemType"] if problemType != solutionProblemType: # find the mismatched items in ProblemType diff --git a/tensilelite/Tensile/Parallel.py b/tensilelite/Tensile/Parallel.py index 99950aba04..b0d44ef7b9 100644 --- a/tensilelite/Tensile/Parallel.py +++ b/tensilelite/Tensile/Parallel.py @@ -207,7 +207,7 @@ def ParallelMap2(function, objects, message="", enable=True, multiArg=True, retu pargs = zip(objects, itertools.repeat(globalParameters)) 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(verbose=1000, n_jobs=threadCount,timeout=99999, return_as=return_as)(delayed(pcall)(function, a, params) for a, params in pargs) else: rv = Parallel(n_jobs=threadCount,timeout=99999)(delayed(pcall)(function, a, params) for a, params in pargs) diff --git a/tensilelite/Tensile/TensileCreateLibrary/Run.py b/tensilelite/Tensile/TensileCreateLibrary/Run.py index 718d1e6680..8da668f0ff 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Run.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Run.py @@ -63,6 +63,26 @@ def wrapper(*args, **kwargs): return wrapper +def copyStaticFiles(outputPath=None): + if outputPath is None: + outputPath = globalParameters["WorkingPath"] + libraryStaticFiles = [ + "TensileTypes.h", + "tensile_bfloat16.h", + "tensile_float8_bfloat8.h", + "hip_f8_impl.h", + "KernelHeader.h", + "ReductionTemplate.h", + "memory_gfx.h" ] + + for fileName in libraryStaticFiles: + # copy file + shutil.copy( os.path.join(globalParameters["SourcePath"], fileName), \ + outputPath ) + + return libraryStaticFiles + + class KernelCodeGenResult(NamedTuple): err: int src: str @@ -129,14 +149,16 @@ def removeInvalidSolutionsAndKernels(results, kernels, solutions, errorTolerant, 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" + path = Path(asmPath) / str(os.getpid()) + path.mkdir(exist_ok=True) + filepath = path / f"{result.name}.s" isa = result.isa wfsize = result.wavefrontSize - with open(path, "w", encoding="utf-8") as f: + with open(filepath, "w", encoding="utf-8") as f: f.write(result.src) del result # result.src is very large so let gc know to clean up asap - return path, isa, wfsize + return filepath, isa, wfsize def writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H): @@ -164,6 +186,13 @@ def writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNE kernelHeaderFile.write(HeaderText) +def getKernelWriterAssembly(solutions, assembler): + kernelSerialNaming = Solution.getSerialNaming(solutions) + kernelMinNaming = Solution.getMinNaming(solutions) + kernelWriterAssembly = KernelWriterAssembly(kernelMinNaming, kernelSerialNaming, assembler) + return kernelWriterAssembly + + def writeSolutionsAndKernels(outputPath, asmToolchain, srcToolchain, solutions, kernels, kernelHelperObjs, \ kernelWriterAssembly, errorTolerant=False, generateSourcesAndExit=False, compress=True): codeObjectFiles = [] @@ -209,135 +238,125 @@ def assemble(ret): return codeObjectFiles, numKernels - +@timing def writeSolutionsAndKernelsTCL(outputPath, asmToolchain, srcToolchain, solutions, assembler, compress=True): - pushWorkingPath('build_tmp') pushWorkingPath(os.path.basename(outputPath).upper()) asmPath = ensurePath(os.path.join(globalParameters["WorkingPath"], "assembly")) - kernels, kernelHelperObjs, _ = generateKernelObjectsFromSolutions(solutions) - kernelWriterAssembly, _, _ = getSolutionAndKernelWriters(solutions, kernels, assembler) - - asmKernels = [k for k in kernels if k['KernelLanguage'] == 'Assembly'] - - visited = set() - duplicates = 0 - for k in asmKernels: - base = kernelWriterAssembly.getKernelFileBase(k) - k.duplicate = True if base in visited else False - duplicates += k.duplicate - print2(f"Duplicate: {base}") - visited.add(base) - print1(f"Number of duplicates: {duplicates}") - - uniqueAsmKernels = [k for k in asmKernels if not k.duplicate] - numAsmKernels = len(asmKernels) - numKernels = len(kernels) - assert numKernels == numAsmKernels, "Only assembly kernels are supported in TensileLite" - def assemble(ret): - p, isa, wavefrontsize = ret + asmKernels = [k.getKernels()[0] for k in solutions if k['KernelLanguage'] == 'Assembly'] + #kernelHelperObjs = generateKernelHelperObjects(asmKernels) + kernelWriterAssembly = getKernelWriterAssembly(asmKernels, assembler) + + uniqueAsmKernels = [k for k in asmKernels if "BuildKernel" in k] + unique = set() + for k in uniqueAsmKernels: + name = kernelWriterAssembly.getKernelFileBase(k) + if name not in unique: + unique.add(name) + else: + print1(f"{k['SolutionIndex']} {k['LogicFileName']}") + + pksResults = [processKernelSource(kernelWriterAssembly, TensileInstructions(), k) for k in uniqueAsmKernels] + for p, isa, wavefrontsize in [writeAssembly(asmPath, k) for k in pksResults]: asmToolchain.assemble(str(p), str(p.with_suffix(".o")), getGfxName(isa), wavefrontsize) - unaryProcessKernelSource = functools.partial(processKernelSource, kernelWriterAssembly, TensileInstructions()) - unaryWriteAssembly = functools.partial(writeAssembly, asmPath) - 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) - buildAssemblyCodeObjectFiles(asmToolchain, asmKernels, kernelWriterAssembly, outputPath, compress) + buildAssemblyCodeObjectFiles(asmToolchain, uniqueAsmKernels, kernelWriterAssembly, outputPath, compress) - writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H) - srcKernelFile = Path(outputPath) / "Kernels.cpp" - buildSourceCodeObjectFile(srcToolchain, outputPath, srcKernelFile) + #writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H) + #srcKernelFile = Path(outputPath) / "Kernels.cpp" + #buildSourceCodeObjectFile(srcToolchain, outputPath, srcKernelFile) popWorkingPath() # build_tmp popWorkingPath() # workingDir - return numKernels + return len(uniqueAsmKernels), len(asmKernels) -@timing -def getSolutionAndKernelWriters(solutions, kernels, assembler): +def generateKernelHelperObjects(solutions): + return list(dict.fromkeys([solution.getHelperKernelObjects() for solution in solutions])) - kernelSerialNaming = Solution.getSerialNaming(kernels) - solutionMinNaming = Solution.getMinNaming(solutions) - kernelMinNaming = Solution.getMinNaming(kernels) - kernelWriterAssembly = KernelWriterAssembly(kernelMinNaming, kernelSerialNaming, assembler) - return (kernelWriterAssembly, kernelMinNaming, solutionMinNaming) +@timing +def generateSolutions(args, cxxCompiler, logicDir): + 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 = [] + files = glob.glob(logicDir + "/*.yaml") + for f in files: + solutions.extend(LibraryIO.parseLibraryLogicFile(f, cxxCompiler, archs).solutions) + numSoln = len(solutions) + #solutions = list(dict.fromkeys(solutions).keys()) + return solutions, numSoln, (numSoln-len(solutions)) + + +#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) -@timing -def copyStaticFiles(outputPath=None): - if outputPath is None: - outputPath = globalParameters["WorkingPath"] - libraryStaticFiles = [ - "TensileTypes.h", - "tensile_bfloat16.h", - "tensile_float8_bfloat8.h", - "hip_f8_impl.h", - "KernelHeader.h", - "ReductionTemplate.h", - "memory_gfx.h" ] +def getArchitectures(arguments): + if ";" in arguments["Architecture"]: + archs = arguments["Architecture"].split(";") + else: + archs = arguments["Architecture"].split("_") + + logicArchs = set() + for arch in archs: + if arch in architectureMap: + logicArchs.add(architectureMap[arch]) + else: + printExit("Architecture %s not supported" % arch) + return archs, logicArchs - for fileName in libraryStaticFiles: - # copy file - shutil.copy( os.path.join(globalParameters["SourcePath"], fileName), \ - outputPath ) - return libraryStaticFiles +def getLogicFileList(arguments): + archs, _ = getArchitectures(arguments) + + 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)) + + if not os.path.exists(arguments["LogicPath"]): + printExit(f"LogicPath {arguments['LogicPath']} doesn't exist") + globPattern = os.path.join(arguments["LogicPath"], f"**/{arguments['LogicFilter']}.yaml") + print1(f"# LogicFilter: {globPattern}") + logicFiles = (os.path.join(arguments["LogicPath"], file) for file in glob.iglob(globPattern, recursive=True)) + print1(f"# Experimental: {arguments['Experimental']}") -@timing -def generateKernelObjectsFromSolutions(solutions): - kernels = [] - kernelHelperObjs = [] - kernelNames = set() - kernelHelperNames = set() - - for solution in solutions: - solutionKernels = solution.getKernels() - for kernel in solutionKernels: - kName = Solution.getKeyNoInternalArgs(kernel) - 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 - kernelHelperObjs = list(dict.fromkeys(kernelHelperObjs)) - return (kernels, kernelHelperObjs, kernelHelperNames) + if not arguments["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))] + + print2(f"# LibraryLogicFiles: {len(logicFiles)}") + for logicFile in logicFiles: + print2("# %s" % logicFile) + return list(set(str(Path(l).parent) for l in logicFiles)) -@timing -def generateSolutions(args, cxxCompiler, logicFile): - 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 = LibraryIO.parseLibraryLogicFile(logicFile, cxxCompiler, archs).solutions - numSoln = len(solutions) - solutions = dict.fromkeys(solutions).keys() - print1(f"Number of duplicate solutions: {numSoln - len(solutions)}") - return solutions - - -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) +@profile +def build(arguments, cxxCompiler, assembler, asmToolchain, srcToolchain, logicFiles): + solutions, totalSoln, dupSoln = generateSolutions(arguments, cxxCompiler, logicFiles) + numKernels = writeSolutionsAndKernelsTCL(arguments["OutputPath"], asmToolchain, srcToolchain, solutions, + assembler, compress=arguments["UseCompression"]) + return numKernels, totalSoln, dupSoln ################################################################################ @@ -374,55 +393,45 @@ def run(): asmToolchain = AssemblyToolchain(assembler, offloadBundler, globalParameters["BuildIdKind"], arguments["CodeObjectVersion"]) srcToolchain = SourceToolchain(cxxCompiler, offloadBundler, globalParameters["BuildIdKind"], globalParameters["AsanBuild"], globalParameters["SaveTemps"]) - if not os.path.exists(arguments["LogicPath"]): - printExit(f"LogicPath {arguments['LogicPath']} doesn't exist") - - if ";" in arguments["Architecture"]: - archs = arguments["Architecture"].split(";") - else: - archs = arguments["Architecture"].split("_") - logicArchs = set() - for arch in archs: - if arch in architectureMap: - logicArchs.add(architectureMap[arch]) - else: - printExit("Architecture %s not supported" % arch) - - logicExtFormat = ".yaml" - if arguments["LogicFormat"] == "yaml": - pass - elif arguments["LogicFormat"] == "json": - logicExtFormat = ".json" - else: - printExit("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) - + logicFiles = getLogicFileList(arguments) copyStaticFiles(arguments["OutputPath"]) + unaryBuild = functools.partial(build, arguments, cxxCompiler, assembler, asmToolchain, srcToolchain) + result = ParallelMap2(unaryBuild, logicFiles, "Building Library", multiArg=False, return_as="generator_unordered") + totalKernels = 0 + totalUnique = 0 + totDup = 0 + totSoln = 0 + dupKernels = 0 + + for n, total, dup in result: + totalKernels += n[1] + totalUnique += n[0] + totDup += dup + totSoln += total + + print1("# Tensile Library Writer DONE") + print1(HR) + print1("") + + stop = timer() + + print1(f"Total time (s): {(stop-start):3.2f}") + print1(f"Total kernels processed: {totalUnique}") + print1(f"Total kernels: {totalKernels}") + print1(f"Kernels processed per second: {(totalKernels/(stop-start)):3.2f}") + print1(f"Total solutions processed: {totSoln - totDup}") + print1(f"Duplicate solutions removed: {totDup}") + print1(f"Duplicate kernels removed: {dupKernels}") + + + + + + + + - unaryGenerateSolutions = functools.partial(generateSolutions, arguments, cxxCompiler) - solutions = ParallelMap2(unaryGenerateSolutions, logicFiles, "Reading logic files and generating solutions", multiArg=False) - solutions = (s for solutionList in solutions for s in solutionList) - numKernels = writeSolutionsAndKernelsTCL(arguments["OutputPath"], asmToolchain, srcToolchain, solutions, - assembler, compress=arguments["UseCompression"]) # archs = [getGfxName(arch) for arch in globalParameters['SupportedISA'] \ # if globalParameters["AsmCaps"][arch]["SupportedISA"]] @@ -439,14 +448,4 @@ def validLogicFile(p: Path): # for name, lib in newMasterLibrary.lazyLibraries.items(): # filename = os.path.join(newLibraryDir, name) # lib.applyNaming(kernelMinNaming) - # LibraryIO.write(filename, Utils.state(lib), arguments["LibraryFormat"]) - - 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}") + # LibraryIO.write(filename, Utils.state(lib), arguments["LibraryFormat"]) \ No newline at end of file diff --git a/tensilelite/Tensile/Toolchain/Assembly.py b/tensilelite/Tensile/Toolchain/Assembly.py index 0b44f36dff..d9230908e1 100644 --- a/tensilelite/Tensile/Toolchain/Assembly.py +++ b/tensilelite/Tensile/Toolchain/Assembly.py @@ -169,17 +169,15 @@ def _batchObjectFiles(objFiles: List[str], coPathDest: Union[Path, str], maxObjF def buildAssemblyCodeObjectFiles(toolchain: AssemblyToolchain, kernels, writerAsm, outputPath, compress: bool=True): - isAsm = lambda k: k["KernelLanguage"] == "Assembly" - extObj = ".o" extCo = ".co" extCoRaw = ".co.raw" destDir = Path(ensurePath(os.path.join(outputPath, 'library'))) - asmDir = Path(ensurePath(os.path.join(globalParameters["WorkingPath"], "assembly"))) + asmDir = Path(ensurePath(os.path.join(globalParameters["WorkingPath"], f"assembly/{str(os.getpid())}"))) archKernelMap = collections.defaultdict(list) - for k in filter(isAsm, kernels): + for k in kernels: archKernelMap[tuple(k['ISA'])].append(k) coFiles = [] From 40e9cd2018de530b3f6cb80794ff1eea5224e6fb Mon Sep 17 00:00:00 2001 From: David Dixon Date: Fri, 24 Jan 2025 19:37:04 +0000 Subject: [PATCH 10/25] validMFMA adds several GB to peak --- tensilelite/Tensile/Common.py | 1 - 1 file changed, 1 deletion(-) diff --git a/tensilelite/Tensile/Common.py b/tensilelite/Tensile/Common.py index 467d72c28c..115eeffd8f 100644 --- a/tensilelite/Tensile/Common.py +++ b/tensilelite/Tensile/Common.py @@ -429,7 +429,6 @@ def getArchitectureName(gfxName): validSMFMA["_format9"].append([MI[0],MI[1],MI[2],MI[3],2**bm,tt0,tt1,2**wave_m, 2**wave_n]) validSparseMatrixInstructions = validSMFMA["H"] + validSMFMA["B"] + validSMFMA["4xi8"] validMatrixInstructions = validMatrixInstructions + validSparseMatrixInstructions + validSMFMA["_format9"] -validMatrixInstructions = validMatrixInstructions + validSparseMatrixInstructions From a82ddbfbe72b49776fb28af0ca06b13ced1ae2eb Mon Sep 17 00:00:00 2001 From: David Dixon Date: Tue, 28 Jan 2025 17:29:01 +0000 Subject: [PATCH 11/25] Added very basic loadd balancing which improves build times --- .../Tensile/TensileCreateLibrary/Run.py | 31 +++++++++++++++---- 1 file changed, 25 insertions(+), 6 deletions(-) diff --git a/tensilelite/Tensile/TensileCreateLibrary/Run.py b/tensilelite/Tensile/TensileCreateLibrary/Run.py index 8da668f0ff..9534fb135d 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Run.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Run.py @@ -277,15 +277,16 @@ def generateKernelHelperObjects(solutions): @timing -def generateSolutions(args, cxxCompiler, logicDir): +def generateSolutions(args, cxxCompiler, logicDirs): 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 = [] - files = glob.glob(logicDir + "/*.yaml") - for f in files: - solutions.extend(LibraryIO.parseLibraryLogicFile(f, cxxCompiler, archs).solutions) + for logicDir in logicDirs: + files = glob.glob(logicDir + "/*.yaml") + for f in files: + solutions.extend(LibraryIO.parseLibraryLogicFile(f, cxxCompiler, archs).solutions) numSoln = len(solutions) #solutions = list(dict.fromkeys(solutions).keys()) return solutions, numSoln, (numSoln-len(solutions)) @@ -393,10 +394,28 @@ def run(): asmToolchain = AssemblyToolchain(assembler, offloadBundler, globalParameters["BuildIdKind"], arguments["CodeObjectVersion"]) srcToolchain = SourceToolchain(cxxCompiler, offloadBundler, globalParameters["BuildIdKind"], globalParameters["AsanBuild"], globalParameters["SaveTemps"]) - logicFiles = getLogicFileList(arguments) + def distribute(lst, n): + import heapq + 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) + lists[index].append(f) + heapq.heappush(totals, (total + value, index)) + return lists + + def getSize(logicDirs): + from os.path import getsize + from operator import itemgetter + filesizes = [(sum([int(getsize(f)) for f in glob.glob(d + "/*.yaml")]), d) for d in logicDirs] + return filesizes + + logicDirs = distribute(getSize(getLogicFileList(arguments)), 2*arguments["CpuThreads"]) + copyStaticFiles(arguments["OutputPath"]) unaryBuild = functools.partial(build, arguments, cxxCompiler, assembler, asmToolchain, srcToolchain) - result = ParallelMap2(unaryBuild, logicFiles, "Building Library", multiArg=False, return_as="generator_unordered") + result = ParallelMap2(unaryBuild, logicDirs, "Building Library", multiArg=False, return_as="generator_unordered") totalKernels = 0 totalUnique = 0 totDup = 0 From 18fad184b369a48fc8054748fb1a8471e7b640c6 Mon Sep 17 00:00:00 2001 From: David Dixon Date: Tue, 28 Jan 2025 22:33:26 +0000 Subject: [PATCH 12/25] Add static requireParameters --- .../Tensile/Utilities/RequiredParameters.py | 118 ++++++++++++++++++ 1 file changed, 118 insertions(+) create mode 100644 tensilelite/Tensile/Utilities/RequiredParameters.py 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 From 4da5f392817004e9df2f4e3b91e227cff8300e95 Mon Sep 17 00:00:00 2001 From: David Dixon Date: Mon, 3 Feb 2025 09:42:44 -0600 Subject: [PATCH 13/25] Use new codeObjectFile entry in logic file --- tensilelite/Tensile/CustomYamlLoader.py | 2 +- tensilelite/Tensile/LibraryIO.py | 72 +++--- tensilelite/Tensile/Parallel.py | 2 +- .../Tensile/TensileCreateLibrary/Run.py | 213 ++++++++++++------ 4 files changed, 181 insertions(+), 108 deletions(-) diff --git a/tensilelite/Tensile/CustomYamlLoader.py b/tensilelite/Tensile/CustomYamlLoader.py index bab8c68750..0fcebaaccf 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 = 3 arch = load_yaml_sequence_item(yaml_path, loader_type, GFX_ARCH_IDX) if isinstance(arch, dict): diff --git a/tensilelite/Tensile/LibraryIO.py b/tensilelite/Tensile/LibraryIO.py index 57b0ef8b7c..3c7a9608f4 100644 --- a/tensilelite/Tensile/LibraryIO.py +++ b/tensilelite/Tensile/LibraryIO.py @@ -204,16 +204,16 @@ def parseSolutionsData(data, srcFile, cxxCompiler): printExit("Solution file {} is missing required fields (len = {} < 3" \ .format(srcFile, len(data))) - versionString = data[0]["MinimumRequiredVersion"] + versionString = data[1]["MinimumRequiredVersion"] if not versionIsCompatible(versionString): printWarning("Version = {} in solution file {} does not match Tensile version = {}" \ .format(srcFile, versionString, __version__) ) - if "ProblemSizes" not in data[1]: + if "ProblemSizes" not in data[2]: printExit("Solution file {} doesn't begin with ProblemSizes".format(srcFile)) - problemSizesConfig = data[1]["ProblemSizes"] - solutionStartIdxInData = 2 + problemSizesConfig = data[2]["ProblemSizes"] + solutionStartIdxInData = 3 if (len(data) > solutionStartIdxInData) and "BiasTypeArgs" in data[solutionStartIdxInData]: solutionStartIdxInData += 1 if (len(data) > solutionStartIdxInData) and "ActivationArgs" in data[solutionStartIdxInData]: @@ -323,31 +323,32 @@ def parseLibraryLogicList(data, srcFile="?"): .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["codeObjectFile"] = data[0]["codeObjectFile"] + rv["MinimumRequiredVersion"] = data[1]["MinimumRequiredVersion"] + rv["ScheduleName"] = data[2] + rv["DeviceNames"] = data[4] + rv["ProblemType"] = data[5] + rv["Solutions"] = data[6] + + if type(data[3]) is dict: + rv["ArchitectureName"] = data[3]["Architecture"] + rv["CUCount"] = data[3]["CUCount"] else: - rv["ArchitectureName"] = data[2] + rv["ArchitectureName"] = data[3] rv["CUCount"] = None # TODOBEN: figure out what to do with these... - rv["ExactLogic"] = data[7] - rv["RangeLogic"] = data[8] + rv["ExactLogic"] = data[8] + rv["RangeLogic"] = data[9] # optional fields - if len(data) > 10 and data[10]: - rv["PerfMetric"] = data[10] + if len(data) > 10 and data[11]: + rv["PerfMetric"] = data[11] # library logic fields libraryType = None - if len(data) > 11 and data[11]: - libraryType = data[11] + if len(data) > 12 and data[12]: + libraryType = data[12] else: printExit("Library logic file {} is missing required field matching property." \ .format(srcFile)) @@ -355,13 +356,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[6])] rv["Library"]["distance"] = None else: rv["LibraryType"] = "Matching" rv["Library"] = {} - rv["Library"]["indexOrder"] = data[6] - rv["Library"]["table"] = data[7] + rv["Library"]["indexOrder"] = data[7] + rv["Library"]["table"] = data[8] rv["Library"]["distance"] = libraryType return rv @@ -369,23 +370,24 @@ 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] + codeObjectFile = data[0] + versionString = data[1] + scheduleName = data[2] + architectureName = data[3] + deviceNames = data[5] + problemTypeState = data[5] + solutionStates = data[6] + indexOrder = data[7] + exactLogic = data[8] + rangeLogic = data[9] 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,\ + return (codeObjectFile, versionString, scheduleName, architectureName, deviceNames,\ problemTypeState, solutionStates, indexOrder, exactLogic, rangeLogic, otherFields) diff --git a/tensilelite/Tensile/Parallel.py b/tensilelite/Tensile/Parallel.py index b0d44ef7b9..99950aba04 100644 --- a/tensilelite/Tensile/Parallel.py +++ b/tensilelite/Tensile/Parallel.py @@ -207,7 +207,7 @@ def ParallelMap2(function, objects, message="", enable=True, multiArg=True, retu pargs = zip(objects, itertools.repeat(globalParameters)) if joblibParallelSupportsGenerator(): - rv = Parallel(verbose=1000, 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=return_as)(delayed(pcall)(function, a, params) for a, params in pargs) else: rv = Parallel(n_jobs=threadCount,timeout=99999)(delayed(pcall)(function, a, params) for a, params in pargs) diff --git a/tensilelite/Tensile/TensileCreateLibrary/Run.py b/tensilelite/Tensile/TensileCreateLibrary/Run.py index 9534fb135d..af4c339abb 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Run.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Run.py @@ -27,10 +27,11 @@ import itertools import os import shutil +import subprocess from pathlib import Path from timeit import default_timer as timer -from typing import NamedTuple, List, Optional, Sequence, Union +from typing import Dict, NamedTuple, List, Optional, Union from Tensile import Utils from Tensile.Toolchain.Assembly import AssemblyToolchain, buildAssemblyCodeObjectFiles @@ -43,10 +44,11 @@ from Tensile.KernelWriterAssembly import KernelWriterAssembly from Tensile.KernelWriterBase import KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H from Tensile import LibraryIO -from Tensile.SolutionLibrary import MasterSolutionLibrary from Tensile.SolutionStructs import Solution -from Tensile.CustomYamlLoader import load_logic_gfx_arch +from Tensile.CustomYamlLoader import load_logic_gfx_arch, load_yaml_sequence_item from Tensile.Utilities.Profile import profile +from Tensile.Utilities.RequiredParameters import getRequiredParametersMin +from Tensile.SolutionLibrary import MasterSolutionLibrary from .ParseArguments import parseArguments @@ -258,9 +260,9 @@ def writeSolutionsAndKernelsTCL(outputPath, asmToolchain, srcToolchain, solution print1(f"{k['SolutionIndex']} {k['LogicFileName']}") pksResults = [processKernelSource(kernelWriterAssembly, TensileInstructions(), k) for k in uniqueAsmKernels] - for p, isa, wavefrontsize in [writeAssembly(asmPath, k) for k in pksResults]: - asmToolchain.assemble(str(p), str(p.with_suffix(".o")), getGfxName(isa), wavefrontsize) - buildAssemblyCodeObjectFiles(asmToolchain, uniqueAsmKernels, kernelWriterAssembly, outputPath, compress) + #for p, isa, wavefrontsize in [writeAssembly(asmPath, k) for k in pksResults]: + # asmToolchain.assemble(str(p), str(p.with_suffix(".o")), getGfxName(isa), wavefrontsize) + #buildAssemblyCodeObjectFiles(asmToolchain, uniqueAsmKernels, kernelWriterAssembly, outputPath, compress) #writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H) #srcKernelFile = Path(outputPath) / "Kernels.cpp" @@ -277,19 +279,22 @@ def generateKernelHelperObjects(solutions): @timing -def generateSolutions(args, cxxCompiler, logicDirs): +def generateSolutions(args, cxxCompiler, logicFiles): 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 = [] - for logicDir in logicDirs: - files = glob.glob(logicDir + "/*.yaml") - for f in files: - solutions.extend(LibraryIO.parseLibraryLogicFile(f, cxxCompiler, archs).solutions) + libraries = [] + for logicFileGroup in logicFiles: + for logicFile in logicFileGroup[1]: + libraryLogic = LibraryIO.parseLibraryLogicFile(logicFile, cxxCompiler, archs) + solutions.extend(libraryLogic.solutions) + libraries.append((libraryLogic.architecture, libraryLogic.library)) + numSoln = len(solutions) #solutions = list(dict.fromkeys(solutions).keys()) - return solutions, numSoln, (numSoln-len(solutions)) + return solutions, libraries, numSoln, (numSoln-len(solutions)) #def generateMatchTable(masterSolutionLibraries): @@ -336,9 +341,9 @@ def validLogicFile(p: Path): printExit(f"LogicPath {arguments['LogicPath']} doesn't exist") globPattern = os.path.join(arguments["LogicPath"], f"**/{arguments['LogicFilter']}.yaml") - print1(f"# LogicFilter: {globPattern}") + print1(f"# LogicFilter: {globPattern}") logicFiles = (os.path.join(arguments["LogicPath"], file) for file in glob.iglob(globPattern, recursive=True)) - print1(f"# Experimental: {arguments['Experimental']}") + 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)] @@ -349,15 +354,109 @@ def validLogicFile(p: Path): for logicFile in logicFiles: print2("# %s" % logicFile) - return list(set(str(Path(l).parent) for l in logicFiles)) + return logicFiles + + +def numberOfBuildKernerls(logicFile): + from operator import itemgetter + result = subprocess.run(['/bin/grep', "BuildKernel", logicFile], stderr=subprocess.PIPE, stdout=subprocess.PIPE, check=False) + return int(str(result.stdout).count("BuildKernel")) +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 schedule(logicFiles: list, numberOfTasks: int): + from yaml import Loader + problemMap = {} + for logicFile in logicFiles: + codeObjectFile = load_yaml_sequence_item(logicFile, Loader, 0) + codeObjectFile = codeObjectFile["codeObjectFile"] + if codeObjectFile in problemMap: + problemMap[codeObjectFile].append(logicFile) + else: + problemMap[codeObjectFile] = [logicFile] + + result = [] + for codeObjectFile, logicFiles in problemMap.items(): + count = sum(numberOfBuildKernerls(logicFile) for logicFile in logicFiles) + result.append((count, logicFiles)) + + return distribute(result, numberOfTasks) # need to convert list of list of tuples to list of list of strings + + +def updateParentMasterLibrary( + gfxName: str, + masterLib: MasterSolutionLibrary, + masterLibraries: Dict[str, MasterSolutionLibrary], + nextIdx: Dict[str, int], +) -> None: + if gfxName in masterLibraries: + nextIdx[gfxName] = masterLibraries[gfxName].merge(masterLib, nextIdx[gfxName]) + else: + masterLibraries[gfxName] = masterLib + nextIdx[gfxName] = 0 + +def updateMasterLibrary( + gfxName: str, + currMasterLib: MasterSolutionLibrary, + prevMasterLib: MasterSolutionLibrary, + nextIdx: int, +) -> None: + if prevMasterLib is not None: + nextIdx = prevMasterLib.merge(currMasterLib, nextIdx) + else: + prevMasterLib = currMasterLib + nextIdx = 0 + return prevMasterLib, nextIdx + +#import asyncio +#import aiofiles +#import msgpack + +#async def write_to_file(filename, data): +# # Open the file in asynchronous mode +# async with aiofiles.open(filename, 'w') as file: +# #await file.write(str(data)) +# msgpack.pack(data, file) + @profile def build(arguments, cxxCompiler, assembler, asmToolchain, srcToolchain, logicFiles): - solutions, totalSoln, dupSoln = generateSolutions(arguments, cxxCompiler, logicFiles) + + start = timer() + solutions, libraries, totalSoln, dupSoln = generateSolutions(arguments, cxxCompiler, logicFiles) + + #_masterLib = None + #_nextSolutionIdx = 0 + #if len(libraries) > 0: + # for gfxName, lib in libraries: + # _masterLib, _nextSolutionIdx = updateMasterLibrary(gfxName, lib, _masterLib, _nextSolutionIdx) + # Can we do this asynchronously before the call to writeSolutionsAndKernels? + #newLibraryDir = Path(arguments["OutputPath"]) / "library" + #for name, lib in list(_masterLib.lazyLibraries.items()): + #catalogPath = newLibraryDir / name + #parents = catalogPath.parents[:2] + #print1(f"# LAZY CATALOG: {parents[1].name}/{parents[0].name}/{catalogPath.name}") + #lib.applyNaming(getRequiredParametersMin()) # <-- This should be able to be replaced directly with `name`? + #LibraryIO.write(str(catalogPath), Utils.state(lib), "msgpack") + #asyncio.run(write_to_file(str(catalogPath)+".dat", Utils.state(lib))) + numKernels = writeSolutionsAndKernelsTCL(arguments["OutputPath"], asmToolchain, srcToolchain, solutions, assembler, compress=arguments["UseCompression"]) - return numKernels, totalSoln, dupSoln + stop = timer() + print1(f"Total time (s): {(stop-start):3.2f}") + print1(f"Kernels per second: {(numKernels[0]/(stop-start)):3.2f}") + print1(f" {numKernels[0]} {[l[1] for l in logicFiles]}") + return libraries, numKernels, totalSoln, dupSoln ################################################################################ @@ -389,44 +488,43 @@ def run(): print1(f"# Architecture(s): {arguments['Architecture']}") print1(f"# Library Format: {arguments['LibraryFormat']}") - assignGlobalParameters(arguments, cxxCompiler) + libraryDir = Path(arguments["OutputPath"]) / "library" + libraryDir.mkdir(exist_ok=True) + + unsortedLogic = getLogicFileList(arguments) + logicFiles = list(filter(lambda x: x != [], schedule(unsortedLogic, 2*arguments["CpuThreads"]))) + assignGlobalParameters(arguments, cxxCompiler) asmToolchain = AssemblyToolchain(assembler, offloadBundler, globalParameters["BuildIdKind"], arguments["CodeObjectVersion"]) srcToolchain = SourceToolchain(cxxCompiler, offloadBundler, globalParameters["BuildIdKind"], globalParameters["AsanBuild"], globalParameters["SaveTemps"]) - - def distribute(lst, n): - import heapq - 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) - lists[index].append(f) - heapq.heappush(totals, (total + value, index)) - return lists - - def getSize(logicDirs): - from os.path import getsize - from operator import itemgetter - filesizes = [(sum([int(getsize(f)) for f in glob.glob(d + "/*.yaml")]), d) for d in logicDirs] - return filesizes - - logicDirs = distribute(getSize(getLogicFileList(arguments)), 2*arguments["CpuThreads"]) - copyStaticFiles(arguments["OutputPath"]) unaryBuild = functools.partial(build, arguments, cxxCompiler, assembler, asmToolchain, srcToolchain) - result = ParallelMap2(unaryBuild, logicDirs, "Building Library", multiArg=False, return_as="generator_unordered") + result = ParallelMap2(unaryBuild, logicFiles, "Building Library", multiArg=False, return_as="generator_unordered") + totalKernels = 0 totalUnique = 0 totDup = 0 totSoln = 0 dupKernels = 0 - for n, total, dup in result: - totalKernels += n[1] - totalUnique += n[0] - totDup += dup - totSoln += total + #baseName = "TensileLibrary_" + #masterLibs = {} + #nextIdx = {} + + for library, n, total, dup in result: + #if len(library) > 0: + # for gfxName, lib in library: + # updateParentMasterLibrary(gfxName, lib, masterLibs, nextIdx) + totalKernels += n[1] + totalUnique += n[0] + totDup += dup + totSoln += total + + #for arch, masterLib in masterLibs.items(): + # name = baseName + "lazy_" + arch + # print1(f"WRITING PARENT CATALOG: {name}") + # masterLib.applyNaming(getRequiredParametersMin()) # <-- This should be able to be replaced directly with `name`? + # LibraryIO.write(str(libraryDir / name), Utils.state(masterLib), arguments["LibraryFormat"]) print1("# Tensile Library Writer DONE") print1(HR) @@ -441,30 +539,3 @@ def getSize(logicDirs): print1(f"Total solutions processed: {totSoln - totDup}") print1(f"Duplicate solutions removed: {totDup}") print1(f"Duplicate kernels removed: {dupKernels}") - - - - - - - - - - - - # archs = [getGfxName(arch) for arch in globalParameters['SupportedISA'] \ - # if globalParameters["AsmCaps"][arch]["SupportedISA"]] - # newLibraryDir = ensurePath(os.path.join(arguments["OutputPath"], 'library')) - - # for archName, newMasterLibrary in masterLibraries.items(): - # if archName in archs: - # if globalParameters["LazyLibraryLoading"]: - # masterFile = os.path.join(newLibraryDir, "TensileLibrary_lazy_"+archName) - # else: - # masterFile = os.path.join(newLibraryDir, "TensileLibrary_"+archName) - # newMasterLibrary.applyNaming(kernelMinNaming) - # LibraryIO.write(masterFile, Utils.state(newMasterLibrary), arguments["LibraryFormat"]) - # for name, lib in newMasterLibrary.lazyLibraries.items(): - # filename = os.path.join(newLibraryDir, name) - # lib.applyNaming(kernelMinNaming) - # LibraryIO.write(filename, Utils.state(lib), arguments["LibraryFormat"]) \ No newline at end of file From 9cf1cedcfb58faa73f1bf05d36d5b1d4867d1fb6 Mon Sep 17 00:00:00 2001 From: David Dixon Date: Mon, 3 Feb 2025 20:17:04 +0000 Subject: [PATCH 14/25] Working master solution library --- tensilelite/Tensile/SolutionLibrary.py | 2 +- .../Tensile/TensileCreateLibrary/Run.py | 64 ++++++++----------- 2 files changed, 27 insertions(+), 39 deletions(-) diff --git a/tensilelite/Tensile/SolutionLibrary.py b/tensilelite/Tensile/SolutionLibrary.py index f23a04bcbf..352e8e0c56 100644 --- a/tensilelite/Tensile/SolutionLibrary.py +++ b/tensilelite/Tensile/SolutionLibrary.py @@ -569,7 +569,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/TensileCreateLibrary/Run.py b/tensilelite/Tensile/TensileCreateLibrary/Run.py index af4c339abb..936b563d14 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Run.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Run.py @@ -260,9 +260,9 @@ def writeSolutionsAndKernelsTCL(outputPath, asmToolchain, srcToolchain, solution print1(f"{k['SolutionIndex']} {k['LogicFileName']}") pksResults = [processKernelSource(kernelWriterAssembly, TensileInstructions(), k) for k in uniqueAsmKernels] - #for p, isa, wavefrontsize in [writeAssembly(asmPath, k) for k in pksResults]: - # asmToolchain.assemble(str(p), str(p.with_suffix(".o")), getGfxName(isa), wavefrontsize) - #buildAssemblyCodeObjectFiles(asmToolchain, uniqueAsmKernels, kernelWriterAssembly, outputPath, compress) + for p, isa, wavefrontsize in [writeAssembly(asmPath, k) for k in pksResults]: + asmToolchain.assemble(str(p), str(p.with_suffix(".o")), getGfxName(isa), wavefrontsize) + buildAssemblyCodeObjectFiles(asmToolchain, uniqueAsmKernels, kernelWriterAssembly, outputPath, compress) #writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H) #srcKernelFile = Path(outputPath) / "Kernels.cpp" @@ -401,10 +401,10 @@ def updateParentMasterLibrary( nextIdx: Dict[str, int], ) -> None: if gfxName in masterLibraries: - nextIdx[gfxName] = masterLibraries[gfxName].merge(masterLib, nextIdx[gfxName]) + nextIdx= masterLibraries[gfxName].merge(masterLib, nextIdx) else: masterLibraries[gfxName] = masterLib - nextIdx[gfxName] = 0 + def updateMasterLibrary( gfxName: str, @@ -419,15 +419,6 @@ def updateMasterLibrary( nextIdx = 0 return prevMasterLib, nextIdx -#import asyncio -#import aiofiles -#import msgpack - -#async def write_to_file(filename, data): -# # Open the file in asynchronous mode -# async with aiofiles.open(filename, 'w') as file: -# #await file.write(str(data)) -# msgpack.pack(data, file) @profile def build(arguments, cxxCompiler, assembler, asmToolchain, srcToolchain, logicFiles): @@ -435,20 +426,17 @@ def build(arguments, cxxCompiler, assembler, asmToolchain, srcToolchain, logicFi start = timer() solutions, libraries, totalSoln, dupSoln = generateSolutions(arguments, cxxCompiler, logicFiles) - #_masterLib = None - #_nextSolutionIdx = 0 - #if len(libraries) > 0: - # for gfxName, lib in libraries: - # _masterLib, _nextSolutionIdx = updateMasterLibrary(gfxName, lib, _masterLib, _nextSolutionIdx) + _masterLib = None + _nextSolutionIdx = 0 + if len(libraries) > 0: + for gfxName, lib in libraries: + _masterLib, _nextSolutionIdx = updateMasterLibrary(gfxName, lib, _masterLib, _nextSolutionIdx) # Can we do this asynchronously before the call to writeSolutionsAndKernels? - #newLibraryDir = Path(arguments["OutputPath"]) / "library" - #for name, lib in list(_masterLib.lazyLibraries.items()): - #catalogPath = newLibraryDir / name - #parents = catalogPath.parents[:2] - #print1(f"# LAZY CATALOG: {parents[1].name}/{parents[0].name}/{catalogPath.name}") - #lib.applyNaming(getRequiredParametersMin()) # <-- This should be able to be replaced directly with `name`? - #LibraryIO.write(str(catalogPath), Utils.state(lib), "msgpack") - #asyncio.run(write_to_file(str(catalogPath)+".dat", Utils.state(lib))) + newLibraryDir = Path(arguments["OutputPath"]) / "library" + for name, lib in list(_masterLib.lazyLibraries.items()): + catalogPath = newLibraryDir / name + lib.applyNaming(getRequiredParametersMin()) # <-- This should be able to be replaced directly with `name`? + LibraryIO.write(str(catalogPath), Utils.state(lib), "dat") numKernels = writeSolutionsAndKernelsTCL(arguments["OutputPath"], asmToolchain, srcToolchain, solutions, assembler, compress=arguments["UseCompression"]) @@ -456,6 +444,7 @@ def build(arguments, cxxCompiler, assembler, asmToolchain, srcToolchain, logicFi print1(f"Total time (s): {(stop-start):3.2f}") print1(f"Kernels per second: {(numKernels[0]/(stop-start)):3.2f}") print1(f" {numKernels[0]} {[l[1] for l in logicFiles]}") + return libraries, numKernels, totalSoln, dupSoln @@ -507,24 +496,23 @@ def run(): totSoln = 0 dupKernels = 0 - #baseName = "TensileLibrary_" - #masterLibs = {} - #nextIdx = {} + baseName = "TensileLibrary_" + masterLibs = {} + nextIdx = 0 for library, n, total, dup in result: - #if len(library) > 0: - # for gfxName, lib in library: - # updateParentMasterLibrary(gfxName, lib, masterLibs, nextIdx) + if len(library) > 0: + for gfxName, lib in library: + updateParentMasterLibrary(gfxName, lib, masterLibs, nextIdx) totalKernels += n[1] totalUnique += n[0] totDup += dup totSoln += total - #for arch, masterLib in masterLibs.items(): - # name = baseName + "lazy_" + arch - # print1(f"WRITING PARENT CATALOG: {name}") - # masterLib.applyNaming(getRequiredParametersMin()) # <-- This should be able to be replaced directly with `name`? - # LibraryIO.write(str(libraryDir / name), Utils.state(masterLib), arguments["LibraryFormat"]) + for arch, masterLib in masterLibs.items(): + name = baseName + "lazy_" + arch + masterLib.applyNaming(getRequiredParametersMin()) # <-- This should be able to be replaced directly with `name`? + LibraryIO.write(str(libraryDir / name), Utils.state(masterLib), arguments["LibraryFormat"]) print1("# Tensile Library Writer DONE") print1(HR) From 275df7e67a653b974a990d7fa922ee0d57aa86a7 Mon Sep 17 00:00:00 2001 From: David Dixon Date: Mon, 3 Feb 2025 14:20:38 -0600 Subject: [PATCH 15/25] Use msgpack not dat --- tensilelite/Tensile/TensileCreateLibrary/Run.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensilelite/Tensile/TensileCreateLibrary/Run.py b/tensilelite/Tensile/TensileCreateLibrary/Run.py index 936b563d14..b4f6ead274 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Run.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Run.py @@ -436,7 +436,7 @@ def build(arguments, cxxCompiler, assembler, asmToolchain, srcToolchain, logicFi for name, lib in list(_masterLib.lazyLibraries.items()): catalogPath = newLibraryDir / name lib.applyNaming(getRequiredParametersMin()) # <-- This should be able to be replaced directly with `name`? - LibraryIO.write(str(catalogPath), Utils.state(lib), "dat") + LibraryIO.write(str(catalogPath), Utils.state(lib), "msgpack") numKernels = writeSolutionsAndKernelsTCL(arguments["OutputPath"], asmToolchain, srcToolchain, solutions, assembler, compress=arguments["UseCompression"]) From cd021bcbad00c4d7031aab50cd5c2e41362c3a2f Mon Sep 17 00:00:00 2001 From: David Dixon Date: Fri, 7 Feb 2025 15:00:27 +0000 Subject: [PATCH 16/25] Draft implementation of new build algo working end to end --- .../Tensile/Source/tensile_float8_bfloat8.h | 9 +- .../TensileCreateLibrary/ParseArguments.py | 6 +- .../Tensile/TensileCreateLibrary/Run.py | 206 ++++++++++++------ tensilelite/Tensile/Toolchain/Source.py | 145 ++++-------- tensilelite/Tensile/Toolchain/Validators.py | 33 ++- 5 files changed, 222 insertions(+), 177 deletions(-) diff --git a/tensilelite/Tensile/Source/tensile_float8_bfloat8.h b/tensilelite/Tensile/Source/tensile_float8_bfloat8.h index d8b6efb99b..af01d22a85 100644 --- a/tensilelite/Tensile/Source/tensile_float8_bfloat8.h +++ b/tensilelite/Tensile/Source/tensile_float8_bfloat8.h @@ -166,14 +166,17 @@ static void set_hip_f8_bias_mode_optimal() hip_f8_bias_mode_bit_host = true; } -static inline HIP_HOST_DEVICE bool get_hip_f8_bias_mode() -{ #if defined(__HIP_DEVICE_COMPILE__) +static inline HIP_DEVICE bool get_hip_f8_bias_mode() +{ return hip_f8_bias_mode_bit_device; +} #else +static inline HIP_HOST bool get_hip_f8_bias_mode() +{ return hip_f8_bias_mode_bit_host; -#endif } +#endif static bool isOcpF8; inline bool IsOCPSupported() diff --git a/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py b/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py index 3ecc3f4204..04a04397e5 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py +++ b/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py @@ -53,11 +53,13 @@ def parseArguments(input: Optional[List[str]] = None) -> Dict[str, Any]: help=f"Default: {ToolchainDefaults.CXX_COMPILER}") argParser.add_argument("--c-compiler", dest="CCompiler", action="store", default=ToolchainDefaults.C_COMPILER) argParser.add_argument("--cmake-cxx-compiler", dest="CmakeCxxCompiler", action="store") + 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("--offload-bundler", dest="OffloadBundler", action="store", default=ToolchainDefaults.OFFLOAD_BUNDLER) argParser.add_argument("--assembler", dest="Assembler", action="store", default=ToolchainDefaults.ASSEMBLER) argParser.add_argument("--code-object-version", dest="CodeObjectVersion", choices=["4", "5"], default="4", action="store") argParser.add_argument("--architecture", dest="Architecture", type=str, action="store", default="all", help="Supported archs: " + " ".join(architectureMap.keys())) - argParser.add_argument("--short-file-names", dest="ShortNames", action="store_true", default=False) + argParser.add_argument("--short-file-names", dest="ShortNames", action="store_true", default=False) argParser.add_argument("--no-compress", dest="NoCompress", action="store_true", help="Don't compress assembly code objects.") argParser.add_argument("--experimental", dest="Experimental", action="store_true", help="Include logic files in directories named 'Experimental'.") @@ -117,6 +119,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 b4f6ead274..9e5ba94172 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Run.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Run.py @@ -188,6 +188,44 @@ def writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNE kernelHeaderFile.write(HeaderText) +def writeHelper(outputPath, kernelHelperObj) -> str: + name = kernelHelperObj.getKernelName() + KERNEL_HELPER_FILENAME_CPP = name + ".cpp" + KERNEL_HELPER_FILENAME_H = name + ".h" + kernelSourceFilename = os.path.join(os.path.normcase(outputPath), "Kernels", KERNEL_HELPER_FILENAME_CPP) + kernelHeaderFilename = os.path.join(os.path.normcase(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") + if globalParameters["RuntimeLanguage"] == "HIP": + 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") + 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: + 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 getKernelWriterAssembly(solutions, assembler): kernelSerialNaming = Solution.getSerialNaming(solutions) kernelMinNaming = Solution.getMinNaming(solutions) @@ -247,7 +285,7 @@ def writeSolutionsAndKernelsTCL(outputPath, asmToolchain, srcToolchain, solution asmPath = ensurePath(os.path.join(globalParameters["WorkingPath"], "assembly")) asmKernels = [k.getKernels()[0] for k in solutions if k['KernelLanguage'] == 'Assembly'] - #kernelHelperObjs = generateKernelHelperObjects(asmKernels) + kernelHelperObjs = generateKernelHelperObjects(asmKernels) kernelWriterAssembly = getKernelWriterAssembly(asmKernels, assembler) uniqueAsmKernels = [k for k in asmKernels if "BuildKernel" in k] @@ -262,20 +300,20 @@ def writeSolutionsAndKernelsTCL(outputPath, asmToolchain, srcToolchain, solution pksResults = [processKernelSource(kernelWriterAssembly, TensileInstructions(), k) for k in uniqueAsmKernels] for p, isa, wavefrontsize in [writeAssembly(asmPath, k) for k in pksResults]: asmToolchain.assemble(str(p), str(p.with_suffix(".o")), getGfxName(isa), wavefrontsize) + buildAssemblyCodeObjectFiles(asmToolchain, uniqueAsmKernels, kernelWriterAssembly, outputPath, compress) - #writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H) - #srcKernelFile = Path(outputPath) / "Kernels.cpp" - #buildSourceCodeObjectFile(srcToolchain, outputPath, srcKernelFile) - popWorkingPath() # build_tmp popWorkingPath() # workingDir - return len(uniqueAsmKernels), len(asmKernels) + return kernelHelperObjs, len(uniqueAsmKernels), len(asmKernels) def generateKernelHelperObjects(solutions): - return list(dict.fromkeys([solution.getHelperKernelObjects() for solution in solutions])) + khos = [] + for solution in solutions: + khos.extend(solution.getHelperKernelObjects()) + return list(dict.fromkeys(khos)) @timing @@ -293,7 +331,6 @@ def generateSolutions(args, cxxCompiler, logicFiles): libraries.append((libraryLogic.architecture, libraryLogic.library)) numSoln = len(solutions) - #solutions = list(dict.fromkeys(solutions).keys()) return solutions, libraries, numSoln, (numSoln-len(solutions)) @@ -420,32 +457,77 @@ def updateMasterLibrary( return prevMasterLib, nextIdx +def genLazyMasterSolutionLibrary(libraryPath, libraryFormat, libraries): + _masterLib = None + _nextSolutionIdx = 0 + if len(libraries) > 0: + for gfxName, lib in libraries: + _masterLib, _nextSolutionIdx = updateMasterLibrary(gfxName, lib, _masterLib, _nextSolutionIdx) + # Can we do this asynchronously before the call to writeSolutionsAndKernels? + newLibraryDir = libraryPath + for name, lib in list(_masterLib.lazyLibraries.items()): + catalogPath = newLibraryDir / name + lib.applyNaming(getRequiredParametersMin()) # <-- This should be able to be replaced directly with `name`? + LibraryIO.write(str(catalogPath), Utils.state(lib), libraryFormat) + + @profile def build(arguments, cxxCompiler, assembler, asmToolchain, srcToolchain, logicFiles): start = timer() + libraryPath = Path(arguments["OutputPath"]) / "library" solutions, libraries, totalSoln, dupSoln = generateSolutions(arguments, cxxCompiler, logicFiles) - - _masterLib = None - _nextSolutionIdx = 0 - if len(libraries) > 0: - for gfxName, lib in libraries: - _masterLib, _nextSolutionIdx = updateMasterLibrary(gfxName, lib, _masterLib, _nextSolutionIdx) - # Can we do this asynchronously before the call to writeSolutionsAndKernels? - newLibraryDir = Path(arguments["OutputPath"]) / "library" - for name, lib in list(_masterLib.lazyLibraries.items()): - catalogPath = newLibraryDir / name - lib.applyNaming(getRequiredParametersMin()) # <-- This should be able to be replaced directly with `name`? - LibraryIO.write(str(catalogPath), Utils.state(lib), "msgpack") - - numKernels = writeSolutionsAndKernelsTCL(arguments["OutputPath"], asmToolchain, srcToolchain, solutions, - assembler, compress=arguments["UseCompression"]) + genLazyMasterSolutionLibrary(libraryPath, arguments["LibraryFormat"], libraries) + khos, numUniqueKernels, numKernels = writeSolutionsAndKernelsTCL(arguments["OutputPath"], asmToolchain, srcToolchain, solutions, + assembler, compress=arguments["UseCompression"]) stop = timer() - print1(f"Total time (s): {(stop-start):3.2f}") - print1(f"Kernels per second: {(numKernels[0]/(stop-start)):3.2f}") - print1(f" {numKernels[0]} {[l[1] for l in logicFiles]}") - - return libraries, numKernels, totalSoln, dupSoln + print2(f"Total time (s): {(stop-start):3.2f}") + print2(f"Kernels per second: {(numUniqueKernels/(stop-start)):3.2f}") + print2(f" {numUniqueKernels} {[l[1] for l in logicFiles]}") + + return libraries, khos, numUniqueKernels, numKernels, totalSoln, dupSoln + + +def generateParentLibrary(libraryFormat: str, libraryPath: Union[Path, str], masterLibs: Dict[str, MasterSolutionLibrary]): + for arch, masterLib in masterLibs.items(): + name = "TensileLibrary_" + "lazy_" + arch + masterLib.applyNaming(getRequiredParametersMin()) # <-- This should be able to be replaced directly with `name`? + LibraryIO.write(str(libraryPath / name), Utils.state(masterLib), libraryFormat) + + +def createDirectories(outputPath): + outputPath = Path(outputPath) + buildTmp = outputPath.parent / "build_tmp" / str(outputPath.name).upper() + srcCodeObjectPath = buildTmp / "code_object_tmp" + srcCodeObjectPath.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, srcCodeObjectPath, libraryPath + + +def extracBuildResults(result): + numKernels = 0 + numUniqueKernerls = 0 + numDuplicateKernels = 0 + numSoln = 0 + numDuplicateSoln = 0 + masterLibs = {} + nextIdx = 0 + kho = [] + + for library, khos, uniqueKernels, numKerns, soln, dupSoln in result: + if len(library) > 0: + for gfxName, lib in library: + updateParentMasterLibrary(gfxName, lib, masterLibs, nextIdx) + numUniqueKernerls += uniqueKernels + numKernels += numKerns + numDuplicateSoln += dupSoln + numSoln += soln + kho.extend(khos) + + return kho, masterLibs, numKernels, numUniqueKernerls, numDuplicateKernels, numSoln, numDuplicateSoln ################################################################################ @@ -463,9 +545,10 @@ def run(): arguments = parseArguments() ensurePath(arguments["OutputPath"]) arguments["OutputPath"] = os.path.abspath(arguments["OutputPath"]) + outputPath, buildTmpPath, srcCodeObjectPath, libraryPath = createDirectories(arguments["OutputPath"]) - cxxCompiler, cCompiler, offloadBundler, assembler, hipconfig = validateToolchain( - arguments["CxxCompiler"], arguments["CCompiler"], arguments["OffloadBundler"], arguments["Assembler"], ToolchainDefaults.HIP_CONFIG + cxxCompiler, cCompiler, offloadBundler, rocObjExtract, rocObjLs, assembler, hipconfig = validateToolchain( + arguments["CxxCompiler"], arguments["CCompiler"], arguments["OffloadBundler"], arguments["RocObjExtract"], arguments["RocObjLs"], arguments["Assembler"], ToolchainDefaults.HIP_CONFIG ) print1(f"# HIP Version: {getVersion(hipconfig, regex=r'(.+)')}") @@ -473,46 +556,36 @@ def run(): print1(f"# C Compiler: {cCompiler} (version {getVersion(cCompiler)})") print1(f"# Assembler: {assembler} (version {getVersion(assembler)})") print1(f"# Offload Bundler: {offloadBundler} (version {getVersion(offloadBundler)})") + print1(f"# Object Extractor: {rocObjExtract} (version {getVersion(rocObjLs, '-v')})") + print1(f"# Object Lister: {rocObjLs} (version {getVersion(rocObjLs, '-v')})") print1(f"# Code Object Version: {arguments['CodeObjectVersion']}") print1(f"# Architecture(s): {arguments['Architecture']}") print1(f"# Library Format: {arguments['LibraryFormat']}") - libraryDir = Path(arguments["OutputPath"]) / "library" - libraryDir.mkdir(exist_ok=True) + assignGlobalParameters(arguments, cxxCompiler) - unsortedLogic = getLogicFileList(arguments) - logicFiles = list(filter(lambda x: x != [], schedule(unsortedLogic, 2*arguments["CpuThreads"]))) + if ";" in arguments["Architecture"]: + archs = arguments["Architecture"].split(";") # user arg list format + else: + archs = arguments["Architecture"].split("_") # workaround for cmake list in list issue - assignGlobalParameters(arguments, cxxCompiler) asmToolchain = AssemblyToolchain(assembler, offloadBundler, globalParameters["BuildIdKind"], arguments["CodeObjectVersion"]) - srcToolchain = SourceToolchain(cxxCompiler, offloadBundler, globalParameters["BuildIdKind"], globalParameters["AsanBuild"], globalParameters["SaveTemps"]) - copyStaticFiles(arguments["OutputPath"]) + srcToolchain = SourceToolchain(cxxCompiler, rocObjExtract, rocObjLs, globalParameters["BuildIdKind"], globalParameters["AsanBuild"], globalParameters["SaveTemps"]) + + copyStaticFiles(outputPath) + unsortedLogic = getLogicFileList(arguments) + logicFiles = list(filter(lambda x: x != [], schedule(unsortedLogic, 2*arguments["CpuThreads"]))) unaryBuild = functools.partial(build, arguments, cxxCompiler, assembler, asmToolchain, srcToolchain) result = ParallelMap2(unaryBuild, logicFiles, "Building Library", multiArg=False, return_as="generator_unordered") - totalKernels = 0 - totalUnique = 0 - totDup = 0 - totSoln = 0 - dupKernels = 0 - - baseName = "TensileLibrary_" - masterLibs = {} - nextIdx = 0 - - for library, n, total, dup in result: - if len(library) > 0: - for gfxName, lib in library: - updateParentMasterLibrary(gfxName, lib, masterLibs, nextIdx) - totalKernels += n[1] - totalUnique += n[0] - totDup += dup - totSoln += total - - for arch, masterLib in masterLibs.items(): - name = baseName + "lazy_" + arch - masterLib.applyNaming(getRequiredParametersMin()) # <-- This should be able to be replaced directly with `name`? - LibraryIO.write(str(libraryDir / name), Utils.state(masterLib), arguments["LibraryFormat"]) + kho, masterLibs, numKernels, numUniqueKernels, numDuplicateKernels, numSoln, numDuplicateSoln = extracBuildResults(result) + unaryWriteHelpers = functools.partial(writeHelper, outputPath) + srcFiles = ParallelMap2(unaryWriteHelpers, list(dict.fromkeys(kho)), "Generating Kernels code", multiArg=False, return_as="list") + kernelsLib = str(srcCodeObjectPath / "Kernels.so") + srcToolchain.compile(srcFiles, str(kernelsLib), str(outputPath), archs) + buildSourceCodeObjectFile(srcToolchain, libraryPath, kernelsLib) + + generateParentLibrary(arguments["LibraryFormat"], libraryPath, masterLibs) print1("# Tensile Library Writer DONE") print1(HR) @@ -521,9 +594,10 @@ def run(): stop = timer() print1(f"Total time (s): {(stop-start):3.2f}") - print1(f"Total kernels processed: {totalUnique}") - print1(f"Total kernels: {totalKernels}") - print1(f"Kernels processed per second: {(totalKernels/(stop-start)):3.2f}") - print1(f"Total solutions processed: {totSoln - totDup}") - print1(f"Duplicate solutions removed: {totDup}") - print1(f"Duplicate kernels removed: {dupKernels}") + print1(f"Total kernels: {numKernels}") + print1(f"Total kernels processed: {numUniqueKernels}") + print1(f"Duplicate kernels removed: {numDuplicateKernels}") + print1(f"Kernels processed per second: {(numKernels/(stop-start)):3.2f}") + print1(f"Total solutions processed: {numSoln}") + print1(f"Duplicate solutions: {numDuplicateSoln}") + diff --git a/tensilelite/Tensile/Toolchain/Source.py b/tensilelite/Tensile/Toolchain/Source.py index ab0eba5220..a50223f13a 100644 --- a/tensilelite/Tensile/Toolchain/Source.py +++ b/tensilelite/Tensile/Toolchain/Source.py @@ -32,17 +32,18 @@ from timeit import default_timer as timer from typing import List, Union -from ..Common import globalParameters, print1, print2, ensurePath, splitArchs +from ..Common import globalParameters, print2, printExit class SourceToolchain: - def __init__(self, compiler: str, bundler: str, buildIdKind: str, asanBuild: bool=False, saveTemps: bool=False): + def __init__(self, compiler: str, objDump: str, objLs: str, buildIdKind: str, asanBuild: bool=False, saveTemps: bool=False): self.compiler = compiler - self.bundler = bundler + self.objLs = objLs + self.objDump = objDump self.buildIdKind = buildIdKind self.asanBuild = asanBuild self.saveTemps = saveTemps - def invoke(self, args: List[str], desc: str=""): + def invoke(self, args: List[str], desc: str="", workingDir=None): """Invokes a subprocess with the provided arguments. Args: @@ -54,7 +55,10 @@ def invoke(self, args: List[str], desc: str=""): """ print2(f"{desc}: {' '.join(args)}") try: - out = subprocess.check_output(args, stderr=subprocess.STDOUT) + if workingDir: + out = subprocess.check_output(args, stderr=subprocess.STDOUT, cwd=workingDir) + else: + out = subprocess.check_output(args, stderr=subprocess.STDOUT) except subprocess.CalledProcessError as err: raise RuntimeError( f"Error with {desc}: {err.output}\n" @@ -63,30 +67,22 @@ def invoke(self, args: List[str], desc: str=""): print2(f"Output: {out}") return out - def compile(self, srcPath: str, destPath: str, includePath: str, gfxs: List[str]): - """Compiles a source file into an object file. - - Args: - cmdlineArchs: List of architectures for offloading. - kernelFile: The path to the kernel source file. - buildPath: The build directory path. - objectFilename: The name of the output object file. - outputPath: The output directory path. - globalParameters: A dictionary of global parameters. - - Raises: - RuntimeError: If the compilation command fails. - """ + def compile(self, srcPaths: List[str], destPath: str, includePath: str, gfxs: List[str]): launcher = shlex.split(os.environ.get("Tensile_CXX_COMPILER_LAUNCHER", "")) hipFlags = [ + "-shared", + "-fPIC", + "-fgpu-rdc", + "-Xoffload-linker", + "--lto-partitions=16", "-D__HIP_HCC_COMPAT_MODE__=1", - "--offload-device-only", - "-x", "hip", "-O3", + "-x", "hip", "-O3", "-I", includePath, - "-Xoffload-linker", f"--build-id={self.buildIdKind}", "-std=c++17", + "-parallel-jobs=64" ] + if self.asanBuild: hipFlags.extend(["-fsanitize=address", "-shared-libasan", "-fuse-ld=lld"]) if self.saveTemps: @@ -96,75 +92,40 @@ def compile(self, srcPath: str, destPath: str, includePath: str, gfxs: List[str] archFlags = [f"--offload-arch={gfx}" for gfx in gfxs] - args = [ - *launcher, self.compiler, *hipFlags, *archFlags, srcPath, "-c", "-o", destPath - ] + args = [*launcher, self.compiler, *hipFlags, *archFlags] + args += srcPaths + args +=["-o", destPath] return self.invoke(args, f"Compiling HIP source kernels into objects (.cpp -> .o)") - def targets(self, objFile: str): - """Lists the target triples in an object file. + def list(self, sharedObjFile: str): + """Lists the code objects in shared object. Args: - objFile: The object file path. + sharedObjFile: Name of object file to list. Returns: - List of target triples in the object file. + List of objects embedded in shared object file. """ - args = [self.bundler, "--type=o", f"--input={objFile}", "-list"] - return self.invoke(args, f"Listing target triples in object file").decode().split("\n") + args = [self.objLs, sharedObjFile] + return [e.strip().split()[1:] for e in self.invoke(args, f"Listing code objects in shared object", Path(sharedObjFile).parent).decode().split("\n") if e.strip().split()[1:]] - def unbundle(self, target: str, srcPath: str, destPath: str): - """Unbundles source code object files using the Clang Offload Bundler. + + def extract(self, filename: str): + """Extracts code objects from a shared object. Args: - target: The target triple, see https://llvm.org/docs/AMDGPUUsage.html#target-triples. - infile: The path to the input object file. - outfileRaw: The path to the unbundled code object. + objFile: Name pf object file to extract. - Raises: - RuntimeError: If unbundling the source code object file fails. + Returns: + List of target triples in the object file. """ - args = [ - self.bundler, - "--type=o", - f"--targets={target}", - f"--input={srcPath}", - f"--output={destPath}", - "--unbundle", - ] - - return self.invoke(args, f"Unbundling source code object file") - + args = [self.objDump, filename] + return self.invoke(args, f"Extracting code object.", Path(filename.replace("file:","")).parent) -def _computeSourceCodeObjectFilename(target: str, base: str, buildPath: Union[Path, str], arch: str) -> Union[Path, None]: - """Generates a code object file path using the target, base, and build path. - Args: - target: The target triple. - base: The base name for the output file (name without extension). - buildPath: The build directory path. - - Returns: - Path to the code object file. - """ - coPath = None - buildPath = Path(buildPath) - if "TensileLibrary" in base and "fallback" in base: - coPath = buildPath / "{0}_{1}.hsaco.raw".format(base, arch) - elif "TensileLibrary" in base: - variant = [t for t in ["", "xnack-", "xnack+"] if t in target][-1] - baseVariant = base + "-" + variant if variant else base - if arch in baseVariant: - coPath = buildPath / (baseVariant + ".hsaco.raw") - else: - coPath= buildPath / "{0}.so-000-{1}.hsaco.raw".format(base, arch) - - return coPath - - -def buildSourceCodeObjectFile(toolchain: SourceToolchain, outputPath: Union[Path, str], kernelPath: Union[Path, str]) -> List[str]: +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: @@ -176,40 +137,16 @@ def buildSourceCodeObjectFile(toolchain: SourceToolchain, outputPath: Union[Path Returns: List of paths to the created code objects. """ - start = timer() - - buildPath = Path(ensurePath(os.path.join(globalParameters['WorkingPath'], 'code_object_tmp'))) - destPath = Path(ensurePath(os.path.join(outputPath, 'library'))) - kernelPath = Path(kernelPath) if "CmakeCxxCompiler" in globalParameters and globalParameters["CmakeCxxCompiler"] is not None: os.environ["CMAKE_CXX_COMPILER"] = globalParameters["CmakeCxxCompiler"] - objFilename = kernelPath.stem + '.o' - coPathsRaw = [] - coPaths= [] - - _, cmdlineArchs = splitArchs() - - objPath = str(buildPath / objFilename) - toolchain.compile(str(kernelPath), objPath, str(outputPath), cmdlineArchs) - - for target in toolchain.targets(objPath): + for target, filename in toolchain.list(sharedObjPath): match = re.search("gfx.*$", target) if match: arch = re.sub(":", "-", match.group()) - coPathRaw = _computeSourceCodeObjectFilename(target, kernelPath.stem, buildPath, arch) - if not coPathRaw: continue - toolchain.unbundle(target, objPath, str(coPathRaw)) - - coPath = str(destPath / coPathRaw.stem) - coPathsRaw.append(coPathRaw) - coPaths.append(coPath) - - for src, dst in zip(coPathsRaw, coPaths): + toolchain.extract(filename) + print(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) - - stop = timer() - print1(f"buildSourceCodeObjectFile time (s): {(stop-start):3.2f}") - - return coPaths diff --git a/tensilelite/Tensile/Toolchain/Validators.py b/tensilelite/Tensile/Toolchain/Validators.py index 86c3ff074b..74c45061ce 100644 --- a/tensilelite/Tensile/Toolchain/Validators.py +++ b/tensilelite/Tensile/Toolchain/Validators.py @@ -61,6 +61,8 @@ 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") + 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") ASSEMBLER = osSelect(linux="amdclang++", windows="clang++.exe") HIP_CONFIG = osSelect(linux="hipconfig", windows="hipconfig") @@ -94,11 +96,35 @@ def supportedCxxCompiler(compiler: str) -> bool: return _supportedComponent(compiler, [ToolchainDefaults.CXX_COMPILER]) -def supportedOffloadBundler(bundler: str) -> bool: +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 supportedBundler(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. @@ -141,7 +167,7 @@ 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) + supportedCxxCompiler(file), supportedCCompiler(file), supportedBundler(file), supportedExtract(file), supportedLs(file), supportedHip(file) )): raise ValueError(f"{file} is not a supported toolchain component for OS: {os.name}") @@ -190,4 +216,5 @@ def getVersion(executable: str, versionFlag: str="--version", regex: str=r"versi match = re.search(regex, output, re.IGNORECASE) return match.group(1) if match else "" except Exception as e: + pass raise RuntimeError(f"Failed to get version when calling {args}: {e}") From 79eaa3933ae695b9715577335e088c6da269b899 Mon Sep 17 00:00:00 2001 From: David Dixon Date: Sat, 8 Feb 2025 16:43:10 +0000 Subject: [PATCH 17/25] Use the unique kernel count --- tensilelite/Tensile/TensileCreateLibrary/Run.py | 2 +- tensilelite/Tensile/Toolchain/Assembly.py | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/tensilelite/Tensile/TensileCreateLibrary/Run.py b/tensilelite/Tensile/TensileCreateLibrary/Run.py index 9e5ba94172..c75c02d472 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Run.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Run.py @@ -597,7 +597,7 @@ def run(): print1(f"Total kernels: {numKernels}") print1(f"Total kernels processed: {numUniqueKernels}") print1(f"Duplicate kernels removed: {numDuplicateKernels}") - print1(f"Kernels processed per second: {(numKernels/(stop-start)):3.2f}") + print1(f"Kernels processed per second: {(numUniqueKernels/(stop-start)):3.2f}") print1(f"Total solutions processed: {numSoln}") print1(f"Duplicate solutions: {numDuplicateSoln}") diff --git a/tensilelite/Tensile/Toolchain/Assembly.py b/tensilelite/Tensile/Toolchain/Assembly.py index d9230908e1..72a579b2a5 100644 --- a/tensilelite/Tensile/Toolchain/Assembly.py +++ b/tensilelite/Tensile/Toolchain/Assembly.py @@ -196,7 +196,7 @@ def buildAssemblyCodeObjectFiles(toolchain: AssemblyToolchain, kernels, writerAs if coName: coFileMap[asmDir / (coName + extCoRaw)].append(str(asmDir / (writerAsm.getKernelFileBase(kernel) + extObj))) for coFileRaw, objFiles in coFileMap.items(): - objFiles = _batchObjectFiles(objFiles, coFileRaw) + objFiles = _batchObjectFiles(set(objFiles), coFileRaw) # shouldn't need a set here toolchain.link(objFiles, str(coFileRaw)) coFile = destDir / coFileRaw.name.replace(extCoRaw, extCo) if compress: From a9a6d6f99c67a6eab164213b9927cae79a7fc422 Mon Sep 17 00:00:00 2001 From: David Dixon Date: Sun, 9 Feb 2025 13:26:04 -0600 Subject: [PATCH 18/25] TensileCreateLibrary clean up --- tensilelite/Tensile/BenchmarkProblems.py | 3 +- tensilelite/Tensile/KernelWriter.py | 8 +- tensilelite/Tensile/KernelWriterAssembly.py | 4 +- tensilelite/Tensile/Parallel.py | 33 +- .../Tensile/TensileCreateLibrary/IO.py | 113 +++++ .../Tensile/TensileCreateLibrary/Logic.py | 92 ++++ .../TensileCreateLibrary/ParseArguments.py | 2 +- .../Tensile/TensileCreateLibrary/Run.py | 475 ++++-------------- .../Tensile/TensileCreateLibrary/Tuning.py | 147 ++++++ .../Tensile/TensileCreateLibrary/__init__.py | 2 +- tensilelite/Tensile/Toolchain/Assembly.py | 35 +- tensilelite/Tensile/Toolchain/Source.py | 1 - 12 files changed, 498 insertions(+), 417 deletions(-) create mode 100644 tensilelite/Tensile/TensileCreateLibrary/IO.py create mode 100644 tensilelite/Tensile/TensileCreateLibrary/Logic.py create mode 100644 tensilelite/Tensile/TensileCreateLibrary/Tuning.py diff --git a/tensilelite/Tensile/BenchmarkProblems.py b/tensilelite/Tensile/BenchmarkProblems.py index 1cc2139057..1d95df8a47 100644 --- a/tensilelite/Tensile/BenchmarkProblems.py +++ b/tensilelite/Tensile/BenchmarkProblems.py @@ -40,7 +40,8 @@ printExit, printWarning, ensurePath, startTime, validParameters from .KernelWriterAssembly import KernelWriterAssembly from .SolutionStructs import Solution, ProblemType, ProblemSizes -from .TensileCreateLibrary import copyStaticFiles, writeSolutionsAndKernels +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 diff --git a/tensilelite/Tensile/KernelWriter.py b/tensilelite/Tensile/KernelWriter.py index 378d47cc36..2749f4ecdf 100644 --- a/tensilelite/Tensile/KernelWriter.py +++ b/tensilelite/Tensile/KernelWriter.py @@ -359,9 +359,9 @@ class KernelWriter(metaclass=abc.ABCMeta): ############################################################################## # Init ############################################################################## - def __init__(self, kernelMinNaming, kernelSerialNaming, assembler: str): + def __init__(self, kernelMinNaming, assembler: str): self.kernelMinNaming = kernelMinNaming - self.kernelSerialNaming = kernelSerialNaming + #self.kernelSerialNaming = kernelSerialNaming self.assembler = assembler self.ti = None @@ -5016,8 +5016,8 @@ def _getKernelSource(self, kernel: Solution): def getKernelFileBase(self, kernel): if isCustomKernelConfig(kernel): fileBase = kernel["CustomKernelName"] - elif globalParameters["ShortNames"]: - fileBase = Solution.getNameSerial(kernel, self.kernelSerialNaming) + #elif globalParameters["ShortNames"]: + # fileBase = Solution.getNameSerial(kernel, self.kernelSerialNaming) else: fileBase = self._shortenFileBase(kernel) return fileBase diff --git a/tensilelite/Tensile/KernelWriterAssembly.py b/tensilelite/Tensile/KernelWriterAssembly.py index db429f3236..4148590ec3 100644 --- a/tensilelite/Tensile/KernelWriterAssembly.py +++ b/tensilelite/Tensile/KernelWriterAssembly.py @@ -67,8 +67,8 @@ class KernelWriterAssembly(KernelWriter): ############################################################################## # Init ############################################################################## - def __init__(self, kernelMinNaming, kernelSerialNaming, assembler: str): - super(KernelWriterAssembly, self).__init__(kernelMinNaming, kernelSerialNaming, assembler) + def __init__(self, kernelMinNaming, assembler: str): + super(KernelWriterAssembly, self).__init__(kernelMinNaming, assembler) def getSourceFileString(self, kernel) -> Tuple[int, str]: diff --git a/tensilelite/Tensile/Parallel.py b/tensilelite/Tensile/Parallel.py index 99950aba04..82a3a6e925 100644 --- a/tensilelite/Tensile/Parallel.py +++ b/tensilelite/Tensile/Parallel.py @@ -28,6 +28,7 @@ import time import concurrent.futures +from typing import NamedTuple from joblib import Parallel, delayed def joblibParallelSupportsGenerator(): @@ -173,7 +174,15 @@ def ParallelMapReturnAsGenerator(function, objects, message="", enable=True, mul for result in concurrent.futures.as_completed(resultFutures): yield result.result() -def ParallelMap2(function, objects, message="", enable=True, multiArg=True, return_as="list"): + + +class ParallelMapConfig(NamedTuple): + message: str = "" + enable: bool = True + multiArg: bool = False + return_as: str = "generator_unordered" + +def ParallelMap2(function, config: ParallelMapConfig, objects): """ Generally equivalent to list(map(function, objects)), possibly executing in parallel. @@ -182,36 +191,38 @@ def ParallelMap2(function, objects, message="", enable=True, multiArg=True, retu 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) + if config.return_as in ('generator', 'generator_unordered') and not joblibParallelSupportsGenerator(): + return ParallelMapReturnAsGenerator(function, objects, message, config.enable, config.multiArg) from .Common import globalParameters from . import Utils - threadCount = CPUThreadCount(enable) + threadCount = 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 Utils.tqdm(objects, message)] + return [function(*args) if config.multiArg else function(args) for args in Utils.tqdm(objects, message)] countMessage = "" try: countMessage = " for {} tasks".format(len(objects)) except TypeError: pass - - if message != "": message += ": " - print("{0}Launching {1} threads{2}...".format(message, threadCount, countMessage)) + + message = config.message + if message != "": + message += ": " + print(f"{message} Launching {threadCount} threads{countMessage}...") sys.stdout.flush() currentTime = time.time() - pcall = pcallWithGlobalParamsMultiArg if multiArg else pcallWithGlobalParamsSingleArg + pcall = pcallWithGlobalParamsMultiArg if config.multiArg else pcallWithGlobalParamsSingleArg pargs = zip(objects, itertools.repeat(globalParameters)) 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, params) for a, params in pargs) else: rv = Parallel(n_jobs=threadCount,timeout=99999)(delayed(pcall)(function, a, params) for a, params in pargs) totalTime = time.time() - currentTime - print("{0}Done. ({1:.1f} secs elapsed)".format(message, totalTime)) + print(f"{message}Done. ({totalTime:.1f} secs elapsed)") sys.stdout.flush() return rv diff --git a/tensilelite/Tensile/TensileCreateLibrary/IO.py b/tensilelite/Tensile/TensileCreateLibrary/IO.py new file mode 100644 index 0000000000..e5bc2a0ba9 --- /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 +from Tensile.LibraryIO import parseLibraryLogicFile, write +from Tensile.SolutionLibrary import MasterSolutionLibrary +from Tensile.Utilities.RequiredParameters import getRequiredParametersMin +from Tensile.Utils import state + +from os import getpid +from pathlib import Path +from typing import Dict, Union + + +def generateSolutionsAndLibraries(archs, cxxCompiler, logicFiles): + solutions = [] + libraries = [] + for logicFileGroup in logicFiles: + for logicFile in logicFileGroup[1]: + libraryLogic = parseLibraryLogicFile(logicFile, cxxCompiler, archs) + 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 build 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(getRequiredParametersMin()) # <-- This should be able to be replaced directly with `name`? + 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(getRequiredParametersMin()) # <-- This should be able to be replaced directly with `name`? + 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..6dd8f184f4 --- /dev/null +++ b/tensilelite/Tensile/TensileCreateLibrary/Logic.py @@ -0,0 +1,92 @@ +################################################################################ +# +# 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 printExit, print1, print2 +from Tensile.CustomYamlLoader import load_logic_gfx_arch, load_yaml_sequence_item + +from glob import iglob +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)) + + if not logicPath.exists(): + printExit(f"LogicPath {str(logicPath)} doesn't exist") + + globPattern = str(logicPath / f"**/{logicFilter}.yaml") + print1(f"# LogicFilter: {globPattern}") + logicFiles = (str(logicPath / file) for file in iglob(globPattern, recursive=True)) + print1(f"# Experimental: {experimental}") + + 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))] + + 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 numberOfBuildKernerls(logicFile): + result = run(['/bin/grep', "BuildKernel", logicFile], stderr=PIPE, stdout=PIPE, check=False) + return int(str(result.stdout).count("BuildKernel")) + + +def schedule(logicFiles: list, numberOfTasks: int): + from yaml import Loader + problemMap = {} + for logicFile in logicFiles: + codeObjectFile = load_yaml_sequence_item(logicFile, Loader, 0) + codeObjectFile = codeObjectFile["codeObjectFile"] + if codeObjectFile in problemMap: + problemMap[codeObjectFile].append(logicFile) + else: + problemMap[codeObjectFile] = [logicFile] + + result = [] + for codeObjectFile, logicFiles in problemMap.items(): + count = sum(numberOfBuildKernerls(logicFile) for logicFile in logicFiles) + result.append((count, logicFiles)) + + return 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 04a04397e5..59dbc94cc7 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py +++ b/tensilelite/Tensile/TensileCreateLibrary/ParseArguments.py @@ -87,7 +87,7 @@ def parseArguments(input: Optional[List[str]] = None) -> Dict[str, Any]: argParser.add_argument("--address-sanitizer", dest="AsanBuild", action="store_true", default=False, help="Enable ASAN build.") argParser.add_argument("--keep-build-tmp", dest="KeepBuildTmp", action="store_true", - default=False, help="Do not remove the temporary build directory (may required hundreds of GBs of space)"), + help="Do not remove the temporary build directory (may required hundreds of GBs of space)"), argParser.add_argument("--logic-filter", dest="LogicFilter", action="store", default="*", type=str, help="Cutomsized logic filter, default is *, i.e. all logics." " Example: gfx942/Equality/* for building equality of gfx942 only") diff --git a/tensilelite/Tensile/TensileCreateLibrary/Run.py b/tensilelite/Tensile/TensileCreateLibrary/Run.py index c75c02d472..af2198ac25 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Run.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Run.py @@ -1,6 +1,6 @@ ################################################################################ # -# Copyright (C) 2022-2024 Advanced Micro Devices, Inc. All rights reserved. +# Copyright (C) 2022-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 @@ -23,33 +23,29 @@ ################################################################################ import functools -import glob -import itertools import os import shutil -import subprocess from pathlib import Path from timeit import default_timer as timer -from typing import Dict, NamedTuple, List, Optional, Union +from typing import Dict, NamedTuple, List, Optional -from Tensile import Utils from Tensile.Toolchain.Assembly import AssemblyToolchain, buildAssemblyCodeObjectFiles from Tensile.Toolchain.Source import SourceToolchain, buildSourceCodeObjectFile from Tensile.Toolchain.Validators import validateToolchain, getVersion, ToolchainDefaults from Tensile.TensileInstructions import getGfxName, TensileInstructions -from Tensile.Common import globalParameters, HR, print1, print2, printExit, ensurePath, \ - CHeader, assignGlobalParameters, architectureMap, IsaVersion, pushWorkingPath, \ - popWorkingPath, ParallelMap2 +from Tensile.Common import globalParameters, HR, print1, print2, printWarning, \ + assignGlobalParameters, IsaVersion, ParallelMap2 +from Tensile.Parallel import ParallelMapConfig +from Tensile.Utilities.RequiredParameters import getRequiredParametersMin from Tensile.KernelWriterAssembly import KernelWriterAssembly -from Tensile.KernelWriterBase import KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H -from Tensile import LibraryIO from Tensile.SolutionStructs import Solution -from Tensile.CustomYamlLoader import load_logic_gfx_arch, load_yaml_sequence_item from Tensile.Utilities.Profile import profile -from Tensile.Utilities.RequiredParameters import getRequiredParametersMin from Tensile.SolutionLibrary import MasterSolutionLibrary +from .IO import generateSolutionsAndLibraries, genLazyMasterSolutionLibrary, \ + generateParentLibrary, writeAssembly, writeHelper +from .Logic import logicFileList, schedule from .ParseArguments import parseArguments def timing(func): @@ -84,7 +80,6 @@ def copyStaticFiles(outputPath=None): return libraryStaticFiles - class KernelCodeGenResult(NamedTuple): err: int src: str @@ -95,7 +90,7 @@ class KernelCodeGenResult(NamedTuple): wavefrontSize: int -def processKernelSource(kernelWriterAssembly, ti, kernel) -> KernelCodeGenResult: +def _processKernelSource(kernelWriterAssembly, ti, kernel) -> KernelCodeGenResult: """ Generate source for a single kernel. Returns (error, source, header, kernelName). @@ -110,228 +105,23 @@ def processKernelSource(kernelWriterAssembly, ti, kernel) -> KernelCodeGenResult return KernelCodeGenResult(err, src, header, asmFilename, objFilename, tuple(kernel["ISA"]), kernel["WavefrontSize"]) -def removeInvalidSolutionsAndKernels(results, kernels, solutions, errorTolerant, globalParameters): - removeKernels = [] - removeKernelNames = [] - removeSolutions = [] - removeResults = [] - - for kernIdx, r in Utils.tqdm(enumerate(results)) if globalParameters["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 = Solution.getKeyNoInternalArgs(kernels[kernIdx]) - 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 Utils.tqdm(solutions, "Finding invalid solutions") if globalParameters["PrintLevel"] > 1 else solutions: - solutionKernels = solution.getKernels() - for kernel in solutionKernels: - kName = Solution.getKeyNoInternalArgs(kernel) - if kName in removeKernelNames: - removeSolutions.append(solution) - break - - for solut in removeSolutions: - solutions.remove(solut) - - for rel in removeResults: - results.remove(rel) - - -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) / str(os.getpid()) - path.mkdir(exist_ok=True) - filepath = path / f"{result.name}.s" - isa = result.isa - wfsize = result.wavefrontSize - with open(filepath, "w", encoding="utf-8") as f: - f.write(result.src) - del result # result.src is very large so let gc know to clean up asap - - return filepath, 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") - if globalParameters["RuntimeLanguage"] == "HIP": - 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 writeHelper(outputPath, kernelHelperObj) -> str: - name = kernelHelperObj.getKernelName() - KERNEL_HELPER_FILENAME_CPP = name + ".cpp" - KERNEL_HELPER_FILENAME_H = name + ".h" - kernelSourceFilename = os.path.join(os.path.normcase(outputPath), "Kernels", KERNEL_HELPER_FILENAME_CPP) - kernelHeaderFilename = os.path.join(os.path.normcase(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") - if globalParameters["RuntimeLanguage"] == "HIP": - 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") - 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: - 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 getKernelWriterAssembly(solutions, assembler): - kernelSerialNaming = Solution.getSerialNaming(solutions) - kernelMinNaming = Solution.getMinNaming(solutions) - kernelWriterAssembly = KernelWriterAssembly(kernelMinNaming, kernelSerialNaming, assembler) - return kernelWriterAssembly - - -def writeSolutionsAndKernels(outputPath, asmToolchain, srcToolchain, solutions, kernels, kernelHelperObjs, \ - kernelWriterAssembly, errorTolerant=False, generateSourcesAndExit=False, compress=True): - codeObjectFiles = [] - - pushWorkingPath('build_tmp') - pushWorkingPath(os.path.basename(outputPath).upper()) - asmPath = ensurePath(os.path.join(globalParameters["WorkingPath"], "assembly")) - - asmKernels = [k for k in kernels if k['KernelLanguage'] == 'Assembly'] - - visited = set() - duplicates = 0 - for k in asmKernels: - base = kernelWriterAssembly.getKernelFileBase(k) - k.duplicate = True if base in visited else False - duplicates += k.duplicate - print2(f"Duplicate: {base}") - visited.add(base) - print1(f"Number of duplicates: {duplicates}") - - numAsmKernels = len(asmKernels) - numKernels = len(asmKernels) - assert numKernels == numAsmKernels, "Only assembly kernels are supported in TensileLite" - asmIter = zip(itertools.repeat(kernelWriterAssembly), itertools.repeat(TensileInstructions()), asmKernels) - asmResults = ParallelMap2(processKernelSource, asmIter, "Generating assembly kernels") - removeInvalidSolutionsAndKernels(asmResults, asmKernels, solutions, errorTolerant, globalParameters) - def assemble(ret): - p, isa, wavefrontsize = ret - asmToolchain.assemble(str(p), str(p.with_suffix(".o")), getGfxName(isa), wavefrontsize) - unaryWriteAssembly = functools.partial(writeAssembly, asmPath) - 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, asmKernels, kernelWriterAssembly, outputPath, compress) - buildSourceCodeObjectFile(srcToolchain, outputPath, srcKernelFile) - - popWorkingPath() # build_tmp - popWorkingPath() # workingDir - - return codeObjectFiles, numKernels - @timing -def writeSolutionsAndKernelsTCL(outputPath, asmToolchain, srcToolchain, solutions, assembler, compress=True): - pushWorkingPath('build_tmp') - pushWorkingPath(os.path.basename(outputPath).upper()) - asmPath = ensurePath(os.path.join(globalParameters["WorkingPath"], "assembly")) - - asmKernels = [k.getKernels()[0] for k in solutions if k['KernelLanguage'] == 'Assembly'] - kernelHelperObjs = generateKernelHelperObjects(asmKernels) - kernelWriterAssembly = getKernelWriterAssembly(asmKernels, assembler) - - uniqueAsmKernels = [k for k in asmKernels if "BuildKernel" in k] - unique = set() - for k in uniqueAsmKernels: - name = kernelWriterAssembly.getKernelFileBase(k) - if name not in unique: - unique.add(name) - else: - print1(f"{k['SolutionIndex']} {k['LogicFileName']}") - - pksResults = [processKernelSource(kernelWriterAssembly, TensileInstructions(), k) for k in uniqueAsmKernels] - for p, isa, wavefrontsize in [writeAssembly(asmPath, k) for k in pksResults]: +def buildAssemblyKernels(asmPath: Path, asmToolchain: AssemblyToolchain, kernelWriterAssembly: KernelWriterAssembly, ti: TensileInstructions, removeTemporaries: bool, solnLibs: List[tuple]): + uniqueAsmKernels = [k.getKernels()[0] for k in solnLibs[0] if "BuildKernel" in k] + pksResults = [_processKernelSource(kernelWriterAssembly, ti, k) for k in uniqueAsmKernels] + for p, isa, wavefrontsize in set([writeAssembly(asmPath, k) for k in pksResults]): asmToolchain.assemble(str(p), str(p.with_suffix(".o")), getGfxName(isa), wavefrontsize) + if removeTemporaries: + p.unlink() - buildAssemblyCodeObjectFiles(asmToolchain, uniqueAsmKernels, kernelWriterAssembly, outputPath, compress) - - popWorkingPath() # build_tmp - popWorkingPath() # workingDir - - return kernelHelperObjs, len(uniqueAsmKernels), len(asmKernels) + return uniqueAsmKernels, solnLibs[1] def generateKernelHelperObjects(solutions): khos = [] for solution in solutions: khos.extend(solution.getHelperKernelObjects()) - return list(dict.fromkeys(khos)) - - -@timing -def generateSolutions(args, cxxCompiler, logicFiles): - 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 = [] - libraries = [] - for logicFileGroup in logicFiles: - for logicFile in logicFileGroup[1]: - libraryLogic = LibraryIO.parseLibraryLogicFile(logicFile, cxxCompiler, archs) - solutions.extend(libraryLogic.solutions) - libraries.append((libraryLogic.architecture, libraryLogic.library)) - - numSoln = len(solutions) - return solutions, libraries, numSoln, (numSoln-len(solutions)) + return khos #should we deduplicate here? #def generateMatchTable(masterSolutionLibraries): @@ -351,86 +141,6 @@ def generateSolutions(args, cxxCompiler, logicFiles): # LibraryIO.write("MatchTable", matchTable) -def getArchitectures(arguments): - if ";" in arguments["Architecture"]: - archs = arguments["Architecture"].split(";") - else: - archs = arguments["Architecture"].split("_") - - logicArchs = set() - for arch in archs: - if arch in architectureMap: - logicArchs.add(architectureMap[arch]) - else: - printExit("Architecture %s not supported" % arch) - return archs, logicArchs - - -def getLogicFileList(arguments): - archs, _ = getArchitectures(arguments) - - 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)) - - if not os.path.exists(arguments["LogicPath"]): - printExit(f"LogicPath {arguments['LogicPath']} doesn't exist") - - globPattern = os.path.join(arguments["LogicPath"], f"**/{arguments['LogicFilter']}.yaml") - print1(f"# LogicFilter: {globPattern}") - logicFiles = (os.path.join(arguments["LogicPath"], file) for file in glob.iglob(globPattern, recursive=True)) - 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)] - - logicFiles = [file for file in logicFiles if validLogicFile(Path(file))] - - print2(f"# LibraryLogicFiles: {len(logicFiles)}") - for logicFile in logicFiles: - print2("# %s" % logicFile) - - return logicFiles - - -def numberOfBuildKernerls(logicFile): - from operator import itemgetter - result = subprocess.run(['/bin/grep', "BuildKernel", logicFile], stderr=subprocess.PIPE, stdout=subprocess.PIPE, check=False) - return int(str(result.stdout).count("BuildKernel")) - - -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 schedule(logicFiles: list, numberOfTasks: int): - from yaml import Loader - problemMap = {} - for logicFile in logicFiles: - codeObjectFile = load_yaml_sequence_item(logicFile, Loader, 0) - codeObjectFile = codeObjectFile["codeObjectFile"] - if codeObjectFile in problemMap: - problemMap[codeObjectFile].append(logicFile) - else: - problemMap[codeObjectFile] = [logicFile] - - result = [] - for codeObjectFile, logicFiles in problemMap.items(): - count = sum(numberOfBuildKernerls(logicFile) for logicFile in logicFiles) - result.append((count, logicFiles)) - - return distribute(result, numberOfTasks) # need to convert list of list of tuples to list of list of strings - - def updateParentMasterLibrary( gfxName: str, masterLib: MasterSolutionLibrary, @@ -456,78 +166,42 @@ def updateMasterLibrary( nextIdx = 0 return prevMasterLib, nextIdx - -def genLazyMasterSolutionLibrary(libraryPath, libraryFormat, libraries): - _masterLib = None - _nextSolutionIdx = 0 +def processMsl(libraryPath, libraryFormat, solnLibTup): + solns, libraries = solnLibTup if len(libraries) > 0: - for gfxName, lib in libraries: - _masterLib, _nextSolutionIdx = updateMasterLibrary(gfxName, lib, _masterLib, _nextSolutionIdx) - # Can we do this asynchronously before the call to writeSolutionsAndKernels? - newLibraryDir = libraryPath - for name, lib in list(_masterLib.lazyLibraries.items()): - catalogPath = newLibraryDir / name - lib.applyNaming(getRequiredParametersMin()) # <-- This should be able to be replaced directly with `name`? - LibraryIO.write(str(catalogPath), Utils.state(lib), libraryFormat) - - -@profile -def build(arguments, cxxCompiler, assembler, asmToolchain, srcToolchain, logicFiles): - - start = timer() - libraryPath = Path(arguments["OutputPath"]) / "library" - solutions, libraries, totalSoln, dupSoln = generateSolutions(arguments, cxxCompiler, logicFiles) - genLazyMasterSolutionLibrary(libraryPath, arguments["LibraryFormat"], libraries) - khos, numUniqueKernels, numKernels = writeSolutionsAndKernelsTCL(arguments["OutputPath"], asmToolchain, srcToolchain, solutions, - assembler, compress=arguments["UseCompression"]) - stop = timer() - print2(f"Total time (s): {(stop-start):3.2f}") - print2(f"Kernels per second: {(numUniqueKernels/(stop-start)):3.2f}") - print2(f" {numUniqueKernels} {[l[1] for l in logicFiles]}") + masterLib = None + nextSolutionIdx = 0 + for gfxName, lib in libraries: + masterLib, nextSolutionIdx = updateMasterLibrary(gfxName, lib, masterLib, nextSolutionIdx) + genLazyMasterSolutionLibrary(libraryPath, libraryFormat, masterLib) # can we make this async + return solns, libraries - return libraries, khos, numUniqueKernels, numKernels, totalSoln, dupSoln - -def generateParentLibrary(libraryFormat: str, libraryPath: Union[Path, str], masterLibs: Dict[str, MasterSolutionLibrary]): - for arch, masterLib in masterLibs.items(): - name = "TensileLibrary_" + "lazy_" + arch - masterLib.applyNaming(getRequiredParametersMin()) # <-- This should be able to be replaced directly with `name`? - LibraryIO.write(str(libraryPath / name), Utils.state(masterLib), libraryFormat) - - -def createDirectories(outputPath): +def makeDirectories(outputPath): outputPath = Path(outputPath) + outputPath.mkdir(parents=True, exist_ok=True) buildTmp = outputPath.parent / "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, srcCodeObjectPath, libraryPath + return outputPath, buildTmp, asmPath, srcCodeObjectPath, libraryPath -def extracBuildResults(result): - numKernels = 0 - numUniqueKernerls = 0 - numDuplicateKernels = 0 - numSoln = 0 - numDuplicateSoln = 0 +def extractBuildResults(result): masterLibs = {} nextIdx = 0 - kho = [] - - for library, khos, uniqueKernels, numKerns, soln, dupSoln in result: - if len(library) > 0: - for gfxName, lib in library: - updateParentMasterLibrary(gfxName, lib, masterLibs, nextIdx) - numUniqueKernerls += uniqueKernels - numKernels += numKerns - numDuplicateSoln += dupSoln - numSoln += soln - kho.extend(khos) - - return kho, masterLibs, numKernels, numUniqueKernerls, numDuplicateKernels, numSoln, numDuplicateSoln + 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 ################################################################################ @@ -543,9 +217,8 @@ def run(): print2("") arguments = parseArguments() - ensurePath(arguments["OutputPath"]) - arguments["OutputPath"] = os.path.abspath(arguments["OutputPath"]) - outputPath, buildTmpPath, srcCodeObjectPath, libraryPath = createDirectories(arguments["OutputPath"]) + arguments["OutputPath"] = Path(arguments["OutputPath"]).resolve() + outputPath, buildTmp, assemblyPath, srcCodeObjectPath, libraryPath = makeDirectories(arguments["OutputPath"]) cxxCompiler, cCompiler, offloadBundler, rocObjExtract, rocObjLs, assembler, hipconfig = validateToolchain( arguments["CxxCompiler"], arguments["CCompiler"], arguments["OffloadBundler"], arguments["RocObjExtract"], arguments["RocObjLs"], arguments["Assembler"], ToolchainDefaults.HIP_CONFIG @@ -569,23 +242,60 @@ def run(): else: archs = arguments["Architecture"].split("_") # workaround for cmake list in list issue + # Construct Toolchain asmToolchain = AssemblyToolchain(assembler, offloadBundler, globalParameters["BuildIdKind"], arguments["CodeObjectVersion"]) srcToolchain = SourceToolchain(cxxCompiler, rocObjExtract, rocObjLs, globalParameters["BuildIdKind"], globalParameters["AsanBuild"], globalParameters["SaveTemps"]) - copyStaticFiles(outputPath) - unsortedLogic = getLogicFileList(arguments) + # List logic files + unsortedLogic = logicFileList(archs, Path(arguments["LogicPath"]), arguments["LogicFilter"], arguments["Experimental"]) logicFiles = list(filter(lambda x: x != [], schedule(unsortedLogic, 2*arguments["CpuThreads"]))) - unaryBuild = functools.partial(build, arguments, cxxCompiler, assembler, asmToolchain, srcToolchain) - result = ParallelMap2(unaryBuild, logicFiles, "Building Library", multiArg=False, return_as="generator_unordered") - kho, masterLibs, numKernels, numUniqueKernels, numDuplicateKernels, numSoln, numDuplicateSoln = extracBuildResults(result) + # Phase1: Build assembly and master solution libraries + writerAsm = KernelWriterAssembly(getRequiredParametersMin(), asmToolchain.assembler) + unaryGenSolutions = functools.partial(generateSolutionsAndLibraries, archs, cxxCompiler) + unaryProcessMsl = functools.partial(processMsl, libraryPath, arguments["LibraryFormat"]) + unaryBuildAsmKernels = functools.partial(buildAssemblyKernels, assemblyPath, asmToolchain, writerAsm, TensileInstructions(), not arguments["KeepBuildTmp"]) + unaryBuildCOFile = functools.partial(buildAssemblyCodeObjectFiles, asmToolchain, assemblyPath, libraryPath, writerAsm, arguments["UseCompression"]) + + def buildCoAndHelpers(input): + uniqueAsmKernels, libraries = input + unaryBuildCOFile(uniqueAsmKernels) + return generateKernelHelperObjects(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) + + kho, masterLibs = phase1(logicFiles) + + # Phase3: Build Master Solution Library + generateParentLibrary(arguments["LibraryFormat"], libraryPath, masterLibs, arguments["LazyLibraryLoading"]) + del masterLibs + + # Phase2: Write Kernel helpers and build Kernels.co + copyStaticFiles(outputPath) unaryWriteHelpers = functools.partial(writeHelper, outputPath) - srcFiles = ParallelMap2(unaryWriteHelpers, list(dict.fromkeys(kho)), "Generating Kernels code", multiArg=False, return_as="list") + srcFiles = ParallelMap2(unaryWriteHelpers, ParallelMapConfig(message="Generating Kernels code", return_as="list"), list(dict.fromkeys(kho))) + del kho + kernelsLib = str(srcCodeObjectPath / "Kernels.so") - srcToolchain.compile(srcFiles, str(kernelsLib), str(outputPath), archs) + srcToolchain.compile(srcFiles, kernelsLib, str(outputPath), archs) + del srcFiles + buildSourceCodeObjectFile(srcToolchain, libraryPath, kernelsLib) - generateParentLibrary(arguments["LibraryFormat"], libraryPath, masterLibs) + + if not arguments["KeepBuildTmp"]: + if buildTmp.exists() and buildTmp.is_dir(): + shutil.rmtree(buildTmp) + else: + printWarning(f"Cannot remove build_tmp") print1("# Tensile Library Writer DONE") print1(HR) @@ -594,10 +304,9 @@ def run(): stop = timer() print1(f"Total time (s): {(stop-start):3.2f}") - print1(f"Total kernels: {numKernels}") - print1(f"Total kernels processed: {numUniqueKernels}") - print1(f"Duplicate kernels removed: {numDuplicateKernels}") - print1(f"Kernels processed per second: {(numUniqueKernels/(stop-start)):3.2f}") - print1(f"Total solutions processed: {numSoln}") - print1(f"Duplicate solutions: {numDuplicateSoln}") - + #print1(f"Total kernels: {numKernels}") + #print1(f"Total kernels processed: {numUniqueKernels}") + #print1(f"Duplicate kernels removed: {numDuplicateKernels}") + #print1(f"Kernels processed per second: {(numUniqueKernels/(stop-start)):3.2f}") + #print1(f"Total solutions processed: {numSoln}") + #print1(f"Duplicate solutions: {numDuplicateSoln}") diff --git a/tensilelite/Tensile/TensileCreateLibrary/Tuning.py b/tensilelite/Tensile/TensileCreateLibrary/Tuning.py new file mode 100644 index 0000000000..52d6d55dea --- /dev/null +++ b/tensilelite/Tensile/TensileCreateLibrary/Tuning.py @@ -0,0 +1,147 @@ +################################################################################ +# +# 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, print1, print2, printExit, globalParameters, \ + ParallelMap2, pushWorkingPath, popWorkingPath, ensurePath +from Tensile.TensileInstructions import getGfxName, TensileInstructions +from Tensile.KernelWriterBase import KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H +from Tensile.SolutionStructs import Solution +from Tensile.Toolchain.Assembly import buildAssemblyCodeObjectFiles +from Tensile.Toolchain.Source import buildSourceCodeObjectFile +from Tensile.Utils import tqdm + +from .IO import writeAssembly +from .Run import _processKernelSource + +from functools import partial, reduce +from itertools import repeat +from pathlib import Path +import os + +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, globalParameters): + removeKernels = [] + removeKernelNames = [] + removeSolutions = [] + removeResults = [] + + for kernIdx, r in tqdm(enumerate(results)) if globalParameters["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 = Solution.getKeyNoInternalArgs(kernels[kernIdx]) + 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 globalParameters["PrintLevel"] > 1 else solutions: + solutionKernels = solution.getKernels() + for kernel in solutionKernels: + kName = Solution.getKeyNoInternalArgs(kernel) + if kName in removeKernelNames: + removeSolutions.append(solution) + break + + for solut in removeSolutions: + solutions.remove(solut) + + for rel in removeResults: + results.remove(rel) + + +def writeSolutionsAndKernels(outputPath, asmToolchain, srcToolchain, solutions, kernels, kernelHelperObjs, \ + kernelWriterAssembly, errorTolerant=False, generateSourcesAndExit=False, compress=True): + codeObjectFiles = [] + + pushWorkingPath('build_tmp') + pushWorkingPath(os.path.basename(outputPath).upper()) + asmPath = ensurePath(os.path.join(globalParameters["WorkingPath"], "assembly")) + + asmKernels = [k for k in kernels if k['KernelLanguage'] == 'Assembly'] + + visited = set() + duplicates = 0 + for k in asmKernels: + base = kernelWriterAssembly.getKernelFileBase(k) + k.duplicate = True if base in visited else False + duplicates += k.duplicate + print2(f"Duplicate: {base}") + visited.add(base) + print1(f"Number of duplicates: {duplicates}") + + numAsmKernels = len(asmKernels) + numKernels = len(asmKernels) + assert numKernels == numAsmKernels, "Only assembly kernels are supported in TensileLite" + asmIter = zip(repeat(kernelWriterAssembly), repeat(TensileInstructions()), asmKernels) + asmResults = ParallelMap2(_processKernelSource, asmIter, "Generating assembly kernels") + removeInvalidSolutionsAndKernels(asmResults, asmKernels, solutions, errorTolerant, globalParameters) + def assemble(ret): + p, isa, wavefrontsize = ret + asmToolchain.assemble(str(p), str(p.with_suffix(".o")), getGfxName(isa), wavefrontsize) + unaryWriteAssembly = partial(writeAssembly, asmPath) + compose = lambda *F: reduce(lambda f, g: lambda x: f(g(x)), F) + ret = ParallelMap2(compose(assemble, unaryWriteAssembly), asmResults, "Writing assembly kernels", return_as="generator_unordered", multiArg=False) + + _writeHelpers(outputPath, kernelHelperObjs, KERNEL_HELPER_FILENAME_CPP, KERNEL_HELPER_FILENAME_H) + srcKernelFile = Path(outputPath) / "Kernels.cpp" + + if not generateSourcesAndExit: + codeObjectFiles += buildAssemblyCodeObjectFiles(asmToolchain, asmKernels, kernelWriterAssembly, outputPath, compress) + buildSourceCodeObjectFile(srcToolchain, outputPath, srcKernelFile) + + popWorkingPath() # build_tmp + popWorkingPath() # workingDir + + return codeObjectFiles, numKernels diff --git a/tensilelite/Tensile/TensileCreateLibrary/__init__.py b/tensilelite/Tensile/TensileCreateLibrary/__init__.py index 16a8e3dc98..647b4e759d 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/__init__.py +++ b/tensilelite/Tensile/TensileCreateLibrary/__init__.py @@ -1,3 +1,3 @@ from .Run import run from .Run import copyStaticFiles -from .Run import writeSolutionsAndKernels \ No newline at end of file +from .Tuning import writeSolutionsAndKernels \ No newline at end of file diff --git a/tensilelite/Tensile/Toolchain/Assembly.py b/tensilelite/Tensile/Toolchain/Assembly.py index 72a579b2a5..961c004331 100644 --- a/tensilelite/Tensile/Toolchain/Assembly.py +++ b/tensilelite/Tensile/Toolchain/Assembly.py @@ -28,14 +28,15 @@ import shlex import shutil import subprocess -import warnings from pathlib import Path -from typing import List, Literal, Union, Tuple +from typing import List, Literal, Union -from .. import Utils from ..TensileInstructions import getGfxName -from ..Common import globalParameters, print2, ensurePath +from ..Common import globalParameters, print2 +from Tensile.SolutionStructs import Solution +from Tensile.KernelWriterAssembly import KernelWriterAssembly +from Tensile.Utilities.RequiredParameters import getRequiredParametersMin class AssemblyToolchain: def __init__(self, assembler: str, bundler: str, buildIdKind: str, coVersion: Literal[4, 5]): self.assembler = assembler @@ -167,14 +168,13 @@ def _batchObjectFiles(objFiles: List[str], coPathDest: Union[Path, str], maxObjF return newObjFilesOutput -def buildAssemblyCodeObjectFiles(toolchain: AssemblyToolchain, kernels, writerAsm, outputPath, compress: bool=True): +def buildAssemblyCodeObjectFiles(toolchain: AssemblyToolchain, srcDir, destDir, writerAsm, compress: bool, kernels): extObj = ".o" extCo = ".co" extCoRaw = ".co.raw" - destDir = Path(ensurePath(os.path.join(outputPath, 'library'))) - asmDir = Path(ensurePath(os.path.join(globalParameters["WorkingPath"], f"assembly/{str(os.getpid())}"))) + asmDir = srcDir archKernelMap = collections.defaultdict(list) for k in kernels: @@ -187,16 +187,25 @@ def buildAssemblyCodeObjectFiles(toolchain: AssemblyToolchain, kernels, writerAs gfx = getGfxName(arch) - objectFiles = [str(asmDir / (writerAsm.getKernelFileBase(k) + extObj)) for k in archKernels if 'codeObjectFile' not in k] + ### start of map for separate architectures no lazy loading - the absence of codeObjectFile means no lazy loading + objectFiles = [str(asmDir / (writerAsm.getKernelFileBase(k) + extObj)) for k in archKernels if 'codeObjectFile' not in k or k['codeObjectFile'] == "TensileLibrary"] coFileMap = collections.defaultdict(list) if len(objectFiles): coFileMap[asmDir / ("TensileLibrary_"+ gfx + extCoRaw)] = objectFiles - for kernel in archKernels: - coName = kernel.get("codeObjectFile", None) - if coName: - coFileMap[asmDir / (coName + extCoRaw)].append(str(asmDir / (writerAsm.getKernelFileBase(kernel) + extObj))) + ### end of map for separate architectures no lazy loading + + ### start of map for separate architectures and lazy loading + else: + for kernel in archKernels: + coName = kernel.get("codeObjectFile", None) + if coName: + coFileMap[asmDir / (coName + extCoRaw)].append(str(asmDir / (writerAsm.getKernelFileBase(kernel) + extObj))) + ### end of map for separate architectures and lazy loading + for coFileRaw, objFiles in coFileMap.items(): - objFiles = _batchObjectFiles(set(objFiles), coFileRaw) # shouldn't need a set here + # shouldn't need a set here the fact that we do implies we have duplicates + #objFiles = _batchObjectFiles(set(objFiles), coFileRaw) + objFiles = _batchObjectFiles(objFiles, coFileRaw) toolchain.link(objFiles, str(coFileRaw)) coFile = destDir / coFileRaw.name.replace(extCoRaw, extCo) if compress: diff --git a/tensilelite/Tensile/Toolchain/Source.py b/tensilelite/Tensile/Toolchain/Source.py index a50223f13a..81caedc4a9 100644 --- a/tensilelite/Tensile/Toolchain/Source.py +++ b/tensilelite/Tensile/Toolchain/Source.py @@ -146,7 +146,6 @@ def buildSourceCodeObjectFile(toolchain: SourceToolchain, destPath: Union[Path, if match: arch = re.sub(":", "-", match.group()) toolchain.extract(filename) - print(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) From 488f85ac20b0493b2b4be0c8a72ceb2a9b931f0a Mon Sep 17 00:00:00 2001 From: David Dixon Date: Mon, 3 Mar 2025 21:35:47 +0000 Subject: [PATCH 19/25] Use problem type info to compute co name --- tensilelite/Tensile/CodeObjectName.py | 73 +++++++++++++++++ tensilelite/Tensile/LibraryIO.py | 79 ++++++++++++------- .../Tensile/TensileCreateLibrary/Logic.py | 42 ++++++---- 3 files changed, 152 insertions(+), 42 deletions(-) create mode 100644 tensilelite/Tensile/CodeObjectName.py 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/LibraryIO.py b/tensilelite/Tensile/LibraryIO.py index 3c7a9608f4..6d9985603a 100644 --- a/tensilelite/Tensile/LibraryIO.py +++ b/tensilelite/Tensile/LibraryIO.py @@ -29,8 +29,11 @@ from . import Common from . import SolutionLibrary from .CustomYamlLoader import load_yaml_stream +from Tensile.CodeObjectName import codeObjectFileBaseName +from enum import IntEnum from typing import NamedTuple, List + import os import sys @@ -232,6 +235,24 @@ def parseSolutionsData(data, srcFile, cxxCompiler): return (problemSizes, solutions) +class DataIndex(IntEnum): + CODE_OBEJECT_FILE=0 + MINIMUM_REQUIRED_VERSION=1 + SCHEDULE_NAME=2 + DEVICE_PROPERTIES=3 + DEVICE_NAMES=4 + PROBLEM_TYPE=5 + SOLUTIONS=6 + INDEX_ORDER=7 + EXACT_LOGIC=8 + RANGE_LOGIC=9 + PERF_METRIC=11 + LIBRARY_TYPE=12 + + def __index__(self): + return self.value + + class LibraryLogic(NamedTuple): """Return tuple for parseLibraryLogicData()""" schedule: str @@ -308,7 +329,10 @@ def solutionStateToSolution(solutionState, cxxCompiler) -> Solution: solutions = [solutionStateToSolution(solutionState, cxxCompiler) for solutionState in data["Solutions"]] + n = codeObjectFileBaseName(data) newLibrary, codeObjectFile = SolutionLibrary.MasterSolutionLibrary.FromOriginalState(data, solutions, cxxCompiler) + assert n == codeObjectFile, f"codeObjectFileBaseName computed different name than FromOriginalState \n {n}\n {codeObjectFile}" + for solution in solutions: solution["codeObjectFile"] = codeObjectFile @@ -321,34 +345,33 @@ def parseLibraryLogicList(data, srcFile="?"): if len(data) < 9: printExit("Library logic file {} is missing required fields (len = {} < 9)" \ .format(srcFile, len(data))) - rv = {} - rv["codeObjectFile"] = data[0]["codeObjectFile"] - rv["MinimumRequiredVersion"] = data[1]["MinimumRequiredVersion"] - rv["ScheduleName"] = data[2] - rv["DeviceNames"] = data[4] - rv["ProblemType"] = data[5] - rv["Solutions"] = data[6] + rv["codeObjectFile"] = data[DataIndex.CODE_OBEJECT_FILE]["codeObjectFile"] + 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[3]) is dict: - rv["ArchitectureName"] = data[3]["Architecture"] + rv["ArchitectureName"] = data[DataIndex.DEVICE_PROPERTIES]["Architecture"] rv["CUCount"] = data[3]["CUCount"] else: - rv["ArchitectureName"] = data[3] + rv["ArchitectureName"] = data[DataIndex.DEVICE_PROPERTIES] rv["CUCount"] = None # TODOBEN: figure out what to do with these... - rv["ExactLogic"] = data[8] - rv["RangeLogic"] = data[9] + rv["ExactLogic"] = data[DataIndex.EXACT_LOGIC] + rv["RangeLogic"] = data[DataIndex.RANGE_LOGIC] # optional fields - if len(data) > 10 and data[11]: - rv["PerfMetric"] = data[11] + if len(data) > 10 and data[DataIndex.PERF_METRIC]: + rv["PerfMetric"] = data[DataIndex.PERF_METRIC] # library logic fields libraryType = None - if len(data) > 12 and data[12]: - libraryType = data[12] + if len(data) > 12 and data[DataIndex.LIBRARY_TYPE]: + libraryType = data[DataIndex.LIBRARY_TYPE] else: printExit("Library logic file {} is missing required field matching property." \ .format(srcFile)) @@ -356,13 +379,13 @@ def parseLibraryLogicList(data, srcFile="?"): rv["LibraryType"] = "FreeSize" rv["Library"] = {} rv["Library"]["indexOrder"] = None - rv["Library"]["table"] = [0, len(data[6])] + rv["Library"]["table"] = [0, len(data[DataIndex.SOLUTIONS])] rv["Library"]["distance"] = None else: rv["LibraryType"] = "Matching" rv["Library"] = {} - rv["Library"]["indexOrder"] = data[7] - rv["Library"]["table"] = data[8] + rv["Library"]["indexOrder"] = data[DataIndex.INDEX_ORDER] + rv["Library"]["table"] = data[DataIndex.EXACT_LOGIC] rv["Library"]["distance"] = libraryType return rv @@ -370,16 +393,16 @@ def parseLibraryLogicList(data, srcFile="?"): def rawLibraryLogic(data): """Returns a tuple of the data in a library logic file.""" - codeObjectFile = data[0] - versionString = data[1] - scheduleName = data[2] - architectureName = data[3] - deviceNames = data[5] - problemTypeState = data[5] - solutionStates = data[6] - indexOrder = data[7] - exactLogic = data[8] - rangeLogic = data[9] + codeObjectFile = data[DataIndex.CODE_OBEJECT_FILE] + 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) diff --git a/tensilelite/Tensile/TensileCreateLibrary/Logic.py b/tensilelite/Tensile/TensileCreateLibrary/Logic.py index 6dd8f184f4..a0232c3923 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Logic.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Logic.py @@ -22,8 +22,11 @@ # ################################################################################ -from Tensile.Common import printExit, print1, print2 +from Tensile.Common import print1, print2, ParallelMap2 +from Tensile.LibraryIO import DataIndex +from Tensile.Parallel import ParallelMapConfig from Tensile.CustomYamlLoader import load_logic_gfx_arch, load_yaml_sequence_item +from Tensile.CodeObjectName import codeObjectFileBaseName from glob import iglob from pathlib import Path @@ -36,26 +39,25 @@ 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)) - - if not logicPath.exists(): - printExit(f"LogicPath {str(logicPath)} doesn't exist") + + assert logicPath.exists(), f"LogicPath {str(logicPath)} doesn't exist" globPattern = str(logicPath / f"**/{logicFilter}.yaml") - print1(f"# LogicFilter: {globPattern}") logicFiles = (str(logicPath / file) for file in iglob(globPattern, recursive=True)) - print1(f"# Experimental: {experimental}") - 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))] - + + print1(f"# LogicFilter: {globPattern}") + print1(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)] @@ -73,17 +75,29 @@ def numberOfBuildKernerls(logicFile): return int(str(result.stdout).count("BuildKernel")) -def schedule(logicFiles: list, numberOfTasks: int): +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(logicFiles: list, numberOfTasks: int): problemMap = {} - for logicFile in logicFiles: - codeObjectFile = load_yaml_sequence_item(logicFile, Loader, 0) - codeObjectFile = codeObjectFile["codeObjectFile"] + cofiles = ParallelMap2(getCoFileNames, ParallelMapConfig(message="Scheudling work."), logicFiles) + 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(numberOfBuildKernerls(logicFile) for logicFile in logicFiles) From 895c6a95fe4fcd78b8ca8a13b4bb7dcc80cdd65f Mon Sep 17 00:00:00 2001 From: David Dixon Date: Mon, 3 Mar 2025 22:41:34 +0000 Subject: [PATCH 20/25] Improve kernel helper construction --- .../KernelWriterActivationEnumHeader.py | 10 +- .../Tensile/KernelWriterActivationFunction.py | 11 + .../Tensile/KernelWriterActivationOnly.py | 19 ++ tensilelite/Tensile/KernelWriterBetaOnly.py | 20 ++ tensilelite/Tensile/KernelWriterConversion.py | 58 ++++++ tensilelite/Tensile/KernelWriterReduction.py | 15 ++ tensilelite/Tensile/SolutionStructs.py | 189 +++++++++++++++--- .../Tensile/TensileCreateLibrary/Run.py | 63 ++++-- tensilelite/Tensile/Toolchain/Assembly.py | 2 +- 9 files changed, 341 insertions(+), 46 deletions(-) diff --git a/tensilelite/Tensile/KernelWriterActivationEnumHeader.py b/tensilelite/Tensile/KernelWriterActivationEnumHeader.py index 84ba70c27b..b429268a0b 100644 --- a/tensilelite/Tensile/KernelWriterActivationEnumHeader.py +++ b/tensilelite/Tensile/KernelWriterActivationEnumHeader.py @@ -40,13 +40,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 95b8395e55..26afd92ca7 100644 --- a/tensilelite/Tensile/KernelWriterActivationFunction.py +++ b/tensilelite/Tensile/KernelWriterActivationFunction.py @@ -64,6 +64,17 @@ def __init__(self, state, cxxCompiler: str): 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 1d1053cb28..515e5a12d5 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 = globalParameters["IndexChars"] + # 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 = globalParameters["IndexChars"] diff --git a/tensilelite/Tensile/KernelWriterBetaOnly.py b/tensilelite/Tensile/KernelWriterBetaOnly.py index ccf7886c14..e9f04ba74e 100644 --- a/tensilelite/Tensile/KernelWriterBetaOnly.py +++ b/tensilelite/Tensile/KernelWriterBetaOnly.py @@ -280,6 +280,26 @@ def kernelBodyBetaOnly(self): return kStr + @staticmethod + def _getKernelName(solution, btype=None): + indexChars = globalParameters["IndexChars"] + # 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 = globalParameters["IndexChars"] # C dimensions diff --git a/tensilelite/Tensile/KernelWriterConversion.py b/tensilelite/Tensile/KernelWriterConversion.py index 905d5ad625..4dee87bb37 100644 --- a/tensilelite/Tensile/KernelWriterConversion.py +++ b/tensilelite/Tensile/KernelWriterConversion.py @@ -777,6 +777,64 @@ def kernelBody(self): return kStr + @staticmethod + def _getKernelName(solution, num_elements_load, btype=None): + indexChars = globalParameters["IndexChars"] + # 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 = globalParameters["IndexChars"] # C dimensions diff --git a/tensilelite/Tensile/KernelWriterReduction.py b/tensilelite/Tensile/KernelWriterReduction.py index c24472b807..bdb628ff61 100644 --- a/tensilelite/Tensile/KernelWriterReduction.py +++ b/tensilelite/Tensile/KernelWriterReduction.py @@ -55,6 +55,21 @@ def kernelBody(self): return kStr + @staticmethod + def _getKernelName(solution, btype): + # C dimensions + indexChars = globalParameters["IndexChars"] + 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/SolutionStructs.py b/tensilelite/Tensile/SolutionStructs.py index fb8ad4e7b9..b9cb4227d7 100644 --- a/tensilelite/Tensile/SolutionStructs.py +++ b/tensilelite/Tensile/SolutionStructs.py @@ -55,6 +55,92 @@ import operator import sys + +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 reject message : def reject(state, *args): @@ -1111,7 +1197,7 @@ def __init__(self, config, cxxCompiler: str): Solution.assignDerivedParameters(self._state) self._name = config["CustomKernelName"] if isCustomKernelConfig(config) else None - self.initHelperKernelObjects() + #self.initHelperKernelObjects() # these keys are copied from ProblemType to internal that may be overridden InternalKeys = ["UseSgprForGRO","VectorStore"] @@ -1130,17 +1216,38 @@ def getKernels(self): ######################################## # create Helper Kernels def initHelperKernelObjects(self): - self.initBetaOnlyKernelObjects() - self.initConversionKernelObjects() - self.initActivationEnumHeaderObjects() - self.initActivationFunctionObjects() - self.initActivationOnlyKernelObjects() - self.initReductionKernelObjects() + result = [] + + temp = self.initBetaOnlyKernelObjects() + if temp: + result.extend(temp) + + temp = self.initConversionKernelObjects() + if temp: + result.extend(temp) + + temp = self.initActivationEnumHeaderObjects() + if temp: + result.extend(temp) + + temp = self.initActivationFunctionObjects() + 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"]: @@ -1151,20 +1258,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): - self.conversionKernelObjects = [] + conversionKernelObjects = [] load_vector_width = [1, 2] if self["ProblemType"]["DataType"].isDouble() else [1, 2, 4] genPGRPostKernels = True gsuList = [internalParameters["GlobalSplitUPGR"]] @@ -1190,7 +1298,7 @@ def initConversionKernelObjects(self): state["UnrollOnly"] = unrollOnly state["_GlobalAccumulation"] = self["_GlobalAccumulation"] state["ActivationFused"] = self["ActivationFused"] - self.conversionKernelObjects.append(KernelWriterConversion(state, vw)) + conversionKernelObjects.append(KernelWriterConversion(state, vw)) for btype in typeList: state = {} state["ProblemType"] = deepcopy(self["ProblemType"]) @@ -1203,7 +1311,7 @@ def initConversionKernelObjects(self): state["UnrollOnly"] = unrollOnly state["_GlobalAccumulation"] = self["_GlobalAccumulation"] state["ActivationFused"] = self["ActivationFused"] - self.conversionKernelObjects.append(KernelWriterConversion(state, vw)) + conversionKernelObjects.append(KernelWriterConversion(state, vw)) else: state = {} state["ProblemType"] = deepcopy(self["ProblemType"]) @@ -1214,29 +1322,32 @@ def initConversionKernelObjects(self): state["UnrollOnly"] = unrollOnly state["_GlobalAccumulation"] = self["_GlobalAccumulation"] state["ActivationFused"] = self["ActivationFused"] - self.conversionKernelObjects.append(KernelWriterConversion(state, vw)) + conversionKernelObjects.append(KernelWriterConversion(state, vw)) + 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): - 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"])} - self.activationFunctionObjects.append(KernelWriterActivationFunction(state, self.cxxCompiler)) + activationFunctionObjects.append(KernelWriterActivationFunction(state, self.cxxCompiler)) + return activationFunctionObjects def initActivationOnlyKernelObjects(self): - self.activationOnlyKernelObjects = [] + activationOnlyKernelObjects = [] if (self["ActivationFused"] == False) and (self["ProblemType"]["ActivationType"] != 'none') : state = {} state["ProblemType"] = deepcopy(self["ProblemType"]) @@ -1246,10 +1357,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 = {} @@ -1257,14 +1369,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 ######################################## @@ -4446,3 +4581,5 @@ def __ne__(self, other): if result is NotImplemented: return result return not result + + diff --git a/tensilelite/Tensile/TensileCreateLibrary/Run.py b/tensilelite/Tensile/TensileCreateLibrary/Run.py index af2198ac25..6a3bffcaa1 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Run.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Run.py @@ -29,6 +29,7 @@ from pathlib import Path from timeit import default_timer as timer from typing import Dict, NamedTuple, List, Optional +import yaml from Tensile.Toolchain.Assembly import AssemblyToolchain, buildAssemblyCodeObjectFiles from Tensile.Toolchain.Source import SourceToolchain, buildSourceCodeObjectFile @@ -39,7 +40,7 @@ from Tensile.Parallel import ParallelMapConfig from Tensile.Utilities.RequiredParameters import getRequiredParametersMin from Tensile.KernelWriterAssembly import KernelWriterAssembly -from Tensile.SolutionStructs import Solution +from Tensile.SolutionStructs import kernelObjectNames from Tensile.Utilities.Profile import profile from Tensile.SolutionLibrary import MasterSolutionLibrary @@ -113,15 +114,24 @@ def buildAssemblyKernels(asmPath: Path, asmToolchain: AssemblyToolchain, kernelW asmToolchain.assemble(str(p), str(p.with_suffix(".o")), getGfxName(isa), wavefrontsize) if removeTemporaries: p.unlink() - + return uniqueAsmKernels, solnLibs[1] def generateKernelHelperObjects(solutions): khos = [] + visited = set() for solution in solutions: - khos.extend(solution.getHelperKernelObjects()) - return khos #should we deduplicate here? + 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()) + print(len(khos)) + return list(dict.fromkeys(khos)) #def generateMatchTable(masterSolutionLibraries): @@ -204,6 +214,28 @@ def extractBuildResults(result): return flattenedList, masterLibs + +from Tensile.TensileInstructions import DataType + +def DataType_constructor(loader, node): + value = loader.construct_mapping(node) + return DataType(**value) + +def DataType_representer(dumper, data): + return dumper.represent_mapping('!DataType', {'properties': data.properties, 'value': data.value}) + +def ActivationType_constructor(loader, node): + value = loader.construct_mapping(node) + return DataType(**value) + +def DataType_representer(dumper, data): + return dumper.represent_mapping('!DataType', {'properties': data.properties, 'value': data.value}) + +yaml.add_constructor('!DataType', DataType_constructor) +yaml.add_representer(DataType, DataType_representer) + + + ################################################################################ # Tensile Create Library ################################################################################ @@ -260,7 +292,7 @@ def run(): def buildCoAndHelpers(input): uniqueAsmKernels, libraries = input unaryBuildCOFile(uniqueAsmKernels) - return generateKernelHelperObjects(uniqueAsmKernels), libraries + return uniqueAsmKernels, libraries compose = lambda *F: functools.reduce(lambda f, g: lambda x: f(g(x)), F) if arguments["LazyLibraryLoading"]: @@ -272,8 +304,10 @@ def buildCoAndHelpers(input): parMap = functools.partial(ParallelMap2, assembly, ParallelMapConfig(message="Building Objects")) phase1 = compose(buildCoAndHelpers, extractBuildResults, parMap) - kho, masterLibs = phase1(logicFiles) - + #khos, masterLibs = phase1(logicFiles) + kernels, masterLibs = phase1(logicFiles) + khos = generateKernelHelperObjects(kernels) + print(len(khos)) # Phase3: Build Master Solution Library generateParentLibrary(arguments["LibraryFormat"], libraryPath, masterLibs, arguments["LazyLibraryLoading"]) del masterLibs @@ -281,16 +315,11 @@ def buildCoAndHelpers(input): # Phase2: 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"), list(dict.fromkeys(kho))) - del kho - + srcFiles = ParallelMap2(unaryWriteHelpers, ParallelMapConfig(message="Generating Kernels code", return_as="list"), generateKernelHelperObjects(kernels)) # kernelsLib = str(srcCodeObjectPath / "Kernels.so") srcToolchain.compile(srcFiles, kernelsLib, str(outputPath), archs) - del srcFiles - buildSourceCodeObjectFile(srcToolchain, libraryPath, kernelsLib) - if not arguments["KeepBuildTmp"]: if buildTmp.exists() and buildTmp.is_dir(): shutil.rmtree(buildTmp) @@ -302,11 +331,11 @@ def buildCoAndHelpers(input): print1("") stop = timer() - print1(f"Total time (s): {(stop-start):3.2f}") - #print1(f"Total kernels: {numKernels}") + numKernels = len(kernels) + print1(f"Total kernels: {numKernels}") #print1(f"Total kernels processed: {numUniqueKernels}") #print1(f"Duplicate kernels removed: {numDuplicateKernels}") - #print1(f"Kernels processed per second: {(numUniqueKernels/(stop-start)):3.2f}") + print1(f"Kernels processed per second: {(numKernels/(stop-start)):3.2f}") #print1(f"Total solutions processed: {numSoln}") - #print1(f"Duplicate solutions: {numDuplicateSoln}") + #print1(f"Duplicate solutions: {numDuplicateSoln}") \ No newline at end of file diff --git a/tensilelite/Tensile/Toolchain/Assembly.py b/tensilelite/Tensile/Toolchain/Assembly.py index 961c004331..7bf2b18240 100644 --- a/tensilelite/Tensile/Toolchain/Assembly.py +++ b/tensilelite/Tensile/Toolchain/Assembly.py @@ -206,7 +206,7 @@ def buildAssemblyCodeObjectFiles(toolchain: AssemblyToolchain, srcDir, destDir, # shouldn't need a set here the fact that we do implies we have duplicates #objFiles = _batchObjectFiles(set(objFiles), coFileRaw) objFiles = _batchObjectFiles(objFiles, coFileRaw) - toolchain.link(objFiles, str(coFileRaw)) + toolchain.link(set(objFiles), str(coFileRaw)) coFile = destDir / coFileRaw.name.replace(extCoRaw, extCo) if compress: toolchain.compress(str(coFileRaw), str(coFile), gfx) From dd011ab2960c8e5d83655b3e5f12a4f3f0bc1cc9 Mon Sep 17 00:00:00 2001 From: David Dixon Date: Sat, 15 Mar 2025 20:47:45 +0000 Subject: [PATCH 21/25] Corrections --- tensilelite/Tensile/LibraryIO.py | 13 +------------ tensilelite/Tensile/TensileCreateLibrary/Logic.py | 4 ++-- 2 files changed, 3 insertions(+), 14 deletions(-) diff --git a/tensilelite/Tensile/LibraryIO.py b/tensilelite/Tensile/LibraryIO.py index 813ab3285b..771dd6d7bf 100644 --- a/tensilelite/Tensile/LibraryIO.py +++ b/tensilelite/Tensile/LibraryIO.py @@ -271,17 +271,6 @@ def parseLibraryLogicData(data, srcFile, cxxCompiler, archs=None): if isinstance(data, List): data = parseLibraryLogicList(data, srcFile) - #is_arch_valid = lambda cArch, tArch : (cArch == tArch or cArch == "all") - #if not (archs is None) and "ArchitectureName" in data: - # printExit("Shouldn't be here") - # if isinstance(archs, List): - # if len(archs) > 0 and not archs[0] == "all": - # if not (any(is_arch_valid(arch.split(":")[0], data["ArchitectureName"]) for arch in archs)): - # return LibraryLogic("", "", None, [], [], None, srcFile) - # elif isinstance(archs, str): - # if not is_arch_valid(archs.split(":")[0], data["ArchitectureName"]): - # return LibraryLogic("", "", None, [], [], None, srcFile) - if "CUCount" not in data: data["CUCount"] = None @@ -350,7 +339,7 @@ def parseLibraryLogicList(data, srcFile="?"): rv["ProblemType"] = data[DataIndex.PROBLEM_TYPE] rv["Solutions"] = data[DataIndex.SOLUTIONS] - if type(data[3]) is dict: + if type(data[DataIndex.DEVICE_PROPERTIES]) is dict: rv["ArchitectureName"] = data[DataIndex.DEVICE_PROPERTIES]["Architecture"] rv["CUCount"] = data[DataIndex.DEVICE_PROPERTIES]["CUCount"] else: diff --git a/tensilelite/Tensile/TensileCreateLibrary/Logic.py b/tensilelite/Tensile/TensileCreateLibrary/Logic.py index 4d21f06e14..cfb41ff1b0 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Logic.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Logic.py @@ -70,8 +70,8 @@ def distribute(lst, n): def numberOfBuildKernerls(logicFile): - result = run(['/bin/grep', "SolutionIndex", logicFile], stderr=PIPE, stdout=PIPE, check=False) - return int(str(result.stdout).count("SolutionIndex")) + result = run(['/bin/grep', "BuildKernel", logicFile], stderr=PIPE, stdout=PIPE, check=False) + return int(str(result.stdout).count("BuildKernel")) def getCoFileNames(logicFile): From d57a950ccd6f16ca9e36a04906dacad1d6f8c9ab Mon Sep 17 00:00:00 2001 From: David Dixon Date: Sat, 15 Mar 2025 22:35:15 +0000 Subject: [PATCH 22/25] Need to have code object name in logic files --- tensilelite/Tensile/CustomYamlLoader.py | 2 +- tensilelite/Tensile/LibraryIO.py | 30 ++++++++++--------- .../Tensile/TensileCreateLibrary/Logic.py | 24 ++++++++------- 3 files changed, 30 insertions(+), 26 deletions(-) diff --git a/tensilelite/Tensile/CustomYamlLoader.py b/tensilelite/Tensile/CustomYamlLoader.py index bab8c68750..b82e89132d 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 = 3 # DataIndex.DEVICE_PROPERTIES arch = load_yaml_sequence_item(yaml_path, loader_type, GFX_ARCH_IDX) if isinstance(arch, dict): diff --git a/tensilelite/Tensile/LibraryIO.py b/tensilelite/Tensile/LibraryIO.py index 771dd6d7bf..23d92fb0ba 100644 --- a/tensilelite/Tensile/LibraryIO.py +++ b/tensilelite/Tensile/LibraryIO.py @@ -235,17 +235,19 @@ def parseSolutionsData(data, srcFile, cxxCompiler): 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=0 + MINIMUM_REQUIRED_VERSION=1 + SCHEDULE_NAME=2 + DEVICE_PROPERTIES=3 + DEVICE_NAMES=4 + PROBLEM_TYPE=5 + SOLUTIONS=6 + INDEX_ORDER=7 + EXACT_LOGIC=8 + RANGE_LOGIC=9 + PERF_METRIC=11 + LIBRARY_TYPE=12 + def __index__(self): return self.value @@ -356,7 +358,7 @@ def parseLibraryLogicList(data, srcFile="?"): # library logic fields libraryType = None - if len(data) > 11 and data[DataIndex.LIBRARY_TYPE]: + if len(data) > 12 and data[DataIndex.LIBRARY_TYPE]: libraryType = data[DataIndex.LIBRARY_TYPE] else: printExit("Library logic file {} is missing required field matching property." \ @@ -391,8 +393,8 @@ def rawLibraryLogic(data): 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/TensileCreateLibrary/Logic.py b/tensilelite/Tensile/TensileCreateLibrary/Logic.py index cfb41ff1b0..bf289acbb9 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Logic.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Logic.py @@ -76,17 +76,19 @@ def numberOfBuildKernerls(logicFile): 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 + #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 + coBasename = load_yaml_sequence_item(logicFile, Loader, DataIndex.CODE_OBJECT_NAME) + return coBasename["codeObjectFile"], logicFile def schedule(logicFiles: list, numberOfTasks: int): From 19a54f296d8dde8367b3fa3bb429f3cfaa9d4236 Mon Sep 17 00:00:00 2001 From: David Dixon Date: Wed, 19 Mar 2025 16:24:27 +0000 Subject: [PATCH 23/25] Working 942 build --- tensilelite/Tensile/CustomYamlLoader.py | 2 +- tensilelite/Tensile/LibraryIO.py | 29 +++++++-------- .../Tensile/TensileCreateLibrary/IO.py | 4 +-- .../Tensile/TensileCreateLibrary/Logic.py | 36 +++++++++---------- .../Tensile/TensileCreateLibrary/Run.py | 16 +++++++-- 5 files changed, 47 insertions(+), 40 deletions(-) diff --git a/tensilelite/Tensile/CustomYamlLoader.py b/tensilelite/Tensile/CustomYamlLoader.py index b82e89132d..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 = 3 # DataIndex.DEVICE_PROPERTIES + 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/LibraryIO.py b/tensilelite/Tensile/LibraryIO.py index 23d92fb0ba..eac7e08eea 100644 --- a/tensilelite/Tensile/LibraryIO.py +++ b/tensilelite/Tensile/LibraryIO.py @@ -233,26 +233,23 @@ def parseSolutionsData(data, srcFile, cxxCompiler): problemSizes = ProblemSizes(problemType, problemSizesConfig) return (problemSizes, solutions) - class DataIndex(IntEnum): - CODE_OBJECT_NAME=0 - MINIMUM_REQUIRED_VERSION=1 - SCHEDULE_NAME=2 - DEVICE_PROPERTIES=3 - DEVICE_NAMES=4 - PROBLEM_TYPE=5 - SOLUTIONS=6 - INDEX_ORDER=7 - EXACT_LOGIC=8 - RANGE_LOGIC=9 - PERF_METRIC=11 - LIBRARY_TYPE=12 - + 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()""" schedule: str @@ -358,7 +355,7 @@ def parseLibraryLogicList(data, srcFile="?"): # library logic fields libraryType = None - if len(data) > 12 and data[DataIndex.LIBRARY_TYPE]: + 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." \ diff --git a/tensilelite/Tensile/TensileCreateLibrary/IO.py b/tensilelite/Tensile/TensileCreateLibrary/IO.py index 1328c68e93..a3df1d4721 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/IO.py +++ b/tensilelite/Tensile/TensileCreateLibrary/IO.py @@ -36,7 +36,7 @@ def generateSolutionsAndLibraries(archs, cxxCompiler, logicFiles): solutions = [] libraries = [] for logicFileGroup in logicFiles: - for logicFile in logicFileGroup[1]: # should be logicFileGroup[1] + for logicFile in logicFileGroup[1]: libraryLogic = parseLibraryLogicFile(logicFile, cxxCompiler, archs) solutions.extend(libraryLogic.solutions) libraries.append((libraryLogic.architecture, libraryLogic.library)) @@ -45,7 +45,7 @@ def generateSolutionsAndLibraries(archs, cxxCompiler, logicFiles): def writeAssembly(asmPath: Path, result: tuple): if result[0]: - printExit(f"Failed to build kernel {result[3]} because it has error code {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] diff --git a/tensilelite/Tensile/TensileCreateLibrary/Logic.py b/tensilelite/Tensile/TensileCreateLibrary/Logic.py index bf289acbb9..60e385154d 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Logic.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Logic.py @@ -33,7 +33,6 @@ 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): @@ -76,24 +75,25 @@ def numberOfBuildKernerls(logicFile): 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 - coBasename = load_yaml_sequence_item(logicFile, Loader, DataIndex.CODE_OBJECT_NAME) - return coBasename["codeObjectFile"], logicFile - - -def schedule(logicFiles: list, numberOfTasks: int): + 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 + #coBasename = load_yaml_sequence_item(logicFile, Loader, DataIndex.CODE_OBJECT_NAME) + #coBasename = load_yaml_sequence_item(logicFile, Loader, 0) + #return coBasename["codeObjectFile"], logicFile + + +def schedule(logicFiles: list, numberOfTasks: int, procs: int): problemMap = {} - cofiles = ParallelMap2(getCoFileNames, ParallelMapConfig(message="Scheudling work."), logicFiles) + cofiles = ParallelMap2(getCoFileNames, ParallelMapConfig(message="Scheudling work.", procs=procs), logicFiles) for codeObjectFile, logicFile in cofiles: if codeObjectFile in problemMap: problemMap[codeObjectFile].append(logicFile) diff --git a/tensilelite/Tensile/TensileCreateLibrary/Run.py b/tensilelite/Tensile/TensileCreateLibrary/Run.py index 17d761c3e1..ad580db6a3 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Run.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Run.py @@ -115,9 +115,10 @@ def buildAssemblyKernels(asmPath: Path, asmToolchain: AssemblyToolchain, kernelW uniqueAsmKernels = [k for k in kernels if "BuildKernel" in k] pksResults = [_processKernelSource(kernelWriterAssembly, ti, k) for k in uniqueAsmKernels] for p, isa, wavefrontsize in set([writeAssembly(asmPath, k) for k in pksResults]): + #Cijk_Ailk_Bljk_BSS_BH_UserArgs_MT128x256x64_MI16Vcp76y7W1sPIjpGVm_aV-u3wID5BM2LfzmbkyHs7rjQ=.s asmToolchain.assemble(str(p), str(p.with_suffix(".o")), isaToGfx(isa), wavefrontsize) - if removeTemporaries: - p.unlink() + #if removeTemporaries: + # p.unlink() return uniqueAsmKernels, solnLibs[1] @@ -259,11 +260,20 @@ def run(): # List logic files unsortedLogic = logicFileList(archs, Path(arguments["LogicPath"]), arguments["LogicFilter"], arguments["Experimental"]) - logicFiles = list(filter(lambda x: x != [], schedule(unsortedLogic, 2*arguments["CpuThreads"]))) + logicFiles = list(filter(lambda x: x != [], schedule(unsortedLogic, 2*arguments["CpuThreads"], arguments["CpuThreads"]))) # Phase1: Build assembly and master solution libraries writerAsm = KernelWriterAssembly(getRequiredParametersMin(), getRequiredParametersMin(), asmToolchain.assembler, asmToolchain.assemblerVersion) unaryGenSolutions = functools.partial(generateSolutionsAndLibraries, archs, cxxCompiler) + #result = ParallelMap2(unaryGenSolutions, ParallelMapConfig(message="blah"), logicFiles) + #d = {} + #l_path = Path(arguments["LogicPath"]) + #for s, _, l in result: + # logic_file = Path(l).relative_to(l_path) + # d[str(logic_file)] = s[0]["codeObjectFile"] + #import pprint + #pprint.pprint(d, indent=2) + #printExit("done") unaryProcessMsl = functools.partial(processMsl, libraryPath, arguments["LibraryFormat"]) unaryBuildAsmKernels = functools.partial(buildAssemblyKernels, assemblyPath, asmToolchain, writerAsm, TensileInstructions(), not arguments["KeepBuildTmp"]) unaryBuildCOFile = functools.partial(buildAssemblyCodeObjectFiles, asmToolchain, writerAsm, libraryPath, assemblyPath, arguments["UseCompression"]) From c0508cf812f27f138e67608cf2ba7a6af0395f46 Mon Sep 17 00:00:00 2001 From: David Dixon Date: Wed, 9 Apr 2025 18:19:55 +0000 Subject: [PATCH 24/25] Argument ordering --- tensilelite/Tensile/TensileCreateLibrary/Run.py | 16 +++++++++------- tensilelite/Tensile/Toolchain/Assembly.py | 2 +- 2 files changed, 10 insertions(+), 8 deletions(-) diff --git a/tensilelite/Tensile/TensileCreateLibrary/Run.py b/tensilelite/Tensile/TensileCreateLibrary/Run.py index 45f5076876..6b8b2f72e7 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Run.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Run.py @@ -286,9 +286,9 @@ def run(): offloadBundler, ls, extract, + arguments["CpuThreads"], arguments["AsanBuild"], arguments["BuildIdKind"], - arguments["CpuThreads"], save_temps=False ) @@ -300,9 +300,10 @@ def run(): logicFiles = list(filter(lambda x: x != [], schedule(unsortedLogic, 2*arguments["CpuThreads"], arguments["CpuThreads"]))) # Phase1: Build assembly and master solution libraries + kernelMinNaming = getRequiredParametersMin() writerAsm = KernelWriterAssembly( - getRequiredParametersMin(), - getRequiredParametersMin(), + kernelMinNaming, + kernelMinNaming, asmToolchain.assembler, DebugConfig(), ) @@ -328,7 +329,7 @@ def run(): asmToolchain.assembler, writerAsm, rocisa.rocIsa.getInstance().getData(), - getRequiredParametersMin(), + kernelMinNaming, not arguments["KeepBuildTmp"]) unaryBuildCOFile = functools.partial(buildAssemblyCodeObjectFiles, asmToolchain.linker, @@ -336,7 +337,8 @@ def run(): globalParameters["ROCmLdPath"], libraryPath, assemblyPath, - arguments["UseCompression"]) + arguments["UseCompression"], + kernelMinNaming) def buildCoAndHelpers(input): uniqueAsmKernels, libraries = input unaryBuildCOFile(uniqueAsmKernels) @@ -365,8 +367,8 @@ def buildCoAndHelpers(input): 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) + srcToolchain.compiler(srcFiles, kernelsLib, str(outputPath), archs) + buildSourceCodeObjectFile(srcToolchain, libraryPath, kernelsLib) if not arguments["KeepBuildTmp"]: if buildTmp.exists() and buildTmp.is_dir(): diff --git a/tensilelite/Tensile/Toolchain/Assembly.py b/tensilelite/Tensile/Toolchain/Assembly.py index 7a0a4f4249..9d86e85b4a 100644 --- a/tensilelite/Tensile/Toolchain/Assembly.py +++ b/tensilelite/Tensile/Toolchain/Assembly.py @@ -119,7 +119,7 @@ def buildAssemblyCodeObjectFiles( for kernel in archKernels: coName = kernel.get("codeObjectFile", None) if coName: - coFileMap[asmDir / (coName + extCoRaw)].append(str(asmDir / (basename(kernel) + extObj))) + coFileMap[asmDir / (coName + extCoRaw)].append(str(asmDir / (baseName(kernel) + extObj))) for coFileRaw, objFiles in coFileMap.items(): objFiles = _batchObjectFiles(ldPath, objFiles, coFileRaw) From 86cd83239aa060fb0b4cacf09857db74d69f2784 Mon Sep 17 00:00:00 2001 From: David Dixon Date: Sun, 13 Apr 2025 20:07:43 +0000 Subject: [PATCH 25/25] Use file size to schedule logic file processing --- .../Tensile/TensileCreateLibrary/Logic.py | 23 +++++++------------ .../Tensile/TensileCreateLibrary/Run.py | 22 +++++++++++------- 2 files changed, 22 insertions(+), 23 deletions(-) diff --git a/tensilelite/Tensile/TensileCreateLibrary/Logic.py b/tensilelite/Tensile/TensileCreateLibrary/Logic.py index 7b4f7883f1..5a6ce63dd7 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Logic.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Logic.py @@ -22,12 +22,13 @@ # ################################################################################ -from Tensile.Common import print1, print2, ParallelMap2, ParallelMapConfig +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 @@ -47,8 +48,8 @@ def validLogicFile(p: Path): logicFiles = [file for file in logicFiles if validLogicFile(Path(file))] - print1(f"# LogicFilter: {globPattern}") - print1(f"# Experimental: {experimental}") + print(f"# LogicFilter: {globPattern}") + print(f"# Experimental: {experimental}") print2(f"# LibraryLogicFiles: {len(logicFiles)}") for logicFile in logicFiles: print2("# %s" % logicFile) @@ -68,11 +69,6 @@ def distribute(lst, n): return sorted(list_of_lists, key=lambda x: sum(first for first, _ in x), reverse=True) -def numberOfBuildKernerls(logicFile): - result = run(['/bin/grep', "BuildKernel", logicFile], stderr=PIPE, stdout=PIPE, check=False) - return int(str(result.stdout).count("BuildKernel")) - - def getCoFileNames(logicFile): from yaml import Loader data = {} @@ -86,22 +82,19 @@ def getCoFileNames(logicFile): data["CUCount"] = None data["PerfMetric"] = load_yaml_sequence_item(logicFile, Loader, DataIndex.PERF_METRIC.value) return codeObjectFileBaseName(data), logicFile - #coBasename = load_yaml_sequence_item(logicFile, Loader, DataIndex.CODE_OBJECT_NAME) - #coBasename = load_yaml_sequence_item(logicFile, Loader, 0) - #return coBasename["codeObjectFile"], logicFile -def schedule(logicFiles: list, numberOfTasks: int, procs: int): +def schedule(cofiles: list, numberOfTasks: int, procs: int): problemMap = {} - cofiles = ParallelMap2(getCoFileNames, ParallelMapConfig(message="Scheduling work. ", procs=procs), logicFiles) 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(numberOfBuildKernerls(logicFile) for logicFile in logicFiles) + count = sum(getsize(logicFile) for logicFile in logicFiles) result.append((count, logicFiles)) - return distribute(result, numberOfTasks) # need to convert list of list of tuples to list of list of strings + 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/Run.py b/tensilelite/Tensile/TensileCreateLibrary/Run.py index 0c3a9f30d1..73abdc950c 100644 --- a/tensilelite/Tensile/TensileCreateLibrary/Run.py +++ b/tensilelite/Tensile/TensileCreateLibrary/Run.py @@ -49,7 +49,7 @@ from .IO import generateSolutionsAndLibraries, genLazyMasterSolutionLibrary, \ generateParentLibrary, writeAssembly, writeHelper -from .Logic import logicFileList, schedule +from .Logic import logicFileList, schedule, getCoFileNames from .ParseArguments import parseArguments @@ -120,7 +120,6 @@ def buildAssemblyKernels(asmPath: Path, removeTemporaries: bool, solnLibs: List[tuple]): kernels = [s.getKernels()[0] for s in solnLibs[0]] - #uniqueAsmKernels = [k for k in kernels if "BuildKernel" in k] visited = set() duplicates = 0 @@ -129,16 +128,16 @@ def buildAssemblyKernels(asmPath: Path, 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() + if removeTemporaries: + p.unlink() return uniqueAsmKernels, solnLibs[1] @@ -293,8 +292,15 @@ def run(): print(asmToolchain.bundler) # List logic files - unsortedLogic = logicFileList(archs, Path(arguments["LogicPath"]), arguments["LogicFilter"], arguments["Experimental"]) - logicFiles = list(filter(lambda x: x != [], schedule(unsortedLogic, 2*arguments["CpuThreads"], arguments["CpuThreads"]))) + 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()) @@ -366,7 +372,7 @@ def buildCoAndHelpers(input): generateKernelHelperObjects(kernels, isaInfoMap)) kernelsLib = str(srcCodeObjectPath / "Kernels.so") srcToolchain.compiler(srcFiles, kernelsLib, str(outputPath), archs) - buildSourceCodeObjectFile(srcToolchain, libraryPath, kernelsLib) + #buildSourceCodeObjectFile(srcToolchain, libraryPath, kernelsLib) if not arguments["KeepBuildTmp"]: if buildTmp.exists() and buildTmp.is_dir():