ae8f72fa79
* Use native tool for counter collection
* Add native counter collection tool which uses rocprofiler-sdk C++
library public API to get counter collection data
* This is enabled by default, unless --no-native-tool option is
provided or ROCPROF=rocprofv3 env. var. is provided
* This tool is only supported for ROCm version >=7.x.x
* This tool is not supported for attach/detach scenario
* Build native tool shared object during build time
* If using rocprof-compute without building then runtime compilation of
t push native tool shared object is performed
* rocprofiler-sdk tools is still used for services other than counter
collection and data collected by native tool is merged into the
rocpd/csv output of rocprofiler-sdk tool
* Make `rocpd` choice the default choice for `--format-rocprof-output`
option
* If `rocpd` public API from rocprofiler-sdk library is not present,
then fallback to `csv` choice
* In this case only `pmc_perf.csv` is written in workload folder
instead of multiple `csv` files for each profiling run
* Remove `json` choice from `--format-rocprof-output` option since it
functions identical to `csv` option
* Rename option `--rocprofiler-sdk-library-path` to
`--rocprofiler-sdk-tool-path` since we LD_PRELOAD the
rocprofiler-sdk tool shared object and not the rocprofiler-sdk library
shared object
* Fix the meaning of `--dispatch` option in `profile` mode to mention
dispatch iteration filtering instead of dispatch id filtering
* --dispatch option in analyze mode does dispatch id filtering
* Move standalone binary creation logic from cmake file to docker file
* fix native counter collection tool during attach/detach
* improve logging
* fix attach detach with native tool
* fix attach detach with native tool
* do not support attach/detach in native tool
* Update changelog
* add standalone binary creation functionality in cmake
* address review comments
* address review comments
* fix formatting
* address review comments
* Adding paths for cmake to search. Also updated min. cmake requirement to 3.21 as this was when hip was supported.
Signed-off-by: Carrie Fallows <Carrie.Fallows@amd.com>
* Update hip compiler ID check, sometimes comes up as Clang, sometimes ROCMClang- depends on setup.
Updated formatting.
Signed-off-by: Carrie Fallows <Carrie.Fallows@amd.com>
* RHEL8.10 unable to compile due to defaulting to old c++ version, need to force c++17
Signed-off-by: Carrie Fallows <Carrie.Fallows@amd.com>
* Updating changelog per docs team recommendations
Signed-off-by: Carrie Fallows <Carrie.Fallows@amd.com>
* Apply suggestions from code review to changelog
Co-authored-by: Pratik Basyal <pratik.basyal@amd.com>
* Do not required HIP complier to build native counter collection tool
* fix cmake
* gersemi formatting on latest cmake change
Signed-off-by: Carrie Fallows <Carrie.Fallows@amd.com>
* ex ci updated dependencies to include rocprofiler-sdk, but cmake was still not capturing the path- there was a commit that added to the cmake_prefix_path entry that specified rocprof-sdk's cmake location ut was too specific for the search paths in find_package's config mode.
removing the cmake_prefix_path var and adding hints to find_package call instead, and specifying config mode so it knows how to construct the search paths
Signed-off-by: Carrie Fallows <Carrie.Fallows@amd.com>
* gersemi run for formatting
Signed-off-by: Carrie Fallows <Carrie.Fallows@amd.com>
* Still need prefix path, should not have been removed in last commit but does need to be shortened to just the rocm path to allow for find_package config mode to do the job
Signed-off-by: Carrie Fallows <Carrie.Fallows@amd.com>
* include cstdint for uint32_t
* Run formatting on helper.cpp
Signed-off-by: Carrie Fallows <Carrie.Fallows@amd.com>
* Remove rocm 7.2 release stuff from version and changelog and handle it in separate pr
* fix version
* fix changelog
* fix changelog
* run ruff formatter
Signed-off-by: Carrie Fallows <Carrie.Fallows@amd.com>
* fix rocprofiler-sdk attach so path
---------
Signed-off-by: Carrie Fallows <Carrie.Fallows@amd.com>
Co-authored-by: Carrie Fallows <Carrie.Fallows@amd.com>
Co-authored-by: Pratik Basyal <pratik.basyal@amd.com>
181 lines
6.7 KiB
Python
181 lines
6.7 KiB
Python
##############################################################################
|
|
# MIT License
|
|
#
|
|
# Copyright (c) 2025 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 csv
|
|
import sqlite3
|
|
from contextlib import closing
|
|
from typing import Any
|
|
|
|
import pandas as pd
|
|
|
|
from utils.logger import console_error
|
|
|
|
# From schema definition in source/share/rocprofiler-sdk-rocpd/data_views.sql
|
|
# in rocprofiler-sdk repository
|
|
COUNTERS_COLLECTION_QUERY = """
|
|
SELECT
|
|
agent_id as GPU_ID,
|
|
dispatch_id as Dispatch_ID,
|
|
grid_size as Grid_Size,
|
|
workgroup_size as Workgroup_Size,
|
|
lds_block_size as LDS_Per_Workgroup,
|
|
scratch_size as Scratch_Per_Workitem,
|
|
vgpr_count as Arch_VGPR,
|
|
accum_vgpr_count as Accum_VGPR,
|
|
sgpr_count as SGPR,
|
|
kernel_name as Kernel_Name,
|
|
start as Start_Timestamp,
|
|
end as End_Timestamp,
|
|
kernel_id as Kernel_ID,
|
|
counter_name as Counter_Name,
|
|
value as Counter_Value
|
|
FROM counters_collection
|
|
"""
|
|
ROCPD_PMC_EVENT_TABLE_NAME_PREFIX = "rocpd_pmc_event_"
|
|
TABLE_NAME_PREFIX_QUERY = (
|
|
"SELECT name FROM sqlite_master WHERE type='table' "
|
|
"AND name LIKE '{table_name_prefix}%'"
|
|
)
|
|
INSERT_QUERY = "INSERT INTO {table_name} ({columns}) VALUES ({placeholders})"
|
|
|
|
|
|
def convert_db_to_csv(
|
|
db_path: str,
|
|
csv_file_path: str,
|
|
) -> None:
|
|
"""
|
|
Read rocpd database and write to CSV file
|
|
"""
|
|
# Read counters_collection view from the database and write to CSV
|
|
try:
|
|
with closing(sqlite3.connect(db_path)) as conn:
|
|
with closing(conn.execute(COUNTERS_COLLECTION_QUERY)) as cursor:
|
|
with open(csv_file_path, "w", newline="") as csvfile:
|
|
writer = csv.writer(csvfile)
|
|
writer.writerow([
|
|
description[0] for description in cursor.description
|
|
])
|
|
for row in cursor:
|
|
writer.writerow(row)
|
|
except OSError as e:
|
|
console_error(f"Database error while converting to CSV: {e}")
|
|
except Exception as e:
|
|
console_error(f"Unexpected error converting database to CSV: {e}")
|
|
|
|
|
|
def process_rocpd_csv(df: pd.DataFrame) -> pd.DataFrame:
|
|
"""
|
|
Merge counters across unique dispatches from the
|
|
input dataframe and return processed dataframe.
|
|
"""
|
|
if df.empty:
|
|
return df
|
|
|
|
data: list[dict[str, Any]] = []
|
|
|
|
# Group by unique kernel and merge into a single row
|
|
for _, group_df in df.groupby([
|
|
"Dispatch_ID",
|
|
"Kernel_Name",
|
|
"Grid_Size",
|
|
"Workgroup_Size",
|
|
"LDS_Per_Workgroup",
|
|
]):
|
|
row = {
|
|
"GPU_ID": group_df["GPU_ID"].iloc[0],
|
|
"Grid_Size": group_df["Grid_Size"].iloc[0],
|
|
"Workgroup_Size": group_df["Workgroup_Size"].iloc[0],
|
|
"LDS_Per_Workgroup": group_df["LDS_Per_Workgroup"].iloc[0],
|
|
"Scratch_Per_Workitem": group_df["Scratch_Per_Workitem"].iloc[0],
|
|
"Arch_VGPR": group_df["Arch_VGPR"].iloc[0],
|
|
"Accum_VGPR": group_df["Accum_VGPR"].iloc[0],
|
|
"SGPR": group_df["SGPR"].iloc[0],
|
|
"Kernel_Name": group_df["Kernel_Name"].iloc[0],
|
|
"Kernel_ID": group_df["Kernel_ID"].iloc[0],
|
|
"Start_Timestamp": group_df["Start_Timestamp"].iloc[0],
|
|
"End_Timestamp": group_df["End_Timestamp"].iloc[0],
|
|
}
|
|
# Each counter will become its own column
|
|
row.update(dict(zip(group_df["Counter_Name"], group_df["Counter_Value"])))
|
|
data.append(row)
|
|
df = pd.DataFrame(data)
|
|
# Rank GPU IDs, map lowest number to 0, next to 1, etc.
|
|
df["GPU_ID"] = df["GPU_ID"].rank(method="dense").astype(int) - 1
|
|
# Reset dispatch IDs
|
|
df["Dispatch_ID"] = range(len(df))
|
|
return df
|
|
|
|
|
|
def update_rocpd_pmc_events(counter_info: pd.DataFrame, rocpd_db_path: str) -> None:
|
|
"""Update pmc_event table in the given rocpd database path"""
|
|
try:
|
|
with closing(sqlite3.connect(rocpd_db_path)) as conn:
|
|
# Get pmc_event table name
|
|
with closing(
|
|
conn.execute(
|
|
TABLE_NAME_PREFIX_QUERY.format(
|
|
table_name_prefix=ROCPD_PMC_EVENT_TABLE_NAME_PREFIX
|
|
)
|
|
)
|
|
) as cursor:
|
|
table_name = cursor.fetchone()
|
|
if table_name is None:
|
|
console_error("No pmc_event table found in the rocpd database")
|
|
table_name = table_name[0]
|
|
|
|
# get pmc_event table data
|
|
guid = table_name[len(ROCPD_PMC_EVENT_TABLE_NAME_PREFIX) :].replace(
|
|
"_", "-"
|
|
)
|
|
columns = ("guid", "event_id", "pmc_id", "value")
|
|
values = list(
|
|
zip(
|
|
# guid
|
|
[guid] * len(counter_info),
|
|
# event_id
|
|
counter_info["dispatch_id"],
|
|
# pmc_id
|
|
counter_info["counter_id"],
|
|
# value
|
|
counter_info["counter_value"],
|
|
)
|
|
)
|
|
|
|
# insert into pmc_event table
|
|
with conn:
|
|
placeholders = ", ".join(["?"] * len(columns))
|
|
conn.executemany(
|
|
INSERT_QUERY.format(
|
|
table_name=table_name,
|
|
columns=", ".join(columns),
|
|
placeholders=placeholders,
|
|
),
|
|
values,
|
|
)
|
|
except OSError as e:
|
|
console_error(f"Database error while updating pmc_event table: {e}")
|
|
except Exception as e:
|
|
console_error(f"Unexpected error updating pmc_event table: {e}")
|