Initial commit

src/utils/csv_converter.py

@@ -0,0 +1,264 @@
+################################################################################
+# Copyright (c) 2021 - 2022 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.
+################################################################################
+
+import argparse
+import collections
+import os
+import sys
+import re
+import pandas as pd
+import getpass
+from pymongo import MongoClient
+from tqdm import tqdm
+import shutil
+
+cache = dict()
+supported_arch = {"gfx906": "mi50", "gfx908": "mi100", "gfx90a": "mi200"}
+MAX_SERVER_SEL_DELAY = 5000  # 5 sec connection timeout
+
+
+def kernel_name_shortener(df, cache, level):
+
+    if level >= 5:
+        return df
+
+    columnName = ""
+    if "KernelName" in df:
+        columnName = "KernelName"
+    if "Name" in df:
+        columnName = "Name"
+
+    if columnName == "KernelName" or columnName == "Name":
+        # loop through all indices
+        for index in df.index:
+
+            original_name = df.loc[index, columnName]
+            if original_name in cache:
+                continue
+
+            # cache miss, add the shortened name to the dictionary
+            new_name = ""
+            matches = ""
+
+            names_and_args = re.compile(r"(?P<name>[( )A-Za-z0-9_]+)([ ,*<>()]+)(::)?")
+
+            # works for name Kokkos::namespace::init_lock_array_kernel_threadid(int) [clone .kd]
+            if names_and_args.search(original_name):
+                matches = names_and_args.findall(original_name)
+            else:
+                # Works for first case  '__amd_rocclr_fillBuffer.kd'
+                # remove .kd and then parse through original regex
+                first_case = re.compile(r"([^\s]+)(.kd)")
+                Mod_name_and_args = re.compile(r"(?P<name>[( )A-Za-z0-9_]+)([ ,*<>()]*)")
+                interim_name = first_case.search(original_name).group(1)
+                matches = Mod_name_and_args.findall(interim_name)
+
+            current_level = 0
+            for name in matches:
+                ##can cause errors if a function name or argument is equal to 'clone'
+                if name[0] == "clone":
+                    continue
+                if len(name) == 3:
+                    if name[2] == "::":
+                        continue
+
+                if current_level < level:
+                    new_name += name[0]
+                # closing '>' is to be taken account by the while loop
+                if name[1].count(">") == 0:
+                    if current_level < level:
+                        if not (current_level == level - 1 and name[1].count("<") > 0):
+                            new_name += name[1]
+                    current_level += name[1].count("<")
+
+                curr_index = 0
+                # cases include '>'  '> >, ' have to go in depth here to not lose account of commas and current level
+                while name[1].count(">") > 0 and curr_index < len(name[1]):
+                    if current_level < level:
+                        new_name += name[1][curr_index:]
+                        current_level -= name[1][curr_index:].count(">")
+                        curr_index = len(name[1])
+                    elif name[1][curr_index] == (">"):
+                        current_level -= 1
+                    curr_index += 1
+
+            cache[original_name] = new_name
+            if new_name == None or new_name == "":
+                cache[original_name] = original_name
+
+        df[columnName] = df[columnName].map(cache)
+
+    return df
+
+
+# Verify target directory and setup connection
+def parse(args, profileAndExport):
+    host = args.host
+    port = str(args.port)
+    username = args.username
+    Extractionlvl = args.kernelVerbose
+
+    if profileAndExport:
+        workload = args.workload + "/" + args.target + "/"
+    else:
+        workload = args.workload
+
+    # Verify directory path is valid
+    print("Pulling data from ", workload)
+    if os.path.isdir(workload):
+        print("The directory exists")
+    else:
+        raise argparse.ArgumentTypeError("Directory does not exist")
+
+    sysInfoPath = workload + "/sysinfo.csv"
+    if os.path.isfile(sysInfoPath):
+        print("Found sysinfo file")
+        sysInfo = pd.read_csv(sysInfoPath)
+        # Extract SoC
+        arch = sysInfo["gpu_soc"][0]
+        soc = supported_arch[arch]
+        # Extract name
+        name = sysInfo["workload_name"][0]
+    else:
+        print("Unable to parse SoC or workload name from sysinfo.csv")
+        sys.exit(1)
+
+    db = "omniperf_" + args.team + "_" + name + "_" + soc
+
+    if Extractionlvl >= 5:
+        print("KernelName shortening disabled")
+    else:
+        print("KernelName shortening enabled")
+
+    print("Kernel name verbose level:", Extractionlvl)
+
+    if args.password == "":
+        try:
+            password = getpass.getpass()
+        except Exception as error:
+            print("PASSWORD ERROR", error)
+        else:
+            print("Password recieved")
+    else:
+        password = args.password
+
+    if db.find(".") != -1 or db.find("-") != -1:
+        raise ValueError("'-' and '.' are not permited in workload name", db)
+
+    connectionInfo = {
+        "username": username,
+        "password": password,
+        "host": host,
+        "port": port,
+        "workload": workload,
+        "db": db,
+    }
+
+    return connectionInfo, Extractionlvl
+
+
+def convert_folder(connectionInfo, Extractionlvl):
+    # Test connection
+    connection_str = (
+        "mongodb://"
+        + connectionInfo["username"]
+        + ":"
+        + connectionInfo["password"]
+        + "@"
+        + connectionInfo["host"]
+        + ":"
+        + connectionInfo["port"]
+        + "/?authSource=admin"
+    )
+    client = MongoClient(connection_str, serverSelectionTimeoutMS=MAX_SERVER_SEL_DELAY)
+    try:
+        client.server_info()
+    except:
+        print("ERROR: Unable to connect to the server")
+        sys.exit(1)
+    # Set up directories
+    if Extractionlvl < 5:
+        newfilepath = connectionInfo["workload"]
+        newfilepath_h = newfilepath + "/renamedFiles/"
+        if not os.path.exists(newfilepath_h):
+            os.mkdir(newfilepath_h)
+        newfilepath = newfilepath_h + connectionInfo["db"] + "/"
+        if not os.path.exists(newfilepath):
+            os.mkdir(newfilepath)
+    # Upload files
+    i = 0
+    file = "blank"
+    for file in tqdm(os.listdir(connectionInfo["workload"])):
+        if file.endswith(".csv"):
+            print(connectionInfo["workload"] + "/" + file)
+            try:
+                fileName = file[0 : file.find(".")]
+                # Only shorten KernelNames if instructed to
+                if Extractionlvl < 5:
+                    t1 = pd.read_csv(
+                        connectionInfo["workload"] + "/" + file,
+                        on_bad_lines="skip",
+                        engine="python",
+                    )
+
+                    t2 = kernel_name_shortener(t1, cache, level=Extractionlvl)
+                    df_saved_file = t2.to_csv(newfilepath + file)
+
+                    cmd = (
+                        'mongoimport --quiet --uri mongodb://{}:{}@{}:{} --authenticationDatabase "admin" --file {} -d {} -c {} --drop --type csv --headerline'
+                    ).format(
+                        connectionInfo["username"],
+                        connectionInfo["password"],
+                        connectionInfo["host"],
+                        connectionInfo["port"],
+                        newfilepath + file,
+                        connectionInfo["db"],
+                        fileName,
+                    )
+                    os.system(cmd)
+                else:
+                    cmd = (
+                        'mongoimport --quiet --uri mongodb://{}:{}@{}:{} --authenticationDatabase "admin" --file {} -d {} -c {} --drop --type csv --headerline'
+                    ).format(
+                        connectionInfo["username"],
+                        connectionInfo["password"],
+                        connectionInfo["host"],
+                        connectionInfo["port"],
+                        connectionInfo["workload"] + "/" + file,
+                        connectionInfo["db"],
+                        fileName,
+                    )
+                    os.system(cmd)
+                i += 1
+            except pd.errors.EmptyDataError:
+                print("Skipping empty csv " + file)
+
+    mydb = client["workload_names"]
+    mycol = mydb["names"]
+    value = {"name": connectionInfo["db"]}
+    newValue = {"name": connectionInfo["db"]}
+    mycol.replace_one(value, newValue, upsert=True)
+    # Remove tmp directory if we shortened KernelNames
+    if Extractionlvl < 5:
+        shutil.rmtree(newfilepath_h)
+    print("{} collections added.".format(i))
+    print("Workload name uploaded")

