diff --git a/source/lib/output/buffered_output.hpp b/source/lib/output/buffered_output.hpp index efeb409820..efe238eb7e 100644 --- a/source/lib/output/buffered_output.hpp +++ b/source/lib/output/buffered_output.hpp @@ -163,7 +163,7 @@ buffered_output::get_num_bytes() const } using hip_buffered_output_t = - buffered_output; + buffered_output; using hsa_buffered_output_t = buffered_output; using marker_buffered_output_t = diff --git a/source/lib/output/generateCSV.cpp b/source/lib/output/generateCSV.cpp index 974d2aa670..68363c9d7f 100644 --- a/source/lib/output/generateCSV.cpp +++ b/source/lib/output/generateCSV.cpp @@ -327,10 +327,10 @@ generate_csv(const output_config& } void -generate_csv(const output_config& cfg, - const metadata& tool_metadata, - const generator& data, - const stats_entry_t& stats) +generate_csv(const output_config& cfg, + const metadata& tool_metadata, + const generator& data, + const stats_entry_t& stats) { if(data.empty()) return; diff --git a/source/lib/output/generateCSV.hpp b/source/lib/output/generateCSV.hpp index b747d919bc..0f75f48784 100644 --- a/source/lib/output/generateCSV.hpp +++ b/source/lib/output/generateCSV.hpp @@ -46,10 +46,10 @@ generate_csv(const output_config& const stats_entry_t& stats); void -generate_csv(const output_config& cfg, - const metadata& tool_metadata, - const generator& data, - const stats_entry_t& stats); +generate_csv(const output_config& cfg, + const metadata& tool_metadata, + const generator& data, + const stats_entry_t& stats); void generate_csv(const output_config& cfg, diff --git a/source/lib/output/generateJSON.cpp b/source/lib/output/generateJSON.cpp index 2f6f046c06..7eb8b243ff 100644 --- a/source/lib/output/generateJSON.cpp +++ b/source/lib/output/generateJSON.cpp @@ -186,7 +186,7 @@ write_json( const output_config& /*cfg*/, const metadata& /*tool_metadata*/, const domain_stats_vec_t& domain_stats, - const generator& hip_api_gen, + const generator& hip_api_gen, const generator& hsa_api_gen, const generator& kernel_dispatch_gen, const generator& memory_copy_gen, diff --git a/source/lib/output/generateJSON.hpp b/source/lib/output/generateJSON.hpp index 8bd2954344..4ffc78bd58 100644 --- a/source/lib/output/generateJSON.hpp +++ b/source/lib/output/generateJSON.hpp @@ -86,7 +86,7 @@ write_json( const output_config& cfg, const metadata& tool_metadata, const domain_stats_vec_t& domain_stats, - const generator& hip_api_gen, + const generator& hip_api_gen, const generator& hsa_api_gen, const generator& kernel_dispatch_gen, const generator& memory_copy_gen, diff --git a/source/lib/output/generateOTF2.cpp b/source/lib/output/generateOTF2.cpp index 0147c07685..f0ace1ff5c 100644 --- a/source/lib/output/generateOTF2.cpp +++ b/source/lib/output/generateOTF2.cpp @@ -359,7 +359,7 @@ write_otf2(const output_config& cfg, const metadata& tool_metadata, uint64_t pid, const std::vector& agent_data, - std::deque* hip_api_data, + std::deque* hip_api_data, std::deque* hsa_api_data, std::deque* kernel_dispatch_data, std::deque* memory_copy_data, diff --git a/source/lib/output/generateOTF2.hpp b/source/lib/output/generateOTF2.hpp index 84894a14cf..1c3b94b03b 100644 --- a/source/lib/output/generateOTF2.hpp +++ b/source/lib/output/generateOTF2.hpp @@ -39,7 +39,7 @@ write_otf2(const output_config& cfg, const metadata& tool_metadata, uint64_t pid, const std::vector& agent_data, - std::deque* hip_api_data, + std::deque* hip_api_data, std::deque* hsa_api_data, std::deque* kernel_dispatch_data, std::deque* memory_copy_data, diff --git a/source/lib/output/generatePerfetto.cpp b/source/lib/output/generatePerfetto.cpp index a24087c420..b807bfd09e 100644 --- a/source/lib/output/generatePerfetto.cpp +++ b/source/lib/output/generatePerfetto.cpp @@ -68,7 +68,7 @@ write_perfetto( const output_config& ocfg, const metadata& tool_metadata, std::vector agent_data, - const generator& hip_api_gen, + const generator& hip_api_gen, const generator& hsa_api_gen, const generator& kernel_dispatch_gen, const generator& memory_copy_gen, @@ -142,9 +142,8 @@ write_perfetto( auto agent_thread_ids_alloc = std::unordered_map>{}; auto agent_queue_ids = std::unordered_map>{}; - auto agent_stream_ids = - std::unordered_map>{}; - auto thread_indexes = std::unordered_map{}; + auto agent_stream_ids = std::unordered_set{}; + auto thread_indexes = std::unordered_map{}; auto thread_tracks = std::unordered_map{}; auto agent_thread_tracks = @@ -190,7 +189,7 @@ write_perfetto( for(auto itr : memory_copy_gen.get(ditr)) { tids.emplace(itr.thread_id); - agent_stream_ids[itr.dst_agent_id].emplace(itr.stream_id); + agent_stream_ids.emplace(itr.stream_id); if(group_by_queue) { agent_thread_ids[itr.dst_agent_id].emplace(itr.thread_id); @@ -208,7 +207,7 @@ write_perfetto( for(auto itr : kernel_dispatch_gen.get(ditr)) { tids.emplace(itr.thread_id); - agent_stream_ids[itr.dispatch_info.agent_id].emplace(itr.stream_id); + agent_stream_ids.emplace(itr.stream_id); if(group_by_queue) { agent_queue_ids[itr.dispatch_info.agent_id].emplace(itr.dispatch_info.queue_id); @@ -290,24 +289,21 @@ write_perfetto( } } - for(const auto& aitr : agent_stream_ids) + for(const auto& sitr : agent_stream_ids) { - for(auto sitr : aitr.second) + const auto stream_id = sitr.handle; + { - const auto stream_id = sitr.handle; + auto _namess = std::stringstream{}; + _namess << fmt::format("STREAM [\" {} \"] ", stream_id); - { - auto _namess = std::stringstream{}; - _namess << fmt::format("STREAM [\" {} \"] ", stream_id); + auto _track = ::perfetto::Track{get_hash_id(_namess.str())}; + auto _desc = _track.Serialize(); + _desc.set_name(_namess.str()); - auto _track = ::perfetto::Track{get_hash_id(_namess.str())}; - auto _desc = _track.Serialize(); - _desc.set_name(_namess.str()); + perfetto::TrackEvent::SetTrackDescriptor(_track, _desc); - perfetto::TrackEvent::SetTrackDescriptor(_track, _desc); - - stream_tracks.emplace(sitr, _track); - } + stream_tracks.emplace(sitr, _track); } } @@ -383,7 +379,9 @@ write_perfetto( "corr_id", itr.correlation_id.internal, "ancestor_id", - itr.correlation_id.ancestor); + itr.correlation_id.ancestor, + "stream_ID", + itr.stream_id.handle); TRACE_EVENT_END( sdk::perfetto_category::name, track, itr.end_timestamp); @@ -575,7 +573,9 @@ write_perfetto( "corr_id", itr.correlation_id.internal, "tid", - itr.thread_id); + itr.thread_id, + "stream_ID", + itr.stream_id.handle); TRACE_EVENT_END(sdk::perfetto_category::name, *_track, itr.end_timestamp); @@ -636,14 +636,15 @@ write_perfetto( auto name = std::string_view{sym->kernel_name}; - ::perfetto::Track* _track = nullptr; + ::perfetto::Track* _track = nullptr; + auto stream_id = (*it)->stream_id; if(group_by_queue) { _track = &agent_queue_tracks.at(info.agent_id).at(info.queue_id); } else { - _track = &stream_tracks.at((*it)->stream_id); + _track = &stream_tracks.at(stream_id); } // Temporary fix until timestamp issues are resolved: Set timestamps to be @@ -674,6 +675,8 @@ write_perfetto( { demangled.emplace(name, common::cxx_demangle(name)); } + // Queue IDs are 1 higher than the track name. Subtracting 1 for consistency + auto queue_id = info.queue_id.handle > 0 ? info.queue_id.handle - 1 : 0; TRACE_EVENT_BEGIN( sdk::perfetto_category::name, @@ -702,19 +705,27 @@ write_perfetto( "corr_id", current.correlation_id.internal, "queue", - info.queue_id.handle, + queue_id, "tid", current.thread_id, "kernel_id", info.kernel_id, - "private_segment_size", + "Scratch_Size", info.private_segment_size, - "group_segment_size", + "LDS_Block_Size", info.group_segment_size, + "VGPR_Count", + sym->arch_vgpr_count, + "Accum_VGPR_Count", + sym->accum_vgpr_count, + "SGPR_Count", + sym->sgpr_count, "workgroup_size", info.workgroup_size.x * info.workgroup_size.y * info.workgroup_size.z, "grid_size", info.grid_size.x * info.grid_size.y * info.grid_size.z, + "stream_ID", + stream_id.handle, [&](::perfetto::EventContext ctx) { auto corr_id = current.correlation_id.internal; auto counter_it = dispatch_counter_id_value.find(corr_id); diff --git a/source/lib/output/generatePerfetto.hpp b/source/lib/output/generatePerfetto.hpp index 18863a53f6..15fa43a7ae 100644 --- a/source/lib/output/generatePerfetto.hpp +++ b/source/lib/output/generatePerfetto.hpp @@ -40,7 +40,7 @@ write_perfetto( const output_config& cfg, const metadata& tool_metadata, std::vector agent_data, - const generator& hip_api_gen, + const generator& hip_api_gen, const generator& hsa_api_gen, const generator& kernel_dispatch_gen, const generator& memory_copy_gen, diff --git a/source/lib/output/generateRocpd.cpp b/source/lib/output/generateRocpd.cpp index 8f601790b8..6567767f45 100644 --- a/source/lib/output/generateRocpd.cpp +++ b/source/lib/output/generateRocpd.cpp @@ -551,7 +551,7 @@ write_rocpd( const output_config& cfg, const metadata& tool_metadata, const std::vector& agent_data, - const generator& hip_api_gen, + const generator& hip_api_gen, const generator& hsa_api_gen, const generator& kernel_dispatch_gen, const generator& memory_copy_gen, diff --git a/source/lib/output/generateRocpd.hpp b/source/lib/output/generateRocpd.hpp index f12f7d2d42..1ed8fac9d1 100644 --- a/source/lib/output/generateRocpd.hpp +++ b/source/lib/output/generateRocpd.hpp @@ -40,7 +40,7 @@ write_rocpd( const output_config& cfg, const metadata& tool_metadata, const std::vector& agent_data, - const generator& hip_api_gen, + const generator& hip_api_gen, const generator& hsa_api_gen, const generator& kernel_dispatch_gen, const generator& memory_copy_gen, diff --git a/source/lib/output/generateStats.cpp b/source/lib/output/generateStats.cpp index e99eee7ce7..13193c9a5f 100644 --- a/source/lib/output/generateStats.cpp +++ b/source/lib/output/generateStats.cpp @@ -83,8 +83,8 @@ generate_stats(const output_config& /*cfg*/, stats_entry_t generate_stats(const output_config& /*cfg*/, - const metadata& tool_metadata, - const generator& data) + const metadata& tool_metadata, + const generator& data) { auto hip_stats = stats_map_t{}; for(auto ditr : data) diff --git a/source/lib/output/generateStats.hpp b/source/lib/output/generateStats.hpp index 801af0c931..a75095c6f4 100644 --- a/source/lib/output/generateStats.hpp +++ b/source/lib/output/generateStats.hpp @@ -37,9 +37,9 @@ generate_stats(const output_config& const generator& data); stats_entry_t -generate_stats(const output_config& cfg, - const metadata& tool_metadata, - const generator& data); +generate_stats(const output_config& cfg, + const metadata& tool_metadata, + const generator& data); stats_entry_t generate_stats(const output_config& cfg, diff --git a/source/lib/output/stream_info.hpp b/source/lib/output/stream_info.hpp index 15b0dfe42d..990c586c0e 100644 --- a/source/lib/output/stream_info.hpp +++ b/source/lib/output/stream_info.hpp @@ -40,8 +40,8 @@ struct tool_buffer_tracing_kernel_dispatch_ext_record_t { using base_type = rocprofiler_buffer_tracing_kernel_dispatch_record_t; - tool_buffer_tracing_kernel_dispatch_ext_record_t(const base_type& _base, - const rocprofiler_stream_id_t& _stream_id) + tool_buffer_tracing_kernel_dispatch_ext_record_t(const base_type& _base, + const rocprofiler_stream_id_t _stream_id) : base_type{_base} , stream_id{_stream_id} {} @@ -65,8 +65,8 @@ struct tool_buffer_tracing_memory_copy_ext_record_t { using base_type = rocprofiler_buffer_tracing_memory_copy_record_t; - tool_buffer_tracing_memory_copy_ext_record_t(const base_type& _base, - const rocprofiler_stream_id_t& _stream_id) + tool_buffer_tracing_memory_copy_ext_record_t(const base_type& _base, + const rocprofiler_stream_id_t _stream_id) : base_type{_base} , stream_id{_stream_id} {} @@ -90,8 +90,8 @@ struct tool_buffer_tracing_memory_allocation_ext_record_t { using base_type = rocprofiler_buffer_tracing_memory_allocation_record_t; - tool_buffer_tracing_memory_allocation_ext_record_t(const base_type& _base, - const rocprofiler_stream_id_t& _stream_id) + tool_buffer_tracing_memory_allocation_ext_record_t(const base_type& _base, + const rocprofiler_stream_id_t _stream_id) : base_type{_base} , stream_id{_stream_id} {} @@ -110,6 +110,30 @@ struct tool_buffer_tracing_memory_allocation_ext_record_t rocprofiler_stream_id_t stream_id = {}; }; +struct tool_buffer_tracing_hip_api_ext_record_t : rocprofiler_buffer_tracing_hip_api_ext_record_t +{ + using base_type = rocprofiler_buffer_tracing_hip_api_ext_record_t; + + tool_buffer_tracing_hip_api_ext_record_t(const base_type& _base, + const rocprofiler_stream_id_t _stream_id) + : base_type{_base} + , stream_id{_stream_id} + {} + + tool_buffer_tracing_hip_api_ext_record_t() = delete; + ~tool_buffer_tracing_hip_api_ext_record_t() = default; + tool_buffer_tracing_hip_api_ext_record_t(const tool_buffer_tracing_hip_api_ext_record_t&) = + default; + tool_buffer_tracing_hip_api_ext_record_t(tool_buffer_tracing_hip_api_ext_record_t&&) noexcept = + default; + tool_buffer_tracing_hip_api_ext_record_t& operator =( + const tool_buffer_tracing_hip_api_ext_record_t&) = default; + tool_buffer_tracing_hip_api_ext_record_t& operator =( + tool_buffer_tracing_hip_api_ext_record_t&&) noexcept = default; + + rocprofiler_stream_id_t stream_id = {}; +}; + } // namespace tool } // namespace rocprofiler @@ -144,5 +168,13 @@ save(ArchiveT& SAVE_DATA_FIELD(stream_id); } +template +void +save(ArchiveT& ar, const ::rocprofiler::tool::tool_buffer_tracing_hip_api_ext_record_t& data) +{ + cereal::save(ar, static_cast(data)); + SAVE_DATA_FIELD(stream_id); +} + #undef SAVE_DATA_FIELD } // namespace cereal diff --git a/source/lib/python/rocpd/source/perfetto.cpp b/source/lib/python/rocpd/source/perfetto.cpp index 9908941fbb..95d08b7719 100644 --- a/source/lib/python/rocpd/source/perfetto.cpp +++ b/source/lib/python/rocpd/source/perfetto.cpp @@ -212,9 +212,8 @@ write_perfetto( auto agent_thread_ids_alloc = std::unordered_map>{}; auto agent_queue_ids = std::unordered_map>{}; - auto agent_stream_ids = - std::unordered_map>{}; - auto thread_indexes = std::unordered_map{}; + auto agent_stream_ids = std::unordered_set{}; + auto thread_indexes = std::unordered_map{}; auto thread_tracks = std::unordered_map{}; auto agent_thread_tracks = @@ -222,16 +221,14 @@ write_perfetto( auto agent_queue_tracks = std::unordered_map>{}; - auto agent_stream_tracks = - std::unordered_map>{}; + auto stream_tracks = std::unordered_map{}; { for(auto ditr : memory_copy_gen) for(const auto& itr : memory_copy_gen.get(ditr)) { auto stream_id = rocprofiler_stream_id_t{.handle = itr.stream_id}; - agent_stream_ids[itr.dst_agent_abs_index].emplace(stream_id); + agent_stream_ids.emplace(stream_id); if(ocfg.group_by_queue) { agent_thread_ids[itr.dst_agent_abs_index].emplace(itr.tid); @@ -251,7 +248,7 @@ write_perfetto( { auto stream_id = rocprofiler_stream_id_t{.handle = itr.stream_id}; auto queue_id = rocprofiler_queue_id_t{.handle = itr.queue_id}; - agent_stream_ids[itr.agent_abs_index].emplace(stream_id); + agent_stream_ids.emplace(stream_id); if(ocfg.group_by_queue) { agent_queue_ids[itr.agent_abs_index].emplace(queue_id); @@ -335,32 +332,19 @@ write_perfetto( } } - for(const auto& [abs_index, stream_ids] : agent_stream_ids) + for(const auto& sitr : agent_stream_ids) { - const auto _agent = agent_data.at(abs_index).first; - // auto agent_index_info = agent_data.at(abs_index).second; - for(auto sitr : stream_ids) - { - const auto stream_id = sitr.handle; + const auto stream_id = sitr.handle; - auto _name = - fmt::format("COMPUTE AGENT [{}] STREAM [{}]", _agent.logical_node_id, stream_id); + auto _name = fmt::format("STREAM [{}]", stream_id); - if(_agent.type == "CPU") - _name = fmt::format("{} (CPU)", _name); - else if(_agent.type == "GPU") - _name = fmt::format("{} (GPU)", _name); - else - _name = fmt::format("{} (UNK)", _name); + auto _track = ::perfetto::Track{get_hash_id(_name), this_pid_track}; + auto _desc = _track.Serialize(); + _desc.set_name(_name); - auto _track = ::perfetto::Track{get_hash_id(_name), this_pid_track}; - auto _desc = _track.Serialize(); - _desc.set_name(_name); + ::perfetto::TrackEvent::SetTrackDescriptor(_track, _desc); - ::perfetto::TrackEvent::SetTrackDescriptor(_track, _desc); - - agent_stream_tracks[abs_index].emplace(sitr, _track); - } + stream_tracks.emplace(sitr, _track); } // Fetch counter values @@ -482,7 +466,7 @@ write_perfetto( else { auto stream_id = rocprofiler_stream_id_t{.handle = itr.stream_id}; - _track = &agent_stream_tracks.at(itr.dst_agent_abs_index).at(stream_id); + _track = &stream_tracks.at(stream_id); } auto src_agent_index = agent_data.at(itr.src_agent_abs_index).second; @@ -511,7 +495,9 @@ write_perfetto( "corr_id", itr.stack_id, "tid", - itr.tid); + itr.tid, + "stream_id", + itr.stream_id); TRACE_EVENT_END( sdk::perfetto_category::name, *_track, itr.end); } @@ -535,7 +521,7 @@ write_perfetto( } else { - _track = &agent_stream_tracks.at(agent_id).at(stream_id); + _track = &stream_tracks.at(stream_id); } // Temporary fix until timestamp issues are resolved: Set timestamps to be @@ -590,14 +576,16 @@ write_perfetto( current.tid, "kernel_id", current.kernel_id, - "private_segment_size", + "Scratch_Size", current.scratch_size, - "group_segment_size", + "LDS_Block_Size", current.lds_size, "workgroup_size", to_string(current.workgroup_size), "grid_size", to_string(current.grid_size), + "stream_id", + current.stream_id, [&](::perfetto::EventContext ctx) { for(auto& [counter_id, counter_value] : counter_id_value) { diff --git a/source/lib/rocprofiler-sdk-tool/tool.cpp b/source/lib/rocprofiler-sdk-tool/tool.cpp index d1a789c571..a694286b8c 100644 --- a/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -615,6 +615,7 @@ hip_stream_display_callback(rocprofiler_callback_tracing_record_t record, auto* stream_handle_data = static_cast(record.payload); auto stream_id = stream_handle_data->stream_id; + // STREAM_HANDLE_CREATE and DESTROY are no-ops if(record.operation == ROCPROFILER_HIP_STREAM_CREATE) { @@ -1067,7 +1068,10 @@ buffered_tracing_callback(rocprofiler_context_id_t /*context*/, auto* record = static_cast(header->payload); - tool::write_ring_buffer(*record, domain_type::HIP); + auto stream_id = get_stream_id(record); + tool::write_ring_buffer( + tool::tool_buffer_tracing_hip_api_ext_record_t{*record, stream_id}, + domain_type::HIP); } else if(header->kind == ROCPROFILER_BUFFER_TRACING_RCCL_API) { @@ -2023,10 +2027,11 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) tool::get_config().benchmark_mode != tool::config::benchmark::execution_profile) { auto external_corr_id_request_kinds = - std::array{ + std::array{ ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_KERNEL_DISPATCH, ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_MEMORY_COPY, - ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_MEMORY_ALLOCATION}; + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_MEMORY_ALLOCATION, + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_HIP_RUNTIME_API}; ROCPROFILER_CALL(rocprofiler_configure_external_correlation_id_request_service( get_client_ctx(),