From af1a6a6b37f3dc2b46b4fa2691a92fb67bfbfa3b Mon Sep 17 00:00:00 2001 From: 7b30f3f5e26d48061f873d04cc7e1d1f_amdeng <103e5765c35bad2b9d03171132741341@amd.com> Date: Wed, 21 May 2025 11:55:48 -0400 Subject: [PATCH] Add VGPR and SGPR counts to kernel trace file (#371) Signed-off-by: Tim Co-authored-by: sonadeem [ROCm/rocprofiler-sdk commit: 856491944b730cde7ea3a302ae1e2b0101b4e75d] --- .../rocprofiler-sdk/source/lib/output/csv.hpp | 2 +- .../source/lib/output/generateCSV.cpp | 21 ++++++++++++++----- 2 files changed, 17 insertions(+), 6 deletions(-) diff --git a/projects/rocprofiler-sdk/source/lib/output/csv.hpp b/projects/rocprofiler-sdk/source/lib/output/csv.hpp index 2ea467b5d6..4bd2d1b45e 100644 --- a/projects/rocprofiler-sdk/source/lib/output/csv.hpp +++ b/projects/rocprofiler-sdk/source/lib/output/csv.hpp @@ -109,7 +109,7 @@ using list_derived_metrics_csv_encoder = csv_encoder<5>; using scratch_memory_encoder = csv_encoder<8>; using stats_csv_encoder = csv_encoder<8>; 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 pc_sampling_stochastic_csv_encoder = csv_encoder<10>; } // namespace csv diff --git a/projects/rocprofiler-sdk/source/lib/output/generateCSV.cpp b/projects/rocprofiler-sdk/source/lib/output/generateCSV.cpp index ee39f11ddb..974d2aa670 100644 --- a/projects/rocprofiler-sdk/source/lib/output/generateCSV.cpp +++ b/projects/rocprofiler-sdk/source/lib/output/generateCSV.cpp @@ -272,8 +272,11 @@ generate_csv(const output_config& "Correlation_Id", "Start_Timestamp", "End_Timestamp", - "Private_Segment_Size", - "Group_Segment_Size", + "LDS_Block_Size", + "Scratch_Size", + "VGPR_Count", + "Accum_VGPR_Count", + "SGPR_Count", "Workgroup_Size_X", "Workgroup_Size_Y", "Workgroup_Size_Z", @@ -285,9 +288,14 @@ generate_csv(const output_config& { for(auto record : data.get(ditr)) { - auto row_ss = std::stringstream{}; - auto kernel_name = tool_metadata.get_kernel_name(record.dispatch_info.kernel_id, + auto row_ss = std::stringstream{}; + auto kernel_name = tool_metadata.get_kernel_name(record.dispatch_info.kernel_id, 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( row_ss, tool_metadata.get_kind_name(record.kind), @@ -302,8 +310,11 @@ generate_csv(const output_config& record.correlation_id.internal, record.start_timestamp, record.end_timestamp, + lds_block_size_v, 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.y, record.dispatch_info.workgroup_size.z,