src/utils/gfx_perfmon_builder.py

@@ -0,0 +1,402 @@
+################################################################################
+# Copyright (c) 2021 - 2022 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.
+################################################################################
+
+from PyQt5.QtWidgets import (
+    QMainWindow,
+    QApplication,
+    QTreeView,
+    QTableWidget,
+    QTableWidgetItem,
+)
+from PyQt5.QtWidgets import (
+    QHBoxLayout,
+    QWidget,
+    QAction,
+    QFileDialog,
+    QAbstractItemView,
+    qApp,
+)
+from PyQt5.QtGui import QStandardItemModel, QStandardItem
+from lxml import html
+import sys
+
+
+# class view(QWidget):
+class mainWindow(QMainWindow):
+    def __init__(self):
+        super(QMainWindow, self).__init__()
+
+        ###############################################################################
+        # SOC Parameters
+        ##############################################################################
+
+        # Per IP block max number of simulutaneous counters
+        # GFX IP Blocks
+        self.perfmon_config = {
+            "SQ": 8,
+            "TA": 2,
+            "TD": 2,
+            "TCP": 4,
+            "TCC": 4,
+            "CPC": 2,
+            "CPF": 2,
+            "SPI": 2,
+            "GRBM": 2,
+            "GDS": 4,
+        }
+
+        # GFX Architectures
+        self.soc_arch_list = ["gfx906", "gfx908", "gfx90a"]
+
+        ###############################################################################
+        # Window layout Design
+        ##############################################################################
+
+        self.block_list = []
+        self.nodes_dict = {}  # list of QStandardItem
+
+        self.tree = QTreeView(self)
+        self.table = QTableWidget()
+
+        # Tree setup
+        self.tree.header().setDefaultSectionSize(180)
+        self.model = QStandardItemModel()
+        self.model.setHorizontalHeaderLabels(["Metric", "Block", "Event", "Definition"])
+
+        self.tree.setModel(self.model)
+        # self.importData(data)
+        self.tree.setEditTriggers(QAbstractItemView.NoEditTriggers)
+
+        # Set up click processing
+        self.tree.clicked.connect(self.pmc_select)
+        # self.tree.expandAll()
+
+        # Table setup
+        tableHeader = list(self.perfmon_config.keys())
+        self.table.setColumnCount(len(tableHeader))
+        self.table.setHorizontalHeaderLabels(tableHeader)
+        self.table.setEditTriggers(QAbstractItemView.NoEditTriggers)
+        self.table.showGrid()
+
+        self.setWindowTitle("GFX Perfmon Builder")
+        # layout: lhs: metrics; rhs: selected perfmon
+        layout = QHBoxLayout(self)
+        layout.addWidget(self.tree)
+        layout.addWidget(self.table)
+
+        widget = QWidget()
+        widget.setLayout(layout)
+        self.setCentralWidget(widget)
+
+        # Add Status
+        self.statusBar()
+
+        ###############################################################################
+        # Window Menu Design
+        ##############################################################################
+
+        # Setup file menu
+        menuBar = self.menuBar()
+        menuBar.setNativeMenuBar(False)
+
+        openAction = QAction("&Open", self)
+        openAction.setShortcut("Ctrl+O")
+        openAction.setStatusTip("Open GFX Metrics file")
+        openAction.triggered.connect(self.openGFXDialog)
+
+        saveAction = QAction("&Save", self)
+        saveAction.setShortcut("Ctrl+S")
+        saveAction.setStatusTip("Save to PMC file")
+        saveAction.triggered.connect(self.exportGFXDialog)
+
+        exitAction = QAction("&Exit", self)
+        exitAction.setShortcut("Ctrl+Q")
+        exitAction.setStatusTip("Exit")
+        exitAction.triggered.connect(self.close)
+
+        # Create new action
+        fileMenu = menuBar.addMenu("&File")
+        fileMenu.addActions([openAction, saveAction])
+        fileMenu.addSeparator()
+        fileMenu.addActions([exitAction])
+
+    def openGFXDialog(self):
+        options = QFileDialog.Options()
+        options |= QFileDialog.DontUseNativeDialog
+        fileName, _ = QFileDialog.getOpenFileName(
+            self, "Open GFX Metrics", "", "XML Files (*.xml)", "XML(*.xml)"
+        )
+
+        # Parse the xml
+        if fileName:
+            xmlparsed = html.parse(fileName)
+            self.importData(xmlparsed)
+
+    def exportGFXDialog(self):
+        options = QFileDialog.Options()
+        options |= QFileDialog.DontUseNativeDialog
+        fileName, _ = QFileDialog.getSaveFileName(
+            self, "Export PMC Counters", "", "Text File (*.txt)", "Text File(*.txt)"
+        )
+
+        # Parse the xml
+        if fileName:
+            self.exportPMCCounters(fileName)
+
+    def exportPMCCounters(self, fileName):
+
+        f = open(fileName, "w")
+
+        total_IP_blocks = len(list(self.perfmon_config.keys()))
+        for row in range(self.table.rowCount()):
+            pmc_str = "pmc: "
+            for col in range(total_IP_blocks):
+                cell = self.table.item(row, col)
+
+                if cell:
+                    pmc_str = pmc_str + "  ".join(cell.text().split("\n")) + "  "
+
+            f.write(pmc_str + "\n")
+
+        # Add standard lines
+        f.write("\n\n")
+        f.write("gpu: \n")
+        f.write("dispatch: \n")
+        f.write("kernel: \n")
+
+        f.close()
+
+        return
+
+    def pmc_metric_selected(self, metric_name, col):
+
+        # check if the metric already exists
+        metric_selected = False
+
+        for row in range(self.table.rowCount()):
+            entry = self.table.item(row, col)
+            if entry:
+                pmc_list = entry.text().split(sep="\n")
+                if metric_name in pmc_list:
+                    metric_selected = True
+                    break
+
+        return metric_selected
+
+    def pmc_remove_metric(self, metric_name, IP_block):
+        # Remove the metric to pmc table, if it is selected
+
+        # Map SQC to SQ, since they share the same Perfmon block
+        if IP_block == "SQC":
+            IP_block = "SQ"
+
+        # not action if it is for a ghost IP!
+        if not IP_block in list(self.perfmon_config.keys()):
+            return
+
+        # This is the column we need to add/remove perfmon counters
+        col = list(self.perfmon_config.keys()).index(IP_block)
+
+        if not self.pmc_metric_selected(metric_name, col):
+            return
+
+        pmc_list = []
+        for row in range(self.table.rowCount()):
+            entry = self.table.item(row, col)
+
+            if entry:
+                pmc_list = pmc_list + entry.text().split(sep="\n")
+                # clear the cell, we will re-allocate the pmc
+
+            self.table.takeItem(row, col)
+
+        # allowed PMC counters per batch
+        max_pmc_num = self.perfmon_config[IP_block]
+
+        # remote this metric and re-segment the list and refill all rows in this column
+        pmc_list.remove(metric_name)
+
+        # We are empty now, do nothing
+        if len(pmc_list) == 0:
+            return
+
+        for row in range((len(pmc_list) + max_pmc_num - 1) // max_pmc_num):
+            start_index = row * max_pmc_num
+            pmc_str = "\n".join(pmc_list[start_index : start_index + max_pmc_num])
+            self.table.setItem(row, col, QTableWidgetItem(pmc_str))
+
+        # Remove last row, if empty
+        last_row = self.table.rowCount() - 1
+        empty_row = True
+        total_cols = len(list(self.perfmon_config.keys()))
+        for cindex in range(total_cols):
+            x = self.table.item(last_row, cindex)
+
+            if x and x.text():
+                empty_row = False
+                break
+
+        if empty_row:
+            self.table.removeRow(last_row)
+
+    def pmc_add_metric(self, metric_name, IP_block):
+        # Add the metric to pmc table, if not there yet
+
+        # Map SQC to SQ, since they share the same Perfmon block
+        if IP_block == "SQC":
+            IP_block = "SQ"
+
+        if not IP_block in list(self.perfmon_config.keys()):
+            return
+
+        # This is the column we need to add/remove perfmon counters
+        col = list(self.perfmon_config.keys()).index(IP_block)
+
+        # check if the metric already exists
+        if self.pmc_metric_selected(metric_name, col):
+            return
+
+        # metric is not bucket yet, add it!
+        if self.table.rowCount() == 0:
+            # starting from scratch!
+            self.table.insertRow(0)
+            self.table.setItem(0, col, QTableWidgetItem(metric_name))
+            return
+
+        # find the row to insert
+        for row in range(self.table.rowCount()):
+            entry = self.table.item(row, col)
+            if not entry:
+                # print("search insert pos, row:", row, ", cell empty")
+                break
+
+            if len(entry.text().split(sep="\n")) < self.perfmon_config[IP_block]:
+                # print("found")
+                break
+
+        entry = self.table.item(row, col)
+        if not entry:
+            # put it into the empty cell
+            self.table.setItem(row, col, QTableWidgetItem(metric_name))
+            return
+
+        pmc_list = entry.text().split(sep="\n")
+
+        if len(pmc_list) < self.perfmon_config[IP_block]:
+            # we still have hit per-IP HW counters limit, add it to the last row
+            pmc_list.append(metric_name)
+            pmc_str = "\n".join(pmc_list)
+            self.table.setItem(row, col, QTableWidgetItem(pmc_str))
+            self.table.resizeRowsToContents()
+        else:
+            # Start a new row
+            row = row + 1
+            self.table.insertRow(row)
+            self.table.setItem(row, col, QTableWidgetItem(metric_name))
+
+    def pmc_select(self, item):
+
+        metric_name = item.data()
+        if (
+            not metric_name in self.nodes_dict
+            or not self.nodes_dict[metric_name].isCheckable()
+        ):
+            return
+
+        # only proper metrics check/uncheck is processed here.
+        IP_block = item.data().split(sep="_")[0]
+
+        if self.nodes_dict[metric_name].checkState() == 0:
+            # unselect the metric in the table if it is currently selected
+            self.pmc_remove_metric(metric_name, IP_block)
+
+        elif self.nodes_dict[metric_name].checkState() == 2:
+            self.pmc_add_metric(metric_name, IP_block)
+
+    # Function to save populate treeview with a dictionary
+    def importData(self, xmlparsed, root=None):
+
+        self.model.setRowCount(0)
+        if root is None:
+            root = self.model.invisibleRootItem()
+
+        for x in xmlparsed.getiterator():
+
+            # Add SoC node to Root
+            if x.tag in self.soc_arch_list:
+                parent = root
+                parent.appendRow([QStandardItem(x.tag)])
+                self.nodes_dict[x.tag] = parent.child(parent.rowCount() - 1)
+
+            # check all metrics in an SoC family
+            if x.tag == "metric" and x.getparent().tag in self.soc_arch_list:
+
+                # New IP block (e.g., SQ), detected, create a new hierarchy for the block
+                if not x.attrib["block"] in self.block_list:
+
+                    self.block_list.append(x.attrib["block"])
+                    parent = self.nodes_dict[x.getparent().tag]  # the SoC node
+                    parent.appendRow(
+                        [
+                            QStandardItem(x.attrib["block"]),
+                            QStandardItem(""),
+                            QStandardItem(""),
+                            QStandardItem(""),
+                        ]
+                    )
+
+                    # record the tree node for the block
+                    self.nodes_dict[x.attrib["block"]] = parent.child(
+                        parent.rowCount() - 1
+                    )
+
+                # Add metric node to the Block node
+                parent = self.nodes_dict[x.attrib["block"]]
+                metric_name = QStandardItem(x.attrib["name"])
+                metric_name.setCheckable(True)
+                parent.appendRow(
+                    [
+                        metric_name,
+                        QStandardItem(x.attrib["block"]),
+                        QStandardItem(x.attrib["event"]),
+                        QStandardItem(x.attrib["descr"]),
+                    ]
+                )
+
+                self.nodes_dict[x.attrib["name"]] = parent.child(parent.rowCount() - 1)
+
+
+if __name__ == "__main__":
+
+    app = QApplication(sys.argv)
+    app.setStyle("Fusion")
+
+    # populate the view with GFX metrics.xml
+    window = mainWindow()
+
+    # show the view
+    window.setGeometry(300, 100, 600, 300)
+    # view.setWindowTitle('GFX Perfmon Counters')
+    window.show()
+
+    # start the application
+    sys.exit(app.exec_())

src/utils/perfagg.py

@@ -0,0 +1,293 @@
+################################################################################
+# Copyright (c) 2021 - 2022 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.
+################################################################################
+
+import sys, os, pathlib, shutil, subprocess, argparse, glob, re
+import numpy as np
+import math
+
+prog = "omniperf"
+
+# Per IP block max number of simulutaneous counters
+# GFX IP Blocks
+perfmon_config = {
+    "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 perfmon_coalesce(pmc_files_list, workload_dir, soc):
+
+    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()
+            if "SQ_ACCUM_PREV_HIRES" in counters:
+                # save  all level counters separately
+
+                nindex = counters.index("SQ_ACCUM_PREV_HIRES")
+                level_counter = counters[nindex - 1]
+
+                # 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()
+
+                continue
+
+            # 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]
+
+    # 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, workload_dir, soc):
+
+    workload_perfmon_dir = workload_dir + "/perfmon"
+
+    # 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
+    fd = open(workload_perfmon_dir + "/pmc_perf.txt", "w")
+
+    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)
+        fd.write(line + "\n")
+
+    fd.write("\ngpu:\n")
+    fd.write("range:\n")
+    fd.write("kernel:\n")
+    fd.close()
+
+
+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, workload_dir, soc)
+    perfmon_emit(pmc_list, workload_dir, soc)
+
+
+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, workload_dir, soc)
+    perfmon_emit(pmc_list, workload_dir, soc)

