Skip to content

[rocprofiler-compute] Fix for multi process workload profiling #2418

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
jharryma merged 5 commits into from
Dec 23, 2025
Merged

[rocprofiler-compute] Fix for multi process workload profiling #2418

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
c1bbc46
Fix for multi process workload profiling
vedithal-amd Dec 19, 2025
3528f22
Add test case for multiprocess profiling using rocflop workload
vedithal-amd Dec 20, 2025
e626d56
Add rocflop
vedithal-amd Dec 22, 2025
22901c8
Fix native counter csv to rocprofv3 csv conversion
vedithal-amd Dec 22, 2025
6220049
python formatting using ruff 0.14 instead of 0.13
vedithal-amd Dec 22, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions projects/rocprofiler-compute/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -749,6 +749,7 @@ if(INSTALL_TESTS)
tests/hip_dynamic_shared
tests/laplace_eqn
tests/mat_mul_max
tests/rocflop
DESTINATION ${CMAKE_INSTALL_LIBEXECDIR}/${PROJECT_NAME}/tests
COMPONENT tests
)
Expand Down
54 changes: 25 additions & 29 deletions projects/rocprofiler-compute/src/lib/rocprofiler_compute_tool.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -77,11 +77,11 @@ for the agent and returns a pointer to it.
#include <iostream>
#include <memory>
#include <mutex>
#include <random>
#include <set>
#include <shared_mutex>
#include <sstream>
#include <string>
#include <unistd.h>
#include <unordered_map>
#include <vector>

Expand Down Expand Up @@ -148,7 +148,7 @@ struct counter_info_record_t {
// Tool data struct, now includes a vector of counter_info_record_t
struct tool_data_t {
std::mutex mut{};
std::unique_ptr<std::ostream> output_stream{nullptr};
std::string output_filename{};
std::unordered_map<uint64_t, std::string> counter_id_name_map{};
std::string requested_counters{};
std::string kernel_filter_include_regex{};
Expand Down Expand Up @@ -614,14 +614,28 @@ void generate_output(tool_data_t *tool_data) {
}),
tool_data->counter_records.end());
}

if (tool_data->counter_records.empty()) {
return;
}
// Write collected counter records and clean up
if (auto &os = tool_data->output_stream) {
if (!tool_data->output_filename.empty()) {
std::ofstream ofs(tool_data->output_filename);
if (!ofs.is_open()) {
std::cerr << "Failed to open output file: " << tool_data->output_filename
<< std::endl;
return;
}
// Write header at the beginning of the file
ofs << "dispatch_id,gpu_id,kernel_id,lds_per_workgroup,"
"counter_id,counter_name,counter_value\n";
for (const auto &r : tool_data->counter_records)
*os << r.dispatch_id << ',' << r.agent_id << "," << r.kernel_id << ','
ofs << r.dispatch_id << ',' << r.agent_id << "," << r.kernel_id << ','
<< r.LDS_memory_size << ',' << r.counter_id << ',' << r.counter_name
<< ',' << r.counter_value << '\n';
os->flush();
ofs.flush();
std::clog << "[rocprofiler-compute] [" << __FUNCTION__
<< "] Counter collection data has been written to: "
<< tool_data->output_filename << std::endl;
}
}

Expand All @@ -638,18 +652,13 @@ void tool_fini(void *user_data) {

} // namespace

