Initial overhaul of Profile mode
Signed-off-by: colramos-amd <colramos@amd.com>
[ROCm/rocprofiler-compute commit: 87388227ca]
This commit is contained in:
1 والد
e8fa3e3510
کامیت
ec13e780d9
421 files changed
+16949
-1912
No files matched your search
@@ -12,3 +12,4 @@ dash-svg
|
||||
dash-bootstrap-components
|
||||
kaleido
|
||||
setuptools
|
||||
pyfiglet
|
||||
+138
-149
@@ -22,36 +22,23 @@
|
||||
# SOFTWARE.
|
||||
##############################################################################el
|
||||
|
||||
import os
|
||||
import argparse
|
||||
import shutil
|
||||
import os
|
||||
|
||||
|
||||
from common import (
|
||||
OMNIPERF_HOME,
|
||||
PROG,
|
||||
SOC_LIST,
|
||||
) # Import global variables
|
||||
from common import getVersion, getVersionDisplay
|
||||
|
||||
|
||||
def parse(my_parser):
|
||||
# versioning info
|
||||
vData = getVersion()
|
||||
versionString = getVersionDisplay(vData["version"], vData["sha"], vData["mode"])
|
||||
|
||||
def omniarg_parser(parser, omniperf_home, omniperf_version):
|
||||
# -----------------------------------------
|
||||
# Parse arguments (dependent on mode)
|
||||
# -----------------------------------------
|
||||
|
||||
## General Command Line Options
|
||||
## ----------------------------
|
||||
general_group = my_parser.add_argument_group("General Options")
|
||||
my_parser._positionals.title = "Modes"
|
||||
my_parser._optionals.title = "Help"
|
||||
general_group.add_argument("-v", "--version", action="version", version=versionString)
|
||||
general_group = parser.add_argument_group("General Options")
|
||||
parser._positionals.title = "Modes"
|
||||
parser._optionals.title = "Help"
|
||||
general_group.add_argument("-v", "--version", action="version", version=omniperf_version["ver_pretty"])
|
||||
|
||||
subparsers = my_parser.add_subparsers(
|
||||
subparsers = parser.add_subparsers(
|
||||
dest="mode", help="Select mode of interaction with the target application:"
|
||||
)
|
||||
|
||||
@@ -61,17 +48,17 @@ def parse(my_parser):
|
||||
"profile",
|
||||
help="Profile the target application",
|
||||
usage="""
|
||||
\nomniperf profile --name <workload_name> [profile options] [roofline options] -- <profile_cmd>
|
||||
\nomniperf profile --name <workload_name> [profile options] [roofline options] -- <profile_cmd>
|
||||
|
||||
\n\n-------------------------------------------------------------------------------
|
||||
\nExamples:
|
||||
\n\tomniperf profile -n vcopy_all -- ./vcopy 1048576 256
|
||||
\n\tomniperf profile -n vcopy_SPI_TCC -b SQ TCC -- ./vcopy 1048576 256
|
||||
\n\tomniperf profile -n vcopy_kernel -k vecCopy -- ./vcopy 1048576 256
|
||||
\n\tomniperf profile -n vcopy_disp -d 0 -- ./vcopy 1048576 256
|
||||
\n\tomniperf profile -n vcopy_roof --roof-only -- ./vcopy 1048576 256
|
||||
\n-------------------------------------------------------------------------------\n
|
||||
""",
|
||||
\n\n-------------------------------------------------------------------------------
|
||||
\nExamples:
|
||||
\n\tomniperf profile -n vcopy_all -- ./vcopy 1048576 256
|
||||
\n\tomniperf profile -n vcopy_SPI_TCC -b SQ TCC -- ./vcopy 1048576 256
|
||||
\n\tomniperf profile -n vcopy_kernel -k vecCopy -- ./vcopy 1048576 256
|
||||
\n\tomniperf profile -n vcopy_disp -d 0 -- ./vcopy 1048576 256
|
||||
\n\tomniperf profile -n vcopy_roof --roof-only -- ./vcopy 1048576 256
|
||||
\n-------------------------------------------------------------------------------\n
|
||||
""",
|
||||
prog="tool",
|
||||
allow_abbrev=False,
|
||||
formatter_class=lambda prog: argparse.RawTextHelpFormatter(
|
||||
@@ -79,14 +66,16 @@ def parse(my_parser):
|
||||
),
|
||||
)
|
||||
profile_parser._optionals.title = "Help"
|
||||
|
||||
general_group = profile_parser.add_argument_group("General Options")
|
||||
profile_group = profile_parser.add_argument_group("Profile Options")
|
||||
roofline_group = profile_parser.add_argument_group("Standalone Roofline Options")
|
||||
|
||||
general_group.add_argument("-v", "--version", action="version", version=versionString)
|
||||
general_group.add_argument("-v", "--version", action="version", version=omniperf_version["ver_pretty"])
|
||||
general_group.add_argument(
|
||||
"-V", "--verbose", help="Increase output verbosity", action="count", default=0
|
||||
)
|
||||
|
||||
profile_group.add_argument(
|
||||
"-n",
|
||||
"--name",
|
||||
@@ -103,7 +92,7 @@ def parse(my_parser):
|
||||
metavar="",
|
||||
type=str,
|
||||
dest="path",
|
||||
default=os.getcwd() + "/workloads",
|
||||
default=os.path.join(os.getcwd(), "workloads"),
|
||||
required=False,
|
||||
help="\t\t\tSpecify path to save workload.\n\t\t\t(DEFAULT: {}/workloads/<name>)".format(
|
||||
os.getcwd()
|
||||
@@ -120,6 +109,27 @@ def parse(my_parser):
|
||||
default=None,
|
||||
help="\t\t\tKernel filtering.",
|
||||
)
|
||||
profile_group.add_argument(
|
||||
"-d",
|
||||
"--dispatch",
|
||||
type=str,
|
||||
metavar="",
|
||||
nargs="+",
|
||||
dest="dispatch",
|
||||
required=False,
|
||||
help="\t\t\tDispatch ID filtering.",
|
||||
)
|
||||
profile_group.add_argument(
|
||||
"-b",
|
||||
"--ipblocks",
|
||||
type=str,
|
||||
dest="ipblocks",
|
||||
metavar="",
|
||||
nargs="+",
|
||||
required=False,
|
||||
choices=["SQ", "SQC", "TA", "TD", "TCP", "TCC", "SPI", "CPC", "CPF"],
|
||||
help="\t\t\tIP block filtering:\n\t\t\t SQ\n\t\t\t SQC\n\t\t\t TA\n\t\t\t TD\n\t\t\t TCP\n\t\t\t TCC\n\t\t\t SPI\n\t\t\t CPC\n\t\t\t CPF",
|
||||
)
|
||||
|
||||
result = shutil.which("rocscope")
|
||||
if result:
|
||||
@@ -160,28 +170,6 @@ def parse(my_parser):
|
||||
profile_group.add_argument(
|
||||
"--kernel-summaries", default=False, dest="summaries", help=argparse.SUPPRESS
|
||||
)
|
||||
|
||||
profile_group.add_argument(
|
||||
"-b",
|
||||
"--ipblocks",
|
||||
type=str,
|
||||
dest="ipblocks",
|
||||
metavar="",
|
||||
nargs="+",
|
||||
required=False,
|
||||
choices=["SQ", "SQC", "TA", "TD", "TCP", "TCC", "SPI", "CPC", "CPF"],
|
||||
help="\t\t\tIP block filtering:\n\t\t\t SQ\n\t\t\t SQC\n\t\t\t TA\n\t\t\t TD\n\t\t\t TCP\n\t\t\t TCC\n\t\t\t SPI\n\t\t\t CPC\n\t\t\t CPF",
|
||||
)
|
||||
profile_group.add_argument(
|
||||
"-d",
|
||||
"--dispatch",
|
||||
type=str,
|
||||
metavar="",
|
||||
nargs="+",
|
||||
dest="dispatch",
|
||||
required=False,
|
||||
help="\t\t\tDispatch ID filtering.",
|
||||
)
|
||||
profile_group.add_argument(
|
||||
"--join-type",
|
||||
metavar="",
|
||||
@@ -205,7 +193,7 @@ def parse(my_parser):
|
||||
help="\t\t\tProvide command for profiling after double dash.",
|
||||
)
|
||||
profile_group.add_argument(
|
||||
"--kernelVerbose",
|
||||
"--kernel-verbose",
|
||||
required=False,
|
||||
metavar="",
|
||||
help="\t\t\tSpecify Kernel Name verbose level 1-5. Lower the level, shorter the kernel name. (DEFAULT: 2) (DISABLE: 5)",
|
||||
@@ -268,14 +256,14 @@ def parse(my_parser):
|
||||
"database",
|
||||
help="Interact with Omniperf database",
|
||||
usage="""
|
||||
\nomniperf database <interaction type> [connection options]
|
||||
\nomniperf database <interaction type> [connection options]
|
||||
|
||||
\n\n-------------------------------------------------------------------------------
|
||||
\nExamples:
|
||||
\n\tomniperf database --import -H pavii1 -u temp -t asw -w workloads/vcopy/mi200/
|
||||
\n\tomniperf database --remove -H pavii1 -u temp -w omniperf_asw_sample_mi200
|
||||
\n-------------------------------------------------------------------------------\n
|
||||
""",
|
||||
\n\n-------------------------------------------------------------------------------
|
||||
\nExamples:
|
||||
\n\tomniperf database --import -H pavii1 -u temp -t asw -w workloads/vcopy/mi200/
|
||||
\n\tomniperf database --remove -H pavii1 -u temp -w omniperf_asw_sample_mi200
|
||||
\n-------------------------------------------------------------------------------\n
|
||||
""",
|
||||
prog="tool",
|
||||
allow_abbrev=False,
|
||||
formatter_class=lambda prog: argparse.RawTextHelpFormatter(
|
||||
@@ -288,7 +276,7 @@ def parse(my_parser):
|
||||
interaction_group = db_parser.add_argument_group("Interaction Type")
|
||||
connection_group = db_parser.add_argument_group("Connection Options")
|
||||
|
||||
general_group.add_argument("-v", "--version", action="version", version=versionString)
|
||||
general_group.add_argument("-v", "--version", action="version", version=omniperf_version["ver_pretty"])
|
||||
general_group.add_argument(
|
||||
"-V", "--verbose", help="Increase output verbosity", action="count", default=0
|
||||
)
|
||||
@@ -357,15 +345,15 @@ def parse(my_parser):
|
||||
"analyze",
|
||||
help="Analyze existing profiling results at command line",
|
||||
usage="""
|
||||
\nomniperf analyze --path <workload_path> [analyze options]
|
||||
\nomniperf analyze --path <workload_path> [analyze options]
|
||||
|
||||
\n\n-------------------------------------------------------------------------------
|
||||
\nExamples:
|
||||
\n\tomniperf analyze -p workloads/vcopy/mi200/ --list-metrics gfx90a
|
||||
\n\tomniperf analyze -p workloads/mixbench/mi200/ --filter-dispatch-ids 12 34 --decimal 3
|
||||
\n\tomniperf analyze -p workloads/mixbench/mi200/ --gui
|
||||
\n-------------------------------------------------------------------------------\n
|
||||
""",
|
||||
\n\n-------------------------------------------------------------------------------
|
||||
\nExamples:
|
||||
\n\tomniperf analyze -p workloads/vcopy/mi200/ --list-metrics gfx90a
|
||||
\n\tomniperf analyze -p workloads/mixbench/mi200/ --dispatch 12 34 --decimal 3
|
||||
\n\tomniperf analyze -p workloads/mixbench/mi200/ --gui
|
||||
\n-------------------------------------------------------------------------------\n
|
||||
""",
|
||||
prog="tool",
|
||||
allow_abbrev=False,
|
||||
formatter_class=lambda prog: argparse.RawTextHelpFormatter(
|
||||
@@ -376,8 +364,9 @@ def parse(my_parser):
|
||||
|
||||
general_group = analyze_parser.add_argument_group("General Options")
|
||||
analyze_group = analyze_parser.add_argument_group("Analyze Options")
|
||||
analyze_advanced_group = analyze_parser.add_argument_group("Advanced Options")
|
||||
|
||||
general_group.add_argument("-v", "--version", action="version", version=versionString)
|
||||
general_group.add_argument("-v", "--version", action="version", version=omniperf_version["ver_pretty"])
|
||||
general_group.add_argument(
|
||||
"-V", "--verbose", help="Increase output verbosity", action="count", default=0
|
||||
)
|
||||
@@ -392,31 +381,16 @@ def parse(my_parser):
|
||||
action="append",
|
||||
help="\t\tSpecify the raw data root dirs or desired results directory.",
|
||||
)
|
||||
analyze_group.add_argument(
|
||||
"-o",
|
||||
"--output",
|
||||
metavar="",
|
||||
dest="output_file",
|
||||
help="\t\tSpecify the output file.",
|
||||
)
|
||||
analyze_group.add_argument(
|
||||
"--list-kernels",
|
||||
action="store_true",
|
||||
help="\t\tList kernels. Top 10 kernels sorted by duration (descending order).",
|
||||
help="\t\tList all detected kernels. Sorted by duration (descending order).",
|
||||
)
|
||||
analyze_group.add_argument(
|
||||
"--list-metrics",
|
||||
metavar="",
|
||||
choices=["gfx906", "gfx908", "gfx90a"],
|
||||
help="\t\tList metrics can be customized to analyze on specific arch:\n\t\t gfx906\n\t\t gfx908\n\t\t gfx90a",
|
||||
)
|
||||
analyze_group.add_argument(
|
||||
"-b",
|
||||
"--metric",
|
||||
dest="filter_metrics",
|
||||
metavar="",
|
||||
nargs="+",
|
||||
help="\t\tSpecify IP block/metric id(s) from --list-metrics for filtering.",
|
||||
help="\t\tList all availible metrics for analysis on specified arch:\n\t\t gfx906\n\t\t gfx908\n\t\t gfx90a",
|
||||
)
|
||||
analyze_group.add_argument(
|
||||
"-k",
|
||||
@@ -437,6 +411,14 @@ def parse(my_parser):
|
||||
action="append",
|
||||
help="\t\tSpecify dispatch id(s) for filtering.",
|
||||
)
|
||||
analyze_group.add_argument(
|
||||
"-b",
|
||||
"--metric",
|
||||
dest="filter_metrics",
|
||||
metavar="",
|
||||
nargs="+",
|
||||
help="\t\tSpecify IP block/metric id(s) from --list-metrics for filtering.",
|
||||
)
|
||||
analyze_group.add_argument(
|
||||
"--gpu-id",
|
||||
dest="gpu_id",
|
||||
@@ -445,62 +427,11 @@ def parse(my_parser):
|
||||
help="\t\tSpecify GPU id(s) for filtering.",
|
||||
)
|
||||
analyze_group.add_argument(
|
||||
"-n",
|
||||
"--normal-unit",
|
||||
dest="normal_unit",
|
||||
"-o",
|
||||
"--output",
|
||||
metavar="",
|
||||
default="per_wave",
|
||||
choices=["per_wave", "per_cycle", "per_second", "per_kernel"],
|
||||
help="\t\tSpecify the normalization unit: (DEFAULT: per_wave)\n\t\t per_wave\n\t\t per_cycle\n\t\t per_second\n\t\t per_kernel",
|
||||
)
|
||||
analyze_group.add_argument(
|
||||
"--max-kernel-num",
|
||||
dest="max_kernel_num",
|
||||
metavar="",
|
||||
type=int,
|
||||
default=10,
|
||||
help="\t\tSpecify the maximum number of kernels shown (DEFAULT: 10)",
|
||||
)
|
||||
analyze_group.add_argument(
|
||||
"--config-dir",
|
||||
dest="config_dir",
|
||||
metavar="",
|
||||
help="\t\tSpecify the directory of customized configs.",
|
||||
default=OMNIPERF_HOME.joinpath("omniperf_analyze/configs"),
|
||||
)
|
||||
analyze_group.add_argument(
|
||||
"-t",
|
||||
"--time-unit",
|
||||
dest="time_unit",
|
||||
metavar="",
|
||||
default="ns",
|
||||
choices=["s", "ms", "us", "ns"],
|
||||
help="\t\tSpecify display time unit in kernel top stats: (DEFAULT: ns)\n\t\t s\n\t\t ms\n\t\t us\n\t\t ns",
|
||||
)
|
||||
analyze_group.add_argument(
|
||||
"--decimal",
|
||||
type=int,
|
||||
metavar="",
|
||||
default=2,
|
||||
help="\t\tSpecify the decimal to display. (DEFAULT: 2)",
|
||||
)
|
||||
analyze_group.add_argument(
|
||||
"--save-dfs",
|
||||
dest="df_file_dir",
|
||||
metavar="",
|
||||
help="\t\tSpecify the dirctory to save analysis dataframe csv files.",
|
||||
)
|
||||
analyze_group.add_argument(
|
||||
"--cols",
|
||||
type=int,
|
||||
dest="cols",
|
||||
metavar="",
|
||||
nargs="+",
|
||||
help="\t\tSpecify column indices to display.",
|
||||
)
|
||||
analyze_group.add_argument("-g", action="store_true", help="\t\tDebug single metric.")
|
||||
analyze_group.add_argument(
|
||||
"--dependency", action="store_true", help="\t\tList the installation dependency."
|
||||
dest="output_file",
|
||||
help="\t\tSpecify an output file to save analysis results.",
|
||||
)
|
||||
analyze_group.add_argument(
|
||||
"--gui",
|
||||
@@ -509,19 +440,77 @@ def parse(my_parser):
|
||||
const=8050,
|
||||
help="\t\tActivate a GUI to interate with Omniperf metrics.\n\t\tOptionally, specify port to launch application (DEFAULT: 8050)",
|
||||
)
|
||||
analyze_group.add_argument(
|
||||
analyze_advanced_group.add_argument(
|
||||
"--random-port",
|
||||
action="store_true",
|
||||
help="\t\tRandomly generate a port to launch GUI application.\n\t\tRegistered Ports range inclusive (1024-49151).",
|
||||
)
|
||||
analyze_group.add_argument(
|
||||
"--kernelVerbose",
|
||||
analyze_advanced_group.add_argument(
|
||||
"--max-kernel-num",
|
||||
dest="max_kernel_num",
|
||||
metavar="",
|
||||
type=int,
|
||||
default=10,
|
||||
help="\t\tSpecify the maximum number of kernels shown in \"Top Stats\" table (DEFAULT: 10)",
|
||||
)
|
||||
analyze_advanced_group.add_argument(
|
||||
"-n",
|
||||
"--normal-unit",
|
||||
dest="normal_unit",
|
||||
metavar="",
|
||||
default="per_wave",
|
||||
choices=["per_wave", "per_cycle", "per_second", "per_kernel"],
|
||||
help="\t\tSpecify the normalization unit: (DEFAULT: per_wave)\n\t\t per_wave\n\t\t per_cycle\n\t\t per_second\n\t\t per_kernel",
|
||||
)
|
||||
analyze_advanced_group.add_argument(
|
||||
"-t",
|
||||
"--time-unit",
|
||||
dest="time_unit",
|
||||
metavar="",
|
||||
default="ns",
|
||||
choices=["s", "ms", "us", "ns"],
|
||||
help="\t\tSpecify display time unit in kernel top stats: (DEFAULT: ns)\n\t\t s\n\t\t ms\n\t\t us\n\t\t ns",
|
||||
)
|
||||
analyze_advanced_group.add_argument(
|
||||
"--decimal",
|
||||
type=int,
|
||||
metavar="",
|
||||
default=2,
|
||||
help="\t\tSpecify desired decimal precision of analysis results. (DEFAULT: 2)",
|
||||
)
|
||||
analyze_advanced_group.add_argument(
|
||||
"--config-dir",
|
||||
dest="config_dir",
|
||||
metavar="",
|
||||
help="\t\tSpecify the directory of customized configs.",
|
||||
default=omniperf_home.joinpath("omniperf_analyze/configs"),
|
||||
)
|
||||
analyze_advanced_group.add_argument(
|
||||
"--save-dfs",
|
||||
dest="df_file_dir",
|
||||
metavar="",
|
||||
help="\t\tSpecify the dirctory to save analysis dataframe csv files.",
|
||||
)
|
||||
analyze_advanced_group.add_argument(
|
||||
"--cols",
|
||||
type=int,
|
||||
dest="cols",
|
||||
metavar="",
|
||||
nargs="+",
|
||||
help="\t\tSpecify column indices to display.",
|
||||
)
|
||||
analyze_advanced_group.add_argument("-g", action="store_true", help="\t\tDebug single metric.")
|
||||
analyze_advanced_group.add_argument(
|
||||
"--dependency", action="store_true", help="\t\tList the installation dependency."
|
||||
)
|
||||
analyze_advanced_group.add_argument(
|
||||
"--kernel-verbose",
|
||||
required=False,
|
||||
metavar="",
|
||||
help="\t\tSpecify Kernel Name verbose level 1-5. Lower the level, shorter the kernel name. (DEFAULT: 5) (DISABLE: 5)",
|
||||
default=5,
|
||||
type=int,
|
||||
)
|
||||
analyze_group.add_argument(
|
||||
analyze_advanced_group.add_argument(
|
||||
"--report-diff", default=0, nargs="?", type=int, help=argparse.SUPPRESS
|
||||
)
|
||||
)
|
||||
@@ -1,110 +0,0 @@
|
||||
##############################################################################bl
|
||||
# MIT License
|
||||
#
|
||||
# Copyright (c) 2021 - 2023 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.
|
||||
##############################################################################el
|
||||
|
||||
import os
|
||||
import sys
|
||||
import io
|
||||
from pathlib import Path
|
||||
import subprocess
|
||||
import shutil
|
||||
|
||||
OMNIPERF_HOME = Path(__file__).resolve().parent
|
||||
|
||||
# OMNIPERF INFO
|
||||
PROG = "omniperf"
|
||||
SOC_LIST = ["mi50", "mi100", "mi200", "vega10"]
|
||||
DISTRO_MAP = {"platform:el8": "rhel8", "15.3": "sle15sp3", "20.04": "ubuntu20_04"}
|
||||
|
||||
|
||||
def resolve_rocprof(returnPath=False):
|
||||
# ROCPROF INFO
|
||||
if not "ROCPROF" in os.environ.keys():
|
||||
rocprof_cmd = "rocprof"
|
||||
else:
|
||||
rocprof_cmd = os.environ["ROCPROF"]
|
||||
|
||||
rocprof_path = shutil.which(rocprof_cmd)
|
||||
|
||||
if not rocprof_path:
|
||||
print("\nError: Unable to resolve path to %s binary" % rocprof_cmd)
|
||||
print(
|
||||
"Please verify installation or set ROCPROF environment variable with full path."
|
||||
)
|
||||
sys.exit(1)
|
||||
else:
|
||||
# Resolve any sym links in file path
|
||||
rocprof_path = os.path.realpath(rocprof_path.rstrip("\n"))
|
||||
print("ROC Profiler: ", rocprof_path)
|
||||
if returnPath:
|
||||
return rocprof_path
|
||||
else:
|
||||
return rocprof_cmd
|
||||
|
||||
|
||||
def getVersion():
|
||||
# symantic version info
|
||||
version = os.path.join(OMNIPERF_HOME.parent, "VERSION")
|
||||
try:
|
||||
with open(version, "r") as file:
|
||||
VER = file.read().replace("\n", "")
|
||||
except EnvironmentError:
|
||||
print("ERROR: Cannot find VERSION file at {}".format(version))
|
||||
sys.exit(1)
|
||||
|
||||
# git version info
|
||||
gitDir = os.path.join(OMNIPERF_HOME.parent, ".git")
|
||||
if (shutil.which("git") is not None) and os.path.exists(gitDir):
|
||||
gitQuery = subprocess.run(
|
||||
["git", "log", "--pretty=format:%h", "-n", "1"],
|
||||
stdout=subprocess.PIPE,
|
||||
stderr=subprocess.DEVNULL,
|
||||
)
|
||||
if gitQuery.returncode != 0:
|
||||
SHA = "unknown"
|
||||
MODE = "unknown"
|
||||
else:
|
||||
SHA = gitQuery.stdout.decode("utf-8")
|
||||
MODE = "dev"
|
||||
else:
|
||||
shaFile = os.path.join(OMNIPERF_HOME.parent, "VERSION.sha")
|
||||
try:
|
||||
with open(shaFile, "r") as file:
|
||||
SHA = file.read().replace("\n", "")
|
||||
except EnvironmentError:
|
||||
print("ERROR: Cannot find VERSION.sha file at {}".format(shaFile))
|
||||
sys.exit(1)
|
||||
|
||||
MODE = "release"
|
||||
|
||||
versionData = {"version": VER, "sha": SHA, "mode": MODE}
|
||||
return versionData
|
||||
|
||||
|
||||
def getVersionDisplay(version, sha, mode):
|
||||
buf = io.StringIO()
|
||||
print("-" * 40, file=buf)
|
||||
print("Omniperf version: %s (%s)" % (version, mode), file=buf)
|
||||
print("Git revision: %s" % sha, file=buf)
|
||||
print("-" * 40, file=buf)
|
||||
return buf.getvalue()
|
||||
+4
-18
@@ -22,22 +22,8 @@
|
||||
# SOFTWARE.
|
||||
##############################################################################el
|
||||
|
||||
from abc import ABC, abstractmethod
|
||||
|
||||
class OmniSoC_Base():
|
||||
def __init__(self,args):
|
||||
self.__args = args
|
||||
|
||||
# Required methods to be implemented by child classes
|
||||
@abstractmethod
|
||||
def profiling_setup(self):
|
||||
"""Perform any SoC-specific setup prior to profiling.
|
||||
"""
|
||||
pass
|
||||
|
||||
@abstractmethod
|
||||
def analysis_setup(self):
|
||||
"""Perform any SoC-specific setup prior to analysis.
|
||||
"""
|
||||
pass
|
||||
from pathlib import Path
|
||||
|
||||
# NB: Creating a new module to share global vars across modules
|
||||
omniperf_home = Path(__file__).resolve().parent
|
||||
prog = "omniperf"
|
||||
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
@@ -24,866 +24,24 @@
|
||||
# SOFTWARE.
|
||||
##############################################################################el
|
||||
|
||||
import sys
|
||||
import os
|
||||
import io
|
||||
import selectors
|
||||
import argparse
|
||||
import subprocess
|
||||
import glob
|
||||
import pandas as pd
|
||||
from datetime import datetime
|
||||
from pathlib import Path as path
|
||||
import warnings
|
||||
import shutil
|
||||
|
||||
from parser import parse
|
||||
from utils import specs
|
||||
from utils.perfagg import perfmon_filter, pmc_filter, pmc_perf_split, join_prof
|
||||
from utils import remove_workload
|
||||
from utils import csv_processor # Import workload
|
||||
from omniperf_analyze.omniperf_analyze import roofline_only # Standalone roofline
|
||||
from omniperf_analyze.omniperf_analyze import analyze # CLI analysis
|
||||
from common import resolve_rocprof
|
||||
|
||||
from common import (
|
||||
OMNIPERF_HOME,
|
||||
PROG,
|
||||
SOC_LIST,
|
||||
DISTRO_MAP,
|
||||
) # Import global variables
|
||||
|
||||
from common import getVersion
|
||||
|
||||
################################################
|
||||
# Helper Functions
|
||||
################################################
|
||||
def run_subprocess(cmd):
|
||||
subprocess.run(cmd, check=True)
|
||||
|
||||
def capture_subprocess_output(subprocess_args):
|
||||
# Start subprocess
|
||||
# bufsize = 1 means output is line buffered
|
||||
# universal_newlines = True is required for line buffering
|
||||
process = subprocess.Popen(subprocess_args,
|
||||
bufsize=1,
|
||||
stdout=subprocess.PIPE,
|
||||
stderr=subprocess.STDOUT,
|
||||
universal_newlines=True)
|
||||
|
||||
# Create callback function for process output
|
||||
buf = io.StringIO()
|
||||
def handle_output(stream, mask):
|
||||
# Because the process' output is line buffered, there's only ever one
|
||||
# line to read when this function is called
|
||||
line = stream.readline()
|
||||
buf.write(line)
|
||||
sys.stdout.write(line)
|
||||
|
||||
# Register callback for an "available for read" event from subprocess' stdout stream
|
||||
selector = selectors.DefaultSelector()
|
||||
selector.register(process.stdout, selectors.EVENT_READ, handle_output)
|
||||
|
||||
# Loop until subprocess is terminated
|
||||
while process.poll() is None:
|
||||
# Wait for events and handle them with their registered callbacks
|
||||
events = selector.select()
|
||||
for key, mask in events:
|
||||
callback = key.data
|
||||
callback(key.fileobj, mask)
|
||||
|
||||
# Get process return code
|
||||
return_code = process.wait()
|
||||
selector.close()
|
||||
|
||||
success = (return_code == 0)
|
||||
|
||||
# Store buffered output
|
||||
output = buf.getvalue()
|
||||
buf.close()
|
||||
|
||||
return (success, output)
|
||||
|
||||
|
||||
def get_soc():
|
||||
mspec = specs.get_machine_specs(0)
|
||||
|
||||
target = ""
|
||||
if mspec.GPU == "gfx906":
|
||||
target = "mi50"
|
||||
elif mspec.GPU == "gfx908":
|
||||
target = "mi100"
|
||||
elif mspec.GPU == "gfx90a":
|
||||
target = "mi200"
|
||||
elif mspec.GPU == "gfx900":
|
||||
target = "vega10"
|
||||
else:
|
||||
print("\nInvalid SoC")
|
||||
sys.exit(0)
|
||||
return target
|
||||
|
||||
|
||||
def throw_parse_error(my_parser, message):
|
||||
my_parser.print_help(sys.stderr)
|
||||
print("\n\n")
|
||||
my_parser.error(message)
|
||||
|
||||
|
||||
def isWorkloadEmpty(my_parser, path):
|
||||
if os.path.isfile(path + "/pmc_perf.csv"):
|
||||
temp_df = pd.read_csv(path + "/pmc_perf.csv")
|
||||
if temp_df.dropna().empty:
|
||||
print("Profiling Error: Found empty cells. Profiling data could be corrupt.")
|
||||
sys.exit(0)
|
||||
|
||||
else:
|
||||
throw_parse_error(
|
||||
my_parser, "Profling Error: Cannot find pmc_perf.csv in {}".format(path)
|
||||
)
|
||||
|
||||
|
||||
def replace_timestamps(workload_dir, log_file):
|
||||
df_stamps = pd.read_csv(workload_dir + "/timestamps.csv")
|
||||
if "BeginNs" in df_stamps.columns and "EndNs" in df_stamps.columns:
|
||||
# Update timestamps for all *.csv output files
|
||||
for fname in glob.glob(workload_dir + "/" + "*.csv"):
|
||||
df_pmc_perf = pd.read_csv(fname)
|
||||
|
||||
df_pmc_perf["BeginNs"] = df_stamps["BeginNs"]
|
||||
df_pmc_perf["EndNs"] = df_stamps["EndNs"]
|
||||
df_pmc_perf.to_csv(fname, index=False)
|
||||
else:
|
||||
warning = "WARNING: Incomplete profiling data detected. Unable to update timestamps."
|
||||
warnings.warn(
|
||||
warning
|
||||
)
|
||||
log_file.write(warning + "\n")
|
||||
|
||||
|
||||
def gen_sysinfo(workload_name, workload_dir, ip_blocks, app_cmd, skip_roof):
|
||||
# Record system information
|
||||
mspec = specs.get_machine_specs(0)
|
||||
sysinfo = open(workload_dir + "/" + "sysinfo.csv", "w")
|
||||
|
||||
# write header
|
||||
header = "workload_name,"
|
||||
header += "command,"
|
||||
header += "host_name,host_cpu,host_distro,host_kernel,host_rocmver,date,"
|
||||
header += "gpu_soc,numSE,numCU,numSIMD,waveSize,maxWavesPerCU,maxWorkgroupSize,"
|
||||
header += "L1,L2,sclk,mclk,cur_sclk,cur_mclk,L2Banks,LDSBanks,name,numSQC,hbmBW,"
|
||||
header += "ip_blocks\n"
|
||||
sysinfo.write(header)
|
||||
|
||||
# timestamp
|
||||
now = datetime.now()
|
||||
local_now = now.astimezone()
|
||||
local_tz = local_now.tzinfo
|
||||
local_tzname = local_tz.tzname(local_now)
|
||||
timestamp = now.strftime("%c") + " (" + local_tzname + ")"
|
||||
# host info
|
||||
param = [workload_name]
|
||||
param += ['"' + app_cmd + '"']
|
||||
param += [
|
||||
mspec.hostname,
|
||||
mspec.cpu,
|
||||
mspec.distro,
|
||||
mspec.kernel,
|
||||
mspec.rocmversion,
|
||||
timestamp,
|
||||
]
|
||||
|
||||
# GPU info
|
||||
param += [
|
||||
mspec.GPU,
|
||||
mspec.SE,
|
||||
mspec.CU,
|
||||
mspec.SIMD,
|
||||
mspec.wave_size,
|
||||
mspec.wave_occu,
|
||||
mspec.workgroup_size,
|
||||
]
|
||||
param += [
|
||||
mspec.L1,
|
||||
mspec.L2,
|
||||
mspec.SCLK,
|
||||
mspec.cur_MCLK,
|
||||
mspec.cur_SCLK,
|
||||
mspec.cur_MCLK,
|
||||
]
|
||||
|
||||
blocks = []
|
||||
hbmBW = int(mspec.cur_MCLK) / 1000 * 4096 / 8 * 2
|
||||
if mspec.GPU == "gfx906":
|
||||
param += ["16", "32", "mi50", str(int(mspec.CU) // 4), str(hbmBW)]
|
||||
elif mspec.GPU == "gfx908":
|
||||
param += ["32", "32", "mi100", "48", str(hbmBW)]
|
||||
elif mspec.GPU == "gfx90a":
|
||||
param += ["32", "32", "mi200", "56", str(hbmBW)]
|
||||
if not skip_roof:
|
||||
blocks.append("roofline")
|
||||
|
||||
# ip block info
|
||||
if ip_blocks == None:
|
||||
t = ["SQ", "LDS", "SQC", "TA", "TD", "TCP", "TCC", "SPI", "CPC", "CPF"]
|
||||
blocks += t
|
||||
else:
|
||||
blocks += ip_blocks
|
||||
param.append("|".join(blocks))
|
||||
|
||||
sysinfo.write(",".join(param))
|
||||
sysinfo.close()
|
||||
|
||||
|
||||
def mongo_import(args, profileAndImport):
|
||||
# Validate target directory
|
||||
connectionInfo = csv_processor.parse(args, profileAndImport)
|
||||
# Convert and upload data
|
||||
print("-- Conversion & Upload in Progress --")
|
||||
csv_processor.convert_folder(connectionInfo)
|
||||
print("-- Complete! --")
|
||||
|
||||
################################################
|
||||
# Roofline Helpers
|
||||
################################################
|
||||
def roof_setup(args, my_parser, VER):
|
||||
if args.path == os.getcwd() + "/workloads":
|
||||
args.path += "/" + args.name + "/" + args.target
|
||||
|
||||
# Do we need a new directory for roofline?
|
||||
if not os.path.isdir(args.path):
|
||||
os.makedirs(args.path)
|
||||
|
||||
# Does roof data exist?
|
||||
print("Checking for roofline.csv in ", args.path)
|
||||
roof_path = args.path + "/roofline.csv"
|
||||
roofline_exists = os.path.isfile(roof_path)
|
||||
if not roofline_exists:
|
||||
if get_soc() != "mi200":
|
||||
throw_parse_error(
|
||||
my_parser, "Invalid SoC.\nRoofline only availible on MI200."
|
||||
)
|
||||
mibench(args)
|
||||
|
||||
# Does sysinfo exist?
|
||||
print("Checking for sysinfo.csv in ", args.path)
|
||||
sysinfo_path = args.path + "/sysinfo.csv"
|
||||
sysinfo_exists = os.path.isfile(sysinfo_path)
|
||||
if not sysinfo_exists:
|
||||
print("sysinfo not found")
|
||||
gen_sysinfo(args.name, args.path, [], args.remaining, False)
|
||||
|
||||
# Does app data exist?
|
||||
print("Checking for pmc_perf.csv in ", args.path)
|
||||
app_path = args.path + "/pmc_perf.csv"
|
||||
app_exists = os.path.isfile(app_path)
|
||||
if not app_exists:
|
||||
if get_soc() != "mi200":
|
||||
throw_parse_error(
|
||||
my_parser, "Invalid SoC.\nRoofline only availible on MI200."
|
||||
)
|
||||
if not args.remaining:
|
||||
throw_parse_error(
|
||||
my_parser,
|
||||
"Cannot find existing application data.\nAttempting to generate application data from -- <app_cmd>.\n-- <app_cmd> option is required to generate application data.",
|
||||
)
|
||||
else:
|
||||
characterize_app(args, VER)
|
||||
|
||||
|
||||
def detect_roofline():
|
||||
mspec = specs.get_machine_specs(0)
|
||||
rocm_ver = mspec.rocmversion[:1]
|
||||
|
||||
os_release = path("/etc/os-release").read_text()
|
||||
ubuntu_distro = specs.search(r'VERSION_ID="(.*?)"', os_release)
|
||||
rhel_distro = specs.search(r'PLATFORM_ID="(.*?)"', os_release)
|
||||
sles_distro = specs.search(r'VERSION_ID="(.*?)"', os_release)
|
||||
|
||||
if "ROOFLINE_BIN" in os.environ.keys():
|
||||
rooflineBinary = os.environ["ROOFLINE_BIN"]
|
||||
if os.path.exists(rooflineBinary):
|
||||
print("Detected user-supplied binary")
|
||||
return {"rocm_ver": "override", "distro": "override", "path": rooflineBinary}
|
||||
else:
|
||||
print("ROOFLINE ERROR: user-supplied path to binary not accessible")
|
||||
print("--> ROOFLINE_BIN = %s\n" % target_binary)
|
||||
sys.exit(1)
|
||||
elif rhel_distro == "platform:el8":
|
||||
# Must be a valid RHEL machine
|
||||
distro = rhel_distro
|
||||
elif (
|
||||
(type(sles_distro) == str and len(sles_distro) >= 3) and # confirm string and len
|
||||
sles_distro[:2] == "15" and int(sles_distro[3]) >= 3 # SLES15 and SP >= 3
|
||||
):
|
||||
# Must be a valid SLES machine
|
||||
# Use SP3 binary for all forward compatible service pack versions
|
||||
distro = "15.3"
|
||||
elif ubuntu_distro == "20.04":
|
||||
# Must be a valid Ubuntu machine
|
||||
distro = ubuntu_distro
|
||||
else:
|
||||
print("ROOFLINE ERROR: Cannot find a valid binary for your operating system")
|
||||
sys.exit(1)
|
||||
|
||||
target_binary = {"rocm_ver": rocm_ver, "distro": distro}
|
||||
return target_binary
|
||||
|
||||
|
||||
def mibench(args):
|
||||
print("No roofline data found. Generating...")
|
||||
|
||||
target_binary = detect_roofline()
|
||||
if target_binary["rocm_ver"] == "override":
|
||||
path_to_binary = target_binary["path"]
|
||||
else:
|
||||
path_to_binary = (
|
||||
str(OMNIPERF_HOME)
|
||||
+ "/utils/rooflines/roofline"
|
||||
+ "-"
|
||||
+ DISTRO_MAP[target_binary["distro"]]
|
||||
+ "-"
|
||||
+ args.target.lower()
|
||||
+ "-rocm"
|
||||
+ target_binary["rocm_ver"]
|
||||
)
|
||||
|
||||
# Distro is valid but cant find rocm ver
|
||||
if not os.path.exists(path_to_binary):
|
||||
print("ROOFLINE ERROR: Unable to locate expected binary (%s)." % path_to_binary)
|
||||
sys.exit(1)
|
||||
|
||||
run_subprocess(
|
||||
[
|
||||
path_to_binary,
|
||||
"-o",
|
||||
args.path + "/" + "roofline.csv",
|
||||
"-d",
|
||||
str(args.device),
|
||||
]
|
||||
)
|
||||
|
||||
|
||||
def characterize_app(args, VER):
|
||||
# Basic Info
|
||||
print("\n", PROG, "ver: ", VER)
|
||||
print("Path: ", args.path)
|
||||
print("Target: ", args.target)
|
||||
print("Command: ", args.remaining)
|
||||
print("Kernel Selection: ", args.kernel)
|
||||
print("Dispatch Selection: ", args.dispatch)
|
||||
|
||||
perfmon_dir = str(OMNIPERF_HOME) + "/perfmon_pub"
|
||||
print("permon dir is ", os.path.abspath(perfmon_dir))
|
||||
app_cmd = args.remaining
|
||||
workload_dir = args.path
|
||||
|
||||
# Perfmon filtering
|
||||
pmc_filter(workload_dir, perfmon_dir, args.target)
|
||||
|
||||
# Separate pmc_perf runs
|
||||
pmc_perf_split(workload_dir)
|
||||
|
||||
# Set up a log file
|
||||
log = open(workload_dir + "/log.txt", "w")
|
||||
print("Log: ", workload_dir + "/log.txt\n")
|
||||
|
||||
# Workload profiling
|
||||
for fname in glob.glob(workload_dir + "/perfmon/*.txt"):
|
||||
# Kernel filtering (in-place replacement)
|
||||
if not args.kernel == None:
|
||||
success, output = capture_subprocess_output(
|
||||
[
|
||||
"sed",
|
||||
"-i",
|
||||
"-r",
|
||||
"s%^(kernel:).*%" + "kernel: " + ",".join(args.kernel) + "%g",
|
||||
fname,
|
||||
]
|
||||
)
|
||||
log.write(output)
|
||||
|
||||
# Dispatch filtering (inplace replacement)
|
||||
if not args.dispatch == None:
|
||||
success, output = capture_subprocess_output(
|
||||
[
|
||||
"sed",
|
||||
"-i",
|
||||
"-r",
|
||||
"s%^(range:).*%" + "range: " + " ".join(args.dispatch) + "%g",
|
||||
fname,
|
||||
]
|
||||
)
|
||||
log.write(output)
|
||||
print(fname)
|
||||
if args.use_rocscope == True:
|
||||
run_rocscope(args, fname)
|
||||
else:
|
||||
run_prof(fname, workload_dir, perfmon_dir, app_cmd, args.target, log, args.verbose)
|
||||
|
||||
# Update timestamps
|
||||
replace_timestamps(workload_dir, log)
|
||||
|
||||
if args.use_rocscope == False:
|
||||
# Manually join each pmc_perf*.csv output
|
||||
join_prof(workload_dir, args.join_type, log, args.verbose)
|
||||
# Demangle and overwrite original KernelNames
|
||||
csv_processor.kernel_name_shortener(workload_dir, args.kernelVerbose)
|
||||
|
||||
log.close()
|
||||
|
||||
|
||||
################################################
|
||||
# Profiling Helpers
|
||||
################################################
|
||||
|
||||
def run_rocscope(args, fname):
|
||||
# profile the app
|
||||
if args.use_rocscope == True:
|
||||
result = shutil.which("rocscope")
|
||||
if result:
|
||||
rs_cmd = [
|
||||
result.stdout.decode("ascii").strip(),
|
||||
"metrics",
|
||||
"-p",
|
||||
args.path,
|
||||
"-n",
|
||||
args.name,
|
||||
"-t",
|
||||
fname,
|
||||
"--",
|
||||
]
|
||||
for i in args.remaining.split():
|
||||
rs_cmd.append(i)
|
||||
print(rs_cmd)
|
||||
result = run_subprocess(
|
||||
rs_cmd
|
||||
) # , stdout=subprocess.PIPE, stderr=subprocess.PIPE)
|
||||
if result.returncode != 0:
|
||||
print(result.stderr.decode("ascii"))
|
||||
sys.exit(1)
|
||||
|
||||
|
||||
def run_prof(fname, workload_dir, perfmon_dir, cmd, target, log_file, verbose):
|
||||
global rocprof_cmd
|
||||
|
||||
fbase = os.path.splitext(os.path.basename(fname))[0]
|
||||
|
||||
if verbose:
|
||||
print("pmc file:", os.path.basename(fname))
|
||||
|
||||
# profile the app (run w/ custom config files for mi100)
|
||||
if target == "mi100":
|
||||
print("RUNNING WITH CUSTOM METRICS")
|
||||
success, output = capture_subprocess_output(
|
||||
[
|
||||
rocprof_cmd,
|
||||
"-i",
|
||||
fname,
|
||||
"-m",
|
||||
perfmon_dir + "/" + "metrics.xml",
|
||||
"--timestamp",
|
||||
"on",
|
||||
"-o",
|
||||
workload_dir + "/" + fbase + ".csv",
|
||||
'"' + cmd + '"',
|
||||
]
|
||||
)
|
||||
else:
|
||||
success, output = capture_subprocess_output(
|
||||
[
|
||||
rocprof_cmd,
|
||||
"-i",
|
||||
fname,
|
||||
"--timestamp",
|
||||
"on",
|
||||
"-o",
|
||||
workload_dir + "/" + fbase + ".csv",
|
||||
'"' + cmd + '"',
|
||||
]
|
||||
)
|
||||
# Write output to log
|
||||
log_file.write(output)
|
||||
|
||||
|
||||
|
||||
def omniperf_profile(args, VER):
|
||||
# Verify valid name
|
||||
if args.name.find(".") != -1 or args.name.find("-") != -1:
|
||||
raise ValueError("'-' and '.' are not permited in workload name", args.name)
|
||||
|
||||
# Basic Info
|
||||
print(PROG, "ver: ", VER)
|
||||
print("Path: ", args.path)
|
||||
print("Target: ", args.target)
|
||||
print("Command: ", args.remaining)
|
||||
print("Kernel Selection: ", args.kernel)
|
||||
print("Dispatch Selection: ", args.dispatch)
|
||||
|
||||
if args.ipblocks == None:
|
||||
print("IP Blocks: All")
|
||||
else:
|
||||
print("IP Blocks: ", args.ipblocks)
|
||||
if args.kernelVerbose > 5:
|
||||
print("KernelName verbose level: DISABLED")
|
||||
else:
|
||||
print("KernelName verbose level: ", str(args.kernelVerbose))
|
||||
|
||||
# Set up directories
|
||||
if args.path == os.getcwd() + "/workloads":
|
||||
workload_dir = args.path + "/" + args.name + "/" + args.target
|
||||
else:
|
||||
workload_dir = args.path
|
||||
perfmon_dir = str(OMNIPERF_HOME) + "/perfmon_pub"
|
||||
|
||||
# Perfmon filtering
|
||||
perfmon_filter(workload_dir, perfmon_dir, args)
|
||||
|
||||
# Separate pmc_perf runs
|
||||
pmc_perf_split(workload_dir)
|
||||
|
||||
# Set up a log file
|
||||
log = open(workload_dir + "/log.txt", "w")
|
||||
print("Log: ", workload_dir + "/log.txt\n")
|
||||
|
||||
if not args.lucky == None and args.lucky == True:
|
||||
print("You're feeling lucky - only profiling top N kernels")
|
||||
# look for whether workload_dir exists - create if not
|
||||
try:
|
||||
os.makedirs(workload_dir, exist_ok=True)
|
||||
except Exception as e:
|
||||
print("Unable to create workload directory: ", workload_dir)
|
||||
print(e)
|
||||
sys.exit(1)
|
||||
|
||||
result = shutil.which("rocscope")
|
||||
if result:
|
||||
rs_cmd = [
|
||||
result.stdout.decode("ascii").strip(),
|
||||
"top10",
|
||||
"-p",
|
||||
args.path,
|
||||
"-n",
|
||||
args.name,
|
||||
"--",
|
||||
]
|
||||
for i in args.remaining.split():
|
||||
rs_cmd.append(i)
|
||||
print(rs_cmd)
|
||||
result = run_subprocess(
|
||||
rs_cmd
|
||||
) # , stdout=subprocess.PIPE, stderr=subprocess.PIPE)
|
||||
if result.returncode != 0:
|
||||
print(result.stderr.decode("ascii"))
|
||||
else:
|
||||
print("rocscope must be in the PATH")
|
||||
sys.exit(1)
|
||||
elif not args.summaries == None and args.summaries == True:
|
||||
print("creating kernel summaries")
|
||||
# look for whether workload_dir exists - create if not
|
||||
try:
|
||||
os.makedirs(workload_dir, exist_ok=True)
|
||||
except Exception as e:
|
||||
print("Unable to create workload directory: ", workload_dir)
|
||||
print(e)
|
||||
sys.exit(1)
|
||||
|
||||
result = shutil.which("rocscope")
|
||||
if result:
|
||||
rs_cmd = [
|
||||
result.stdout.decode("ascii").strip(),
|
||||
"summary",
|
||||
"-p",
|
||||
args.path,
|
||||
"-n",
|
||||
args.name,
|
||||
"--",
|
||||
]
|
||||
for i in args.remaining.split():
|
||||
rs_cmd.append(i)
|
||||
print(rs_cmd)
|
||||
result = run_subprocess(
|
||||
rs_cmd
|
||||
) # , stdout=subprocess.PIPE, stderr=subprocess.PIPE)
|
||||
if result.returncode != 0:
|
||||
print(result.stderr.decode("ascii"))
|
||||
else:
|
||||
print("rocscope must be in the PATH")
|
||||
sys.exit(1)
|
||||
|
||||
else:
|
||||
for fname in glob.glob(workload_dir + "/perfmon/*.txt"):
|
||||
# Kernel filtering (in-place replacement)
|
||||
if not args.kernel == None:
|
||||
success, output = capture_subprocess_output(
|
||||
[
|
||||
"sed",
|
||||
"-i",
|
||||
"-r",
|
||||
"s%^(kernel:).*%" + "kernel: " + ",".join(args.kernel) + "%g",
|
||||
fname,
|
||||
]
|
||||
)
|
||||
log.write(output)
|
||||
|
||||
# Dispatch filtering (inplace replacement)
|
||||
if not args.dispatch == None:
|
||||
success, output = capture_subprocess_output(
|
||||
[
|
||||
"sed",
|
||||
"-i",
|
||||
"-r",
|
||||
"s%^(range:).*%" + "range: " + " ".join(args.dispatch) + "%g",
|
||||
fname,
|
||||
]
|
||||
)
|
||||
log.write(output)
|
||||
print(fname)
|
||||
if args.use_rocscope == True:
|
||||
run_rocscope(args, fname)
|
||||
else:
|
||||
run_prof(fname, workload_dir, perfmon_dir, args.remaining, args.target, log, args.verbose)
|
||||
|
||||
|
||||
# Update timestamps
|
||||
replace_timestamps(workload_dir, log)
|
||||
|
||||
if args.use_rocscope == False:
|
||||
# Manually join each pmc_perf*.csv output
|
||||
join_prof(workload_dir, args.join_type, log, args.verbose)
|
||||
# Demangle and overwrite original KernelNames
|
||||
csv_processor.kernel_name_shortener(workload_dir, args.kernelVerbose)
|
||||
|
||||
# Generate sysinfo
|
||||
gen_sysinfo(args.name, workload_dir, args.ipblocks, args.remaining, args.no_roof)
|
||||
|
||||
# Add tracing & roofline metrics (mi200 only)
|
||||
if args.target.lower() == "mi200":
|
||||
# Skip roofline if --no-roof is set.
|
||||
if not args.no_roof:
|
||||
target_binary = detect_roofline()
|
||||
if target_binary["rocm_ver"] == "override":
|
||||
path_to_binary = target_binary["path"]
|
||||
else:
|
||||
path_to_binary = (
|
||||
str(OMNIPERF_HOME)
|
||||
+ "/utils/rooflines/roofline"
|
||||
+ "-"
|
||||
+ DISTRO_MAP[target_binary["distro"]]
|
||||
+ "-"
|
||||
+ args.target.lower()
|
||||
+ "-rocm"
|
||||
+ target_binary["rocm_ver"]
|
||||
)
|
||||
# Distro is valid but cant find valid binary
|
||||
if not os.path.exists(path_to_binary):
|
||||
print(
|
||||
"ROOFLINE ERROR: Unable to locate expected binary (%s))."
|
||||
% path_to_binary
|
||||
)
|
||||
sys.exit(1)
|
||||
|
||||
run_subprocess(
|
||||
[
|
||||
path_to_binary,
|
||||
"-o",
|
||||
workload_dir + "/" + "roofline.csv",
|
||||
"-d",
|
||||
str(args.device),
|
||||
]
|
||||
)
|
||||
# Close log
|
||||
log.close()
|
||||
|
||||
|
||||
################################################
|
||||
# MAIN
|
||||
################################################
|
||||
|
||||
from omniperf_base import Omniperf
|
||||
|
||||
def main():
|
||||
my_parser = argparse.ArgumentParser(
|
||||
description="Command line interface for AMD's GPU profiler, Omniperf",
|
||||
prog="tool",
|
||||
formatter_class=lambda prog: argparse.RawTextHelpFormatter(
|
||||
prog, max_help_position=30
|
||||
),
|
||||
usage="omniperf [mode] [options]",
|
||||
)
|
||||
parse(my_parser)
|
||||
args = my_parser.parse_args()
|
||||
|
||||
vData = getVersion()
|
||||
VER = vData["version"]
|
||||
omniperf = Omniperf()
|
||||
omniperf.parse_args() #TODO: We already do this in the __init__ for Omniperf(). Change that?
|
||||
|
||||
if args.mode == None:
|
||||
throw_parse_error(
|
||||
my_parser,
|
||||
"Omniperf requires you pass a valid mode. Please see documentation.",
|
||||
)
|
||||
##############
|
||||
# PROFILE MODE
|
||||
##############
|
||||
if args.mode == "profile":
|
||||
Extractionlvl = args.kernelVerbose
|
||||
print("Resolving rocprof")
|
||||
global rocprof_cmd
|
||||
rocprof_cmd = resolve_rocprof()
|
||||
# Cannot access parent directories
|
||||
if ".." in str(args.path):
|
||||
throw_parse_error(
|
||||
my_parser, "Access denied. Cannot access parent directories in path ../"
|
||||
)
|
||||
|
||||
# Must have a valid soc in profile mode
|
||||
args.target = get_soc()
|
||||
|
||||
# Verify correct command formatting
|
||||
args.remaining = args.remaining[1:]
|
||||
if args.remaining:
|
||||
if not os.path.isfile(args.remaining[0]):
|
||||
throw_parse_error(
|
||||
my_parser,
|
||||
'Your command "{}" doesn\'t point to a file. Try again.'.format(
|
||||
args.remaining[0]
|
||||
),
|
||||
)
|
||||
args.remaining = " ".join(args.remaining)
|
||||
else:
|
||||
throw_parse_error(
|
||||
my_parser,
|
||||
"Profiling command required. Pass application executable after -- at the end of options.\n\ti.e. omniperf profile -n vcopy -- ./vcopy 1048576 256",
|
||||
)
|
||||
# Name cannot exceed MongoDB max len
|
||||
if len(args.name) > 35:
|
||||
throw_parse_error(my_parser, "--name exceeds 35 character limit. Try again.")
|
||||
|
||||
elif args.roof_only:
|
||||
print("\n--------\nRoofline only\n--------\n")
|
||||
# Setup prerequisits for roofline
|
||||
roof_setup(args, my_parser, VER)
|
||||
# Generate roofline
|
||||
roofline_only(args.path, args.device, args.sort, args.mem_level, args.kernel_names, args.verbose)
|
||||
|
||||
# Profile only
|
||||
else:
|
||||
print("\n-------------\nProfile only\n-------------\n")
|
||||
omniperf_profile(args, VER)
|
||||
|
||||
##############
|
||||
# DATABASE MODE
|
||||
##############
|
||||
if args.mode == "database":
|
||||
# Remove a workload
|
||||
if args.remove and not args.upload:
|
||||
print("\n--------\nRemove workload\n--------\n")
|
||||
fullWorkloadName = args.workload.count("_") >= 3
|
||||
if not fullWorkloadName:
|
||||
throw_parse_error(
|
||||
my_parser,
|
||||
"--workload is not valid. Please use full workload name as seen in GUI when removing (i.e. omniperf_asw_vcopy_mi200)",
|
||||
)
|
||||
if args.host == None or args.username == None:
|
||||
throw_parse_error(
|
||||
my_parser, "--host and --username are required when --remove is set."
|
||||
)
|
||||
remove_workload.remove_workload(args)
|
||||
# Import a workload
|
||||
elif args.upload and not args.remove:
|
||||
print("\n--------\nImport Profiling Results\n--------\n")
|
||||
if (
|
||||
args.host == None
|
||||
or args.team == None
|
||||
or args.username == None
|
||||
or args.workload == None
|
||||
):
|
||||
throw_parse_error(
|
||||
my_parser,
|
||||
"--host, --workload, --username, and --team are all required when --import is set.",
|
||||
)
|
||||
|
||||
if os.path.isdir(os.path.abspath(args.workload)):
|
||||
isWorkloadEmpty(
|
||||
my_parser, args.workload
|
||||
) # Throw warning if workload is empty
|
||||
else:
|
||||
throw_parse_error(
|
||||
my_parser,
|
||||
"--workload is invalid. Please pass path to a valid directory.",
|
||||
)
|
||||
|
||||
if len(args.team) > 13:
|
||||
throw_parse_error(
|
||||
my_parser, "--team exceeds 13 character limit. Try again."
|
||||
)
|
||||
|
||||
args.workload = os.path.abspath(args.workload) # Format path properly
|
||||
|
||||
mongo_import(args, False)
|
||||
else:
|
||||
throw_parse_error(
|
||||
my_parser, "Pass either -i/--import or -r/--remove when import mode"
|
||||
)
|
||||
##############
|
||||
# ANALYZE MODE
|
||||
##############
|
||||
if args.mode == "analyze":
|
||||
if args.list_metrics:
|
||||
analyze(args)
|
||||
else:
|
||||
if args.path:
|
||||
if ".." in str(args.path):
|
||||
throw_parse_error(
|
||||
my_parser,
|
||||
"Access denied. Cannot access parent directories in path ../",
|
||||
)
|
||||
if args.filter_metrics and args.gui:
|
||||
throw_parse_error(
|
||||
my_parser,
|
||||
"""
|
||||
omniperf analyze --path <workload_path> [analyze options]
|
||||
\n\n-------------------------------------------------------------------------------
|
||||
\nExamples:
|
||||
\n\tomniperf analyze -p workloads/vcopy/mi200/ --list-metrics gfx90a
|
||||
\n\tomniperf analyze -p workloads/mixbench/mi200/ --filter-dispatch-ids 12 34 --decimal 3
|
||||
\n\tomniperf analyze -p workloads/mixbench/mi200/ --gui
|
||||
\n-------------------------------------------------------------------------------\n
|
||||
\ntool: error: --gui cannot be used in combination with: -b/--metric
|
||||
"""
|
||||
)
|
||||
print("\n--------\nAnalyze\n--------\n")
|
||||
# Ensure absolute path
|
||||
for dir in args.path:
|
||||
full_path = os.path.abspath(dir[0])
|
||||
dir[0] = full_path
|
||||
if not os.path.isdir(dir[0]):
|
||||
throw_parse_error(
|
||||
my_parser,
|
||||
"Error: invalid directory {}\nPlease try again.".format(
|
||||
dir[0]
|
||||
),
|
||||
)
|
||||
isWorkloadEmpty(
|
||||
my_parser, dir[0]
|
||||
) # Verify workload is valid before analyzing
|
||||
|
||||
analyze(args)
|
||||
else:
|
||||
throw_parse_error(
|
||||
my_parser,
|
||||
"""
|
||||
omniperf analyze --path <workload_path> [analyze options]
|
||||
\n\n-------------------------------------------------------------------------------
|
||||
\nExamples:
|
||||
\n\tomniperf analyze -p workloads/vcopy/mi200/ --list-metrics gfx90a
|
||||
\n\tomniperf analyze -p workloads/mixbench/mi200/ --filter-dispatch-ids 12 34 --decimal 3
|
||||
\n\tomniperf analyze -p workloads/mixbench/mi200/ --gui
|
||||
\n-------------------------------------------------------------------------------\n
|
||||
\ntool: error: the following arguments are required: -p/--path
|
||||
""",
|
||||
)
|
||||
|
||||
sys.exit(0) # Indicate successful on exit
|
||||
mode = omniperf.get_mode()
|
||||
|
||||
# major omniperf execution modes
|
||||
if mode == "profile":
|
||||
omniperf.run_profiler()
|
||||
elif mode == "database":
|
||||
omniperf.update_DB()
|
||||
elif mode == "analyze":
|
||||
omniperf.run_analysis()
|
||||
else:
|
||||
omniperf.error("Unsupported execution mode")
|
||||
|
||||
if __name__ == "__main__":
|
||||
main()
|
||||
@@ -1,47 +0,0 @@
|
||||
#!/usr/bin/env python3
|
||||
|
||||
##############################################################################bl
|
||||
# MIT License
|
||||
#
|
||||
# Copyright (c) 2021 - 2023 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.
|
||||
##############################################################################el
|
||||
|
||||
from omniperf_base import Omniperf
|
||||
|
||||
def main():
|
||||
|
||||
omniperf2 = Omniperf()
|
||||
omniperf2.parseArgs()
|
||||
|
||||
mode = omniperf2.get_mode()
|
||||
|
||||
# major omniperf execution modes
|
||||
if mode == "profile":
|
||||
omniperf2.run_profiler()
|
||||
elif mode == "database":
|
||||
omniperf2.update_DB()
|
||||
elif mode == "analyze":
|
||||
omniperf2.run_analysis()
|
||||
else:
|
||||
omniperf2.error("Unsupported execution mode")
|
||||
|
||||
if __name__ == "__main__":
|
||||
main()
|
||||
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
File renamed without changes.
Loaded 100 of 421 files, more files were not shown because too many files have changed in this diff.
Show more
Reference in new issue
Block a user