AI/rocm-systems
2
0
Fork 0
Código Issues Pull Requests Ações Pacotes Projetos Versões Wiki Atividade

Modified perfetto output for HIP stream display (#431)

Ver Código-Fonte 
* Modified perfetto output for HIP stream display

* Moved stream_map file location and changed perfetto output names Private_Segment_Size and Group_Segment_Size to Scratch_Size and LDS_Block_Size respectively

* Used const_cast to remove const modifier on void*

* Reverted stream_map changes, now using tool_metadata map to track mapping between stream ptrs and stream IDs

* Removed buffer tracing args in perfetto, added tool_...hip buffer record struct that stores the HIP stream ID for display purposes

* Updated rocpd perfetto.cpp to reflect stream changes. Still need to add vgpr values and stream ID for HIP API

* Changes pass-by const reference to pass-by const value
Esse commit está contido em:
Trowbridge, Ian
2025-06-26 14:22:50 -05:00
commit de GitHub
pai 8a461afe20
commit 1f8b8c5e9f
16 arquivos alterados com 126 adições e 90 exclusões
Baixar Arquivo de Patch Baixar Arquivo de Diff Expandir todos os arquivos Colapsar todos os arquivos
source/lib/output/buffered_output.hpp
+1 -1
Ver Arquivo
@@ -163,7 +163,7 @@ buffered_output<Tp, DomainT>::get_num_bytes() const
}
using hip_buffered_output_t =
buffered_output<rocprofiler_buffer_tracing_hip_api_ext_record_t, domain_type::HIP>;
buffered_output<tool_buffer_tracing_hip_api_ext_record_t, domain_type::HIP>;
using hsa_buffered_output_t =
buffered_output<rocprofiler_buffer_tracing_hsa_api_record_t, domain_type::HSA>;
using marker_buffered_output_t =
source/lib/output/generateCSV.cpp
+4 -4
Ver Arquivo
@@ -327,10 +327,10 @@ generate_csv(const output_config&
}
void
generate_csv(const output_config& cfg,
const metadata& tool_metadata,
const generator<rocprofiler_buffer_tracing_hip_api_ext_record_t>& data,
const stats_entry_t& stats)
generate_csv(const output_config& cfg,
const metadata& tool_metadata,
const generator<tool_buffer_tracing_hip_api_ext_record_t>& data,
const stats_entry_t& stats)
{
if(data.empty()) return;
source/lib/output/generateCSV.hpp
+4 -4
Ver Arquivo
@@ -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<rocprofiler_buffer_tracing_hip_api_ext_record_t>& data,
const stats_entry_t& stats);
generate_csv(const output_config& cfg,
const metadata& tool_metadata,
const generator<tool_buffer_tracing_hip_api_ext_record_t>& data,
const stats_entry_t& stats);
void
generate_csv(const output_config& cfg,
source/lib/output/generateJSON.cpp
+1 -1
Ver Arquivo
@@ -186,7 +186,7 @@ write_json(
const output_config& /*cfg*/,
const metadata& /*tool_metadata*/,
const domain_stats_vec_t& domain_stats,
const generator<rocprofiler_buffer_tracing_hip_api_ext_record_t>& hip_api_gen,
const generator<tool_buffer_tracing_hip_api_ext_record_t>& hip_api_gen,
const generator<rocprofiler_buffer_tracing_hsa_api_record_t>& hsa_api_gen,
const generator<tool_buffer_tracing_kernel_dispatch_ext_record_t>& kernel_dispatch_gen,
const generator<tool_buffer_tracing_memory_copy_ext_record_t>& memory_copy_gen,
source/lib/output/generateJSON.hpp
+1 -1
Ver Arquivo
@@ -86,7 +86,7 @@ write_json(
const output_config& cfg,
const metadata& tool_metadata,
const domain_stats_vec_t& domain_stats,
const generator<rocprofiler_buffer_tracing_hip_api_ext_record_t>& hip_api_gen,
const generator<tool_buffer_tracing_hip_api_ext_record_t>& hip_api_gen,
const generator<rocprofiler_buffer_tracing_hsa_api_record_t>& hsa_api_gen,
const generator<tool_buffer_tracing_kernel_dispatch_ext_record_t>& kernel_dispatch_gen,
const generator<tool_buffer_tracing_memory_copy_ext_record_t>& memory_copy_gen,
source/lib/output/generateOTF2.cpp
+1 -1
Ver Arquivo
@@ -359,7 +359,7 @@ write_otf2(const output_config& cfg,
const metadata& tool_metadata,
uint64_t pid,
const std::vector<agent_info>& agent_data,
std::deque<rocprofiler_buffer_tracing_hip_api_ext_record_t>* hip_api_data,
std::deque<tool_buffer_tracing_hip_api_ext_record_t>* hip_api_data,
std::deque<rocprofiler_buffer_tracing_hsa_api_record_t>* hsa_api_data,
std::deque<tool_buffer_tracing_kernel_dispatch_ext_record_t>* kernel_dispatch_data,
std::deque<tool_buffer_tracing_memory_copy_ext_record_t>* memory_copy_data,
source/lib/output/generateOTF2.hpp
+1 -1
Ver Arquivo
@@ -39,7 +39,7 @@ write_otf2(const output_config& cfg,
const metadata& tool_metadata,
uint64_t pid,
const std::vector<agent_info>& agent_data,
std::deque<rocprofiler_buffer_tracing_hip_api_ext_record_t>* hip_api_data,
std::deque<tool_buffer_tracing_hip_api_ext_record_t>* hip_api_data,
std::deque<rocprofiler_buffer_tracing_hsa_api_record_t>* hsa_api_data,
std::deque<tool_buffer_tracing_kernel_dispatch_ext_record_t>* kernel_dispatch_data,
std::deque<tool_buffer_tracing_memory_copy_ext_record_t>* memory_copy_data,
source/lib/output/generatePerfetto.cpp
+37 -26
Ver Arquivo
@@ -68,7 +68,7 @@ write_perfetto(
const output_config& ocfg,
const metadata& tool_metadata,
std::vector<agent_info> agent_data,
const generator<rocprofiler_buffer_tracing_hip_api_ext_record_t>& hip_api_gen,
const generator<tool_buffer_tracing_hip_api_ext_record_t>& hip_api_gen,
const generator<rocprofiler_buffer_tracing_hsa_api_record_t>& hsa_api_gen,
const generator<tool_buffer_tracing_kernel_dispatch_ext_record_t>& kernel_dispatch_gen,
const generator<tool_buffer_tracing_memory_copy_ext_record_t>& memory_copy_gen,
@@ -142,9 +142,8 @@ write_perfetto(
auto agent_thread_ids_alloc = std::unordered_map<rocprofiler_agent_id_t, std::set<uint64_t>>{};
auto agent_queue_ids =
std::unordered_map<rocprofiler_agent_id_t, std::unordered_set<rocprofiler_queue_id_t>>{};
auto agent_stream_ids =
std::unordered_map<rocprofiler_agent_id_t, std::unordered_set<rocprofiler_stream_id_t>>{};
auto thread_indexes = std::unordered_map<rocprofiler_thread_id_t, uint64_t>{};
auto agent_stream_ids = std::unordered_set<rocprofiler_stream_id_t>{};
auto thread_indexes = std::unordered_map<rocprofiler_thread_id_t, uint64_t>{};
auto thread_tracks = std::unordered_map<rocprofiler_thread_id_t, ::perfetto::Track>{};
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<sdk::category::hip_api>::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<sdk::category::memory_copy>::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<sdk::category::kernel_dispatch>::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);
source/lib/output/generatePerfetto.hpp
+1 -1
Ver Arquivo
@@ -40,7 +40,7 @@ write_perfetto(
const output_config& cfg,
const metadata& tool_metadata,
std::vector<agent_info> agent_data,
const generator<rocprofiler_buffer_tracing_hip_api_ext_record_t>& hip_api_gen,
const generator<tool_buffer_tracing_hip_api_ext_record_t>& hip_api_gen,
const generator<rocprofiler_buffer_tracing_hsa_api_record_t>& hsa_api_gen,
const generator<tool_buffer_tracing_kernel_dispatch_ext_record_t>& kernel_dispatch_gen,
const generator<tool_buffer_tracing_memory_copy_ext_record_t>& memory_copy_gen,
source/lib/output/generateRocpd.cpp
+1 -1
Ver Arquivo
@@ -551,7 +551,7 @@ write_rocpd(
const output_config& cfg,
const metadata& tool_metadata,
const std::vector<agent_info>& agent_data,
const generator<rocprofiler_buffer_tracing_hip_api_ext_record_t>& hip_api_gen,
const generator<tool_buffer_tracing_hip_api_ext_record_t>& hip_api_gen,
const generator<rocprofiler_buffer_tracing_hsa_api_record_t>& hsa_api_gen,
const generator<tool_buffer_tracing_kernel_dispatch_ext_record_t>& kernel_dispatch_gen,
const generator<tool_buffer_tracing_memory_copy_ext_record_t>& memory_copy_gen,
source/lib/output/generateRocpd.hpp
+1 -1
Ver Arquivo
@@ -40,7 +40,7 @@ write_rocpd(
const output_config& cfg,
const metadata& tool_metadata,
const std::vector<agent_info>& agent_data,
const generator<rocprofiler_buffer_tracing_hip_api_ext_record_t>& hip_api_gen,
const generator<tool_buffer_tracing_hip_api_ext_record_t>& hip_api_gen,
const generator<rocprofiler_buffer_tracing_hsa_api_record_t>& hsa_api_gen,
const generator<tool_buffer_tracing_kernel_dispatch_ext_record_t>& kernel_dispatch_gen,
const generator<tool_buffer_tracing_memory_copy_ext_record_t>& memory_copy_gen,
source/lib/output/generateStats.cpp
+2 -2
Ver Arquivo
@@ -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<rocprofiler_buffer_tracing_hip_api_ext_record_t>& data)
const metadata& tool_metadata,
const generator<tool_buffer_tracing_hip_api_ext_record_t>& data)
{
auto hip_stats = stats_map_t{};
for(auto ditr : data)
source/lib/output/generateStats.hpp
+3 -3
Ver Arquivo
@@ -37,9 +37,9 @@ generate_stats(const output_config&
const generator<tool_buffer_tracing_kernel_dispatch_ext_record_t>& data);
stats_entry_t
generate_stats(const output_config& cfg,
const metadata& tool_metadata,
const generator<rocprofiler_buffer_tracing_hip_api_ext_record_t>& data);
generate_stats(const output_config& cfg,
const metadata& tool_metadata,
const generator<tool_buffer_tracing_hip_api_ext_record_t>& data);
stats_entry_t
generate_stats(const output_config& cfg,
source/lib/output/stream_info.hpp
+38 -6
Ver Arquivo
@@ -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 <typename ArchiveT>
void
save(ArchiveT& ar, const ::rocprofiler::tool::tool_buffer_tracing_hip_api_ext_record_t& data)
{
cereal::save(ar, static_cast<const rocprofiler_buffer_tracing_hip_api_ext_record_t&>(data));
SAVE_DATA_FIELD(stream_id);
}
#undef SAVE_DATA_FIELD
} // namespace cereal
source/lib/python/rocpd/source/perfetto.cpp
+22 -34
Ver Arquivo
@@ -212,9 +212,8 @@ write_perfetto(
auto agent_thread_ids_alloc = std::unordered_map<uint64_t, std::set<uint64_t>>{};
auto agent_queue_ids =
std::unordered_map<uint64_t, std::unordered_set<rocprofiler_queue_id_t>>{};
auto agent_stream_ids =
std::unordered_map<uint64_t, std::unordered_set<rocprofiler_stream_id_t>>{};
auto thread_indexes = std::unordered_map<uint64_t, uint64_t>{};
auto agent_stream_ids = std::unordered_set<rocprofiler_stream_id_t>{};
auto thread_indexes = std::unordered_map<uint64_t, uint64_t>{};
auto thread_tracks = std::unordered_map<uint64_t, ::perfetto::Track>{};
auto agent_thread_tracks =
@@ -222,16 +221,14 @@ write_perfetto(
auto agent_queue_tracks =
std::unordered_map<uint64_t,
std::unordered_map<rocprofiler_queue_id_t, ::perfetto::Track>>{};
auto agent_stream_tracks =
std::unordered_map<uint64_t,
std::unordered_map<rocprofiler_stream_id_t, ::perfetto::Track>>{};
auto stream_tracks = std::unordered_map<rocprofiler_stream_id_t, ::perfetto::Track>{};
{
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<sdk::category::memory_copy>::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)
{
source/lib/rocprofiler-sdk-tool/tool.cpp
+8 -3
Ver Arquivo
@@ -615,6 +615,7 @@ hip_stream_display_callback(rocprofiler_callback_tracing_record_t record,
auto* stream_handle_data =
static_cast<rocprofiler_callback_tracing_hip_stream_data_t*>(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<rocprofiler_buffer_tracing_hip_api_ext_record_t*>(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<rocprofiler_external_correlation_id_request_kind_t, 3>{
std::array<rocprofiler_external_correlation_id_request_kind_t, 4>{
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(),