std::unique_ptr<tool_data_t> create_tool_data(rocprofiler_client_id_t *id) {
std::unique_ptr<tool_data_t>
create_tool_data(rocprofiler_client_id_t * /*id*/) {
Comment thread
vedithal-amd marked this conversation as resolved.
Resolved
auto tool_data = std::make_unique<tool_data_t>();

// Generate a unique output filename using a random hex string (no libuuid
// dependency)
std::random_device rd;
std::mt19937 gen(rd());
std::uniform_int_distribution<uint32_t> dis(0, 0xFFFFFFFF);
std::stringstream filename_ss;
filename_ss << std::hex << dis(gen);
// Generate a unique output filename using the process ID
std::string base_filename =
"counter_collection_" + filename_ss.str().substr(0, 8) + ".csv";
std::to_string(getpid()) + "_native_counter_collection.csv";

// Require ROCPROF_OUTPUT_PATH to be set, otherwise error out
std::string filename;
Expand All @@ -664,20 +673,7 @@ std::unique_ptr<tool_data_t> create_tool_data(rocprofiler_client_id_t *id) {
// Use the generated base filename along with ROCPROF_OUTPUT_PATH
filename += base_filename;

// Set output stream to file
auto ofs = std::make_unique<std::ofstream>(filename);
if (!ofs->is_open()) {
throw std::runtime_error("Failed to open output file: " + filename);
}
tool_data->output_stream = std::move(ofs);
// Write header at the beginning of the file
*tool_data->output_stream << "dispatch_id,gpu_id,kernel_id,lds_per_workgroup,"
"counter_id,counter_name,counter_value\n";
tool_data->output_stream->flush();

// Write to clog the path of the logging file
std::clog << id->name << " [" << __FUNCTION__
<< "] Logging counter collection to: " << filename << std::endl;
tool_data->output_filename = filename;

// Store ROCPROF env. vars. in tool_data

Expand Down
31 changes: 18 additions & 13 deletions projects/rocprofiler-compute/src/utils/rocpd_data.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,6 +38,7 @@
SELECT
agent_id as GPU_ID,
dispatch_id as Dispatch_ID,
pid as PID,
grid_size as Grid_Size,
workgroup_size as Workgroup_Size,
lds_block_size as LDS_Per_Workgroup,
Expand All @@ -61,24 +62,28 @@
INSERT_QUERY = "INSERT INTO {table_name} ({columns}) VALUES ({placeholders})"


def convert_db_to_csv(
db_path: str,
def convert_dbs_to_csv(
db_paths: list[str],
csv_file_path: str,
) -> None:
"""
Read rocpd database and write to CSV file
Read rocpd databases and write to CSV file
"""
# Read counters_collection view from the database and write to CSV
# Read counters_collection view from the databases 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)
with open(csv_file_path, "w", newline="") as csvfile:
writer = csv.writer(csvfile)
header_written = False
for db_path in db_paths:
with closing(sqlite3.connect(db_path)) as conn:
with closing(conn.execute(COUNTERS_COLLECTION_QUERY)) as cursor:
if not header_written:
writer.writerow([
description[0] for description in cursor.description
])
header_written = True
for row in cursor:
writer.writerow(row)
Comment thread
vedithal-amd marked this conversation as resolved.
Resolved
except OSError as e:
console_error(f"Database error while converting to CSV: {e}")
except Exception as e:
Expand Down
204 changes: 115 additions & 89 deletions projects/rocprofiler-compute/src/utils/utils.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -885,24 +885,48 @@ def temporary_env(env_vars: dict[str, str]) -> Generator[None, None, None]:
rocprof_cmd == "rocprofiler-sdk"
and options["ROCPROF_COUNTER_COLLECTION"] == "0"
):
# Update rocpd database with counter csv created by native tool
rocpd_data.update_rocpd_pmc_events(
pd.read_csv(glob.glob(workload_dir + "/out/pmc_1/*.csv")[0]),
glob.glob(workload_dir + "/out/pmc_1/*/*.db")[0],
)
for db_name in glob.glob(workload_dir + "/out/pmc_1/*/*.db"):
pid = Path(db_name).stem.split("_")[0]
rocpd_data.update_rocpd_pmc_events(
pd.read_csv(
f"{workload_dir}/out/pmc_1/{pid}_native_counter_collection.csv"
),
Comment thread
vedithal-amd marked this conversation as resolved.
Resolved
db_name,
)
console_debug(f"Updated rocpd db {db_name} with native tool counters.")
# Write results_fbase.csv
rocpd_data.convert_db_to_csv(
glob.glob(workload_dir + "/out/pmc_1/*/*.db")[0],
rocpd_data.convert_dbs_to_csv(
glob.glob(workload_dir + "/out/pmc_1/*/*.db"),
workload_dir + f"/results_{fbase}.csv",
)
Comment thread
vedithal-amd marked this conversation as resolved.
Resolved
combined_df = pd.read_csv(workload_dir + f"/results_{fbase}.csv")
# Reset Dispatch_ID based on PID, Kernel_Name, Grid_Size,
# Workgroup_Size, LDS_Per_Workgroup
combined_df["Dispatch_ID"] = combined_df.groupby(
["PID", "Kernel_Name", "Grid_Size", "Workgroup_Size", "LDS_Per_Workgroup"],
sort=False,
).ngroup()
# Reset Kernel_ID based on Kernel_Name, Grid_Size,
# Workgroup_Size, LDS_Per_Workgroup
combined_df["Kernel_ID"] = combined_df.groupby(
["Kernel_Name", "Grid_Size", "Workgroup_Size", "LDS_Per_Workgroup"],
sort=False,
).ngroup()
Comment thread
vedithal-amd marked this conversation as resolved.
Resolved
# Drop PID since its not required
Comment thread
vedithal-amd marked this conversation as resolved.
Resolved
combined_df = combined_df.drop(columns=["PID"])
combined_df.to_csv(workload_dir + f"/results_{fbase}.csv", index=False)

if retain_rocpd_output:
shutil.copyfile(
glob.glob(workload_dir + "/out/pmc_1/*/*.db")[0],
workload_dir + "/" + fbase + ".db",
)
console_warning(
f"Retaining large raw rocpd database: {workload_dir}/{fbase}.db"
)
for db_path in glob.glob(workload_dir + "/out/pmc_1/*/*.db"):
pid = Path(db_path).stem.split("_")[0]
shutil.copyfile(
db_path,
workload_dir + f"/{fbase}_{pid}.db",
)
console_warning(
f"Retaining large raw rocpd database: "
f"{workload_dir}/{fbase}_{pid}.db"
)
# Remove temp directory
shutil.rmtree(workload_dir + "/" + "out")
return
Expand Down Expand Up @@ -1064,81 +1088,83 @@ def convert_native_counter_collection_csv(workload_dir: str) -> None:
trace to write counter collection csv in rocprofiler-sdk format
for further processing to pmc_perf.csv file
"""
counter_data = pd.read_csv(
glob.glob(f"{workload_dir}/out/pmc_1/*.csv")[0], index_col=False
)
# Group by on counter_data based on dispatch_id and
# counter_id and sum the counter_value
counter_data = counter_data.groupby(
["dispatch_id", "counter_name"], as_index=False
).agg({"counter_value": "sum"})
kernel_data_filename = glob.glob(f"{workload_dir}/out/pmc_1/*/*_kernel_trace.csv")[
0
]
kernel_data = pd.read_csv(kernel_data_filename)
rocprofv3_counter_data = pd.DataFrame({
"Correlation_Id": counter_data["dispatch_id"],
"Dispatch_Id": counter_data["dispatch_id"],
"Agent_Id": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"Agent_Id"
].values,
"Queue_Id": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"Queue_Id"
].values,
"Process_Id": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"Thread_Id"
].values,
"Thread_Id": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"Thread_Id"
].values,
"Grid_Size": (
kernel_data.iloc[counter_data["dispatch_id"] - 1][
["Grid_Size_X", "Grid_Size_Y", "Grid_Size_Z"]
]
.prod(axis=1)
.values
),
"Kernel_Id": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"Kernel_Id"
].values,
"Kernel_Name": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"Kernel_Name"
].values,
"Workgroup_Size": (
kernel_data.iloc[counter_data["dispatch_id"] - 1][
["Workgroup_Size_X", "Workgroup_Size_Y", "Workgroup_Size_Z"]
]
.prod(axis=1)
.values
),
"LDS_Block_Size": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"LDS_Block_Size"
].values,
"Scratch_Size": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"Scratch_Size"
].values,
"VGPR_Count": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"VGPR_Count"
].values,
"Accum_VGPR_Count": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"Accum_VGPR_Count"
].values,
"SGPR_Count": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"SGPR_Count"
].values,
"Counter_Name": counter_data["counter_name"],
"Counter_Value": counter_data["counter_value"],
"Start_Timestamp": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"Start_Timestamp"
].values,
"End_Timestamp": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"End_Timestamp"
].values,
})
rocprofv3_counter_data.to_csv(
kernel_data_filename.replace("kernel_trace", "counter_collection"),
index=False,
)
for native_filename in glob.glob(
f"{workload_dir}/out/pmc_1/*_native_counter_collection.csv"
):
pid = Path(native_filename).stem.split("_")[0]
counter_data = pd.read_csv(native_filename, index_col=False)
# Group by on counter_data based on dispatch_id and
# counter_id and sum the counter_value
counter_data = counter_data.groupby(
["dispatch_id", "counter_name"], as_index=False
).agg({"counter_value": "sum"})
kernel_data_filename = glob.glob(
f"{workload_dir}/out/pmc_1/*/{pid}_kernel_trace.csv"
)[0]
Comment thread
vedithal-amd marked this conversation as resolved.
Resolved
kernel_data = pd.read_csv(kernel_data_filename)
rocprofv3_counter_data = pd.DataFrame({
"Correlation_Id": counter_data["dispatch_id"],
"Dispatch_Id": counter_data["dispatch_id"],
"Agent_Id": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"Agent_Id"
].values,
"Queue_Id": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"Queue_Id"
].values,
"Process_Id": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"Thread_Id"
].values,
"Thread_Id": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"Thread_Id"
].values,
"Grid_Size": (
kernel_data.iloc[counter_data["dispatch_id"] - 1][
["Grid_Size_X", "Grid_Size_Y", "Grid_Size_Z"]
]
.prod(axis=1)
.values
),
"Kernel_Id": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"Kernel_Id"
].values,
"Kernel_Name": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"Kernel_Name"
].values,
"Workgroup_Size": (
kernel_data.iloc[counter_data["dispatch_id"] - 1][
["Workgroup_Size_X", "Workgroup_Size_Y", "Workgroup_Size_Z"]
]
.prod(axis=1)
.values
),
"LDS_Block_Size": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"LDS_Block_Size"
].values,
"Scratch_Size": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"Scratch_Size"
].values,
"VGPR_Count": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"VGPR_Count"
].values,
"Accum_VGPR_Count": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"Accum_VGPR_Count"
].values,
"SGPR_Count": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"SGPR_Count"
].values,
"Counter_Name": counter_data["counter_name"],
"Counter_Value": counter_data["counter_value"],
"Start_Timestamp": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"Start_Timestamp"
].values,
"End_Timestamp": kernel_data.iloc[counter_data["dispatch_id"] - 1][
"End_Timestamp"
].values,
})
Comment thread
vedithal-amd marked this conversation as resolved.
Resolved
rocprofv3_counter_data.to_csv(
kernel_data_filename.replace("kernel_trace", "counter_collection"),
index=False,
)


def process_rocprofv3_output(workload_dir: str, using_native_tool: bool) -> list[str]:
Expand Down
8 changes: 8 additions & 0 deletions projects/rocprofiler-compute/tests/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -67,3 +67,11 @@ set_target_properties(
laplace_eqn
PROPERTIES RUNTIME_OUTPUT_DIRECTORY ${CMAKE_SOURCE_DIR}/tests
)

set(ROCFLOP_SOURCES ../sample/rocflop.cpp)
set_source_files_properties(${ROCFLOP_SOURCES} PROPERTIES LANGUAGE HIP)
add_executable(rocflop ${ROCFLOP_SOURCES})
set_target_properties(
rocflop
PROPERTIES RUNTIME_OUTPUT_DIRECTORY ${CMAKE_SOURCE_DIR}/tests
)
Loading
Loading