From 4b7e5f51da1da81f2fa5d12cf76cd4fae2013a77 Mon Sep 17 00:00:00 2001 From: Evgeny Date: Fri, 2 Feb 2018 15:38:28 -0600 Subject: [PATCH] read API [ROCm/rocprofiler commit: ff1e5b67a73dd43861407e0d7750e09b85b29b3a] --- projects/rocprofiler/src/core/context.h | 14 +++++++++-- projects/rocprofiler/src/core/profile.h | 25 ++++++++++++++++++- projects/rocprofiler/src/core/rocprofiler.cpp | 15 +++++++++++ .../rocprofiler/src/util/hsa_rsrc_factory.cpp | 4 +-- .../test/util/hsa_rsrc_factory.cpp | 4 +-- 5 files changed, 55 insertions(+), 7 deletions(-) diff --git a/projects/rocprofiler/src/core/context.h b/projects/rocprofiler/src/core/context.h index 59e24ed70c..1ad4761baf 100644 --- a/projects/rocprofiler/src/core/context.h +++ b/projects/rocprofiler/src/core/context.h @@ -94,9 +94,9 @@ class Group { } hsa_status_t Finalize() { - hsa_status_t status = pmc_profile_.Finalize(start_vector_, stop_vector_); + hsa_status_t status = pmc_profile_.Finalize(start_vector_, stop_vector_, read_vector_); if (status == HSA_STATUS_SUCCESS) { - status = sqtt_profile_.Finalize(start_vector_, stop_vector_); + status = sqtt_profile_.Finalize(start_vector_, stop_vector_, read_vector_); } if (status == HSA_STATUS_SUCCESS) { if (!pmc_profile_.Empty()) ++n_profiles_; @@ -115,6 +115,7 @@ class Group { info_vector_t& GetInfoVector() { return info_vector_; } const pkt_vector_t& GetStartVector() const { return start_vector_; } const pkt_vector_t& GetStopVector() const { return stop_vector_; } + const pkt_vector_t& GetReadVector() const { return read_vector_; } Context* GetContext() { return context_; } uint32_t GetIndex() const { return index_; } @@ -133,6 +134,7 @@ class Group { info_vector_t info_vector_; pkt_vector_t start_vector_; pkt_vector_t stop_vector_; + pkt_vector_t read_vector_; uint32_t n_profiles_; uint32_t refs_; Context* const context_; @@ -286,6 +288,9 @@ class Context { const pkt_vector_t& StopPackets(const uint32_t& group_index) const { return set_[group_index].GetStopVector(); } + const pkt_vector_t& ReadPackets(const uint32_t& group_index) const { + return set_[group_index].GetReadVector(); + } void Start(const uint32_t& group_index, Queue* const queue = NULL) { const pkt_vector_t& start_packets = StartPackets(group_index); @@ -297,6 +302,11 @@ class Context { Queue* const submit_queue = (queue != NULL) ? queue : queue_; submit_queue->Submit(&stop_packets[0], stop_packets.size()); } + void Read(const uint32_t& group_index, Queue* const queue = NULL) { + const pkt_vector_t& read_packets = StopPackets(group_index); + Queue* const submit_queue = (queue != NULL) ? queue : queue_; + submit_queue->Submit(&read_packets[0], read_packets.size()); + } void Submit(const uint32_t& group_index, const packet_t* packet, Queue* const queue = NULL) { Queue* const submit_queue = (queue != NULL) ? queue : queue_; Start(group_index, submit_queue); diff --git a/projects/rocprofiler/src/core/profile.h b/projects/rocprofiler/src/core/profile.h index 6002ba2ef3..53705a9fa3 100644 --- a/projects/rocprofiler/src/core/profile.h +++ b/projects/rocprofiler/src/core/profile.h @@ -10,6 +10,10 @@ #include "util/exception.h" #include "util/hsa_rsrc_factory.h" +#ifndef AQL_PROFILE_READ_API_ENABLE +#define AQL_PROFILE_READ_API_ENABLE 0 +#endif + namespace rocprofiler { struct profile_info_t { const event_t* event; @@ -96,7 +100,7 @@ class Profile { virtual void Insert(const profile_info_t& info) { info_vector_.push_back(info.rinfo); } - hsa_status_t Finalize(pkt_vector_t& start_vector, pkt_vector_t& stop_vector) { + hsa_status_t Finalize(pkt_vector_t& start_vector, pkt_vector_t& stop_vector, pkt_vector_t& read_vector) { hsa_status_t status = HSA_STATUS_SUCCESS; if (!info_vector_.empty()) { @@ -104,6 +108,7 @@ class Profile { const pfn_t* api = rsrc->AqlProfileApi(); packet_t start{}; packet_t stop{}; + packet_t read{}; // Check the profile buffer sizes status = api->hsa_ven_amd_aqlprofile_start(&profile_, NULL); @@ -114,6 +119,12 @@ class Profile { if (status != HSA_STATUS_SUCCESS) AQL_EXC_RAISING(status, "aqlprofile_start"); status = api->hsa_ven_amd_aqlprofile_stop(&profile_, &stop); if (status != HSA_STATUS_SUCCESS) AQL_EXC_RAISING(status, "aqlprofile_stop"); + +#if AQL_PROFILE_READ_API_ENABLE + status = api->hsa_ven_amd_aqlprofile_read(&profile_, &read); + if (status != HSA_STATUS_SUCCESS) AQL_EXC_RAISING(status, "aqlprofile_read"); +#endif // AQL_PROFILE_READ_API_ENABLE + // Set completion signals hsa_signal_t dummy_signal{}; dummy_signal.handle = 0; @@ -122,6 +133,7 @@ class Profile { status = hsa_signal_create(1, 0, NULL, &post_signal); if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "signal_create " << std::hex << status); stop.completion_signal = post_signal; + read.completion_signal = post_signal; completion_signal_ = post_signal; if (is_legacy_) { @@ -130,6 +142,7 @@ class Profile { start_vector.insert(start_vector.end(), LEGACY_SLOT_SIZE_PKT, packet_t{}); stop_vector.insert(stop_vector.end(), LEGACY_SLOT_SIZE_PKT, packet_t{}); + status = api->hsa_ven_amd_aqlprofile_legacy_get_pm4( &start, reinterpret_cast(&start_vector[start_index])); if (status != HSA_STATUS_SUCCESS) @@ -138,9 +151,19 @@ class Profile { &stop, reinterpret_cast(&stop_vector[stop_index])); if (status != HSA_STATUS_SUCCESS) AQL_EXC_RAISING(status, "hsa_ven_amd_aqlprofile_legacy_get_pm4"); + +#if AQL_PROFILE_READ_API_ENABLE + const uint32_t read_index = read_vector.size(); + read_vector.insert(read_vector.end(), LEGACY_SLOT_SIZE_PKT, packet_t{}); + status = api->hsa_ven_amd_aqlprofile_legacy_get_pm4( + &read, reinterpret_cast(&read_vector[read_index])); + if (status != HSA_STATUS_SUCCESS) + AQL_EXC_RAISING(status, "hsa_ven_amd_aqlprofile_legacy_get_pm4"); +#endif // AQL_PROFILE_READ_API_ENABLE } else { start_vector.push_back(start); stop_vector.push_back(stop); + read_vector.push_back(read); } } diff --git a/projects/rocprofiler/src/core/rocprofiler.cpp b/projects/rocprofiler/src/core/rocprofiler.cpp index 9a678a776e..0a627e4826 100644 --- a/projects/rocprofiler/src/core/rocprofiler.cpp +++ b/projects/rocprofiler/src/core/rocprofiler.cpp @@ -241,6 +241,14 @@ PUBLIC_API hsa_status_t rocprofiler_stop(rocprofiler_t* handle, uint32_t group_i API_METHOD_SUFFIX } +// Read profiling +PUBLIC_API hsa_status_t rocprofiler_read(rocprofiler_t* handle, uint32_t group_index) { + API_METHOD_PREFIX + rocprofiler::Context* context = reinterpret_cast(handle); + context->Read(group_index); + API_METHOD_SUFFIX +} + // Get profiling data PUBLIC_API hsa_status_t rocprofiler_get_data(rocprofiler_t* handle, uint32_t group_index) { API_METHOD_PREFIX @@ -263,6 +271,13 @@ PUBLIC_API hsa_status_t rocprofiler_group_stop(rocprofiler_group_t* group) { API_METHOD_SUFFIX } +// Read profiling +PUBLIC_API hsa_status_t rocprofiler_group_read(rocprofiler_group_t* group) { + API_METHOD_PREFIX + rocprofiler_read(group->context, group->index); + API_METHOD_SUFFIX +} + // Get profiling data PUBLIC_API hsa_status_t rocprofiler_group_get_data(rocprofiler_group_t* group) { API_METHOD_PREFIX diff --git a/projects/rocprofiler/src/util/hsa_rsrc_factory.cpp b/projects/rocprofiler/src/util/hsa_rsrc_factory.cpp index f513d7f976..57cdac4b1b 100644 --- a/projects/rocprofiler/src/util/hsa_rsrc_factory.cpp +++ b/projects/rocprofiler/src/util/hsa_rsrc_factory.cpp @@ -107,8 +107,8 @@ HsaRsrcFactory::HsaRsrcFactory() { // Destructor of the class HsaRsrcFactory::~HsaRsrcFactory() { - for (auto p : cpu_list_) free(const_cast(p)); - for (auto p : gpu_list_) free(const_cast(p)); + for (auto p : cpu_list_) delete p; + for (auto p : gpu_list_) delete p; printf("HSA shutdown\n"); hsa_status_t status = hsa_shut_down(); diff --git a/projects/rocprofiler/test/util/hsa_rsrc_factory.cpp b/projects/rocprofiler/test/util/hsa_rsrc_factory.cpp index c32e33c561..ffe43a4296 100644 --- a/projects/rocprofiler/test/util/hsa_rsrc_factory.cpp +++ b/projects/rocprofiler/test/util/hsa_rsrc_factory.cpp @@ -104,8 +104,8 @@ HsaRsrcFactory::HsaRsrcFactory() { // Destructor of the class HsaRsrcFactory::~HsaRsrcFactory() { - for (auto p : cpu_list_) free(const_cast(p)); - for (auto p : gpu_list_) free(const_cast(p)); + for (auto p : cpu_list_) delete p; + for (auto p : gpu_list_) delete p; printf("HSA shutdown\n"); hsa_status_t status = hsa_shut_down();