Merge pull request #120 from guacamoleo/develop

ready for release 3.2

Jenkinsfile

@@ -186,39 +186,23 @@ def docker_build_inside_image( def build_image, compiler_data compiler_args, doc
               tensile ${rel_path_to_src}/Tensile/Configs/test_dgemm_defaults.yaml dgemm_defaults
 
               # thorough tests
-              tensile --runtime-language=HIP --kernel-language=HIP ${rel_path_to_src}/Tensile/Configs/test_hgemm.yaml hgemm
-              tensile --runtime-language=HIP --kernel-language=HIP ${rel_path_to_src}/Tensile/Configs/test_sgemm.yaml sgemm
+              tensile ${rel_path_to_src}/Tensile/Configs/test_hgemm.yaml hgemm
+              tensile ${rel_path_to_src}/Tensile/Configs/test_sgemm.yaml sgemm
 
               # vectors
-              tensile --runtime-language=HIP --kernel-language=HIP ${rel_path_to_src}/Tensile/Configs/test_hgemm_vectors.yaml hgemm_vectors
-              tensile --runtime-language=HIP --kernel-language=HIP ${rel_path_to_src}/Tensile/Configs/test_sgemm_vectors.yaml sgemm_vectors
+              tensile ${rel_path_to_src}/Tensile/Configs/test_hgemm_vectors.yaml hgemm_vectors
+              tensile ${rel_path_to_src}/Tensile/Configs/test_sgemm_vectors.yaml sgemm_vectors
 
               # tensor contractions
-              tensile --runtime-language=HIP --kernel-language=HIP ${rel_path_to_src}/Tensile/Configs/tensor_contraction.yaml tensor_contraction
-          """
+              tensile ${rel_path_to_src}/Tensile/Configs/tensor_contraction.yaml tensor_contraction
 
-          // assembly
-          if( env.NODE_LABELS ==~ /.*gfx803.*/ )
-          {
-            sh  """#!/usr/bin/env bash
-                set -x
-                cd ${paths.project_build_prefix}
-                tensile ${rel_path_to_src}/Tensile/Configs/sgemm_gfx803.yaml sgemm_gfx803
-            """
-          }
-          // assembly
-          if( env.NODE_LABELS ==~ /.*gfx900.*/ )
-          {
-            sh  """#!/usr/bin/env bash
-                set -x
-                cd ${paths.project_build_prefix}
-                tensile ${rel_path_to_src}/Tensile/Configs/sgemm_gfx900.yaml sgemm_gfx900
-            """
-          }
+              # assembly
+              tensile ${rel_path_to_src}/Tensile/Configs/sgemm_asm.yaml sgemm_asm
+          """
 
           // TODO re-enable when jenkins supports opencl
-          //sh "tensile --runtime-language=OCL --kernel-language=OCL ../../Tensile/Configs/test_sgemm_vectors.yaml sgemm_vectors"
-          //sh "tensile --runtime-language=OCL --kernel-language=OCL ${rel_path_to_src}/Tensile/Configs/convolution.yaml tensor_contraction"
+          //sh "tensile --runtime-language=OCL ../../Tensile/Configs/test_sgemm_vectors.yaml sgemm_vectors"
+          //sh "tensile --runtime-language=OCL ${rel_path_to_src}/Tensile/Configs/convolution.yaml tensor_contraction"
 
           archiveArtifacts artifacts: "${paths.project_build_prefix}/dist/*.egg", fingerprint: true
         }

Tensile/BenchmarkProblems.py

@@ -28,7 +28,7 @@
 import time
 
 from BenchmarkStructs import BenchmarkProcess
-from Common import globalParameters, HR, pushWorkingPath, popWorkingPath, print1, print2, printExit, printWarning, ensurePath, startTime, ProgressBar, kernelLanguageIsSource
+from Common import globalParameters, HR, pushWorkingPath, popWorkingPath, print1, print2, printExit, printWarning, ensurePath, startTime, ProgressBar
 from SolutionStructs import Solution, ProblemType
 from SolutionWriter import SolutionWriter
 from KernelWriterSource import KernelWriterSource
@@ -372,11 +372,16 @@ def writeBenchmarkFiles(solutions, problemSizes, stepName, filesToCopy):
   # Min Naming
   ##############################################################################
   kernels = []
+  kernelsBetaOnly = []
   for solution in solutions:
     solutionKernels = solution.getKernels()
     for kernel in solutionKernels:
       if kernel not in kernels:
         kernels.append(kernel)