src/utils/plot_roofline.py

@@ -0,0 +1,661 @@
+################################################################################
+# Copyright (c) 2021 - 2022 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.
+################################################################################
+
+from linecache import cache
+import subprocess
+from operator import sub
+import os
+import sys
+from pathlib import Path
+
+import numpy
+import matplotlib.pyplot as plt
+from matplotlib.pyplot import get, text
+from math import log, pi, sqrt
+import pandas as pd
+import pylab
+
+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"]
+
+################################################
+# 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
+    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 == "curr_ai_l1":
+        return "green"
+    elif catagory == "curr_ai_l2":
+        return "blue"
+    elif catagory == "curr_ai_hbm":
+        return "red"
+    else:
+        raise RuntimeError("Invalid catagory passed to get_color()")
+
+
+# -------------------------------------------------------------------------------------
+#                           Plot BW at each cache level
+# -------------------------------------------------------------------------------------
+def plot_roof(inputs, roof_data):
+    cacheHierarchy = []
+    if inputs["mem"] == "ALL":
+        cacheHierarchy += ["HBM", "L2", "L1", "LDS"]
+    else:
+        cacheHierarchy.append(inputs["mem"])
+    targ_dtype = (
+        "FP32"
+        if float(roof_data["FP32Flops"][0]) > float(roof_data["FP64Flops"][0])
+        else "FP64"
+    )
+    print("Dtype: ", targ_dtype)
+    print(inputs["mem"])
+    x1 = y1 = x2 = y2 = -1
+    x1_mfma = y1_mfma = x2_mfma = y2_mfma = -1
+    target_precision = targ_dtype[2:]
+
+    peakOps = float(roof_data[targ_dtype + "Flops"][0])
+    for i in range(0, len(cacheHierarchy)):
+        # Plot BW line
+        # print("Current cache level: {}".format(cacheHierarchy[i]))
+        curr_bw = cacheHierarchy[i] + "Bw"
+        peakBw = float(roof_data[curr_bw][0])
+
+        peakMFMA = float(roof_data["MFMAF{}Flops".format(target_precision)][0])
+
+        x1 = float(XMIN)
+        y1 = float(XMIN) * peakBw
+
+        x2 = peakOps / peakBw
+        y2 = peakOps
+
+        plt.plot([x1, x2], [y1, y2], color="magenta")
+        # print("Mem Points: [{}, {}], [{}, {}]".format(x1, x2, y1, y2))
+
+        # Plot MFMA lines (NOTE: Assuming MI200 soc)
+        x1_mfma = peakOps / peakBw
+        y1_mfma = peakOps
+
+        x2_mfma = peakMFMA / peakBw
+        y2_mfma = peakMFMA
+
+        plt.plot([x1_mfma, x2_mfma], [y1_mfma, y2_mfma], color="blue")
+        # print("Extend BW Points: [{}, {}], [{}, {}]".format(x1_mfma, x2_mfma, y1_mfma, y2_mfma))
+
+        # These are the points to use:
+        # print("x = [{}, {}]".format(x1,x2_mfma))
+        # print("y = [{}, {}]".format(y1, y2_mfma))
+
+        # Plot BW label
+        x1log = log(x1) / log(10)
+        x2log = log(x2) / log(10)
+        y1log = log(y1) / log(10)
+        y2log = log(y2) / log(10)
+        x_text = 10 ** ((x1log + x2log) / 2)
+        y_text = 10 ** ((y1log + y2log) / 2)
+
+        fig = plt.gcf()
+        size = fig.get_size_inches() * fig.dpi
+        fig_x, fig_y = size
+
+        # dx = log(x2) - log(x1)
+        # dy = log(y2) - log(y1)
+        # x_min, x_max = plt.xlim()
+        # y_min, y_max = plt.ylim()
+        # Dx = dx * fig_x / (log(x_max) - log(x_min))
+        # Dy = dy * fig_y / (log(y_max) - log(y_min))
+        # #fdiv = 0.7 #TODO: improve accuracy of text angle (tilt)
+        # angle = (180.0 / pi) * numpy.arctan(Dy / Dx )#/fdiv)
+
+        dx = abs(log(x2) - log(x1))
+        dy = abs(log(y2) - log(y1))
+        angle = (180.0 / pi) * numpy.arctan(dy / dx)
+        # If user isn't zooming in, print bw labels normally
+        if not inputs["axes"]:
+            text(
+                x_text,
+                y_text,
+                "{} vL1D GB/s".format(int(peakBw))
+                if cacheHierarchy[i].upper() == "L1"
+                else "{} {} GB/s".format(int(peakBw), cacheHierarchy[i].upper()),
+                rotation=angle,
+                rotation_mode="anchor",
+                **get_font(),
+            )
+        else:
+            # if bw line isn't being cut out then plot bw
+            print("if {} < {}".format(inputs["axes"][0], 10**x2log))
+            if inputs["axes"][0] < 10**x2log:
+                text(
+                    10**x2log,
+                    10**y2log,
+                    "{} {} GB/s".format(int(peakBw), cacheHierarchy[i].upper()),
+                    rotation=angle,
+                    rotation_mode="anchor",
+                    **get_font(),
+                )
+
+    # -------------------------------------------------------------------------------------
+    #                                     Plot computing roof
+    # -------------------------------------------------------------------------------------
+    # Plot FMA roof
+    x0 = XMAX
+    if x2 < x0:
+        x0 = x2
+
+    temp_label = "{} VALU GFLOP/sec".format(int(peakOps))
+    plt.plot([x0, XMAX], [peakOps, peakOps], color="magenta")
+    # print("FMA Points: [{}, {}], [{},{}]".format(x0, XMAX, peakOps, peakOps))
+    text(
+        XMAX if not inputs["axes"] else inputs["axes"][1],
+        peakOps - 4000,  # should i keep this fixed at 4000?
+        temp_label,
+        horizontalalignment="right",
+        **get_font(),
+    )
+
+    # Plot MFMA roof
+    if x1_mfma != -1:  # assert that mfma has been assigned
+        x0_mfma = XMAX
+        if x2_mfma < x0_mfma:
+            x0_mfma = x2_mfma
+
+        peakMFMA = float(roof_data["MFMAF{}Flops".format(target_precision)][0])
+        temp_label = "{} MFMA GFLOP/sec".format(int(peakMFMA))
+        plt.plot([x0_mfma, XMAX], [peakMFMA, peakMFMA], color="blue")
+        # print("MFMA Points: [{}, {}], [{},{}]".format(x0_mfma, XMAX, peakMFMA, peakMFMA))
+        text(
+            XMAX if not inputs["axes"] else inputs["axes"][1],
+            peakMFMA + 1000,
+            temp_label,
+            horizontalalignment="right",
+            **get_font(),
+        )
+
+    return targ_dtype
+
+
+# -------------------------------------------------------------------------------------
+#                              Overlay application performance
+# -------------------------------------------------------------------------------------
+# Calculate relevent metrics for ai calculation
+def plot_application(inputs, verbose):
+
+    df = pd.read_csv(inputs["path"] + "/pmc_perf.csv")
+    # 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 = []
+    for index, row in df.iterrows():
+        # CASE: Top kernels
+        if inputs["sort"] == "kernels" and (
+            (row["KernelName"] != kernelName and kernelName != "")
+            or index == df.shape[0] - 1
+        ):
+            if df.shape[0] - 1 == index:
+                calls += 1
+            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,
+                    lds_data / calls,
+                    L1cache_data / calls,
+                    L2cache_data / calls,
+                    hbm_data / calls,
+                    totalDuration,
+                    avgDuration / calls,
+                )
+            )
+            if verbose:
+                print(
+                    "Just added {} to AI_Data at index {}. # of calls: {}".format(
+                        kernelName, index, 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
+
+        kernelName = row["KernelName"]
+        try:
+            total_flops += (
+                (
+                    64
+                    * (
+                        row["SQ_INSTS_VALU_ADD_F16"]
+                        + row["SQ_INSTS_VALU_MUL_F16"]
+                        + (2 * row["SQ_INSTS_VALU_FMA_F16"])
+                        + row["SQ_INSTS_VALU_TRANS_F16"]
+                    )
+                )
+                + (
+                    64
+                    * (
+                        row["SQ_INSTS_VALU_ADD_F32"]
+                        + row["SQ_INSTS_VALU_MUL_F32"]
+                        + (2 * row["SQ_INSTS_VALU_FMA_F32"])
+                        + row["SQ_INSTS_VALU_TRANS_F32"]
+                    )
+                )
+                + (
+                    64
+                    * (
+                        row["SQ_INSTS_VALU_ADD_F64"]
+                        + row["SQ_INSTS_VALU_MUL_F64"]
+                        + (2 * row["SQ_INSTS_VALU_FMA_F64"])
+                        + row["SQ_INSTS_VALU_TRANS_F64"]
+                    )
+                )
+                + (row["SQ_INSTS_VALU_MFMA_MOPS_F16"] * 512)
+                + (row["SQ_INSTS_VALU_MFMA_MOPS_BF16"] * 512)
+                + (row["SQ_INSTS_VALU_MFMA_MOPS_F32"] * 512)
+                + (row["SQ_INSTS_VALU_MFMA_MOPS_F64"] * 512)
+            )
+        except KeyError:
+            if verbose:
+                print("Skipped total_flops at index {}".format(index))
+            pass
+        try:
+            valu_flops += (
+                64
+                * (
+                    row["SQ_INSTS_VALU_ADD_F16"]
+                    + row["SQ_INSTS_VALU_MUL_F16"]
+                    + (2 * row["SQ_INSTS_VALU_FMA_F16"])
+                    + row["SQ_INSTS_VALU_TRANS_F16"]
+                )
+                + 64
+                * (
+                    row["SQ_INSTS_VALU_ADD_F32"]
+                    + row["SQ_INSTS_VALU_MUL_F32"]
+                    + (2 * row["SQ_INSTS_VALU_FMA_F32"])
+                    + row["SQ_INSTS_VALU_TRANS_F32"]
+                )
+                + 64
+                * (
+                    row["SQ_INSTS_VALU_ADD_F64"]
+                    + row["SQ_INSTS_VALU_MUL_F64"]
+                    + (2 * row["SQ_INSTS_VALU_FMA_F64"])
+                    + row["SQ_INSTS_VALU_TRANS_F64"]
+                )
+            )
+        except KeyError:
+            if verbose:
+                print("Skipped valu_flops at index {}".format(index))
+            pass
+
+        try:
+            mfma_flops_f16 += row["SQ_INSTS_VALU_MFMA_MOPS_F16"] * 512
+            mfma_flops_bf16 += row["SQ_INSTS_VALU_MFMA_MOPS_BF16"] * 512
+            mfma_flops_f32 += row["SQ_INSTS_VALU_MFMA_MOPS_F32"] * 512
+            mfma_flops_f64 += row["SQ_INSTS_VALU_MFMA_MOPS_F64"] * 512
+            mfma_iops_i8 += row["SQ_INSTS_VALU_MFMA_MOPS_I8"] * 512
+        except KeyError:
+            if verbose:
+                print("Skipped mfma ops at index {}".format(index))
+            pass
+
+        try:
+            lds_data += (
+                (row["SQ_LDS_IDX_ACTIVE"] - row["SQ_LDS_BANK_CONFLICT"]) * 4 * L2_BANKS
+            )  # L2_BANKS = 32 (since assuming mi200)
+        except KeyError:
+            if verbose:
+                print("Skipped lds_data at index {}".format(index))
+            pass
+
+        try:
+            L1cache_data += row["TCP_TOTAL_CACHE_ACCESSES_sum"] * 64
+        except KeyError:
+            if verbose:
+                print("Skipped L1cache_data at index {}".format(index))
+            pass
+
+        try:
+            L2cache_data += (
+                row["TCP_TCC_WRITE_REQ_sum"] * 64
+                + row["TCP_TCC_ATOMIC_WITH_RET_REQ_sum"] * 64
+                + row["TCP_TCC_ATOMIC_WITHOUT_RET_REQ_sum"] * 64
+                + row["TCP_TCC_READ_REQ_sum"] * 64
+            )
+        except KeyError:
+            if verbose:
+                print("Skipped L2cache_data at index {}".format(index))
+            pass
+        try:
+            hbm_data += (
+                (row["TCC_EA_RDREQ_32B_sum"] * 32)
+                + ((row["TCC_EA_RDREQ_sum"] - row["TCC_EA_RDREQ_32B_sum"]) * 64)
+                + (row["TCC_EA_WRREQ_64B_sum"] * 64)
+                + ((row["TCC_EA_WRREQ_sum"] - row["TCC_EA_WRREQ_64B_sum"]) * 32)
+            )
+        except KeyError:
+            if verbose:
+                print("Skipped hbm_data at index {}".format(index))
+            pass
+
+        totalDuration += row["EndNs"] - row["BeginNs"]
+
+        avgDuration += row["EndNs"] - row["BeginNs"]
+
+        calls += 1
+        if inputs["sort"] == "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 10 intensities ('{}')...".format(inputs["sort"]))
+    intensities = {"curr_ai_l1": [], "curr_ai_l2": [], "curr_ai_hbm": []}
+    curr_perf = []
+    i = 0
+    # Create list of top 5 intensities
+    while i <= 9 and i != len(myList):
+        intensities["curr_ai_l1"].append(
+            myList[i].total_flops / myList[i].L1cache_data
+        ) if myList[i].L1cache_data else intensities["curr_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["curr_ai_l2"].append(
+            myList[i].total_flops / myList[i].L2cache_data
+        ) if myList[i].L2cache_data else intensities["curr_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["curr_ai_hbm"].append(
+            myList[i].total_flops / myList[i].hbm_data
+        ) if myList[i].hbm_data else intensities["curr_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
+
+    print(intensities)
+
+    # fig, ax = plt.subplots()
+
+    plotted_spots = []
+    labels = []
+    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])
+        myScatter = plt.scatter(x, y, c=color, marker="o")
+        plotted_spots.append(myScatter)
+        label = i
+        labels.append(label)
+
+    try:
+        pylab.legend(
+            plotted_spots,
+            labels,
+            prop={"size": (FONT_SIZE - 2)},
+            bbox_to_anchor=(1.04, 1),
+            loc="upper left",
+            title="Top {}".format(inputs["sort"]),
+            title_fontsize=FONT_SIZE,
+        )
+    except Exception as e:
+        sys.stderr.write(f"{e}\n")
+        pylab.legend(
+            plotted_spots,
+            labels,
+            prop={"size": (FONT_SIZE - 2)},
+        )
+
+
+def empirical_roof(args):
+    soc = args.target
+    inputs = {
+        "path": str,
+        "cmd": str,
+        "sort": str,
+        "mem": str,
+        "axes": list,
+        "device": int,
+        # "workgroups": int,
+        # "wsize": int,
+        # "dataset": int,
+        # "experiments": int,
+        # "iter": int
+    }
+
+    inputs["sort"] = args.sort.lower()
+    inputs["mem"] = args.mem_level.upper()
+
+    if inputs["sort"] != "kernels" and inputs["sort"] != "dispatches":
+        sys.exit("Invalid sort. Must be either 'kernels' or 'dispatches'")
+    if (
+        inputs["mem"] != "HBM"
+        and inputs["mem"] != "VL1D"
+        and inputs["mem"] != "L2"
+        and inputs["mem"] != "LDS"
+        and inputs["mem"] != "ALL"
+    ):
+        sys.exit(
+            "Invalid mem-level. Must be one of these option 'LDS', 'L2', 'vL1D', or 'HBM'"
+        )
+    if inputs["mem"] == "VL1D":
+        inputs["mem"] = "L1"
+
+    inputs["device"] = int(args.device)
+    # inputs["workgroups"] = int(args.workgroups)
+    # inputs["wsize"] = int(args.wsize)
+    # inputs["dataset"] = int(args.dataset)
+    # inputs["experiments"] = int(args.experiments)
+    # inputs["iter"] = int(args.iter)
+    inputs["path"] = args.path
+    inputs["cmd"] = args.remaining
+    inputs["axes"] = args.axes
+
+    # device_list = [int(item) for item in args.device.split(',')]
+
+    if soc not in SUPPORTED_SOC:
+        sys.exit("SoC not yet supported for Roofline Analysis")
+
+    # Basic Info
+    print("Path: ", inputs["path"])
+    print("Target: ", soc)
+    print("Memory Level: ", inputs["mem"])
+
+    roofPath = inputs["path"] + "/roofline.csv"
+    # -----------------------------------------------------
+    # Initialize roofline data dictionary from roofline.csv
+    # -----------------------------------------------------
+    roof_data = (
+        {}
+    )  # TODO: consider changing this to an ordered dict for consistency over py versions
+    headers = []
+    with open(roofPath, "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:
+                    roof_data[i] = []
+            else:
+                for i, key in enumerate(headers):
+                    roof_data[key].append(row[i])
+
+            rowCount += 1
+    csvfile.close()
+
+    # Initalize plot
+    f = plt.figure(figsize=(1600 / 100, 1200 / 100), dpi=100)
+    f.add_subplot(111)
+
+    _title_font = get_font()
+    _title_font["size"] += 8
+
+    plt.title("Empirical Roofline", **_title_font)
+    plt.xlabel("Arithmetic Intensity (FLOP/Byte)", **get_font())
+    plt.ylabel("Performance (GFLOP/sec)", **get_font())
+    plt.grid(True, which="major", ls="--", lw=1)
+    plt.grid(True, which="minor", ls="--", lw=0.5)
+    plt.yscale("log")
+    plt.xscale("log")
+    # Adjust axes if instructed
+    if inputs["axes"]:
+        plt.xlim(inputs["axes"][0], inputs["axes"][1])
+        plt.ylim(inputs["axes"][2], inputs["axes"][3])
+
+    # ------------------
+    #  Generate Roofline
+    # ------------------
+    dtype = plot_roof(inputs, roof_data)  # Also returns chosen dtype
+    plot_application(inputs, args.verbose)
+
+    filename = IMGNAME + "_gpu-" + str(inputs["device"]) + "_{}".format(dtype) + ".pdf"
+
+    full_path = os.path.abspath(inputs["path"])
+    path_to_output = full_path + "/" + filename
+
+    print('Saving plot: "{}"...'.format(filename))
+    plt.savefig(path_to_output, bbox_inches="tight", format="pdf")
+    print('File saved to: "{}"'.format(path_to_output))
+    plt.close()

src/utils/remove_workload.py

@@ -0,0 +1,92 @@
+#!/usr/bin/env python3
+
+################################################################################
+# Copyright (c) 2021 - 2022 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.
+################################################################################
+
+import argparse
+import getpass
+from pymongo import MongoClient
+
+# Verify target directory and setup connection
+def remove_workload(args):
+    # parser = argparse.ArgumentParser(description='Remove a workload from an Omniperf Instance')
+
+    # parser.add_argument('-H', '--host',required=True, help="Name or IP address of the server host")
+    # parser.add_argument('-P', '--port', required=False, help="TCP/IP Port (defaults to 27018)", default=27018)
+    # parser.add_argument('-u', '--username', required=True, help="Name of the user to connect with")
+    # parser.add_argument('-p', '--password', required=False, help="The user's password (will be requested later if it's not set)", default='')
+    # parser.add_argument('-w', '--workload', required=True, help="Specify the full workload name to delete")
+
+    host = args.host
+    port = str(args.port)
+    username = args.username
+
+    if args.password == "":
+        try:
+            password = getpass.getpass()
+        except Exception as error:
+            print("PASSWORD ERROR", error)
+        else:
+            print("Password recieved")
+    else:
+        password = args.password
+    workload = args.workload
+
+    # Verify workload  is valid
+    if workload == "admin" or workload == "local":
+        print("You cannot remove this database. Try again.")
+    else:
+        print("Attempting to remove " + workload)
+
+    connection_info = {
+        "username": username,
+        "password": password,
+        "host": host,
+        "port": port,
+        "workload": workload,
+    }
+
+    client = MongoClient(
+        "mongodb://"
+        + connection_info["username"]
+        + ":"
+        + connection_info["password"]
+        + "@"
+        + connection_info["host"]
+        + ":"
+        + connection_info["port"]
+        + "/?authSource=admin"
+    )
+
+    dbToRemove = client[workload]
+
+    # check the collection names on the database
+    col_list = dbToRemove.list_collection_names()
+    # if not col_list:
+    #     print("This workload appears to be empty... Try again.")
+    # else:
+    client.drop_database(dbToRemove)
+    db = client["workload_names"]
+    col = db["names"]
+    col.delete_many({"name": workload})
+
+    print("Done!")

src/utils/resources.py

@@ -0,0 +1,33 @@
+################################################################################
+# Copyright (c) 2021 - 2022 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.
+################################################################################
+
+import pathlib
+import sys
+
+
+def get_resource_path(relative_path):
+    rel_path = pathlib.Path(relative_path)
+    dev_base_path = (
+        pathlib.Path(__file__).resolve().parent.parent.parent.parent
+    )  # omniperf dir (whereas equiv comd in omniperf.py returns src dir)
+    base_path = getattr(sys, "_MEIPASS", dev_base_path)
+    return dev_base_path

src/utils/specs.py

@@ -0,0 +1,241 @@
+"""Get host/gpu specs."""
+
+################################################################################
+# Copyright (c) 2021 - 2022 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.
+################################################################################
+
+import os
+import re
+import socket
+import subprocess
+
+from dataclasses import dataclass
+from pathlib import Path as path
+from textwrap import dedent
+
+gpu_list = {"gfx906", "gfx908", "gfx90a"}
+
+
+@dataclass
+class MachineSpecs:
+    hostname: str
+    cpu: str
+    kernel: str
+    ram: str
+    distro: str
+    rocmversion: str
+    GPU: str
+    L1: str
+    L2: str
+    SCLK: str
+    CU: str
+    SIMD: str
+    SE: str
+    wave_size: str
+    workgroup_size: str
+    cur_SCLK: str
+    cur_MCLK: str
+    wave_occu: str
+
+    def __str__(self):
+        return dedent(
+            f"""\
+        Host info:
+            hostname:       {self.hostname}
+            cpu info:       {self.cpu}
+            ram:            {self.ram}
+            distro:         {self.distro}
+            kernel version: {self.kernel}
+            rocm version:   {self.rocmversion}
+        Device info:
+            GPU:                {self.GPU}
+            L1:                 {self.L1}
+            L2:                 {self.L2}
+            Max SCLK:           {self.SCLK}MHz
+            Current SCLK:       {self.cur_SCLK}MHz
+            Current MCLK:       {self.cur_MCLK}MHz
+            CU:                 {self.CU}
+            SIMD:               {self.SIMD}
+            SE:                 {self.SE}
+            Wave Size:          {self.wave_size}
+            Workgroup Max Size: {self.workgroup_size}
+            Max Wave Occupancy Per CU: {self.wave_occu}
+        """
+        )
+
+
+def gpuinfo():
+
+    rocminfo = run(["rocminfo"]).split("\n")
+
+    for idx1, linetext in enumerate(rocminfo):
+        gpu_id = search(r"^\s*Name\s*:\s+ ([a-zA-Z0-9]+)\s*$", linetext)
+        if gpu_id in gpu_list:
+            break
+
+    if not gpu_id in gpu_list:
+        return None, None, None, None, None, None, None, None, None, None
+
+    L1, L2 = "", ""
+    for idx2, linetext in enumerate(rocminfo[idx1 + 1 :]):
+
+        key = search(r"^\s*L1:\s+ ([a-zA-Z0-9]+)\s*", linetext)
+        if key != None:
+            L1 = key
+            continue
+
+        key = search(r"^\s*L2:\s+ ([a-zA-Z0-9]+)\s*", linetext)
+        if key != None:
+            L2 = key
+            continue
+
+        key = search(r"^\s*Max Clock Freq\. \(MHz\):\s+([0-9]+)", linetext)
+        if key != None:
+            sclk = key
+            continue
+
+        key = search(r"^\s*Compute Unit:\s+ ([a-zA-Z0-9]+)\s*", linetext)
+        if key != None:
+            num_CU = key
+            continue
+
+        key = search(r"^\s*SIMDs per CU:\s+ ([a-zA-Z0-9]+)\s*", linetext)
+        if key != None:
+            num_SIMD = key
+            continue
+
+        key = search(r"^\s*Shader Engines:\s+ ([a-zA-Z0-9]+)\s*", linetext)
+        if key != None:
+            num_SE = key
+            continue
+
+        key = search(r"^\s*Wavefront Size:\s+ ([a-zA-Z0-9]+)\s*", linetext)
+        if key != None:
+            wave_size = key
+            continue
+
+        key = search(r"^\s*Workgroup Max Size:\s+ ([a-zA-Z0-9]+)\s*", linetext)
+        if key != None:
+            grp_size = key
+            continue
+
+        key = search(r"^\s*Max Waves Per CU:\s+ ([a-zA-Z0-9]+)\s*", linetext)
+        if key != None:
+            wave_occu = key
+            break
+
+    return gpu_id, L1, L2, sclk, num_CU, num_SIMD, num_SE, wave_size, grp_size, wave_occu
+
+
+def run(cmd):
+    p = subprocess.run(cmd, stdout=subprocess.PIPE, stderr=subprocess.PIPE)
+    return p.stdout.decode("ascii")
+
+
+def search(pattern, string):
+    m = re.search(pattern, string, re.MULTILINE)
+    if m is not None:
+        return m.group(1)
+    return None
+
+
+def get_machine_specs(devicenum):
+
+    cpuinfo = path("/proc/cpuinfo").read_text()
+    meminfo = path("/proc/meminfo").read_text()
+    version = path("/proc/version").read_text()
+    os_release = path("/etc/os-release").read_text()
+
+    version_loc = [
+        "version",
+        "version-dev",
+        "version-hip-libraries",
+        "version-hiprt",
+        "version-hiprt-devel",
+        "version-hip-sdk",
+        "version-libs",
+        "version-utils",
+    ]
+    for itr in version_loc:
+        _path = os.path.join(os.getenv("ROCM_PATH", "/opt/rocm"), ".info", itr)
+        if os.path.exists(_path):
+            rocm_ver = path(_path).read_text()
+            break
+
+    (
+        gpu_id,
+        L1,
+        L2,
+        sclk,
+        num_CU,
+        num_SIMD,
+        num_SE,
+        wave_size,
+        grp_size,
+        wave_occu,
+    ) = gpuinfo()
+    rocm_smi = run(["rocm-smi"])
+
+    device = rf"^\s*{devicenum}(.*)"
+
+    hostname = socket.gethostname()
+    cpu = search(r"^model name\s*: (.*?)$", cpuinfo)
+    kernel = search(r"version (\S*)", version)
+    ram = search(r"MemTotal:\s*(\S*)", meminfo)
+    distro = search(r'PRETTY_NAME="(.*?)"', os_release)
+    if distro is None:
+        distro = ""
+
+    rocmversion = rocm_ver.strip()
+
+    freq = search(device, rocm_smi).split()
+    cur_sclk = search(r"([0-9]+)", freq[2])
+    if cur_sclk is None:
+        cur_sclk = ""
+
+    cur_mclk = search(r"([0-9]+)", freq[3])
+    if cur_mclk is None:
+        cur_mclk = ""
+
+    return MachineSpecs(
+        hostname,
+        cpu,
+        kernel,
+        ram,
+        distro,
+        rocmversion,
+        gpu_id,
+        L1,
+        L2,
+        sclk,
+        num_CU,
+        num_SIMD,
+        num_SE,
+        wave_size,
+        grp_size,
+        cur_sclk,
+        cur_mclk,
+        wave_occu,
+    )
+
+
+if __name__ == "__main__":
+    print(get_machine_specs(0))