From e21609c80e965ef17f808880ab8d4ca190830bc0 Mon Sep 17 00:00:00 2001 From: Benjamin Welton Date: Mon, 29 Apr 2024 15:48:50 -0700 Subject: [PATCH] Code migration and include fixes (#830) * Migrating code/reformatting * Remove unnecessary includes --------- Co-authored-by: Benjamin Welton --- .../rocprofiler-sdk/aql/packet_construct.cpp | 3 +- .../lib/rocprofiler-sdk/context/context.hpp | 21 +- .../rocprofiler-sdk/counters/CMakeLists.txt | 4 +- .../rocprofiler-sdk/counters/controller.cpp | 145 +++++++ .../rocprofiler-sdk/counters/controller.hpp | 107 ++++++ source/lib/rocprofiler-sdk/counters/core.cpp | 358 +----------------- source/lib/rocprofiler-sdk/counters/core.hpp | 59 +-- .../rocprofiler-sdk/counters/dimensions.cpp | 6 +- .../rocprofiler-sdk/counters/dimensions.hpp | 2 - .../counters/dispatch_handlers.cpp | 283 ++++++++++++++ .../counters/dispatch_handlers.hpp | 57 +++ .../rocprofiler-sdk/counters/evaluate_ast.cpp | 1 - .../rocprofiler-sdk/counters/evaluate_ast.hpp | 3 - .../rocprofiler-sdk/counters/id_decode.cpp | 1 - .../rocprofiler-sdk/counters/id_decode.hpp | 1 - .../lib/rocprofiler-sdk/counters/metrics.cpp | 5 - .../lib/rocprofiler-sdk/counters/metrics.hpp | 6 +- .../rocprofiler-sdk/counters/tests/core.cpp | 8 +- .../counters/tests/dimension.cpp | 5 - .../counters/tests/evaluate_ast_test.cpp | 2 +- .../counters/tests/evaluate_ast_test.hpp | 2 - .../counters/tests/init_order.cpp | 4 - .../counters/tests/metrics_test.cpp | 1 + .../lib/rocprofiler-sdk/hsa/agent_cache.cpp | 9 +- .../lib/rocprofiler-sdk/hsa/agent_cache.hpp | 7 - source/lib/rocprofiler-sdk/hsa/aql_packet.cpp | 1 - source/lib/rocprofiler-sdk/hsa/aql_packet.hpp | 1 - source/lib/rocprofiler-sdk/hsa/async_copy.cpp | 4 - source/lib/rocprofiler-sdk/hsa/hsa.def.cpp | 1 - .../lib/rocprofiler-sdk/hsa/pc_sampling.cpp | 6 - .../hsa/profile_serializer.hpp | 5 +- source/lib/rocprofiler-sdk/hsa/queue.cpp | 4 - source/lib/rocprofiler-sdk/hsa/queue.hpp | 4 - .../rocprofiler-sdk/hsa/scratch_memory.cpp | 1 - source/lib/rocprofiler-sdk/hsa/utils.hpp | 4 - 35 files changed, 622 insertions(+), 509 deletions(-) create mode 100644 source/lib/rocprofiler-sdk/counters/controller.cpp create mode 100644 source/lib/rocprofiler-sdk/counters/controller.hpp create mode 100644 source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp create mode 100644 source/lib/rocprofiler-sdk/counters/dispatch_handlers.hpp diff --git a/source/lib/rocprofiler-sdk/aql/packet_construct.cpp b/source/lib/rocprofiler-sdk/aql/packet_construct.cpp index c76db36d93..9b149a8f0d 100644 --- a/source/lib/rocprofiler-sdk/aql/packet_construct.cpp +++ b/source/lib/rocprofiler-sdk/aql/packet_construct.cpp @@ -21,6 +21,7 @@ // SOFTWARE. #include "lib/rocprofiler-sdk/aql/packet_construct.hpp" +#include "lib/common/logging.hpp" #include "lib/rocprofiler-sdk/hsa/details/fmt.hpp" #include @@ -32,7 +33,7 @@ auto status = (fn); \ if(status != HSA_STATUS_SUCCESS) \ { \ - std::cerr << "HSA Err: " << status << "\n"; \ + ROCP_FATAL << "HSA Err: " << status << "\n"; \ exit(1); \ } \ } diff --git a/source/lib/rocprofiler-sdk/context/context.hpp b/source/lib/rocprofiler-sdk/context/context.hpp index 6152185182..81c2e813c6 100644 --- a/source/lib/rocprofiler-sdk/context/context.hpp +++ b/source/lib/rocprofiler-sdk/context/context.hpp @@ -27,9 +27,7 @@ #include #include "lib/common/container/small_vector.hpp" -#include "lib/common/container/stable_vector.hpp" #include "lib/common/synchronized.hpp" -#include "lib/rocprofiler-sdk/allocator.hpp" #include "lib/rocprofiler-sdk/context/correlation_id.hpp" #include "lib/rocprofiler-sdk/context/domain.hpp" #include "lib/rocprofiler-sdk/counters/core.hpp" @@ -37,7 +35,6 @@ #include "lib/rocprofiler-sdk/thread_trace/att_core.hpp" #include -#include #include #include #include @@ -74,7 +71,7 @@ struct buffer_tracing_service buffer_array_t buffer_data = {}; }; -struct counter_collection_service +struct dispatch_counter_collection_service { // Contains a vector of counter collection instances associated with this context. // Each instance is assocated with an agent and a counter collection profile. @@ -90,14 +87,14 @@ struct counter_collection_service struct context { // size is used to ensure that we never read past the end of the version - size_t size = 0; - uint64_t context_idx = 0; // context id - uint32_t client_idx = 0; // tool id - correlation_tracing_service correlation_tracer = {}; - std::unique_ptr callback_tracer = {}; - std::unique_ptr buffered_tracer = {}; - std::unique_ptr counter_collection = {}; - std::shared_ptr thread_trace = {}; + size_t size = 0; + uint64_t context_idx = 0; // context id + uint32_t client_idx = 0; // tool id + correlation_tracing_service correlation_tracer = {}; + std::unique_ptr callback_tracer = {}; + std::unique_ptr buffered_tracer = {}; + std::unique_ptr counter_collection = {}; + std::shared_ptr thread_trace = {}; }; // set the client index needs to be called before allocate_context() diff --git a/source/lib/rocprofiler-sdk/counters/CMakeLists.txt b/source/lib/rocprofiler-sdk/counters/CMakeLists.txt index b8858a5142..30d9abb67b 100644 --- a/source/lib/rocprofiler-sdk/counters/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk/counters/CMakeLists.txt @@ -1,7 +1,7 @@ set(ROCPROFILER_LIB_COUNTERS_SOURCES metrics.cpp dimensions.cpp evaluate_ast.cpp core.cpp - id_decode.cpp) + id_decode.cpp dispatch_handlers.cpp controller.cpp) set(ROCPROFILER_LIB_COUNTERS_HEADERS metrics.hpp dimensions.hpp evaluate_ast.hpp core.hpp - id_decode.hpp) + id_decode.hpp dispatch_handlers.hpp controller.hpp) target_sources(rocprofiler-object-library PRIVATE ${ROCPROFILER_LIB_COUNTERS_SOURCES} ${ROCPROFILER_LIB_COUNTERS_HEADERS}) diff --git a/source/lib/rocprofiler-sdk/counters/controller.cpp b/source/lib/rocprofiler-sdk/counters/controller.cpp new file mode 100644 index 0000000000..a162b1eca2 --- /dev/null +++ b/source/lib/rocprofiler-sdk/counters/controller.cpp @@ -0,0 +1,145 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include "lib/rocprofiler-sdk/counters/controller.hpp" + +#include +#include +#include +#include + +#include "lib/rocprofiler-sdk/context/context.hpp" + +namespace rocprofiler +{ +namespace counters +{ +CounterController::CounterController() +{ + // Pre-read metrics map file to catch faliures during initial setup. + rocprofiler::counters::getMetricIdMap(); +} + +// Adds a counter collection profile to our global cache. +// Note: these profiles can be used across multiple contexts +// and are independent of the context. +uint64_t +CounterController::add_profile(std::shared_ptr&& config) +{ + static std::atomic profile_val = 1; + uint64_t ret = 0; + _configs.wlock([&](auto& data) { + config->id = rocprofiler_profile_config_id_t{.handle = profile_val}; + data.emplace(profile_val, std::move(config)); + ret = profile_val; + profile_val++; + }); + return ret; +} + +void +CounterController::destroy_profile(uint64_t id) +{ + _configs.wlock([&](auto& data) { data.erase(id); }); +} + +// Setup the counter collection service. counter_callback_info is created here +// to contain the counters that need to be collected (specified in profile_id) and +// the AQL packet generator for injecting packets. Note: the service is created +// in the stop state. +bool +CounterController::configure_dispatch( + rocprofiler_context_id_t context_id, + rocprofiler_buffer_id_t buffer, + rocprofiler_profile_counting_dispatch_callback_t callback, + void* callback_args, + rocprofiler_profile_counting_record_callback_t record_callback, + void* record_callback_args) +{ + auto* ctx_p = rocprofiler::context::get_mutable_registered_context(context_id); + if(!ctx_p) return false; + + auto& ctx = *ctx_p; + + if(!ctx.counter_collection) + { + ctx.counter_collection = + std::make_unique(); + } + + auto& cb = + *ctx.counter_collection->callbacks.emplace_back(std::make_shared()); + + cb.user_cb = callback; + cb.callback_args = callback_args; + cb.context = context_id; + if(buffer.handle != 0) + { + cb.buffer = buffer; + } + cb.internal_context = ctx_p; + cb.record_callback = record_callback; + cb.record_callback_args = record_callback_args; + + return true; +} + +std::shared_ptr +CounterController::get_profile_cfg(rocprofiler_profile_config_id_t id) +{ + std::shared_ptr cfg; + _configs.rlock([&](const auto& map) { cfg = map.at(id.handle); }); + return cfg; +} + +CounterController& +get_controller() +{ + static CounterController controller; + return controller; +} + +uint64_t +create_counter_profile(std::shared_ptr&& config) +{ + return get_controller().add_profile(std::move(config)); +} + +void +destroy_counter_profile(uint64_t id) +{ + get_controller().destroy_profile(id); +} + +std::shared_ptr +get_profile_config(rocprofiler_profile_config_id_t id) +{ + try + { + return get_controller().get_profile_cfg(id); + } catch(std::out_of_range&) + { + return nullptr; + } +} +} // namespace counters +} // namespace rocprofiler \ No newline at end of file diff --git a/source/lib/rocprofiler-sdk/counters/controller.hpp b/source/lib/rocprofiler-sdk/counters/controller.hpp new file mode 100644 index 0000000000..e8b496b4ce --- /dev/null +++ b/source/lib/rocprofiler-sdk/counters/controller.hpp @@ -0,0 +1,107 @@ + +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#pragma once + +#include "lib/common/synchronized.hpp" +#include "lib/rocprofiler-sdk/aql/packet_construct.hpp" +#include "lib/rocprofiler-sdk/counters/evaluate_ast.hpp" +#include "lib/rocprofiler-sdk/counters/metrics.hpp" + +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace counters +{ +// Stores counter profiling information such as the agent +// to collect counters on, the metrics to collect, the hw +// counters needed to evaluate the metrics, and the ASTs. +// This profile can be shared among many rocprof contexts. +struct profile_config +{ + const rocprofiler_agent_t* agent = nullptr; + std::vector metrics{}; + // HW counters that must be collected to compute the above + // metrics (derived metrics are broken down into hw counters + // in this vector). + std::set reqired_hw_counters{}; + // Counters that are not hardware based but based on either a + // static value (such as those in agent) + std::set required_special_counters{}; + // ASTs to evaluate + std::vector asts{}; + rocprofiler_profile_config_id_t id{.handle = 0}; + // Packet generator to create AQL packets for insertion + std::unique_ptr pkt_generator{nullptr}; + // A packet cache of AQL packets. This allows reuse of AQL packets (preventing costly + // allocation of new packets/destruction). + rocprofiler::common::Synchronized>> + packets{}; +}; + +class CounterController +{ +public: + CounterController(); + + // Adds a counter collection profile to our global cache. + // Note: these profiles can be used across multiple contexts + // and are independent of the context. + uint64_t add_profile(std::shared_ptr&& config); + + void destroy_profile(uint64_t id); + // Setup the counter collection service. counter_callback_info is created here + // to contain the counters that need to be collected (specified in profile_id) and + // the AQL packet generator for injecting packets. Note: the service is created + // in the stop state. + static bool configure_dispatch(rocprofiler_context_id_t context_id, + rocprofiler_buffer_id_t buffer, + rocprofiler_profile_counting_dispatch_callback_t callback, + void* callback_args, + rocprofiler_profile_counting_record_callback_t record_callback, + void* record_callback_args); + std::shared_ptr get_profile_cfg(rocprofiler_profile_config_id_t id); + +private: + rocprofiler::common::Synchronized>> + _configs; +}; + +CounterController& +get_controller(); + +uint64_t +create_counter_profile(std::shared_ptr&& config); + +void +destroy_counter_profile(uint64_t id); + +std::shared_ptr +get_profile_config(rocprofiler_profile_config_id_t id); + +} // namespace counters +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/counters/core.cpp b/source/lib/rocprofiler-sdk/counters/core.cpp index 2e6f9414e0..452db26b93 100644 --- a/source/lib/rocprofiler-sdk/counters/core.cpp +++ b/source/lib/rocprofiler-sdk/counters/core.cpp @@ -25,13 +25,10 @@ #include "lib/common/container/small_vector.hpp" #include "lib/common/synchronized.hpp" #include "lib/common/utility.hpp" -#include "lib/rocprofiler-sdk/agent.hpp" -#include "lib/rocprofiler-sdk/aql/helpers.hpp" #include "lib/rocprofiler-sdk/aql/packet_construct.hpp" -#include "lib/rocprofiler-sdk/buffer.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" +#include "lib/rocprofiler-sdk/counters/dispatch_handlers.hpp" #include "lib/rocprofiler-sdk/hsa/queue_controller.hpp" -#include "lib/rocprofiler-sdk/registration.hpp" #include #include @@ -40,118 +37,6 @@ namespace rocprofiler { namespace counters { -class CounterController -{ -public: - CounterController() - { - // Pre-read metrics map file to catch faliures during initial setup. - rocprofiler::counters::getMetricIdMap(); - } - - // Adds a counter collection profile to our global cache. - // Note: these profiles can be used across multiple contexts - // and are independent of the context. - uint64_t add_profile(std::shared_ptr&& config) - { - static std::atomic profile_val = 1; - uint64_t ret = 0; - _configs.wlock([&](auto& data) { - config->id = rocprofiler_profile_config_id_t{.handle = profile_val}; - data.emplace(profile_val, std::move(config)); - ret = profile_val; - profile_val++; - }); - return ret; - } - - void destroy_profile(uint64_t id) - { - _configs.wlock([&](auto& data) { data.erase(id); }); - } - - // Setup the counter collection service. counter_callback_info is created here - // to contain the counters that need to be collected (specified in profile_id) and - // the AQL packet generator for injecting packets. Note: the service is created - // in the stop state. - static bool configure_dispatch(rocprofiler_context_id_t context_id, - rocprofiler_buffer_id_t buffer, - rocprofiler_profile_counting_dispatch_callback_t callback, - void* callback_args, - rocprofiler_profile_counting_record_callback_t record_callback, - void* record_callback_args) - { - auto* ctx_p = rocprofiler::context::get_mutable_registered_context(context_id); - if(!ctx_p) return false; - - auto& ctx = *ctx_p; - - if(!ctx.counter_collection) - { - ctx.counter_collection = - std::make_unique(); - } - - auto& cb = *ctx.counter_collection->callbacks.emplace_back( - std::make_shared()); - - cb.user_cb = callback; - cb.callback_args = callback_args; - cb.context = context_id; - if(buffer.handle != 0) - { - cb.buffer = buffer; - } - cb.internal_context = ctx_p; - cb.record_callback = record_callback; - cb.record_callback_args = record_callback_args; - - return true; - } - - std::shared_ptr get_profile_cfg(rocprofiler_profile_config_id_t id) - { - std::shared_ptr cfg; - _configs.rlock([&](const auto& map) { cfg = map.at(id.handle); }); - return cfg; - } - -private: - rocprofiler::common::Synchronized>> - _configs; -}; - -CounterController& -get_controller() -{ - static CounterController controller; - return controller; -} - -uint64_t -create_counter_profile(std::shared_ptr&& config) -{ - return get_controller().add_profile(std::move(config)); -} - -void -destroy_counter_profile(uint64_t id) -{ - get_controller().destroy_profile(id); -} - -std::shared_ptr -get_profile_config(rocprofiler_profile_config_id_t id) -{ - try - { - return get_controller().get_profile_cfg(id); - } catch(std::out_of_range&) - { - return nullptr; - } -} - rocprofiler_status_t counter_callback_info::setup_profile_config(const hsa::AgentCache& agent, std::shared_ptr& profile) @@ -256,247 +141,6 @@ counter_callback_info::get_packet(std::unique_ptr& return ROCPROFILER_STATUS_SUCCESS; } -/** - * Callback we get from HSA interceptor when a kernel packet is being enqueued. - * - * We return an AQLPacket containing the start/stop/read packets for injection. - */ -std::unique_ptr -queue_cb(const context::context* ctx, - const std::shared_ptr& info, - const hsa::Queue& queue, - const hsa::rocprofiler_packet& pkt, - rocprofiler_kernel_id_t kernel_id, - rocprofiler_dispatch_id_t dispatch_id, - rocprofiler_user_data_t* user_data, - const hsa::Queue::queue_info_session_t::external_corr_id_map_t& extern_corr_ids, - const context::correlation_id* correlation_id) -{ - CHECK(info && ctx); - - // Maybe adds serialization packets to the AQLPacket (if serializer is enabled) - // and maybe adds barrier packets if the state is transitioning from serialized <-> - // unserialized - auto maybe_add_serialization = [&](auto& gen_pkt) { - CHECK_NOTNULL(hsa::get_queue_controller())->serializer().rlock([&](const auto& serializer) { - for(auto& s_pkt : serializer.kernel_dispatch(queue)) - { - gen_pkt->before_krn_pkt.push_back(s_pkt.ext_amd_aql_pm4); - } - }); - }; - - // Packet generated when no instrumentation is performed. May contain serialization - // packets/barrier packets (and can be empty). - auto no_instrumentation = [&]() { - auto ret_pkt = std::make_unique(nullptr); - // If we have a counter collection context but it is not enabled, we still might need - // to add barrier packets to transition from serialized -> unserialized execution. This - // transition is coordinated by the serializer. - maybe_add_serialization(ret_pkt); - info->packet_return_map.wlock([&](auto& data) { data.emplace(ret_pkt.get(), nullptr); }); - return ret_pkt; - }; - - if(!ctx || !ctx->counter_collection) return nullptr; - - bool is_enabled = false; - - ctx->counter_collection->enabled.rlock( - [&](const auto& collect_ctx) { is_enabled = collect_ctx; }); - - if(!is_enabled || !info->user_cb) - { - return no_instrumentation(); - } - - auto _corr_id_v = - rocprofiler_correlation_id_t{.internal = 0, .external = context::null_user_data}; - if(const auto* _corr_id = correlation_id) - { - _corr_id_v.internal = _corr_id->internal; - if(const auto* external = - rocprofiler::common::get_val(extern_corr_ids, info->internal_context)) - { - _corr_id_v.external = *external; - } - } - - auto req_profile = rocprofiler_profile_config_id_t{.handle = 0}; - auto dispatch_data = - common::init_public_api_struct(rocprofiler_profile_counting_dispatch_data_t{}); - - dispatch_data.correlation_id = _corr_id_v; - { - auto dispatch_info = common::init_public_api_struct(rocprofiler_kernel_dispatch_info_t{}); - dispatch_info.kernel_id = kernel_id; - dispatch_info.dispatch_id = dispatch_id; - dispatch_info.agent_id = CHECK_NOTNULL(queue.get_agent().get_rocp_agent())->id; - dispatch_info.queue_id = queue.get_id(); - dispatch_info.private_segment_size = pkt.kernel_dispatch.private_segment_size; - dispatch_info.group_segment_size = pkt.kernel_dispatch.group_segment_size; - dispatch_info.workgroup_size = {pkt.kernel_dispatch.workgroup_size_x, - pkt.kernel_dispatch.workgroup_size_y, - pkt.kernel_dispatch.workgroup_size_z}; - dispatch_info.grid_size = {pkt.kernel_dispatch.grid_size_x, - pkt.kernel_dispatch.grid_size_y, - pkt.kernel_dispatch.grid_size_z}; - dispatch_data.dispatch_info = dispatch_info; - } - - info->user_cb(dispatch_data, &req_profile, user_data, info->callback_args); - - if(req_profile.handle == 0) - { - return no_instrumentation(); - } - - auto prof_config = get_controller().get_profile_cfg(req_profile); - CHECK(prof_config); - - std::unique_ptr ret_pkt; - auto status = info->get_packet(ret_pkt, queue.get_agent(), prof_config); - CHECK_EQ(status, ROCPROFILER_STATUS_SUCCESS) << rocprofiler_get_status_string(status); - - maybe_add_serialization(ret_pkt); - if(ret_pkt->empty) - { - return ret_pkt; - } - - ret_pkt->before_krn_pkt.push_back(ret_pkt->start); - ret_pkt->after_krn_pkt.push_back(ret_pkt->stop); - ret_pkt->after_krn_pkt.push_back(ret_pkt->read); - for(auto& aql_pkt : ret_pkt->after_krn_pkt) - { - aql_pkt.completion_signal.handle = 0; - } - - return ret_pkt; -} - -/** - * Callback called by HSA interceptor when the kernel has completed processing. - */ -void -completed_cb(const context::context* ctx, - const std::shared_ptr& info, - const hsa::Queue& /*queue*/, - hsa::rocprofiler_packet, - const hsa::Queue::queue_info_session_t& session, - inst_pkt_t& pkts) -{ - CHECK(info && ctx); - - std::shared_ptr prof_config; - // Get the Profile Config - std::unique_ptr pkt = nullptr; - info->packet_return_map.wlock([&](auto& data) { - for(auto& [aql_pkt, _] : pkts) - { - const auto* profile = rocprofiler::common::get_val(data, aql_pkt.get()); - if(profile) - { - prof_config = *profile; - data.erase(aql_pkt.get()); - pkt = std::move(aql_pkt); - return; - } - } - }); - - if(!pkt) return; - - CHECK_NOTNULL(hsa::get_queue_controller())->serializer().wlock([&](auto& serializer) { - serializer.kernel_completion_signal(session.queue); - }); - - // We have no profile config, nothing to output. - if(!prof_config) return; - - auto decoded_pkt = EvaluateAST::read_pkt(prof_config->pkt_generator.get(), *pkt); - EvaluateAST::read_special_counters( - *prof_config->agent, prof_config->required_special_counters, decoded_pkt); - - prof_config->packets.wlock([&](auto& pkt_vector) { - if(pkt) - { - pkt_vector.emplace_back(std::move(pkt)); - } - }); - - common::container::small_vector out; - rocprofiler::buffer::instance* buf = nullptr; - - if(info->buffer) - { - buf = CHECK_NOTNULL(buffer::get_buffer(info->buffer->handle)); - } - - auto _corr_id_v = - rocprofiler_correlation_id_t{.internal = 0, .external = context::null_user_data}; - if(const auto* _corr_id = session.correlation_id) - { - _corr_id_v.internal = _corr_id->internal; - if(const auto* external = rocprofiler::common::get_val( - session.tracing_data.external_correlation_ids, info->internal_context)) - { - _corr_id_v.external = *external; - } - } - - auto _dispatch_id = session.callback_record.dispatch_info.dispatch_id; - for(auto& ast : prof_config->asts) - { - std::vector>> cache; - auto* ret = ast.evaluate(decoded_pkt, cache); - CHECK(ret); - ast.set_out_id(*ret); - - out.reserve(out.size() + ret->size()); - for(auto& val : *ret) - { - val.dispatch_id = _dispatch_id; - out.emplace_back(val); - } - } - - if(!out.empty()) - { - if(buf) - { - auto _header = - common::init_public_api_struct(rocprofiler_profile_counting_dispatch_record_t{}); - _header.num_records = out.size(); - _header.correlation_id = _corr_id_v; - _header.dispatch_info = session.callback_record.dispatch_info; - buf->emplace(ROCPROFILER_BUFFER_CATEGORY_COUNTERS, - ROCPROFILER_COUNTER_RECORD_PROFILE_COUNTING_DISPATCH_HEADER, - _header); - - for(auto itr : out) - buf->emplace( - ROCPROFILER_BUFFER_CATEGORY_COUNTERS, ROCPROFILER_COUNTER_RECORD_VALUE, itr); - } - else - { - CHECK(info->record_callback); - - auto dispatch_data = - common::init_public_api_struct(rocprofiler_profile_counting_dispatch_data_t{}); - - dispatch_data.dispatch_info = session.callback_record.dispatch_info; - dispatch_data.correlation_id = _corr_id_v; - - info->record_callback(dispatch_data, - out.data(), - out.size(), - session.user_data, - info->record_callback_args); - } - } -} - void start_context(const context::context* ctx) { diff --git a/source/lib/rocprofiler-sdk/counters/core.hpp b/source/lib/rocprofiler-sdk/counters/core.hpp index e5b07f9f43..d86903efad 100644 --- a/source/lib/rocprofiler-sdk/counters/core.hpp +++ b/source/lib/rocprofiler-sdk/counters/core.hpp @@ -25,12 +25,12 @@ #include #include #include +#include -#include "lib/rocprofiler-sdk/aql/helpers.hpp" -#include "lib/rocprofiler-sdk/aql/packet_construct.hpp" -#include "lib/rocprofiler-sdk/counters/evaluate_ast.hpp" -#include "lib/rocprofiler-sdk/counters/metrics.hpp" +#include "lib/common/synchronized.hpp" +#include "lib/rocprofiler-sdk/counters/controller.hpp" #include "lib/rocprofiler-sdk/hsa/agent_cache.hpp" +#include "lib/rocprofiler-sdk/hsa/rocprofiler_packet.hpp" namespace rocprofiler { @@ -40,32 +40,6 @@ struct context; } namespace counters { -// Stores counter profiling information such as the agent -// to collect counters on, the metrics to collect, the hw -// counters needed to evaluate the metrics, and the ASTs. -// This profile can be shared among many rocprof contexts. -struct profile_config -{ - const rocprofiler_agent_t* agent = nullptr; - std::vector metrics{}; - // HW counters that must be collected to compute the above - // metrics (derived metrics are broken down into hw counters - // in this vector). - std::set reqired_hw_counters{}; - // Counters that are not hardware based but based on either a - // static value (such as those in agent) - std::set required_special_counters{}; - // ASTs to evaluate - std::vector asts{}; - rocprofiler_profile_config_id_t id{.handle = 0}; - // Packet generator to create AQL packets for insertion - std::unique_ptr pkt_generator{nullptr}; - // A packet cache of AQL packets. This allows reuse of AQL packets (preventing costly - // allocation of new packets/destruction). - rocprofiler::common::Synchronized>> - packets{}; -}; - // Internal counter struct that stores the state needed to handle an intercepted // HSA kernel packet. struct counter_callback_info @@ -124,30 +98,5 @@ start_context(const context::context*); void stop_context(const context::context*); - -std::unique_ptr -queue_cb(const context::context* ctx, - const std::shared_ptr& info, - const hsa::Queue& queue, - const hsa::rocprofiler_packet& pkt, - rocprofiler_kernel_id_t kernel_id, - rocprofiler_dispatch_id_t dispatch_id, - rocprofiler_user_data_t* user_data, - const hsa::Queue::queue_info_session_t::external_corr_id_map_t& extern_corr_ids, - const context::correlation_id* correlation_id); - -using ClientID = int64_t; -using inst_pkt_t = common::container:: - small_vector, ClientID>, 4>; - -void -completed_cb(const context::context* ctx, - const std::shared_ptr& info, - const hsa::Queue& queue, - hsa::rocprofiler_packet, - const hsa::Queue::queue_info_session_t& session, - inst_pkt_t& pkts); - -std::shared_ptr get_profile_config(rocprofiler_profile_config_id_t); } // namespace counters } // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/counters/dimensions.cpp b/source/lib/rocprofiler-sdk/counters/dimensions.cpp index d367219e28..1ee3c7fb44 100644 --- a/source/lib/rocprofiler-sdk/counters/dimensions.cpp +++ b/source/lib/rocprofiler-sdk/counters/dimensions.cpp @@ -23,19 +23,19 @@ #include "dimensions.hpp" #include -#include #include #include #include +#include +#include + #include "lib/common/static_object.hpp" -#include "lib/common/synchronized.hpp" #include "lib/common/utility.hpp" #include "lib/rocprofiler-sdk/aql/helpers.hpp" #include "lib/rocprofiler-sdk/aql/packet_construct.hpp" #include "lib/rocprofiler-sdk/counters/evaluate_ast.hpp" -#include "lib/rocprofiler-sdk/hsa/queue_controller.hpp" namespace rocprofiler { diff --git a/source/lib/rocprofiler-sdk/counters/dimensions.hpp b/source/lib/rocprofiler-sdk/counters/dimensions.hpp index c7b437d0d8..ee88d03225 100644 --- a/source/lib/rocprofiler-sdk/counters/dimensions.hpp +++ b/source/lib/rocprofiler-sdk/counters/dimensions.hpp @@ -22,9 +22,7 @@ #pragma once -#include #include -#include #include #include #include diff --git a/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp new file mode 100644 index 0000000000..0dfcdaca40 --- /dev/null +++ b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp @@ -0,0 +1,283 @@ + + +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include "lib/rocprofiler-sdk/counters/dispatch_handlers.hpp" + +#include "lib/common/container/small_vector.hpp" +#include "lib/common/synchronized.hpp" +#include "lib/common/utility.hpp" +#include "lib/rocprofiler-sdk/buffer.hpp" +#include "lib/rocprofiler-sdk/context/context.hpp" +#include "lib/rocprofiler-sdk/counters/core.hpp" +#include "lib/rocprofiler-sdk/hsa/queue_controller.hpp" + +#include +#include + +namespace rocprofiler +{ +namespace counters +{ +/** + * Callback we get from HSA interceptor when a kernel packet is being enqueued. + * + * We return an AQLPacket containing the start/stop/read packets for injection. + */ +std::unique_ptr +queue_cb(const context::context* ctx, + const std::shared_ptr& info, + const hsa::Queue& queue, + const hsa::rocprofiler_packet& pkt, + rocprofiler_kernel_id_t kernel_id, + rocprofiler_dispatch_id_t dispatch_id, + rocprofiler_user_data_t* user_data, + const hsa::Queue::queue_info_session_t::external_corr_id_map_t& extern_corr_ids, + const context::correlation_id* correlation_id) +{ + CHECK(info && ctx); + + // Maybe adds serialization packets to the AQLPacket (if serializer is enabled) + // and maybe adds barrier packets if the state is transitioning from serialized <-> + // unserialized + auto maybe_add_serialization = [&](auto& gen_pkt) { + CHECK_NOTNULL(hsa::get_queue_controller())->serializer().rlock([&](const auto& serializer) { + for(auto& s_pkt : serializer.kernel_dispatch(queue)) + { + gen_pkt->before_krn_pkt.push_back(s_pkt.ext_amd_aql_pm4); + } + }); + }; + + // Packet generated when no instrumentation is performed. May contain serialization + // packets/barrier packets (and can be empty). + auto no_instrumentation = [&]() { + auto ret_pkt = std::make_unique(nullptr); + // If we have a counter collection context but it is not enabled, we still might need + // to add barrier packets to transition from serialized -> unserialized execution. This + // transition is coordinated by the serializer. + maybe_add_serialization(ret_pkt); + info->packet_return_map.wlock([&](auto& data) { data.emplace(ret_pkt.get(), nullptr); }); + return ret_pkt; + }; + + if(!ctx || !ctx->counter_collection) return nullptr; + + bool is_enabled = false; + + ctx->counter_collection->enabled.rlock( + [&](const auto& collect_ctx) { is_enabled = collect_ctx; }); + + if(!is_enabled || !info->user_cb) + { + return no_instrumentation(); + } + + auto _corr_id_v = + rocprofiler_correlation_id_t{.internal = 0, .external = context::null_user_data}; + if(const auto* _corr_id = correlation_id) + { + _corr_id_v.internal = _corr_id->internal; + if(const auto* external = + rocprofiler::common::get_val(extern_corr_ids, info->internal_context)) + { + _corr_id_v.external = *external; + } + } + + auto req_profile = rocprofiler_profile_config_id_t{.handle = 0}; + auto dispatch_data = + common::init_public_api_struct(rocprofiler_profile_counting_dispatch_data_t{}); + + dispatch_data.correlation_id = _corr_id_v; + { + auto dispatch_info = common::init_public_api_struct(rocprofiler_kernel_dispatch_info_t{}); + dispatch_info.kernel_id = kernel_id; + dispatch_info.dispatch_id = dispatch_id; + dispatch_info.agent_id = CHECK_NOTNULL(queue.get_agent().get_rocp_agent())->id; + dispatch_info.queue_id = queue.get_id(); + dispatch_info.private_segment_size = pkt.kernel_dispatch.private_segment_size; + dispatch_info.group_segment_size = pkt.kernel_dispatch.group_segment_size; + dispatch_info.workgroup_size = {pkt.kernel_dispatch.workgroup_size_x, + pkt.kernel_dispatch.workgroup_size_y, + pkt.kernel_dispatch.workgroup_size_z}; + dispatch_info.grid_size = {pkt.kernel_dispatch.grid_size_x, + pkt.kernel_dispatch.grid_size_y, + pkt.kernel_dispatch.grid_size_z}; + dispatch_data.dispatch_info = dispatch_info; + } + + info->user_cb(dispatch_data, &req_profile, user_data, info->callback_args); + + if(req_profile.handle == 0) + { + return no_instrumentation(); + } + + auto prof_config = get_controller().get_profile_cfg(req_profile); + CHECK(prof_config); + + std::unique_ptr ret_pkt; + auto status = info->get_packet(ret_pkt, queue.get_agent(), prof_config); + CHECK_EQ(status, ROCPROFILER_STATUS_SUCCESS) << rocprofiler_get_status_string(status); + + maybe_add_serialization(ret_pkt); + if(ret_pkt->empty) + { + return ret_pkt; + } + + ret_pkt->before_krn_pkt.push_back(ret_pkt->start); + ret_pkt->after_krn_pkt.push_back(ret_pkt->stop); + ret_pkt->after_krn_pkt.push_back(ret_pkt->read); + for(auto& aql_pkt : ret_pkt->after_krn_pkt) + { + aql_pkt.completion_signal.handle = 0; + } + + return ret_pkt; +} + +/** + * Callback called by HSA interceptor when the kernel has completed processing. + */ +void +completed_cb(const context::context* ctx, + const std::shared_ptr& info, + const hsa::Queue& /*queue*/, + hsa::rocprofiler_packet, + const hsa::Queue::queue_info_session_t& session, + inst_pkt_t& pkts) +{ + CHECK(info && ctx); + + std::shared_ptr prof_config; + // Get the Profile Config + std::unique_ptr pkt = nullptr; + info->packet_return_map.wlock([&](auto& data) { + for(auto& [aql_pkt, _] : pkts) + { + const auto* profile = rocprofiler::common::get_val(data, aql_pkt.get()); + if(profile) + { + prof_config = *profile; + data.erase(aql_pkt.get()); + pkt = std::move(aql_pkt); + return; + } + } + }); + + if(!pkt) return; + + CHECK_NOTNULL(hsa::get_queue_controller())->serializer().wlock([&](auto& serializer) { + serializer.kernel_completion_signal(session.queue); + }); + + // We have no profile config, nothing to output. + if(!prof_config) return; + + auto decoded_pkt = EvaluateAST::read_pkt(prof_config->pkt_generator.get(), *pkt); + EvaluateAST::read_special_counters( + *prof_config->agent, prof_config->required_special_counters, decoded_pkt); + + prof_config->packets.wlock([&](auto& pkt_vector) { + if(pkt) + { + pkt_vector.emplace_back(std::move(pkt)); + } + }); + + common::container::small_vector out; + rocprofiler::buffer::instance* buf = nullptr; + + if(info->buffer) + { + buf = CHECK_NOTNULL(buffer::get_buffer(info->buffer->handle)); + } + + auto _corr_id_v = + rocprofiler_correlation_id_t{.internal = 0, .external = context::null_user_data}; + if(const auto* _corr_id = session.correlation_id) + { + _corr_id_v.internal = _corr_id->internal; + if(const auto* external = rocprofiler::common::get_val( + session.tracing_data.external_correlation_ids, info->internal_context)) + { + _corr_id_v.external = *external; + } + } + + auto _dispatch_id = session.callback_record.dispatch_info.dispatch_id; + for(auto& ast : prof_config->asts) + { + std::vector>> cache; + auto* ret = ast.evaluate(decoded_pkt, cache); + CHECK(ret); + ast.set_out_id(*ret); + + out.reserve(out.size() + ret->size()); + for(auto& val : *ret) + { + val.dispatch_id = _dispatch_id; + out.emplace_back(val); + } + } + + if(!out.empty()) + { + if(buf) + { + auto _header = + common::init_public_api_struct(rocprofiler_profile_counting_dispatch_record_t{}); + _header.num_records = out.size(); + _header.correlation_id = _corr_id_v; + _header.dispatch_info = session.callback_record.dispatch_info; + buf->emplace(ROCPROFILER_BUFFER_CATEGORY_COUNTERS, + ROCPROFILER_COUNTER_RECORD_PROFILE_COUNTING_DISPATCH_HEADER, + _header); + + for(auto itr : out) + buf->emplace( + ROCPROFILER_BUFFER_CATEGORY_COUNTERS, ROCPROFILER_COUNTER_RECORD_VALUE, itr); + } + else + { + CHECK(info->record_callback); + + auto dispatch_data = + common::init_public_api_struct(rocprofiler_profile_counting_dispatch_data_t{}); + + dispatch_data.dispatch_info = session.callback_record.dispatch_info; + dispatch_data.correlation_id = _corr_id_v; + + info->record_callback(dispatch_data, + out.data(), + out.size(), + session.user_data, + info->record_callback_args); + } + } +} +} // namespace counters +} // namespace rocprofiler \ No newline at end of file diff --git a/source/lib/rocprofiler-sdk/counters/dispatch_handlers.hpp b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.hpp new file mode 100644 index 0000000000..894091ddc6 --- /dev/null +++ b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.hpp @@ -0,0 +1,57 @@ + +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#pragma once + +#include "lib/rocprofiler-sdk/context/context.hpp" +#include "lib/rocprofiler-sdk/hsa/aql_packet.hpp" + +namespace rocprofiler +{ +namespace counters +{ +using ClientID = int64_t; +using inst_pkt_t = common::container:: + small_vector, ClientID>, 4>; + +void +completed_cb(const context::context* ctx, + const std::shared_ptr& info, + const hsa::Queue& /*queue*/, + hsa::rocprofiler_packet, + const hsa::Queue::queue_info_session_t& session, + inst_pkt_t& pkts); + +std::unique_ptr +queue_cb(const context::context* ctx, + const std::shared_ptr& info, + const hsa::Queue& queue, + const hsa::rocprofiler_packet& pkt, + rocprofiler_kernel_id_t kernel_id, + rocprofiler_dispatch_id_t dispatch_id, + rocprofiler_user_data_t* user_data, + const hsa::Queue::queue_info_session_t::external_corr_id_map_t& extern_corr_ids, + const context::correlation_id* correlation_id); + +} // namespace counters +} // namespace rocprofiler \ No newline at end of file diff --git a/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp b/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp index 42398529bc..8fdcf787ae 100644 --- a/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp +++ b/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp @@ -31,7 +31,6 @@ #include #include -#include "lib/common/synchronized.hpp" #include "lib/common/utility.hpp" #include "lib/rocprofiler-sdk/counters/dimensions.hpp" #include "lib/rocprofiler-sdk/counters/parser/reader.hpp" diff --git a/source/lib/rocprofiler-sdk/counters/evaluate_ast.hpp b/source/lib/rocprofiler-sdk/counters/evaluate_ast.hpp index 95ccf1ed0c..d574bccb09 100644 --- a/source/lib/rocprofiler-sdk/counters/evaluate_ast.hpp +++ b/source/lib/rocprofiler-sdk/counters/evaluate_ast.hpp @@ -22,12 +22,9 @@ #pragma once -#include -#include #include #include -#include "lib/common/utility.hpp" #include "lib/rocprofiler-sdk/aql/packet_construct.hpp" #include "lib/rocprofiler-sdk/counters/dimensions.hpp" #include "lib/rocprofiler-sdk/counters/metrics.hpp" diff --git a/source/lib/rocprofiler-sdk/counters/id_decode.cpp b/source/lib/rocprofiler-sdk/counters/id_decode.cpp index a26d6703e9..bd9bb02493 100644 --- a/source/lib/rocprofiler-sdk/counters/id_decode.cpp +++ b/source/lib/rocprofiler-sdk/counters/id_decode.cpp @@ -23,7 +23,6 @@ #include "lib/rocprofiler-sdk/counters/id_decode.hpp" #include -#include #include #include "lib/common/static_object.hpp" diff --git a/source/lib/rocprofiler-sdk/counters/id_decode.hpp b/source/lib/rocprofiler-sdk/counters/id_decode.hpp index 2c51ec1f20..d48abd5235 100644 --- a/source/lib/rocprofiler-sdk/counters/id_decode.hpp +++ b/source/lib/rocprofiler-sdk/counters/id_decode.hpp @@ -22,7 +22,6 @@ #pragma once -#include #include #include diff --git a/source/lib/rocprofiler-sdk/counters/metrics.cpp b/source/lib/rocprofiler-sdk/counters/metrics.cpp index 7c4bd43be8..3a1657cfc2 100644 --- a/source/lib/rocprofiler-sdk/counters/metrics.cpp +++ b/source/lib/rocprofiler-sdk/counters/metrics.cpp @@ -24,22 +24,17 @@ #include -#include "lib/common/defines.hpp" #include "lib/common/filesystem.hpp" #include "lib/common/static_object.hpp" -#include "lib/common/synchronized.hpp" #include "lib/common/utility.hpp" #include "lib/common/xml.hpp" #include "lib/rocprofiler-sdk/agent.hpp" -#include "dimensions.hpp" #include "glog/logging.h" #include // for dladdr -#include #include #include -#include namespace rocprofiler { diff --git a/source/lib/rocprofiler-sdk/counters/metrics.hpp b/source/lib/rocprofiler-sdk/counters/metrics.hpp index 2a7f85ce6c..b6c1e7add9 100644 --- a/source/lib/rocprofiler-sdk/counters/metrics.hpp +++ b/source/lib/rocprofiler-sdk/counters/metrics.hpp @@ -23,16 +23,14 @@ #pragma once #include -#include #include #include #include +#include +#include #include -#include "fmt/core.h" -#include "fmt/ranges.h" - namespace rocprofiler { namespace counters diff --git a/source/lib/rocprofiler-sdk/counters/tests/core.cpp b/source/lib/rocprofiler-sdk/counters/tests/core.cpp index 17d37ec3ca..694c79a578 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/core.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/core.cpp @@ -21,12 +21,10 @@ // SOFTWARE. #include "lib/rocprofiler-sdk/counters/core.hpp" -#include "lib/common/static_object.hpp" #include "lib/common/utility.hpp" #include "lib/rocprofiler-sdk/agent.hpp" -#include "lib/rocprofiler-sdk/buffer.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" -#include "lib/rocprofiler-sdk/counters/id_decode.hpp" +#include "lib/rocprofiler-sdk/counters/dispatch_handlers.hpp" #include "lib/rocprofiler-sdk/counters/metrics.hpp" #include "lib/rocprofiler-sdk/hsa/agent_cache.hpp" #include "lib/rocprofiler-sdk/hsa/queue.hpp" @@ -44,7 +42,6 @@ #include #include -#include #include #include #include @@ -427,7 +424,8 @@ TEST(core, check_callbacks) ROCPROFILER_CALL(rocprofiler_create_context(&get_client_ctx()), "context creation failed"); context::context ctx; - ctx.counter_collection = std::make_unique(); + ctx.counter_collection = + std::make_unique(); ctx.counter_collection->enabled.wlock([](auto& data) { data = true; }); ASSERT_TRUE(hsa::get_queue_controller() != nullptr); diff --git a/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp b/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp index dfd6bf7d0d..6cca31e169 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp @@ -20,18 +20,13 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE // SOFTWARE. -#include "lib/common/static_object.hpp" #include "lib/common/utility.hpp" #include "lib/rocprofiler-sdk/agent.hpp" #include "lib/rocprofiler-sdk/aql/packet_construct.hpp" -#include "lib/rocprofiler-sdk/buffer.hpp" -#include "lib/rocprofiler-sdk/context/context.hpp" -#include "lib/rocprofiler-sdk/counters/core.hpp" #include "lib/rocprofiler-sdk/counters/dimensions.hpp" #include "lib/rocprofiler-sdk/counters/id_decode.hpp" #include "lib/rocprofiler-sdk/counters/metrics.hpp" #include "lib/rocprofiler-sdk/hsa/agent_cache.hpp" -#include "lib/rocprofiler-sdk/hsa/queue.hpp" #include "lib/rocprofiler-sdk/hsa/queue_controller.hpp" #include "lib/rocprofiler-sdk/registration.hpp" diff --git a/source/lib/rocprofiler-sdk/counters/tests/evaluate_ast_test.cpp b/source/lib/rocprofiler-sdk/counters/tests/evaluate_ast_test.cpp index 0839dbf525..012e6d0554 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/evaluate_ast_test.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/evaluate_ast_test.cpp @@ -27,8 +27,8 @@ #include #include -#include "evaluate_ast_test.hpp" #include "lib/rocprofiler-sdk/agent.hpp" +#include "lib/rocprofiler-sdk/counters/evaluate_ast.hpp" #include "lib/rocprofiler-sdk/counters/parser/reader.hpp" namespace diff --git a/source/lib/rocprofiler-sdk/counters/tests/evaluate_ast_test.hpp b/source/lib/rocprofiler-sdk/counters/tests/evaluate_ast_test.hpp index e96d71aa32..5af84ed3ce 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/evaluate_ast_test.hpp +++ b/source/lib/rocprofiler-sdk/counters/tests/evaluate_ast_test.hpp @@ -23,9 +23,7 @@ #pragma once #include -#include #include -#include "lib/rocprofiler-sdk/counters/evaluate_ast.hpp" struct test_data { diff --git a/source/lib/rocprofiler-sdk/counters/tests/init_order.cpp b/source/lib/rocprofiler-sdk/counters/tests/init_order.cpp index 8539e024a1..06f297a219 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/init_order.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/init_order.cpp @@ -22,9 +22,7 @@ #include "lib/common/static_object.hpp" #include "lib/common/utility.hpp" -#include "lib/rocprofiler-sdk/buffer.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" -#include "lib/rocprofiler-sdk/counters/id_decode.hpp" #include "lib/rocprofiler-sdk/counters/metrics.hpp" #include "lib/rocprofiler-sdk/registration.hpp" @@ -35,10 +33,8 @@ #include #include -#include #include #include -#include using namespace rocprofiler::counters; diff --git a/source/lib/rocprofiler-sdk/counters/tests/metrics_test.cpp b/source/lib/rocprofiler-sdk/counters/tests/metrics_test.cpp index eb51cf8356..8ff888ac9d 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/metrics_test.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/metrics_test.cpp @@ -29,6 +29,7 @@ #include +#include "lib/common/logging.hpp" #include "lib/rocprofiler-sdk/agent.hpp" #include "lib/rocprofiler-sdk/counters/metrics.hpp" diff --git a/source/lib/rocprofiler-sdk/hsa/agent_cache.cpp b/source/lib/rocprofiler-sdk/hsa/agent_cache.cpp index 67d6f052b7..dd8af53426 100644 --- a/source/lib/rocprofiler-sdk/hsa/agent_cache.cpp +++ b/source/lib/rocprofiler-sdk/hsa/agent_cache.cpp @@ -22,17 +22,10 @@ #include "agent_cache.hpp" +#include #include -#include -#include -#include #include -#include "lib/common/defines.hpp" -#include "lib/common/filesystem.hpp" -#include "lib/common/synchronized.hpp" -#include "lib/common/utility.hpp" - namespace { // This function checks to see if the provided diff --git a/source/lib/rocprofiler-sdk/hsa/agent_cache.hpp b/source/lib/rocprofiler-sdk/hsa/agent_cache.hpp index 02d5cc182d..2ae4936fe1 100644 --- a/source/lib/rocprofiler-sdk/hsa/agent_cache.hpp +++ b/source/lib/rocprofiler-sdk/hsa/agent_cache.hpp @@ -26,16 +26,9 @@ #include #include -#include "fmt/core.h" -#include "fmt/ranges.h" - -#include #include -#include -#include #include -#include "lib/common/utility.hpp" // Construct const and non-const accessor functions #define CONST_NONCONST_ACCESSOR(RTYPE, NAME, VAL) \ diff --git a/source/lib/rocprofiler-sdk/hsa/aql_packet.cpp b/source/lib/rocprofiler-sdk/hsa/aql_packet.cpp index 2ff93eb961..d43a842cab 100644 --- a/source/lib/rocprofiler-sdk/hsa/aql_packet.cpp +++ b/source/lib/rocprofiler-sdk/hsa/aql_packet.cpp @@ -23,7 +23,6 @@ #include "lib/rocprofiler-sdk/hsa/aql_packet.hpp" #include #include -#include #define CHECK_HSA(fn, message) \ if((fn) != HSA_STATUS_SUCCESS) \ diff --git a/source/lib/rocprofiler-sdk/hsa/aql_packet.hpp b/source/lib/rocprofiler-sdk/hsa/aql_packet.hpp index cb50d44d85..3235a20158 100644 --- a/source/lib/rocprofiler-sdk/hsa/aql_packet.hpp +++ b/source/lib/rocprofiler-sdk/hsa/aql_packet.hpp @@ -27,7 +27,6 @@ #include #include -#include namespace rocprofiler { diff --git a/source/lib/rocprofiler-sdk/hsa/async_copy.cpp b/source/lib/rocprofiler-sdk/hsa/async_copy.cpp index dff29c68dd..f2afb233b4 100644 --- a/source/lib/rocprofiler-sdk/hsa/async_copy.cpp +++ b/source/lib/rocprofiler-sdk/hsa/async_copy.cpp @@ -21,16 +21,12 @@ // THE SOFTWARE. #include "lib/rocprofiler-sdk/hsa/async_copy.hpp" -#include "lib/common/defines.hpp" #include "lib/common/scope_destructor.hpp" #include "lib/common/static_object.hpp" #include "lib/common/utility.hpp" #include "lib/rocprofiler-sdk/agent.hpp" -#include "lib/rocprofiler-sdk/buffer.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" -#include "lib/rocprofiler-sdk/hsa/details/ostream.hpp" #include "lib/rocprofiler-sdk/hsa/hsa.hpp" -#include "lib/rocprofiler-sdk/hsa/utils.hpp" #include "lib/rocprofiler-sdk/registration.hpp" #include "lib/rocprofiler-sdk/tracing/fwd.hpp" #include "lib/rocprofiler-sdk/tracing/tracing.hpp" diff --git a/source/lib/rocprofiler-sdk/hsa/hsa.def.cpp b/source/lib/rocprofiler-sdk/hsa/hsa.def.cpp index 502b16874f..1314a9fd89 100644 --- a/source/lib/rocprofiler-sdk/hsa/hsa.def.cpp +++ b/source/lib/rocprofiler-sdk/hsa/hsa.def.cpp @@ -20,7 +20,6 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN // THE SOFTWARE. -#include "lib/common/defines.hpp" #include "lib/rocprofiler-sdk/hsa/defines.hpp" #include "lib/rocprofiler-sdk/hsa/hsa.hpp" diff --git a/source/lib/rocprofiler-sdk/hsa/pc_sampling.cpp b/source/lib/rocprofiler-sdk/hsa/pc_sampling.cpp index 7992edabed..c73802dcb8 100644 --- a/source/lib/rocprofiler-sdk/hsa/pc_sampling.cpp +++ b/source/lib/rocprofiler-sdk/hsa/pc_sampling.cpp @@ -24,13 +24,8 @@ #if ROCPROFILER_SDK_HSA_PC_SAMPLING > 0 -# include "lib/common/defines.hpp" -# include "lib/common/utility.hpp" -# include "lib/rocprofiler-sdk/buffer.hpp" -# include "lib/rocprofiler-sdk/context/context.hpp" # include "lib/rocprofiler-sdk/hsa/defines.hpp" # include "lib/rocprofiler-sdk/hsa/hsa.hpp" -# include "lib/rocprofiler-sdk/hsa/queue_controller.hpp" # include # include @@ -46,7 +41,6 @@ # include # include # include -# include # include # include diff --git a/source/lib/rocprofiler-sdk/hsa/profile_serializer.hpp b/source/lib/rocprofiler-sdk/hsa/profile_serializer.hpp index 20bb2e8c11..f0bb28c5a3 100644 --- a/source/lib/rocprofiler-sdk/hsa/profile_serializer.hpp +++ b/source/lib/rocprofiler-sdk/hsa/profile_serializer.hpp @@ -28,11 +28,8 @@ #include "lib/rocprofiler-sdk/hsa/hsa_barrier.hpp" #include "lib/rocprofiler-sdk/hsa/queue.hpp" -#include -#include -#include +#include #include -#include namespace rocprofiler { diff --git a/source/lib/rocprofiler-sdk/hsa/queue.cpp b/source/lib/rocprofiler-sdk/hsa/queue.cpp index e77e9e6c18..a28dab0487 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/source/lib/rocprofiler-sdk/hsa/queue.cpp @@ -23,8 +23,6 @@ #include "lib/rocprofiler-sdk/hsa/queue.hpp" #include "lib/common/scope_destructor.hpp" #include "lib/common/utility.hpp" -#include "lib/rocprofiler-sdk/agent.hpp" -#include "lib/rocprofiler-sdk/buffer.hpp" #include "lib/rocprofiler-sdk/code_object/code_object.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" #include "lib/rocprofiler-sdk/hsa/details/fmt.hpp" @@ -43,8 +41,6 @@ #include #include -#include -#include // static assert for rocprofiler_packet ABI compatibility static_assert(sizeof(hsa_ext_amd_aql_pm4_packet_t) == sizeof(hsa_kernel_dispatch_packet_t), diff --git a/source/lib/rocprofiler-sdk/hsa/queue.hpp b/source/lib/rocprofiler-sdk/hsa/queue.hpp index ee2aa72c55..67059fb016 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue.hpp +++ b/source/lib/rocprofiler-sdk/hsa/queue.hpp @@ -29,7 +29,6 @@ #include "lib/common/container/small_vector.hpp" #include "lib/common/synchronized.hpp" -#include "lib/common/utility.hpp" #include "lib/rocprofiler-sdk/hsa/agent_cache.hpp" #include "lib/rocprofiler-sdk/hsa/aql_packet.hpp" #include "lib/rocprofiler-sdk/hsa/queue_info_session.hpp" @@ -45,11 +44,8 @@ #include #include #include -#include #include -#include #include -#include namespace rocprofiler { diff --git a/source/lib/rocprofiler-sdk/hsa/scratch_memory.cpp b/source/lib/rocprofiler-sdk/hsa/scratch_memory.cpp index 60bd00725f..5829d5cb16 100644 --- a/source/lib/rocprofiler-sdk/hsa/scratch_memory.cpp +++ b/source/lib/rocprofiler-sdk/hsa/scratch_memory.cpp @@ -23,7 +23,6 @@ #include "lib/rocprofiler-sdk/hsa/scratch_memory.hpp" #include "lib/common/defines.hpp" #include "lib/common/utility.hpp" -#include "lib/rocprofiler-sdk/buffer.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" #include "lib/rocprofiler-sdk/hsa/defines.hpp" #include "lib/rocprofiler-sdk/hsa/hsa.hpp" diff --git a/source/lib/rocprofiler-sdk/hsa/utils.hpp b/source/lib/rocprofiler-sdk/hsa/utils.hpp index b86a3bb1cc..b6b4764132 100644 --- a/source/lib/rocprofiler-sdk/hsa/utils.hpp +++ b/source/lib/rocprofiler-sdk/hsa/utils.hpp @@ -24,7 +24,6 @@ #include -#include "lib/common/mpl.hpp" #include "lib/common/stringize_arg.hpp" #include @@ -36,11 +35,8 @@ #include #include -#include #include #include -#include -#include #if !defined(ROCPROFILER_HSA_RUNTIME_EXT_AMD_VERSION) # define ROCPROFILER_HSA_RUNTIME_EXT_AMD_VERSION \