+    solutionKernelsBetaOnly = solution.getKernelsBetaOnly()
+    for kernel in solutionKernelsBetaOnly:
+      if kernel not in kernelsBetaOnly:
+        kernelsBetaOnly.append(kernel)
 
   solutionSerialNaming = Solution.getSerialNaming(solutions)
   kernelSerialNaming = Solution.getSerialNaming(kernels)
@@ -385,27 +390,27 @@ def writeBenchmarkFiles(solutions, problemSizes, stepName, filesToCopy):
   solutionWriter = SolutionWriter( \
       solutionMinNaming, solutionSerialNaming, \
       kernelMinNaming, kernelSerialNaming)
-  if kernelLanguageIsSource():
-    kernelWriter = KernelWriterSource( \
-        kernelMinNaming, kernelSerialNaming)
-  else:
-    kernelWriter = KernelWriterAssembly( \
-        kernelMinNaming, kernelSerialNaming)
+  kernelWriterSource = KernelWriterSource( \
+      kernelMinNaming, kernelSerialNaming)
+  kernelWriterAssembly = KernelWriterAssembly( \
+      kernelMinNaming, kernelSerialNaming)
 
   # write solution, kernels and CMake
   writeSolutionsAndKernels( \
-      globalParameters["WorkingPath"], solutions, solutionWriter, kernelWriter)
-
+      globalParameters["WorkingPath"], solutions, kernels, kernelsBetaOnly, \
+      solutionWriter, kernelWriterSource, kernelWriterAssembly )
 
   ##############################################################################
   # Write CMake
   ##############################################################################
 
   clientName = "TensileBenchmark_%s" % stepName
-  writeCMake(globalParameters["WorkingPath"], solutions, filesToCopy, clientName)
+  writeCMake(globalParameters["WorkingPath"], solutions, kernels, filesToCopy, \
+      clientName)
 
   forBenchmark = True
-  writeClientParameters(forBenchmark, solutions, problemSizes, stepName, filesToCopy)
+  writeClientParameters(forBenchmark, solutions, problemSizes, stepName, \
+      filesToCopy)
 
 
 ################################################################################
