Files
rocm-systems/projects/rocprofiler-compute/src/omniperf
T

887 строки
30 KiB
Python
Исходник Обычный вид История

2022-11-04 14:49:36 -05:00
#!/usr/bin/env python3
2023-02-13 09:26:12 -06:00
##############################################################################bl
# MIT License
2023-02-13 14:50:24 -06:00
#
# Copyright (c) 2021 - 2023 Advanced Micro Devices, Inc. All Rights Reserved.
2023-02-13 14:50:24 -06:00
#
2022-11-04 14:49:36 -05:00
# 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:
2023-02-13 14:50:24 -06:00
#
# The above copyright notice and this permission notice shall be included in all
# copies or substantial portions of the Software.
2023-02-13 14:50:24 -06:00
#
2022-11-04 14:49:36 -05:00
# 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
2022-11-04 14:49:36 -05:00
# 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.
2023-02-13 09:26:12 -06:00
##############################################################################el
2022-11-04 14:49:36 -05:00
import sys
import os
2023-04-05 14:28:19 -05:00
import io
import selectors
2022-11-04 14:49:36 -05:00
import argparse
import subprocess
import glob
import pandas as pd
from datetime import datetime
from pathlib import Path as path
import warnings
2023-08-24 15:54:00 -05:00
import shutil
2022-11-04 14:49:36 -05:00
from parser import parse
from utils import specs
2023-05-05 15:07:20 -05:00
from utils.perfagg import perfmon_filter, pmc_filter, pmc_perf_split, join_prof
2022-11-04 14:49:36 -05:00
from utils import remove_workload
from utils import csv_processor # Import workload
2023-02-13 14:50:24 -06:00
from omniperf_analyze.omniperf_analyze import roofline_only # Standalone roofline
from omniperf_analyze.omniperf_analyze import analyze # CLI analysis
from common import resolve_rocprof
2022-11-04 14:49:36 -05:00
from common import (
OMNIPERF_HOME,
PROG,
SOC_LIST,
DISTRO_MAP,
) # Import global variables
from common import getVersion
2022-11-04 14:49:36 -05:00
################################################
# Helper Functions
################################################
2022-11-18 12:50:07 -06:00
def run_subprocess(cmd):
2023-02-13 14:50:24 -06:00
subprocess.run(cmd, check=True)
2023-04-05 14:28:19 -05:00
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)
2022-11-18 12:50:07 -06:00
2022-11-04 14:49:36 -05:00
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"
2022-11-04 14:49:36 -05:00
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)
)
2023-05-16 15:39:45 -05:00
def replace_timestamps(workload_dir, log_file):
2022-11-04 14:49:36 -05:00
df_stamps = pd.read_csv(workload_dir + "/timestamps.csv")
if "BeginNs" in df_stamps.columns and "EndNs" in df_stamps.columns:
2023-05-05 15:07:20 -05:00
# Update timestamps for all *.csv output files
for fname in glob.glob(workload_dir + "/" + "*.csv"):
df_pmc_perf = pd.read_csv(fname)
2022-11-04 14:49:36 -05:00
2023-05-05 15:07:20 -05:00
df_pmc_perf["BeginNs"] = df_stamps["BeginNs"]
df_pmc_perf["EndNs"] = df_stamps["EndNs"]
df_pmc_perf.to_csv(fname, index=False)
else:
2023-05-16 15:39:45 -05:00
warning = "WARNING: Incomplete profiling data detected. Unable to update timestamps."
2023-02-13 14:50:24 -06:00
warnings.warn(
2023-05-16 15:39:45 -05:00
warning
2023-02-13 14:50:24 -06:00
)
2023-05-16 15:39:45 -05:00
log_file.write(warning + "\n")
2022-11-04 14:49:36 -05:00
2022-11-14 09:51:48 -06:00
def gen_sysinfo(workload_name, workload_dir, ip_blocks, app_cmd, skip_roof):
2022-11-04 14:49:36 -05:00
# Record system information
mspec = specs.get_machine_specs(0)
sysinfo = open(workload_dir + "/" + "sysinfo.csv", "w")
# write header
header = "workload_name,"
2022-11-14 09:51:48 -06:00
header += "command,"
2022-11-04 14:49:36 -05:00
header += "host_name,host_cpu,host_distro,host_kernel,host_rocmver,date,"
header += "gpu_soc,numSE,numCU,numSIMD,waveSize,maxWavesPerCU,maxWorkgroupSize,"
2023-05-19 16:08:23 -05:00
header += "L1,L2,sclk,mclk,cur_sclk,cur_mclk,L2Banks,LDSBanks,name,numSQC,hbmBW,"
2022-11-04 14:49:36 -05:00
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]
2023-02-13 14:50:24 -06:00
param += ['"' + app_cmd + '"']
2022-11-04 14:49:36 -05:00
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":
2023-05-19 16:08:23 -05:00
param += ["16", "32", "mi50", str(int(mspec.CU) // 4), str(hbmBW)]
2022-11-04 14:49:36 -05:00
elif mspec.GPU == "gfx908":
2023-05-19 16:08:23 -05:00
param += ["32", "32", "mi100", "48", str(hbmBW)]
2022-11-04 14:49:36 -05:00
elif mspec.GPU == "gfx90a":
2023-05-19 16:08:23 -05:00
param += ["32", "32", "mi200", "56", str(hbmBW)]
2022-11-04 14:49:36 -05:00
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
2023-08-24 18:57:26 -05:00
connectionInfo = csv_processor.parse(args, profileAndImport)
2022-11-04 14:49:36 -05:00
# Convert and upload data
print("-- Conversion & Upload in Progress --")
2023-08-24 18:57:26 -05:00
csv_processor.convert_folder(connectionInfo)
2022-11-04 14:49:36 -05:00
print("-- Complete! --")
################################################
# Roofline Helpers
################################################
2023-02-22 14:37:20 -06:00
def roof_setup(args, my_parser, VER):
if args.path == os.getcwd() + "/workloads":
2023-02-22 14:37:20 -06:00
args.path += "/" + args.name + "/" + args.target
2022-11-04 14:49:36 -05:00
2023-02-22 14:37:20 -06:00
# Do we need a new directory for roofline?
if not os.path.isdir(args.path):
os.makedirs(args.path)
2023-02-22 14:37:20 -06:00
# 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)
2023-02-22 14:37:20 -06:00
# 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)
2023-02-22 14:37:20 -06:00
# 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:
2023-02-22 14:37:20 -06:00
characterize_app(args, VER)
2023-02-13 14:50:24 -06:00
2022-11-04 14:49:36 -05:00
def detect_roofline():
mspec = specs.get_machine_specs(0)
rocm_ver = mspec.rocmversion[:1]
os_release = path("/etc/os-release").read_text()
2022-11-07 10:25:30 -06:00
ubuntu_distro = specs.search(r'VERSION_ID="(.*?)"', os_release)
2022-11-04 14:49:36 -05:00
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}
2022-11-04 14:49:36 -05:00
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
):
2022-11-04 14:49:36 -05:00
# Must be a valid SLES machine
# Use SP3 binary for all forward compatible service pack versions
distro = "15.3"
2022-11-07 10:25:30 -06:00
elif ubuntu_distro == "20.04":
# Must be a valid Ubuntu machine
distro = ubuntu_distro
2022-11-04 14:49:36 -05:00
else:
print("ROOFLINE ERROR: Cannot find a valid binary for your operating system")
sys.exit(1)
target_binary = {"rocm_ver": rocm_ver, "distro": distro}
2022-11-04 14:49:36 -05:00
return target_binary
2023-02-13 14:50:24 -06:00
2022-11-04 14:49:36 -05:00
def mibench(args):
print("No roofline data found. Generating...")
2022-11-04 14:49:36 -05:00
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"]
)
2022-11-04 14:49:36 -05:00
# 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)
2022-11-04 14:49:36 -05:00
sys.exit(1)
2022-11-18 12:50:07 -06:00
run_subprocess(
2022-11-04 14:49:36 -05:00
[
path_to_binary,
"-o",
args.path + "/" + "roofline.csv",
"-d",
str(args.device),
]
)
2023-02-13 14:50:24 -06:00
2023-02-22 14:37:20 -06:00
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)
2023-04-05 14:28:19 -05:00
print("Dispatch Selection: ", args.dispatch)
2023-02-22 14:37:20 -06:00
2022-11-04 14:49:36 -05:00
perfmon_dir = str(OMNIPERF_HOME) + "/perfmon_pub"
print("permon dir is ", os.path.abspath(perfmon_dir))
2023-02-22 14:37:20 -06:00
app_cmd = args.remaining
2023-04-05 14:28:19 -05:00
workload_dir = args.path
2022-11-04 14:49:36 -05:00
# Perfmon filtering
2023-02-22 14:37:20 -06:00
pmc_filter(workload_dir, perfmon_dir, args.target)
2022-11-04 14:49:36 -05:00
2023-05-05 15:07:20 -05:00
# Separate pmc_perf runs
2023-05-08 11:55:45 -05:00
pmc_perf_split(workload_dir)
2023-05-05 15:07:20 -05:00
2023-04-05 14:28:19 -05:00
# Set up a log file
log = open(workload_dir + "/log.txt", "w")
print("Log: ", workload_dir + "/log.txt\n")
2022-11-04 14:49:36 -05:00
# Workload profiling
for fname in glob.glob(workload_dir + "/perfmon/*.txt"):
2023-02-22 14:37:20 -06:00
# Kernel filtering (in-place replacement)
if not args.kernel == None:
2023-05-16 15:39:45 -05:00
success, output = capture_subprocess_output(
2023-02-22 14:37:20 -06:00
[
"sed",
"-i",
"-r",
"s%^(kernel:).*%" + "kernel: " + ",".join(args.kernel) + "%g",
fname,
]
)
2023-05-16 15:39:45 -05:00
log.write(output)
2023-02-22 14:37:20 -06:00
# Dispatch filtering (inplace replacement)
if not args.dispatch == None:
2023-05-16 15:39:45 -05:00
success, output = capture_subprocess_output(
2023-02-22 14:37:20 -06:00
[
"sed",
"-i",
"-r",
"s%^(range:).*%" + "range: " + " ".join(args.dispatch) + "%g",
fname,
]
)
2023-05-16 15:39:45 -05:00
log.write(output)
2022-11-04 14:49:36 -05:00
print(fname)
2023-02-22 14:37:20 -06:00
if args.use_rocscope == True:
run_rocscope(args, fname)
else:
2023-04-05 14:28:19 -05:00
run_prof(fname, workload_dir, perfmon_dir, app_cmd, args.target, log, args.verbose)
2023-06-26 15:30:38 -05:00
# Update timestamps
2023-05-16 15:39:45 -05:00
replace_timestamps(workload_dir, log)
2022-11-04 14:49:36 -05:00
2023-05-05 15:07:20 -05:00
if args.use_rocscope == False:
2023-08-15 14:00:36 -05:00
# Manually join each pmc_perf*.csv output
2023-05-16 15:39:45 -05:00
join_prof(workload_dir, args.join_type, log, args.verbose)
2023-08-15 14:00:36 -05:00
# Demangle and overwrite original KernelNames
csv_processor.kernel_name_shortener(workload_dir, args.kernelVerbose)
2023-08-15 14:00:36 -05:00
2023-05-16 15:39:45 -05:00
log.close()
2023-05-05 15:07:20 -05:00
2023-02-13 14:50:24 -06:00
2022-11-04 14:49:36 -05:00
################################################
# Profiling Helpers
################################################
def run_rocscope(args, fname):
# profile the app
if args.use_rocscope == True:
2023-08-24 15:54:00 -05:00
result = shutil.which("rocscope")
if result:
2023-02-13 14:50:24 -06:00
rs_cmd = [
result.stdout.decode("ascii").strip(),
"metrics",
"-p",
args.path,
"-n",
args.name,
"-t",
fname,
"--",
]
2023-01-25 15:30:04 -06:00
for i in args.remaining.split():
rs_cmd.append(i)
print(rs_cmd)
2023-02-13 14:50:24 -06:00
result = run_subprocess(
rs_cmd
) # , stdout=subprocess.PIPE, stderr=subprocess.PIPE)
if result.returncode != 0:
2023-02-13 14:50:24 -06:00
print(result.stderr.decode("ascii"))
sys.exit(1)
2022-11-04 14:49:36 -05:00
2023-04-05 14:28:19 -05:00
def run_prof(fname, workload_dir, perfmon_dir, cmd, target, log_file, verbose):
2022-11-04 14:49:36 -05:00
global rocprof_cmd
fbase = os.path.splitext(os.path.basename(fname))[0]
if verbose:
print("pmc file:", os.path.basename(fname))
2023-02-15 14:44:29 -06:00
# profile the app (run w/ custom config files for mi100)
if target == "mi100":
print("RUNNING WITH CUSTOM METRICS")
2023-04-05 14:28:19 -05:00
success, output = capture_subprocess_output(
2023-02-15 14:44:29 -06:00
[
rocprof_cmd,
"-i",
fname,
"-m",
perfmon_dir + "/" + "metrics.xml",
"--timestamp",
"on",
"-o",
workload_dir + "/" + fbase + ".csv",
'"' + cmd + '"',
]
)
else:
2023-04-05 14:28:19 -05:00
success, output = capture_subprocess_output(
2023-02-15 14:44:29 -06:00
[
rocprof_cmd,
"-i",
fname,
"--timestamp",
"on",
"-o",
workload_dir + "/" + fbase + ".csv",
'"' + cmd + '"',
]
)
2023-04-05 14:28:19 -05:00
# Write output to log
log_file.write(output)
2022-11-04 14:49:36 -05:00
def omniperf_profile(args, VER):
2023-02-13 14:50:24 -06:00
# Verify valid name
if args.name.find(".") != -1 or args.name.find("-") != -1:
raise ValueError("'-' and '.' are not permited in workload name", args.name)
2022-11-04 14:49:36 -05:00
# 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)
2022-11-04 14:49:36 -05:00
if args.ipblocks == None:
2023-04-05 14:28:19 -05:00
print("IP Blocks: All")
2022-11-04 14:49:36 -05:00
else:
2023-04-05 14:28:19 -05:00
print("IP Blocks: ", args.ipblocks)
2023-08-15 14:00:36 -05:00
if args.kernelVerbose > 5:
print("KernelName verbose level: DISABLED")
else:
print("KernelName verbose level: ", str(args.kernelVerbose))
2022-11-04 14:49:36 -05:00
# Set up directories
workload_dir = args.path + "/" + args.name + "/" + args.target
perfmon_dir = str(OMNIPERF_HOME) + "/perfmon_pub"
2023-02-13 14:50:24 -06:00
2022-11-04 14:49:36 -05:00
# Perfmon filtering
perfmon_filter(workload_dir, perfmon_dir, args)
2023-05-05 15:07:20 -05:00
# Separate pmc_perf runs
pmc_perf_split(workload_dir)
2023-04-05 14:28:19 -05:00
# 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")
2023-02-13 14:50:24 -06:00
# look for whether workload_dir exists - create if not
try:
2023-02-13 14:50:24 -06:00
os.makedirs(workload_dir, exist_ok=True)
except Exception as e:
print("Unable to create workload directory: ", workload_dir)
print(e)
sys.exit(1)
2022-11-04 14:49:36 -05:00
2023-08-24 15:54:00 -05:00
result = shutil.which("rocscope")
if result:
2023-02-13 14:50:24 -06:00
rs_cmd = [
result.stdout.decode("ascii").strip(),
"top10",
"-p",
args.path,
"-n",
args.name,
"--",
]
2023-01-25 15:30:04 -06:00
for i in args.remaining.split():
rs_cmd.append(i)
print(rs_cmd)
2023-02-13 14:50:24 -06:00
result = run_subprocess(
rs_cmd
) # , stdout=subprocess.PIPE, stderr=subprocess.PIPE)
if result.returncode != 0:
2023-02-13 14:50:24 -06:00
print(result.stderr.decode("ascii"))
else:
print("rocscope must be in the PATH")
sys.exit(1)
2022-12-07 23:45:39 +00:00
elif not args.summaries == None and args.summaries == True:
print("creating kernel summaries")
2023-02-13 14:50:24 -06:00
# look for whether workload_dir exists - create if not
2022-12-07 23:45:39 +00:00
try:
2023-02-13 14:50:24 -06:00
os.makedirs(workload_dir, exist_ok=True)
2022-12-07 23:45:39 +00:00
except Exception as e:
print("Unable to create workload directory: ", workload_dir)
print(e)
sys.exit(1)
2022-11-04 14:49:36 -05:00
2023-08-24 15:54:00 -05:00
result = shutil.which("rocscope")
if result:
2023-02-13 14:50:24 -06:00
rs_cmd = [
result.stdout.decode("ascii").strip(),
"summary",
"-p",
args.path,
"-n",
args.name,
"--",
]
2023-01-25 15:30:04 -06:00
for i in args.remaining.split():
rs_cmd.append(i)
2022-12-07 23:45:39 +00:00
print(rs_cmd)
2023-02-13 14:50:24 -06:00
result = run_subprocess(
rs_cmd
) # , stdout=subprocess.PIPE, stderr=subprocess.PIPE)
2022-12-07 23:45:39 +00:00
if result.returncode != 0:
2023-02-13 14:50:24 -06:00
print(result.stderr.decode("ascii"))
2022-12-07 23:45:39 +00:00
else:
print("rocscope must be in the PATH")
sys.exit(1)
2022-11-04 14:49:36 -05:00
else:
for fname in glob.glob(workload_dir + "/perfmon/*.txt"):
# Kernel filtering (in-place replacement)
if not args.kernel == None:
2023-05-16 15:39:45 -05:00
success, output = capture_subprocess_output(
[
"sed",
"-i",
"-r",
"s%^(kernel:).*%" + "kernel: " + ",".join(args.kernel) + "%g",
fname,
]
)
2023-05-16 15:39:45 -05:00
log.write(output)
# Dispatch filtering (inplace replacement)
if not args.dispatch == None:
2023-05-16 15:39:45 -05:00
success, output = capture_subprocess_output(
[
"sed",
"-i",
"-r",
"s%^(range:).*%" + "range: " + " ".join(args.dispatch) + "%g",
fname,
]
)
2023-05-16 15:39:45 -05:00
log.write(output)
2023-02-22 14:37:20 -06:00
print(fname)
if args.use_rocscope == True:
run_rocscope(args, fname)
else:
2023-04-05 14:28:19 -05:00
run_prof(fname, workload_dir, perfmon_dir, args.remaining, args.target, log, args.verbose)
2023-08-10 11:13:27 -05:00
2023-06-26 15:30:38 -05:00
# Update timestamps
2023-05-16 15:39:45 -05:00
replace_timestamps(workload_dir, log)
2023-05-05 15:07:20 -05:00
if args.use_rocscope == False:
2023-08-15 14:00:36 -05:00
# Manually join each pmc_perf*.csv output
2023-05-16 15:39:45 -05:00
join_prof(workload_dir, args.join_type, log, args.verbose)
2023-08-15 14:00:36 -05:00
# Demangle and overwrite original KernelNames
csv_processor.kernel_name_shortener(workload_dir, args.kernelVerbose)
2022-11-04 14:49:36 -05:00
# Generate sysinfo
2022-11-14 09:51:48 -06:00
gen_sysinfo(args.name, workload_dir, args.ipblocks, args.remaining, args.no_roof)
2022-11-04 14:49:36 -05:00
# 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()
2022-11-07 17:04:39 -06:00
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
2022-11-04 14:49:36 -05:00
if not os.path.exists(path_to_binary):
print(
"ROOFLINE ERROR: Unable to locate expected binary (%s))."
% path_to_binary
2022-11-04 14:49:36 -05:00
)
sys.exit(1)
2022-11-18 12:50:07 -06:00
run_subprocess(
2022-11-04 14:49:36 -05:00
[
path_to_binary,
"-o",
workload_dir + "/" + "roofline.csv",
"-d",
str(args.device),
2022-11-04 14:49:36 -05:00
]
)
2023-04-05 14:28:19 -05:00
# Close log
log.close()
2022-11-04 14:49:36 -05:00
2022-11-04 14:49:36 -05:00
################################################
# MAIN
################################################
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"]
2022-11-04 14:49:36 -05:00
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":
2023-08-10 11:13:27 -05:00
Extractionlvl = args.kernelVerbose
print("Resolving rocprof")
global rocprof_cmd
rocprof_cmd = resolve_rocprof()
2023-02-22 14:37:20 -06:00
# Cannot access parent directories
2022-11-04 14:49:36 -05:00
if ".." in str(args.path):
throw_parse_error(
my_parser, "Access denied. Cannot access parent directories in path ../"
)
2023-02-22 14:37:20 -06:00
# Must have a valid soc in profile mode
args.target = get_soc()
2022-11-04 14:49:36 -05:00
# 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,
2023-08-17 09:50:38 -05:00
"Profiling command required. Pass application executable after -- at the end of options.\n\ti.e. omniperf profile -n vcopy -- ./vcopy 1048576 256",
2022-11-04 14:49:36 -05:00
)
# 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
2023-02-22 14:37:20 -06:00
roof_setup(args, my_parser, VER)
2022-11-04 14:49:36 -05:00
# Generate roofline
roofline_only(args.path, args.device, args.sort, args.mem_level, args.kernel_names, args.verbose)
2022-11-04 14:49:36 -05:00
# Profile only
else:
print("\n-------------\nProfile only\n-------------\n")
omniperf_profile(args, VER)
2022-11-04 14:49:36 -05:00
##############
# 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)
2022-11-04 14:49:36 -05:00
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
"""
)
2022-11-04 14:49:36 -05:00
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)
2022-11-04 14:49:36 -05:00
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
""",
2022-11-04 14:49:36 -05:00
)
sys.exit(0) # Indicate successful on exit
if __name__ == "__main__":
main()