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

Add VGPR and SGPR counts to kernel trace file (#371)

Browse Source 
Signed-off-by: Tim <Tim.Gu@Amd.com>
Co-authored-by: sonadeem <sohaib.nadeem@amd.com>
This commit is contained in:
7b30f3f5e26d48061f873d04cc7e1d1f_amdeng
2025-05-21 11:55:48 -04:00
committed by GitHub
parent 14c2dc55ff
commit 856491944b
2 changed files with 17 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files
source/lib/output/csv.hpp
+1 -1
View File
@@ -109,7 +109,7 @@ using list_derived_metrics_csv_encoder = csv_encoder<5>;
using scratch_memory_encoder = csv_encoder<8>; using scratch_memory_encoder = csv_encoder<8>;
using stats_csv_encoder = csv_encoder<8>; using stats_csv_encoder = csv_encoder<8>;
using pc_sampling_host_trap_csv_encoder = csv_encoder<6>; using pc_sampling_host_trap_csv_encoder = csv_encoder<6>;
using kernel_trace_with_stream_csv_encoder = csv_encoder<19>; using kernel_trace_with_stream_csv_encoder = csv_encoder<22>;
using memory_copy_with_stream_csv_encoder = csv_encoder<8>; using memory_copy_with_stream_csv_encoder = csv_encoder<8>;
using pc_sampling_stochastic_csv_encoder = csv_encoder<10>; using pc_sampling_stochastic_csv_encoder = csv_encoder<10>;
} // namespace csv } // namespace csv
source/lib/output/generateCSV.cpp
+16 -5
View File
@@ -272,8 +272,11 @@ generate_csv(const output_config&
"Correlation_Id", "Correlation_Id",
"Start_Timestamp", "Start_Timestamp",
"End_Timestamp", "End_Timestamp",
"Private_Segment_Size", "LDS_Block_Size",
"Group_Segment_Size", "Scratch_Size",
"VGPR_Count",
"Accum_VGPR_Count",
"SGPR_Count",
"Workgroup_Size_X", "Workgroup_Size_X",
"Workgroup_Size_Y", "Workgroup_Size_Y",
"Workgroup_Size_Z", "Workgroup_Size_Z",
@@ -285,9 +288,14 @@ generate_csv(const output_config&
{ {
for(auto record : data.get(ditr)) for(auto record : data.get(ditr))
{ {
auto row_ss = std::stringstream{}; auto row_ss = std::stringstream{};
auto kernel_name = tool_metadata.get_kernel_name(record.dispatch_info.kernel_id, auto kernel_name = tool_metadata.get_kernel_name(record.dispatch_info.kernel_id,
record.correlation_id.external.value); record.correlation_id.external.value);
const auto* kernel_info =
tool_metadata.get_kernel_symbol(record.dispatch_info.kernel_id);
auto lds_block_size_v =
(kernel_info->group_segment_size + (lds_block_size - 1)) & ~(lds_block_size - 1);
rocprofiler::tool::csv::kernel_trace_with_stream_csv_encoder::write_row( rocprofiler::tool::csv::kernel_trace_with_stream_csv_encoder::write_row(
row_ss, row_ss,
tool_metadata.get_kind_name(record.kind), tool_metadata.get_kind_name(record.kind),
@@ -302,8 +310,11 @@ generate_csv(const output_config&
record.correlation_id.internal, record.correlation_id.internal,
record.start_timestamp, record.start_timestamp,
record.end_timestamp, record.end_timestamp,
lds_block_size_v,
record.dispatch_info.private_segment_size, record.dispatch_info.private_segment_size,
record.dispatch_info.group_segment_size, kernel_info->arch_vgpr_count,
kernel_info->accum_vgpr_count,
kernel_info->sgpr_count,
record.dispatch_info.workgroup_size.x, record.dispatch_info.workgroup_size.x,
record.dispatch_info.workgroup_size.y, record.dispatch_info.workgroup_size.y,
record.dispatch_info.workgroup_size.z, record.dispatch_info.workgroup_size.z,