Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Tcl cleanup #1444

Draft
wants to merge 8 commits into
base: develop
Choose a base branch
from
Draft
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
13 changes: 1 addition & 12 deletions install.sh
Original file line number Diff line number Diff line change
Expand Up @@ -395,7 +395,6 @@ tensile_logic=
tensile_cov=
tensile_threads=$(nproc)
tensile_fork=
tensile_merge_files=
tensile_tag=
tensile_test_local_path=
tensile_version=
Expand Down Expand Up @@ -423,7 +422,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,merge-files,no-merge-files,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_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
Expand Down Expand Up @@ -503,12 +502,6 @@ while true; do
-n|--no_tensile|--no-tensile)
build_tensile=false
shift ;;
--merge-files)
tensile_merge_files=true
shift ;;
-no-merge-files)
tensile_merge_files=false
shift 2;;
-u|--use-custom-version)
tensile_version=${2}
shift 2;;
Expand Down Expand Up @@ -769,10 +762,6 @@ pushd .
fi
fi

if [[ "${tensile_merge_files}" == false ]]; then
tensile_opt="${tensile_opt} -DTensile_MERGE_FILES=OFF"
fi

if [[ "${tensile_msgpack_backend}" == true ]]; then
tensile_opt="${tensile_opt} -DTensile_LIBRARY_FORMAT=msgpack"
else
Expand Down
3 changes: 0 additions & 3 deletions tensilelite/Tensile/BenchmarkProblems.py
Original file line number Diff line number Diff line change
Expand Up @@ -112,9 +112,6 @@ def generateCustomKernelSolutions(problemType, customKernels, internalSupportPar
def writeBenchmarkFiles(stepBaseDir, solutions, problemSizes, \
biasTypeArgs, factorDimArgs, activationArgs, icacheFlushArgs, stepName, solutionSummationSizes, cxxCompiler, assembler, offloadBundler):
"""Write all the files needed for a given benchmarking step"""
if not globalParameters["MergeFiles"]:
ensurePath(os.path.join(globalParameters["WorkingPath"], "Solutions"))
ensurePath(os.path.join(globalParameters["WorkingPath"], "Kernels"))

copyStaticFiles()

Expand Down
3 changes: 1 addition & 2 deletions tensilelite/Tensile/BuildCommands/AssemblyCommands.py
Original file line number Diff line number Diff line change
Expand Up @@ -82,8 +82,7 @@ def buildAssemblyCodeObjectFiles(kernels, kernelWriterAssembly, outputPath, asse
continue

gfx = getGfxName(arch)

if globalParameters["MergeFiles"] or globalParameters["NumMergedFiles"] > 1 or globalParameters["LazyLibraryLoading"]:
if globalParameters["LazyLibraryLoading"]:
objectFiles = [str(asmDir / (kernelWriterAssembly.getKernelFileBase(k) + extObj)) for k in archKernels if 'codeObjectFile' not in k]

coFileMap = collections.defaultdict(list)
Expand Down
14 changes: 1 addition & 13 deletions tensilelite/Tensile/ClientWriter.py
Original file line number Diff line number Diff line change
Expand Up @@ -215,12 +215,6 @@ def getBuildClientLibraryScript(buildPath, libraryLogicPath, cxxCompiler):

callCreateLibraryCmd = globalParameters["ScriptPath"] + "/bin/TensileCreateLibrary"


if globalParameters["MergeFiles"]:
callCreateLibraryCmd += " --merge-files"
else:
callCreateLibraryCmd += " --no-merge-files"

if globalParameters["ShortNames"]:
callCreateLibraryCmd += " --short-file-names"
else:
Expand Down Expand Up @@ -727,13 +721,7 @@ def writeClientParameters(forBenchmark, solutions, problemSizes, stepName, \
"""

if forBenchmark:
if globalParameters["MergeFiles"]:
h += "#include \"Solutions.h\"\n"
else:
for solution in solutions:
solutionName = solutionWriter.getSolutionName(solution)
h += "#include \"" + solutionName + ".h\"\n"
h += "#include \"Solutions.h\"\n"
h += "#include \"Solutions.h\"\n"
h += "#include \"ReferenceCPU.h\"\n"
h += "\n"
else:
Expand Down
7 changes: 0 additions & 7 deletions tensilelite/Tensile/Common.py
Original file line number Diff line number Diff line change
Expand Up @@ -218,8 +218,6 @@
globalParameters["DeviceLDS"] = 65536 # LDS bytes per CU, for computing occupancy
globalParameters["MaxLDS"] = 65536 # max LDS a kernel should attempt to use
globalParameters["ShortNames"] = False # on windows kernel names can get too long; =True will convert solution/kernel names to serial ids
globalParameters["MergeFiles"] = True # F=store every solution and kernel in separate file; T=store all solutions in single file
globalParameters["NumMergedFiles"] = 1 # The number of files that kernels should be split between when merging

globalParameters["MaxFileName"] = 64 # If a file name would be longer than this, shorten it with a hash.
globalParameters["SupportedISA"] = [(8,0,3), (9,0,0), (9,0,6), (9,0,8), (9,0,10), (9,4,0), (9,4,1), (9,4,2), (10,1,0), (10,1,1), (10,1,2), (10,3,0), (11,0,0), (11,0,1), (11,0,2), (12,0,0), (12,0,1)] # assembly kernels writer supports these architectures
Expand Down Expand Up @@ -1727,11 +1725,6 @@ def assignGlobalParameters(config, cxxCompiler=None):

validParameters["ISA"] = [(0,0,0), *globalParameters["SupportedISA"]]

if "MergeFiles" in config and "NumMergedFiles" in config:
if not config["MergeFiles"] and config["NumMergedFiles"] > 1:
config["NumMergedFiles"] = 1
printWarning("--num-merged-files and --no-merge-files specified, ignoring --num-merged-files")

# For ubuntu platforms, call dpkg to grep the version of hip-clang. This check is platform specific, and in the future
# additional support for yum, dnf zypper may need to be added. On these other platforms, the default version of
# '0.0.0' will persist
Expand Down
4 changes: 2 additions & 2 deletions tensilelite/Tensile/GenerateSummations.py
Original file line number Diff line number Diff line change
Expand Up @@ -63,8 +63,8 @@ def createLibraryForBenchmark(logicPath, libraryPath, currentPath):

pythonExePath = os.path.join(os.path.dirname(os.path.realpath(__file__)), "bin", "TensileCreateLibrary")
args = [pythonExePath, \
"--merge-files", "--new-client-only", "--no-short-file-names", "--no-library-print-debug", \
"--architecture=all", "--code-object-version=default", "--library-format=yaml", \
"--new-client-only", "--no-short-file-names", "--no-library-print-debug", \
"--architecture=all", "--code-object-version=default", "--cxx-compiler="+globalParameters["CxxCompiler"], "--library-format=yaml", \
logicPath, libraryPath, "HIP"]

try:
Expand Down
2 changes: 0 additions & 2 deletions tensilelite/Tensile/KernelWriter.py
Original file line number Diff line number Diff line change
Expand Up @@ -5209,8 +5209,6 @@ def getSourceFileString(self, kernel):
def getHeaderFileString(self, kernel):
kernelName = self.getKernelName(kernel)
fileString = "" # CHeader
if not globalParameters["MergeFiles"] or globalParameters["NumMergedFiles"] > 1:
fileString += "#pragma once\n\n"
if not globalParameters["CodeFromFiles"]:
fileString += "extern const unsigned char %s_coba[]; // code object byte array\n" % kernelName

Expand Down
4 changes: 0 additions & 4 deletions tensilelite/Tensile/KernelWriterActivationEnumHeader.py
Original file line number Diff line number Diff line change
Expand Up @@ -55,10 +55,6 @@ def getSourceFileString(self):

def getHeaderFileString(self):
fileString = "" # CHeader
if not globalParameters["MergeFiles"]:
fileString += CHeader
fileString += "#pragma once\n\n"

activationCDataType = self.state["ProblemType"]["ActivationComputeDataType"]
supportedBy = ActivationType.SupportedBy.ALL if self.state["ProblemType"]["ActivationType"] == 'all' else ActivationType.SupportedBy.HIPBLASLT
enumName = "%sActivationType_%s"%(self.actGradientPrefix, activationCDataType.toChar())
Expand Down
6 changes: 0 additions & 6 deletions tensilelite/Tensile/KernelWriterActivationFunction.py
Original file line number Diff line number Diff line change
Expand Up @@ -149,12 +149,6 @@ def getHeaderFileString(self):
activation = ActivationInline(activationCDataType, not self.state["ProblemType"]["ActivationNoGuard"])

fileString = "" # CHeader
if not globalParameters["MergeFiles"]:
fileString += CHeader
fileString += "#pragma once\n\n"
fileString += "#include \"Tensile%sActivationEnum_%s.h\"\n"%(self.actGradientPrefix, activationCDataType.toChar())
fileString += "\n"

fileString += "#pragma clang diagnostic push\n"
fileString += "#pragma clang diagnostic ignored \"-Winline-asm\"\n"
fileString += self.functionSignature()
Expand Down
21 changes: 0 additions & 21 deletions tensilelite/Tensile/KernelWriterActivationOnly.py
Original file line number Diff line number Diff line change
Expand Up @@ -229,34 +229,13 @@ def getKernelName(self):

def getSourceFileString(self):
fileString = ""
if not globalParameters["MergeFiles"]:
fileString += "\n"
fileString += "#include \"%s.h\"\n" % self.kernelName
fileString += "\n"

fileString += self.functionSignature()
fileString += self.kernelBody()

return (0, fileString)

def getHeaderFileString(self):
fileString = "" # CHeader
if not globalParameters["MergeFiles"]:
fileString += CHeader
fileString += "#pragma once\n\n"
fileString += "\n"
fileString += "#include <KernelHeader.h>\n\n"
fileString += "#include <hip/hip_runtime.h>\n"
fileString += "#include <hip/hip_fp16.h>\n"
fileString += "\n"
activationCDataType = self.state["ProblemType"]["ActivationComputeDataType"]
if self.state["ProblemType"]["ActivationType"] in ['all', 'hipblaslt_all']:
fileString += "#include \"Tensile%sActivation%s_%s_%s.h\"\n"%(self.actGradientPrefix, \
self.gaurdStr, \
activationCDataType.toChar(), \
self.state["ProblemType"]["ActivationType"])
fileString += "\n"

fileString += self.functionSignature()
fileString += ";\n"

Expand Down
13 changes: 0 additions & 13 deletions tensilelite/Tensile/KernelWriterBetaOnly.py
Original file line number Diff line number Diff line change
Expand Up @@ -301,11 +301,6 @@ def getKernelName(self):
def getSourceFileString(self):
fileString = ""

if not globalParameters["MergeFiles"]:
fileString += "\n"
fileString += "#include \"%s.h\"\n" % self.kernelName
fileString += "\n"

for toggle in [True, False]:
self.state["ProblemType"]["GroupedGemm"] = toggle
self.kernelName = self.getKernelName()
Expand All @@ -316,14 +311,6 @@ def getSourceFileString(self):

def getHeaderFileString(self):
fileString = "" # CHeader
if not globalParameters["MergeFiles"]:
fileString += CHeader
fileString += "#pragma once\n\n"
fileString += "\n"
fileString += "#include <KernelHeader.h>\n\n"
fileString += "#include <hip/hip_runtime.h>\n"
fileString += "#include <hip/hip_fp16.h>\n"
fileString += "\n"

for toggle in [True, False]:
self.state["ProblemType"]["GroupedGemm"] = toggle
Expand Down
21 changes: 0 additions & 21 deletions tensilelite/Tensile/KernelWriterConversion.py
Original file line number Diff line number Diff line change
Expand Up @@ -836,22 +836,6 @@ def getKernelName(self):

def getHeaderFileString(self):
fileString = "" # CHeader
if not globalParameters["MergeFiles"]:
fileString += CHeader
fileString += "#pragma once\n\n"
fileString += "\n"
fileString += "#include <KernelHeader.h>\n\n"
fileString += "#include <hip/hip_runtime.h>\n"
fileString += "#include <hip/hip_fp16.h>\n"
fileString += "\n"
activationCDataType = self.state["ProblemType"]["ActivationComputeDataType"]
if self.state["ProblemType"]["ActivationType"] in ['all', 'hipblaslt_all']:
fileString += "#include \"Tensile%sActivation%s_%s_%s.h\"\n"%(self.actGradientPrefix, \
self.gaurdStr, \
activationCDataType.toChar(), \
self.state["ProblemType"]["ActivationType"])
fileString += "\n"

backupGSU = self.state["GlobalSplitU"]
backupUnroll = self.state["UnrollOnly"]
for gsu in self.gsuKernels:
Expand All @@ -873,11 +857,6 @@ def getHeaderFileString(self):

def getSourceFileString(self):
fileString = ""
if not globalParameters["MergeFiles"]:
fileString += "\n"
fileString += "#include \"%s.h\"\n" % self.kernelName
fileString += "\n"

backupGSU = self.state["GlobalSplitU"]
backupUnroll = self.state["UnrollOnly"]
for gsu in self.gsuKernels:
Expand Down
14 changes: 0 additions & 14 deletions tensilelite/Tensile/KernelWriterReduction.py
Original file line number Diff line number Diff line change
Expand Up @@ -65,15 +65,6 @@ def getKernelName(self):

def getHeaderFileString(self):
fileString = "" # CHeader
if not globalParameters["MergeFiles"]:
fileString += CHeader
fileString += "#pragma once\n\n"
fileString += "\n"
fileString += "#include <KernelHeader.h>\n\n"
fileString += "#include <hip/hip_runtime.h>\n"
fileString += "#include <hip/hip_fp16.h>\n"
fileString += "\n"

indexChars = globalParameters["IndexChars"]
# C dimensions
indicesStr = ""
Expand Down Expand Up @@ -101,11 +92,6 @@ def getHeaderFileString(self):

def getSourceFileString(self):
fileString = ""
if not globalParameters["MergeFiles"]:
fileString += "\n"
fileString += "#include \"%s.h\"\n" % self.kernelName
fileString += "\n"

fileString += self.kernelBody()

return (0, fileString)
4 changes: 0 additions & 4 deletions tensilelite/Tensile/Tensile.py
Original file line number Diff line number Diff line change
Expand Up @@ -119,8 +119,6 @@ def splitExtraParameters(par):
help="set PrintLevel=2 and CMakeBuildType=Debug")
argParser.add_argument("--short-names", dest="shortNames", action="store_true", \
help="use serial kernel and solution names")
argParser.add_argument("--no-merge-files", dest="noMergeFiles", action="store_true", \
help="kernels and solutions written to individual files")
argParser.add_argument("--cxx-compiler", dest="CxxCompiler", choices=[ToolchainDefaults.CXX_COMPILER], \
action="store", default=ToolchainDefaults.CXX_COMPILER, help="select which C++/HIP compiler to use")
argParser.add_argument("--c-compiler", dest="CCompiler", choices=[ToolchainDefaults.C_COMPILER], \
Expand Down Expand Up @@ -167,8 +165,6 @@ def argUpdatedGlobalParameters(args):
rv["CMakeBuildType"] = "Debug"
if args.shortNames:
rv["ShortNames"] = True
if args.noMergeFiles:
rv["MergeFiles"] = False
if args.client_build_path:
rv["ClientBuildPath"] = args.client_build_path
if args.client_lock:
Expand Down
Loading