@@ -467,8 +472,8 @@ def addResults( self, hardcodedParameterList, benchmarkPermutations, \
       winningParameters = {}
       for paramName in benchmarkPermutations[0]:
         winningParameters[paramName] = winningSolution[paramName]
-      print2("HCP[%u] Winner: idx=%u, gflops=%f, param=%s" \
-          % ( hardcodedIdx, winningIdx, winningScore, winningParameters))
+      #print2("HCP[%u] Winner: idx=%u, gflops=%f, param=%s" \
+      #    % ( hardcodedIdx, winningIdx, winningScore, winningParameters))
       matches = WinningParameterDict.get(hardcodedParameters, self.winners)
       if len(matches) != 1:
         printExit("Didn't find exactly 1 match")

Tensile/ClientWriter.py

@@ -140,7 +140,7 @@ def writeRunScript(path, libraryLogicPath, forBenchmark):
   if os.name != "nt":
     runScriptFile.write("#!/bin/sh\n")
   q = "" if os.name == "nt" else "\""
-  runScriptFile.write("%s & echo %s%s%s & echo %s# Configuring CMake for Client%s & echo %s%s%s\n" \
+  runScriptFile.write("%s && echo %s%s%s && echo %s# Configuring CMake for Client%s && echo %s%s%s\n" \
       % (echoLine, q, HR, q, q, q, q, HR, q))
   runScriptFile.write("cmake")
   # runtime and kernel language
@@ -155,8 +155,6 @@ def writeRunScript(path, libraryLogicPath, forBenchmark):
     # for library client
     runScriptFile.write(" -DTensile_ROOT=%s" \
         % os.path.join(globalParameters["ScriptPath"], "..") )
-    runScriptFile.write(" -DTensile_KERNEL_LANGUAGE=%s" \
-        % globalParameters["KernelLanguage"])
     runScriptFile.write(" -DTensile_CLIENT_BENCHMARK=OFF")
     runScriptFile.write(" -DTensile_LOGIC_PATH=%s" % libraryLogicPath)
     runScriptFile.write(" -DTensile_LIBRARY_PRINT_DEBUG=%s" \
@@ -175,7 +173,7 @@ def writeRunScript(path, libraryLogicPath, forBenchmark):
   runScriptFile.write(" -DTensile_MERGE_FILES=%s" \
       % ("ON" if globalParameters["MergeFiles"] else "OFF"))
   runScriptFile.write(" ../source\n")
-  runScriptFile.write("%s & echo %s%s%s & echo %s# Building Client%s & echo %s%s%s\n" \
+  runScriptFile.write("%s && echo %s%s%s && echo %s# Building Client%s && echo %s%s%s\n" \
       % (echoLine, q, HR, q, q, q, q, HR, q))
   runScriptFile.write("cmake --build . --config %s%s\n" \
       % (globalParameters["CMakeBuildType"], " -- -j 8" \
@@ -216,8 +214,8 @@ def writeRunScript(path, libraryLogicPath, forBenchmark):
           "client.exe")
     else:
       executablePath = os.path.join(executablePath, "client")
-    runScriptFile.write("%s & echo %s%s%s & echo %s# Library Client Path:%s & echo %s\n" \
-        % (echoLine, q, HR, q, q, q, executablePath) )
+    runScriptFile.write("%s && echo %s%s%s && echo %s# Library Client:%s && echo %s# %s%s && %s\n" \
+        % (echoLine, q, HR, q, q, q, q, executablePath, q, executablePath) )
   runScriptFile.close()
   if os.name != "nt":
     os.chmod(runScriptName, 0777)

Tensile/Common.py

@@ -22,6 +22,7 @@
 import sys
 from __init__ import __version__
 from collections import OrderedDict
+from subprocess import Popen, PIPE
 import time
 
 startTime = time.time()
@@ -38,10 +39,8 @@
 globalParameters["IndexChars"] =  "IJKLMNOPQRSTUVWXYZ"
 if os.name == "nt":
   globalParameters["RuntimeLanguage"] = "OCL"
-  globalParameters["KernelLanguage"] = "OCL"
 else:
   globalParameters["RuntimeLanguage"] = "HIP"
-  globalParameters["KernelLanguage"] = "HIP"
 # print level
 globalParameters["PrintLevel"] = 1
 globalParameters["LibraryPrintDebug"] = False
@@ -86,6 +85,8 @@
 globalParameters["MaxLDS"] = 32768
 globalParameters["DeviceLDS"] = 32768
 globalParameters["MinimumRequiredVersion"] = "0.0.0"
+globalParameters["SupportedISA"] = [(8,0,3), (9,0,0)]
+globalParameters["CurrentISA"] = (0,0,0)
 
 ################################################################################
 # Default Benchmark Parameters
@@ -108,6 +109,8 @@
 
 validMacroTileSides = [1, 2, 4, 8, 16, 32, 64, 128, 256, 512, 1024, 6, 12, 24, 48, 96, 192, 384, 768 ]
 validMacroTiles = []
+validISA = [(0,0,0)]
+validISA.extend(globalParameters["SupportedISA"])
 for i in validMacroTileSides:
   for j in validMacroTileSides:
     validMacroTiles.append([i, j])
@@ -145,6 +148,8 @@
     "MacroTileShapeMax":          range(1, 64+1),
 
     "EdgeType":                   [ "Branch", "ShiftPtr", "None" ],
+    "KernelLanguage":             [ "Assembly", "Source" ],
+    "ISA":                        validISA,
     "MacroTile":                  validMacroTiles,
 
     }
@@ -153,6 +158,7 @@
     {"LoopDoWhile":               [ False ] },
     {"LoopTail":                  [ True ] },
     {"EdgeType":                  [ "Branch" ] },
+    {"KernelLanguage":            [ "Source" ] },
     {"LdsPad":                    [ 0 ] },
     {"MaxOccupancy":              [ 10 ] },
     {"VectorWidth":               [ 1 ] }, # =2 once fixed
@@ -212,7 +218,7 @@
     "Batched":                  False,
     "IndexAssignmentsA":        [0, 2],
     "IndexAssignmentsB":        [1, 2],
-    "NumDimensionsC":           2,
+    "NumIndicesC":           2,
     "DataType":                 0,
     }
 defaultProblemSizes = [{"Range": [ [2880], 0, 0 ]}]
@@ -224,21 +230,13 @@
 # Default Analysis Parameters
 ################################################################################
 defaultAnalysisParameters = {
-    "ScheduleName":       "Default",
+    "ScheduleName":       "Tensile",
     "DeviceNames":  ["Unspecified"],
     #"BranchPenalty":              0, # microseconds / kernel
-    #"SmoothOutliers":         False, # enforce monotonic data
     "SolutionImportanceMin":      0.01, # = 0.01=1% total time saved by keeping this solution
     }
 
 
-################################################################################
-# Kernel Language Belongs to Source or Assembly?
-################################################################################
-def kernelLanguageIsSource():
-  return globalParameters["KernelLanguage"] \
-      in ["OCL", "HIP"]
-
 ################################################################################
 # Searching Nested Lists / Dictionaries
 ################################################################################
@@ -315,11 +313,32 @@ def printExit(message):
 def assignGlobalParameters( config ):
   global globalParameters
 
+  # read current gfx version
+  if os.name != "nt":
+    process = Popen(["/opt/rocm/bin/rocm_agent_enumerator", "-t", "GPU"], stdout=PIPE)
+    line = process.stdout.readline()
+    while line != "":
+      gfxIdx = line.find("gfx")
+      if gfxIdx >= 0:
+        major = int(line[gfxIdx+3:gfxIdx+4])
+        minor = int(line[gfxIdx+4:gfxIdx+5])
+        step  = int(line[gfxIdx+5:gfxIdx+6])
+        if (major,minor,step) in globalParameters["SupportedISA"]:
+          print1("# Host-Detected: gfx%u%u%u"%(major, minor, step))
+          globalParameters["CurrentISA"] = (major, minor, step)
+        line = process.stdout.readline()
+    if globalParameters["CurrentISA"] == (0,0,0):
+      printWarning("Did not detected ISA: %s" % globalParameters["SupportedISA"])
+    if process.returncode:
+      printWarning("rocm_agen_enumerator exited with code %u" % process.returncode)
+
+  # Minimum Required Version
   if "MinimumRequiredVersion" in config:
     if not versionIsCompatible(config["MinimumRequiredVersion"]):
       printExit("Benchmark.yaml file requires version=%s is not compatible with current Tensile version=%s" \
           % (config["MinimumRequiredVersion"], __version__) )
 
+  # User-specified global parameters
   print2("GlobalParameters:")
   for key in globalParameters:
     defaultValue = globalParameters[key]
@@ -339,6 +358,7 @@ def assignGlobalParameters( config ):
     globalParameters[key] = value
 
 
+
 ################################################################################
 # Assign Parameters
 ################################################################################
@@ -374,8 +394,8 @@ def ensurePath( path ):
 # Is query version compatible with current version
 ################################################################################
 def versionIsCompatible(queryVersionString):
-  (qMajor, qMinor, qPatch) = queryVersionString.split(".")
-  (tMajor, tMinor, tPatch) = __version__.split(".")
+  (qMajor, qMinor, qStep) = queryVersionString.split(".")
+  (tMajor, tMinor, tStep) = __version__.split(".")
 
   # major version must match exactly
   if qMajor != tMajor:
@@ -385,7 +405,7 @@ def versionIsCompatible(queryVersionString):
   if qMinor > tMinor:
     return False
   if qMinor == tMinor:
-    if qPatch > tPatch:
+    if qStep > tStep:
       return False
   return True
 

Tensile/Configs/rocblas_cgemm.yaml

@@ -33,10 +33,7 @@ BenchmarkProblems:
         - ProblemSizes:
           - Exact: [ 2880, 2880, 2, 1536 ]
         - EdgeType: ["ShiftPtr"]
-        - LoopDoWhile: [False]
         - LoopTail: [True]
-        - NumLoadsCoalescedA: [1]
-        - NumLoadsCoalescedB: [1]
       ForkParameters:
         - ThreadTile:
           - [ 4, 8 ]
@@ -75,10 +72,7 @@ BenchmarkProblems:
         - ProblemSizes:
           - Exact: [ 2880, 2880, 2, 1536 ]
         - EdgeType: ["ShiftPtr"]
-        - LoopDoWhile: [False]
         - LoopTail: [True]
-        - NumLoadsCoalescedA: [1]
-        - NumLoadsCoalescedB: [1]
       ForkParameters:
         - ThreadTile:
           - [ 4, 8 ]
@@ -117,10 +111,7 @@ BenchmarkProblems:
         - ProblemSizes:
           - Exact: [ 2880, 2880, 2, 1536 ]
         - EdgeType: ["ShiftPtr"]
-        - LoopDoWhile: [False]
         - LoopTail: [True]
-        - NumLoadsCoalescedA: [1]
-        - NumLoadsCoalescedB: [1]
       ForkParameters:
         - ThreadTile:
           - [ 4, 8 ]
@@ -160,10 +151,7 @@ BenchmarkProblems:
         - ProblemSizes:
           - Exact: [ 2880, 2880, 2, 1536 ]
         - EdgeType: ["ShiftPtr"]
-        - LoopDoWhile: [False]
         - LoopTail: [True]
-        - NumLoadsCoalescedA: [1]
-        - NumLoadsCoalescedB: [1]
       ForkParameters:
         - ThreadTile:
           - [ 4, 8 ]
@@ -188,7 +176,5 @@ BenchmarkProblems:
           - Range: [ [32, 32, 32, 400], [32, 32, 32, 400], [2], [1536] ]
 
 LibraryLogic:
-  ScheduleName: Fiji
-  DeviceNames: ["R9 Nano"]
 
 LibraryClient: