AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Initial overhaul of Profile mode

Browse Source 
Signed-off-by: colramos-amd <colramos@amd.com>
This commit is contained in:
colramos-amd
2023-10-31 15:11:17 -05:00
parent c15ec38eca
commit 87388227ca
421 changed files with 16949 additions and 1912 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/utils/csv_processor.py
+6 -5
View File
@@ -33,7 +33,8 @@ import getpass
from pymongo import MongoClient
from tqdm import tqdm
import glob
from common import resolve_rocprof
import re
import logging
cache = dict()
@@ -129,12 +130,12 @@ def kernel_name_shortener(workload_dir, level):
if level < 5:
cpp_filt = os.path.join("/usr", "bin", "c++filt")
if not os.path.isfile(cpp_filt):
print(
logging.error(
"Error: Could not resolve c++filt in expected directory: {}".format(
cpp_filt
)
)
sys.exit(0)
sys.exit(1)
for fpath in glob.glob(workload_dir + "/*.csv"):
try:
@@ -146,9 +147,9 @@ def kernel_name_shortener(workload_dir, level):
modified_df = shorten_file(orig_df, level)
modified_df.to_csv(fpath, index=False)
except pd.errors.EmptyDataError:
print("Skipping empty csv " + str(fpath))
logging.debug("[profiling] Skipping shortening on empty csv " + str(fpath))
print("KernelName shortening complete!")
logging.info("[profiling] KernelName shortening complete!")
# Verify target directory and setup connection
src/utils/perfagg.py
-555
View File
@@ -1,555 +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 sys, os, shutil, glob, re
import numpy as np
import math
import warnings
import pandas as pd
prog = "omniperf"
# Per IP block max number of simulutaneous counters
# GFX IP Blocks
perfmon_config = {
"vega10": {
"SQ": 8,
"TA": 2,
"TD": 2,
"TCP": 4,
"TCC": 4,
"CPC": 2,
"CPF": 2,
"SPI": 2,
"GRBM": 2,
"GDS": 4,
"TCC_channels": 16,
},
"mi50": {
"SQ": 8,
"TA": 2,
"TD": 2,
"TCP": 4,
"TCC": 4,
"CPC": 2,
"CPF": 2,
"SPI": 2,
"GRBM": 2,
"GDS": 4,
"TCC_channels": 16,
},
"mi100": {
"SQ": 8,
"TA": 2,
"TD": 2,
"TCP": 4,
"TCC": 4,
"CPC": 2,
"CPF": 2,
"SPI": 2,
"GRBM": 2,
"GDS": 4,
"TCC_channels": 32,
},
"mi200": {
"SQ": 8,
"TA": 2,
"TD": 2,
"TCP": 4,
"TCC": 4,
"CPC": 2,
"CPF": 2,
"SPI": 2,
"GRBM": 2,
"GDS": 4,
"TCC_channels": 32,
},
}
def test_df_column_equality(df):
return df.eq(df.iloc[:, 0], axis=0).all(1).all()
# joins disparate runs less dumbly than rocprof
def join_prof(workload_dir, join_type, log_file, verbose, out=None):
# Set default output directory if not specified
if type(workload_dir) == str:
if out is None:
out = workload_dir + "/pmc_perf.csv"
files = glob.glob(workload_dir + "/" + "pmc_perf_*.csv")
elif type(workload_dir) == list:
files = workload_dir
else:
print("ERROR: Invalid workload_dir")
sys.exit(1)
df = None
for i, file in enumerate(files):
_df = pd.read_csv(file) if type(workload_dir) == str else file
if join_type == "kernel":
key = _df.groupby("KernelName").cumcount()
_df["key"] = _df.KernelName + " - " + key.astype(str)
elif join_type == "grid":
key = _df.groupby(["KernelName", "grd"]).cumcount()
_df["key"] = (
_df.KernelName + " - " + _df.grd.astype(str) + " - " + key.astype(str)
)
else:
print("ERROR: Unrecognized --join-type")
sys.exit(1)
if df is None:
df = _df
else:
# join by unique index of kernel
df = pd.merge(df, _df, how="inner", on="key", suffixes=("", f"_{i}"))
# TODO: check for any mismatch in joins
duplicate_cols = {
"gpu": [col for col in df.columns if "gpu" in col],
"grd": [col for col in df.columns if "grd" in col],
"wgr": [col for col in df.columns if "wgr" in col],
"lds": [col for col in df.columns if "lds" in col],
"scr": [col for col in df.columns if "scr" in col],
"spgr": [col for col in df.columns if "sgpr" in col],
}
# Check for vgpr counter in ROCm < 5.3
if "vgpr" in df.columns:
duplicate_cols["vgpr"] = [col for col in df.columns if "vgpr" in col]
# Check for vgpr counter in ROCm >= 5.3
else:
duplicate_cols["arch_vgpr"] = [col for col in df.columns if "arch_vgpr" in col]
duplicate_cols["accum_vgpr"] = [col for col in df.columns if "accum_vgpr" in col]
for key, cols in duplicate_cols.items():
_df = df[cols]
if not test_df_column_equality(_df):
msg = (
"WARNING: Detected differing {} values while joining pmc_perf.csv".format(
key
)
)
warnings.warn(msg)
if log_file:
log_file.write(msg + "\n")
else:
msg = "Successfully joined {} in pmc_perf.csv".format(key)
if log_file:
log_file.write(msg + "\n")
if test_df_column_equality(_df) and verbose:
print(msg)
# now, we can:
#   A) throw away any of the "boring" duplicats
df = df[
[
k
for k in df.keys()
if not any(
check in k
for check in [
# removed merged counters, keep original
"gpu-id_",
"grd_",
"wgr_",
"lds_",
"scr_",
"vgpr_",
"sgpr_",
"Index_",
# un-mergable, remove all
"queue-id",
"queue-index",
"pid",
"tid",
"fbar",
"sig",
"obj",
# rocscope specific merged counters, keep original
"dispatch_",
]
)
]
]
#   B) any timestamps that are _not_ the duration, which is the one we care
#   about
df = df[
[
k
for k in df.keys()
if not any(
check in k
for check in [
"DispatchNs",
"CompleteNs",
# rocscope specific timestamp
"HostDuration",
]
)
]
]
#   C) sanity check the name and key
namekeys = [k for k in df.keys() if "KernelName" in k]
assert len(namekeys)
for k in namekeys[1:]:
assert (df[namekeys[0]] == df[k]).all()
df = df.drop(columns=namekeys[1:])
# now take the median of the durations
bkeys = []
ekeys = []
for k in df.keys():
if "Begin" in k:
bkeys.append(k)
if "End" in k:
ekeys.append(k)
# compute mean begin and end timestamps
endNs = df[ekeys].mean(axis=1)
beginNs = df[bkeys].mean(axis=1)
# and replace
df = df.drop(columns=bkeys)
df = df.drop(columns=ekeys)
df["BeginNs"] = beginNs
df["EndNs"] = endNs
# finally, join the drop key
df = df.drop(columns=["key"])
# save to file and delete old file(s), skip if we're being called outside of Omniperf
if type(workload_dir) == str:
df.to_csv(out, index=False)
if not verbose:
for file in files:
os.remove(file)
else:
return df
def pmc_perf_split(workload_dir):
workload_perfmon_dir = workload_dir + "/perfmon"
lines = open(workload_perfmon_dir + "/pmc_perf.txt", "r").read().splitlines()
# Iterate over each line in pmc_perf.txt
mpattern = r"^pmc:(.*)"
i = 0
for line in lines:
# Verify no comments
stext = line.split("#")[0].strip()
if not stext:
continue
# all pmc counters start with "pmc:"
m = re.match(mpattern, stext)
if m is None:
continue
# Create separate file for each line
fd = open(workload_perfmon_dir + "/pmc_perf_" + str(i) + ".txt", "w")
fd.write(stext + "\n\n")
fd.write("gpu:\n")
fd.write("range:\n")
fd.write("kernel:\n")
fd.close()
i += 1
# Remove old pmc_perf.txt input from perfmon dir
os.remove(workload_perfmon_dir + "/pmc_perf.txt")
def update_pmc_bucket(
counters, save_file, soc, pmc_list=None, stext=None, workload_perfmon_dir=None
):
# Verify inputs.
# If save_file is True, we're being called internally, from perfmon_coalesce
# Else we're being called externally, from rocomni
detected_external_call = False
if save_file and (stext is None or workload_perfmon_dir is None):
raise ValueError(
"stext and workload_perfmon_dir must be specified if save_file is True"
)
if pmc_list is None:
detected_external_call = True
pmc_list = dict(
[
("SQ", []),
("GRBM", []),
("TCP", []),
("TA", []),
("TD", []),
("TCC", []),
("SPI", []),
("CPC", []),
("CPF", []),
("GDS", []),
("TCC2", {}), # per-channel TCC perfmon
]
)
for ch in range(perfmon_config[soc]["TCC_channels"]):
pmc_list["TCC2"][str(ch)] = []
if "SQ_ACCUM_PREV_HIRES" in counters and not detected_external_call:
# save all level counters separately
nindex = counters.index("SQ_ACCUM_PREV_HIRES")
level_counter = counters[nindex - 1]
if save_file:
# Save to level counter file, file name = level counter name
fd = open(workload_perfmon_dir + "/" + level_counter + ".txt", "w")
fd.write(stext + "\n\n")
fd.write("gpu:\n")
fd.write("range:\n")
fd.write("kernel:\n")
fd.close()
return pmc_list
# save normal pmc counters in matching buckets
for counter in counters:
IP_block = counter.split(sep="_")[0].upper()
# SQC and SQ belong to the IP block, coalesce them
if IP_block == "SQC":
IP_block = "SQ"
if IP_block != "TCC":
# Insert unique pmc counters into its bucket
if counter not in pmc_list[IP_block]:
pmc_list[IP_block].append(counter)
else:
# TCC counters processing
m = re.match(r"[\s\S]+\[(\d+)\]", counter)
if m is None:
# Aggregated TCC counters
if counter not in pmc_list[IP_block]:
pmc_list[IP_block].append(counter)
else:
# TCC channel ID
ch = m.group(1)
# fake IP block for per channel TCC
if str(ch) in pmc_list["TCC2"]:
# append unique counter into the channel
if counter not in pmc_list["TCC2"][str(ch)]:
pmc_list["TCC2"][str(ch)].append(counter)
else:
# initial counter in this channel
pmc_list["TCC2"][str(ch)] = [counter]
if detected_external_call:
# sort the per channel counter, so that same counter in all channels can be aligned
for ch in range(perfmon_config[soc]["TCC_channels"]):
pmc_list["TCC2"][str(ch)].sort()
return pmc_list
def perfmon_coalesce(pmc_files_list, soc, workload_dir):
workload_perfmon_dir = workload_dir + "/perfmon"
# match pattern for pmc counters
mpattern = r"^pmc:(.*)"
pmc_list = dict(
[
("SQ", []),
("GRBM", []),
("TCP", []),
("TA", []),
("TD", []),
("TCC", []),
("SPI", []),
("CPC", []),
("CPF", []),
("GDS", []),
("TCC2", {}), # per-channel TCC perfmon
]
)
for ch in range(perfmon_config[soc]["TCC_channels"]):
pmc_list["TCC2"][str(ch)] = []
# Extract all PMC counters and store in separate buckets
for fname in pmc_files_list:
lines = open(fname, "r").read().splitlines()
for line in lines:
# Strip all comements, skip empty lines
stext = line.split("#")[0].strip()
if not stext:
continue
# all pmc counters start with "pmc:"
m = re.match(mpattern, stext)
if m is None:
continue
# we have found all the counters, store them in buckets
counters = m.group(1).split()
# Utilitze helper function once a list of counters has be extracted
save_file = True
pmc_list = update_pmc_bucket(
counters, save_file, soc, pmc_list, stext, workload_perfmon_dir
)
# add a timestamp file
fd = open(workload_perfmon_dir + "/timestamps.txt", "w")
fd.write("pmc:\n\n")
fd.write("gpu:\n")
fd.write("range:\n")
fd.write("kernel:\n")
fd.close()
# sort the per channel counter, so that same counter in all channels can be aligned
for ch in range(perfmon_config[soc]["TCC_channels"]):
pmc_list["TCC2"][str(ch)].sort()
return pmc_list
def perfmon_emit(pmc_list, soc, workload_dir=None):
# Calculate the minimum number of iteration to save the pmc counters
# non-TCC counters
pmc_cnt = [
len(pmc_list[key]) / perfmon_config[soc][key]
for key in pmc_list
if key not in ["TCC", "TCC2"]
]
# TCC counters
tcc_channels = perfmon_config[soc]["TCC_channels"]
tcc_cnt = len(pmc_list["TCC"]) / perfmon_config[soc]["TCC"]
tcc2_cnt = (
np.array([len(pmc_list["TCC2"][str(ch)]) for ch in range(tcc_channels)])
/ perfmon_config[soc]["TCC"]
)
# Total number iterations to write pmc: counters line
niter = max(math.ceil(max(pmc_cnt)), math.ceil(tcc_cnt) + math.ceil(max(tcc2_cnt)))
# Emit PMC counters into pmc config file
if workload_dir:
workload_perfmon_dir = workload_dir + "/perfmon"
fd = open(workload_perfmon_dir + "/pmc_perf.txt", "w")
else:
batches = []
tcc2_index = 0
for iter in range(niter):
# Prefix
line = "pmc: "
# Add all non-TCC counters
for key in pmc_list:
if key not in ["TCC", "TCC2"]:
N = perfmon_config[soc][key]
ip_counters = pmc_list[key][iter * N : iter * N + N]
if ip_counters:
line = line + " " + " ".join(ip_counters)
# Add TCC counters
N = perfmon_config[soc]["TCC"]
tcc_counters = pmc_list["TCC"][iter * N : iter * N + N]
if not tcc_counters:
# TCC per-channel counters
for ch in range(perfmon_config[soc]["TCC_channels"]):
tcc_counters += pmc_list["TCC2"][str(ch)][
tcc2_index * N : tcc2_index * N + N
]
tcc2_index += 1
# TCC aggregated counters
line = line + " " + " ".join(tcc_counters)
if workload_dir:
fd.write(line + "\n")
else:
b = line.split()
b.remove("pmc:")
batches.append(b)
if workload_dir:
fd.write("\ngpu:\n")
fd.write("range:\n")
fd.write("kernel:\n")
fd.close()
else:
return batches
def perfmon_filter(workload_dir, perfmon_dir, args):
workload_perfmon_dir = workload_dir + "/perfmon"
soc = args.target
# Initialize directories
# TODO: Modify this so that data is appended to previous?
if not os.path.isdir(workload_dir):
os.makedirs(workload_dir)
else:
shutil.rmtree(workload_dir)
os.makedirs(workload_perfmon_dir)
ref_pmc_files_list = glob.glob(perfmon_dir + "/" + "pmc_*perf*.txt")
ref_pmc_files_list += glob.glob(perfmon_dir + "/" + soc + "/pmc_*_perf*.txt")
# Perfmon list filtering
if args.ipblocks != None:
for i in range(len(args.ipblocks)):
args.ipblocks[i] = args.ipblocks[i].lower()
mpattern = "pmc_([a-zA-Z0-9_]+)_perf*"
pmc_files_list = []
for fname in ref_pmc_files_list:
fbase = os.path.splitext(os.path.basename(fname))[0]
ip = re.match(mpattern, fbase).group(1)
if ip in args.ipblocks:
pmc_files_list.append(fname)
print("fname: " + fbase + ": Added")
else:
print("fname: " + fbase + ": Skipped")
else:
# default: take all perfmons
pmc_files_list = ref_pmc_files_list
# Coalesce and writeback workload specific perfmon
pmc_list = perfmon_coalesce(pmc_files_list, soc, workload_dir)
perfmon_emit(pmc_list, soc, workload_dir)
def pmc_filter(workload_dir, perfmon_dir, soc):
workload_perfmon_dir = workload_dir + "/perfmon"
if not os.path.isdir(workload_perfmon_dir):
os.makedirs(workload_perfmon_dir)
else:
shutil.rmtree(workload_perfmon_dir)
ref_pmc_files_list = glob.glob(perfmon_dir + "/roofline/" + "pmc_roof_perf.txt")
# ref_pmc_files_list += glob.glob(perfmon_dir + "/" + soc + "/pmc_*_perf*.txt")
pmc_files_list = ref_pmc_files_list
# Coalesce and writeback workload specific perfmon
pmc_list = perfmon_coalesce(pmc_files_list, soc, workload_dir)
perfmon_emit(pmc_list, soc, workload_dir)
src/utils/roofline_calc.py
+533
View File
@@ -0,0 +1,533 @@
##############################################################################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 sys
from dataclasses import dataclass
import csv
################################################
# Global vars
################################################
IMGNAME = "empirRoof"
L2_BANKS = 32 # default assuming mi200
XMIN = 0.01
XMAX = 1000
FONT_SIZE = 16
FONT_COLOR = "black"
FONT_WEIGHT = "bold"
SUPPORTED_SOC = ["mi200"]
TOP_N = 10
################################################
# Helper funcs
################################################
@dataclass
class AI_Data:
KernelName: str
numCalls: float
total_flops: float
valu_flops: float
mfma_flops_f16: float
mfma_flops_bf16: float
mfma_flops_f32: float
mfma_flops_f64: float
mfma_iops_i8: float
lds_data: float
L1cache_data: float
L2cache_data: float
hbm_data: float
totalDuration: float
avgDuration: float
def get_font():
return {
"size": FONT_SIZE,
"color": FONT_COLOR,
"weight": FONT_WEIGHT,
"family": "serif",
}
def get_color(catagory):
if catagory == "ai_l1":
return "green"
elif catagory == "ai_l2":
return "blue"
elif catagory == "ai_hbm":
return "red"
else:
raise RuntimeError("Invalid catagory passed to get_color()")
# -------------------------------------------------------------------------------------
# Plot BW at each cache level
# -------------------------------------------------------------------------------------
def calc_ceilings(roof_specs, benchmark_data, targ_mem_level, verbose):
"""Given benchmarking data, calculate ceilings (or peak performance) for empirical roofline
"""
# TODO: This is where filtering by memory level will need to occur for standalone
graphPoints = {"hbm": [], "l2": [], "l1": [], "lds": [], "valu": [], "mfma": []}
if targ_mem_level == "ALL":
cacheHierarchy = ["HBM", "L2", "L1", "LDS"]
else:
cacheHierarchy = targ_mem_level
x1 = y1 = x2 = y2 = -1
x1_mfma = y1_mfma = x2_mfma = y2_mfma = -1
target_precision = roof_specs["dtype"][2:]
if roof_specs["dtype"] != "FP16" and roof_specs["dtype"] != "I8":
peakOps = float(
benchmark_data[roof_specs["dtype"] + "Flops"][roof_specs["device"]]
)
for i in range(0, len(cacheHierarchy)):
# Plot BW line
if verbose >= 3:
print("Current cache level is ", cacheHierarchy[i])
curr_bw = cacheHierarchy[i] + "Bw"
peakBw = float(benchmark_data[curr_bw][roof_specs["device"]])
if roof_specs["dtype"] == "I8":
peakMFMA = float(benchmark_data["MFMAI8Ops"][roof_specs["device"]])
else:
peakMFMA = float(
benchmark_data["MFMAF{}Flops".format(target_precision)][roof_specs["device"]]
)
x1 = float(XMIN)
y1 = float(XMIN) * peakBw
# Note: No reg peakOps for FP16 or INT8
if roof_specs["dtype"] != "FP16" and roof_specs["dtype"] != "I8":
x2 = peakOps / peakBw
y2 = peakOps
# Plot MFMA lines (NOTE: Assuming MI200 soc)
x1_mfma = peakOps / peakBw
y1_mfma = peakOps
x2_mfma = peakMFMA / peakBw
y2_mfma = peakMFMA
# These are the points to use:
if verbose >= 3:
print("x = [{}, {}]".format(x1, x2_mfma))
print("y = [{}, {}]".format(y1, y2_mfma))
graphPoints[cacheHierarchy[i].lower()].append([x1, x2_mfma])
graphPoints[cacheHierarchy[i].lower()].append([y1, y2_mfma])
graphPoints[cacheHierarchy[i].lower()].append(peakBw)
# -------------------------------------------------------------------------------------
# Plot computing roof
# -------------------------------------------------------------------------------------
# Note: No FMA roof for FP16 or INT8
if roof_specs["dtype"] != "FP16" and roof_specs["dtype"] != "I8":
# Plot FMA roof
x0 = XMAX
if x2 < x0:
x0 = x2
if verbose >= 3:
print("FMA ROOF [{}, {}], [{},{}]".format(x0, XMAX, peakOps, peakOps))
graphPoints["valu"].append([x0, XMAX])
graphPoints["valu"].append([peakOps, peakOps])
graphPoints["valu"].append(peakOps)
# Plot MFMA roof
if (
x1_mfma != -1 or roof_specs["dtype"] == "FP16" or roof_specs["dtype"] == "I8"
): # assert that mfma has been assigned
x0_mfma = XMAX
if x2_mfma < x0_mfma:
x0_mfma = x2_mfma
if verbose >= 3:
print("MFMA ROOF [{}, {}], [{},{}]".format(x0_mfma, XMAX, peakMFMA, peakMFMA))
graphPoints["mfma"].append([x0_mfma, XMAX])
graphPoints["mfma"].append([peakMFMA, peakMFMA])
graphPoints["mfma"].append(peakMFMA)
return graphPoints
# -------------------------------------------------------------------------------------
# Overlay application performance
# -------------------------------------------------------------------------------------
# Calculate relevent metrics for ai calculation
def calc_ai(sort_type, ret_df, verbose):
"""Given counter data, caclulate arithmetic intensity for each kernel in the application.
"""
df = ret_df["pmc_perf"]
# Sort by top kernels or top dispatches?
df = df.sort_values(by=["KernelName"])
df = df.reset_index(drop=True)
total_flops = (
valu_flops
) = (
mfma_flops_bf16
) = (
mfma_flops_f16
) = (
mfma_iops_i8
) = (
mfma_flops_f32
) = (
mfma_flops_f64
) = (
lds_data
) = L1cache_data = L2cache_data = hbm_data = calls = totalDuration = avgDuration = 0.0
kernelName = ""
myList = []
at_end = False
next_kernelName = ""
for idx in df.index:
# CASE: Top kernels
# Calculate + append AI data if
# a) current KernelName is different than previous OR
# b) We've reached the end of list
if idx + 1 == df.shape[0]:
at_end = True
else:
next_kernelName = df["KernelName"][idx + 1]
kernelName = df["KernelName"][idx]
try:
total_flops += (
(
64
* (
df["SQ_INSTS_VALU_ADD_F16"][idx]
+ df["SQ_INSTS_VALU_MUL_F16"][idx]
+ (2 * df["SQ_INSTS_VALU_FMA_F16"][idx])
+ df["SQ_INSTS_VALU_TRANS_F16"][idx]
)
)
+ (
64
* (
df["SQ_INSTS_VALU_ADD_F32"][idx]
+ df["SQ_INSTS_VALU_MUL_F32"][idx]
+ (2 * df["SQ_INSTS_VALU_FMA_F32"][idx])
+ df["SQ_INSTS_VALU_TRANS_F32"][idx]
)
)
+ (
64
* (
df["SQ_INSTS_VALU_ADD_F64"][idx]
+ df["SQ_INSTS_VALU_MUL_F64"][idx]
+ (2 * df["SQ_INSTS_VALU_FMA_F64"][idx])
+ df["SQ_INSTS_VALU_TRANS_F64"][idx]
)
)
+ (df["SQ_INSTS_VALU_MFMA_MOPS_F16"][idx] * 512)
+ (df["SQ_INSTS_VALU_MFMA_MOPS_BF16"][idx] * 512)
+ (df["SQ_INSTS_VALU_MFMA_MOPS_F32"][idx] * 512)
+ (df["SQ_INSTS_VALU_MFMA_MOPS_F64"][idx] * 512)
)
except KeyError:
if verbose >= 3:
print("{}: Skipped total_flops at index {}".format(kernelName[:35], idx))
pass
try:
valu_flops += (
64
* (
df["SQ_INSTS_VALU_ADD_F16"][idx]
+ df["SQ_INSTS_VALU_MUL_F16"][idx]
+ (2 * df["SQ_INSTS_VALU_FMA_F16"][idx])
+ df["SQ_INSTS_VALU_TRANS_F16"][idx]
)
+ 64
* (
df["SQ_INSTS_VALU_ADD_F32"][idx]
+ df["SQ_INSTS_VALU_MUL_F32"][idx]
+ (2 * df["SQ_INSTS_VALU_FMA_F32"][idx])
+ df["SQ_INSTS_VALU_TRANS_F32"][idx]
)
+ 64
* (
df["SQ_INSTS_VALU_ADD_F64"][idx]
+ df["SQ_INSTS_VALU_MUL_F64"][idx]
+ (2 * df["SQ_INSTS_VALU_FMA_F64"][idx])
+ df["SQ_INSTS_VALU_TRANS_F64"][idx]
)
)
except KeyError:
if verbose >= 3:
print("{}: Skipped valu_flops at index {}".format(kernelName[:35], idx))
pass
try:
mfma_flops_f16 += df["SQ_INSTS_VALU_MFMA_MOPS_F16"][idx] * 512
mfma_flops_bf16 += df["SQ_INSTS_VALU_MFMA_MOPS_BF16"][idx] * 512
mfma_flops_f32 += df["SQ_INSTS_VALU_MFMA_MOPS_F32"][idx] * 512
mfma_flops_f64 += df["SQ_INSTS_VALU_MFMA_MOPS_F64"][idx] * 512
mfma_iops_i8 += df["SQ_INSTS_VALU_MFMA_MOPS_I8"][idx] * 512
except KeyError:
if verbose >= 3:
print("{}: Skipped mfma ops at index {}".format(kernelName[:35], idx))
pass
try:
lds_data += (
(df["SQ_LDS_IDX_ACTIVE"][idx] - df["SQ_LDS_BANK_CONFLICT"][idx])
* 4
* L2_BANKS
) # L2_BANKS = 32 (since assuming mi200)
except KeyError:
if verbose >= 3:
print("{}: Skipped lds_data at index {}".format(kernelName[:35], idx))
pass
try:
L1cache_data += df["TCP_TOTAL_CACHE_ACCESSES_sum"][idx] * 64
except KeyError:
if verbose >= 3:
print("{}: Skipped L1cache_data at index {}".format(kernelName[:35], idx))
pass
try:
L2cache_data += (
df["TCP_TCC_WRITE_REQ_sum"][idx] * 64
+ df["TCP_TCC_ATOMIC_WITH_RET_REQ_sum"][idx] * 64
+ df["TCP_TCC_ATOMIC_WITHOUT_RET_REQ_sum"][idx] * 64
+ df["TCP_TCC_READ_REQ_sum"][idx] * 64
)
except KeyError:
if verbose >= 3:
print("{}: Skipped L2cache_data at index {}".format(kernelName[:35], idx))
pass
try:
hbm_data += (
(df["TCC_EA_RDREQ_32B_sum"][idx] * 32)
+ ((df["TCC_EA_RDREQ_sum"][idx] - df["TCC_EA_RDREQ_32B_sum"][idx]) * 64)
+ (df["TCC_EA_WRREQ_64B_sum"][idx] * 64)
+ ((df["TCC_EA_WRREQ_sum"][idx] - df["TCC_EA_WRREQ_64B_sum"][idx]) * 32)
)
except KeyError:
if verbose >= 3:
print("{}: Skipped hbm_data at index {}".format(kernelName[:35], idx))
pass
totalDuration += df["EndNs"][idx] - df["BeginNs"][idx]
avgDuration += df["EndNs"][idx] - df["BeginNs"][idx]
calls += 1
if sort_type == "kernels" and (at_end == True or (kernelName != next_kernelName)):
myList.append(
AI_Data(
kernelName,
calls,
total_flops / calls,
valu_flops / calls,
mfma_flops_f16 / calls,
mfma_flops_bf16 / calls,
mfma_flops_f32 / calls,
mfma_flops_f64 / calls,
mfma_iops_i8 / calls,
lds_data / calls,
L1cache_data / calls,
L2cache_data / calls,
hbm_data / calls,
totalDuration,
avgDuration / calls,
)
)
if verbose >= 2:
print(
"Just added {} to AI_Data at index {}. # of calls: {}".format(
kernelName, idx, calls
)
)
total_flops = (
valu_flops
) = (
mfma_flops_bf16
) = (
mfma_flops_f16
) = (
mfma_iops_i8
) = (
mfma_flops_f32
) = (
mfma_flops_f64
) = (
lds_data
) = (
L1cache_data
) = L2cache_data = hbm_data = calls = totalDuration = avgDuration = 0.0
if sort_type == "dispatches":
myList.append(
AI_Data(
kernelName,
calls,
total_flops,
valu_flops,
mfma_flops_f16,
mfma_flops_bf16,
mfma_flops_f32,
mfma_flops_f64,
mfma_iops_i8,
lds_data,
L1cache_data,
L2cache_data,
hbm_data,
totalDuration,
avgDuration,
)
)
total_flops = (
valu_flops
) = (
mfma_flops_bf16
) = (
mfma_flops_f16
) = (
mfma_iops_i8
) = (
mfma_flops_f32
) = (
mfma_flops_f64
) = (
lds_data
) = (
L1cache_data
) = L2cache_data = hbm_data = calls = totalDuration = avgDuration = 0.0
myList.sort(key=lambda x: x.totalDuration, reverse=True)
# print("Top 5 intensities ('{}')...".format(roof_details["sort"]))
intensities = {"ai_l1": [], "ai_l2": [], "ai_hbm": []}
curr_perf = []
kernelNames = []
i = 0
# Create list of top 5 intensities
while i < TOP_N and i != len(myList):
kernelNames.append(myList[i].KernelName)
intensities["ai_l1"].append(
myList[i].total_flops / myList[i].L1cache_data
) if myList[i].L1cache_data else intensities["ai_l1"].append(0)
# print("cur_ai_L1", myList[i].total_flops/myList[i].L1cache_data) if myList[i].L1cache_data else print("null")
# print()
intensities["ai_l2"].append(
myList[i].total_flops / myList[i].L2cache_data
) if myList[i].L2cache_data else intensities["ai_l2"].append(0)
# print("cur_ai_L2", myList[i].total_flops/myList[i].L2cache_data) if myList[i].L2cache_data else print("null")
# print()
intensities["ai_hbm"].append(
myList[i].total_flops / myList[i].hbm_data
) if myList[i].hbm_data else intensities["ai_hbm"].append(0)
# print("cur_ai_hbm", myList[i].total_flops/myList[i].hbm_data) if myList[i].hbm_data else print("null")
# print()
curr_perf.append(myList[i].total_flops / myList[i].avgDuration) if myList[
i
].avgDuration else curr_perf.append(0)
# print("cur_perf", myList[i].total_flops/myList[i].avgDuration) if myList[i].avgDuration else print("null")
i += 1
intensityPoints = {"ai_l1": [], "ai_l2": [], "ai_hbm": []}
for i in intensities:
values = intensities[i]
color = get_color(i)
x = []
y = []
for entryIndx in range(0, len(values)):
x.append(values[entryIndx])
y.append(curr_perf[entryIndx])
intensityPoints[i].append(x)
intensityPoints[i].append(y)
# Add an entry for kernel names
intensityPoints["kernelNames"] = kernelNames
return intensityPoints
def constuct_roof(roof_specs, targ_mem_level, verbose):
benchmark_results = roof_specs["path"] + "/roofline.csv"
# -----------------------------------------------------
# Initialize roofline data dictionary from roofline.csv
# -----------------------------------------------------
benchmark_data = (
{}
) # TODO: consider changing this to an ordered dict for consistency over py versions
headers = []
try:
with open(benchmark_results, "r") as csvfile:
csvReader = csv.reader(csvfile, delimiter=",")
rowCount = 0
for row in csvReader:
row.pop(0) # remove devID
if rowCount == 0:
headers = row
for i in headers:
benchmark_data[i] = []
else:
for i, key in enumerate(headers):
benchmark_data[key].append(row[i])
rowCount += 1
csvfile.close()
except:
graphPoints = {
"hbm": [None, None, None],
"l2": [None, None, None],
"l1": [None, None, None],
"lds": [None, None, None],
"valu": [None, None, None],
"mfma": [None, None, None],
}
return graphPoints
# ------------------
# Generate Roofline
# ------------------
results = calc_ceilings(roof_specs, benchmark_data, targ_mem_level, verbose)
# for key in results:
# print(key, "->", results[key])
return results
src/utils/utils.py
+367 -1
View File
@@ -23,6 +23,21 @@
##############################################################################el
import logging
import sys
import subprocess
import shutil
import os
import io
import selectors
import pandas as pd
import glob
from utils import specs
from datetime import datetime
from pathlib import Path as path
import config
rocprof_cmd = ""
def demarcate(function):
def wrap_function(*args, **kwargs):
@@ -33,4 +48,355 @@ def demarcate(function):
return wrap_function
def trace_logger(message, *args, **kwargs):
logging.log(logging.TRACE, message, *args, **kwargs)
logging.log(logging.TRACE, message, *args, **kwargs)
def get_version(omniperf_home) -> dict:
"""Return Omniperf versioning info
"""
# 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:
logging.critical("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:
logging.error("ERROR: Cannot find VERSION.sha file at {}".format(shaFile))
sys.exit(1)
MODE = "release"
versionData = {"version": VER, "sha": SHA, "mode": MODE}
return versionData
def get_version_display(version, sha, mode):
"""Pretty print versioning info
"""
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()
def detect_rocprof():
"""Detect loaded rocprof version. Resolve path and set cmd globally.
"""
global rocprof_cmd
# rocprof info
if not "ROCPROF" in os.environ.keys():
rocprof_cmd = "rocprofv2"
else:
rocprof_cmd = os.environ["ROCPROF"]
rocprof_path = shutil.which(rocprof_cmd)
# TODO: this could be more elegant, clean code later
if not rocprof_path:
rocprof_cmd = "rocprof"
rocprof_path = shutil.which(rocprof_cmd)
if not rocprof_path:
logging.error("\nError: Unable to resolve path to %s binary" % rocprof_cmd)
logging.error(
"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"))
logging.info("ROC Profiler: " + str(rocprof_path))
return rocprof_cmd #TODO: Do we still need to return this? It's not being used in the function call
def capture_subprocess_output(subprocess_args):
"""Run specified subprocess and concurrently capture output
"""
# 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 run_prof(fname, workload_dir, perfmon_dir, cmd, target, verbose):
fbase = os.path.splitext(os.path.basename(fname))[0]
logging.debug("pmc file:", os.path.basename(fname))
# profile the app (run w/ custom config files for mi100)
if target == "mi100":
logging.info("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 rocprof output to logging
logging.info(output)
def replace_timestamps(workload_dir):
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."
logging.warning(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 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):
logging._SysExcInfoType("Detected user-supplied binary")
return {"rocm_ver": "override", "distro": "override", "path": rooflineBinary}
else:
logging.error("ROOFLINE ERROR: user-supplied path to binary not accessible")
logging.error("--> 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:
logging.error("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 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)
logging.info(rs_cmd)
success, output = capture_subprocess_output(
rs_cmd
)
if not success:
logging.error(result.stderr.decode("ascii"))
sys.exit(1)
def mibench(args):
"""Run roofline microbenchmark to generate peak BW and FLOP measurements.
"""
logging.info("[roofline] No roofline data found. Generating...")
distro_map = {"platform:el8": "rhel8", "15.3": "sle15sp3", "20.04": "ubuntu20_04"}
target_binary = detect_roofline()
if target_binary["rocm_ver"] == "override":
path_to_binary = target_binary["path"]
else:
path_to_binary = (
str(config.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):
logging.error("ROOFLINE ERROR: Unable to locate expected binary (%s)." % path_to_binary)
sys.exit(1)
subprocess.run(
[
path_to_binary,
"-o",
args.path + "/" + "roofline.csv",
"-d",
str(args.device),
],
check=True
)