From 8e620cc11db015e8941cfc748b285fc514e67953 Mon Sep 17 00:00:00 2001 From: Benjamin Welton Date: Wed, 7 Feb 2024 20:03:21 -0800 Subject: [PATCH] Add support for AQL dimensions (#262) * Add support for AQL dimension changes Adds support for returning dimensions from AQLProfile through rocprofiler to tools. Includes a much larger expanded test suite that covers nearly all files in counter collection. Specific changes below: samples/counter_collection/print_functional_counters: Modified to check the validity of dimensions returned in comparison to the actual underlying data obtained from a kernel execution. rocprofiler-sdk/aql/helpers: adds function calls to support fetching dimension information from AQLProfile. rocprofiler-sdk/aql/packet_construct: modified to allow for events to be exported to aid evaluate_ast in decoding the output buffer. lib/rocprofiler-sdk/counters: Instance count now derived from dimension sizes. rocprofiler_query_counter_dimensions now moved to a callback format to improve usability. rocprofiler-sdk/counters/core: Code migrations and exports of functions for testing. rocprofiler-sdk/counters/dimensions: Generates a dimension cache to be used when querying dimension information for a counter id. rocprofiler-sdk/counters/evaluate_ast: Modified to pass back correct dimension information and to check/determine output dimensions for derived counters. rocprofiler-sdk/counters/id_decode: Modified to have a map between dimension name -> dimension along with a conversion from the aql profile id for a dimension (string) -> integer based id (happens only once during init). rocprofiler-sdk/hsa/queue: Modified to allow for making testing easier. Specifically to allow Queue to now be mocked in unit tests for counter collection. * Merge with changes for serialization * Added suggestions * source formatting (clang-format v11) (#457) Co-authored-by: bwelton * Minor fix * Test change --------- Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Co-authored-by: bwelton [ROCm/rocprofiler-sdk commit: 3eb6a27bc6069504ae125de24850010cca992b17] --- .../samples/counter_collection/main.cpp | 2 +- .../print_functional_counters.cpp | 140 +++- .../source/include/rocprofiler-sdk/counters.h | 36 +- .../rocprofiler-sdk/dispatch_profile.h | 7 + .../source/include/rocprofiler-sdk/fwd.h | 11 +- .../lib/rocprofiler-sdk/aql/CMakeLists.txt | 2 +- .../lib/rocprofiler-sdk/aql/aql_profile_v2.h | 33 + .../lib/rocprofiler-sdk/aql/helpers.cpp | 55 +- .../lib/rocprofiler-sdk/aql/helpers.hpp | 18 + .../rocprofiler-sdk/aql/packet_construct.cpp | 13 + .../rocprofiler-sdk/aql/packet_construct.hpp | 5 + .../source/lib/rocprofiler-sdk/counters.cpp | 95 ++- .../lib/rocprofiler-sdk/counters/core.cpp | 194 +++--- .../lib/rocprofiler-sdk/counters/core.hpp | 28 + .../rocprofiler-sdk/counters/dimensions.cpp | 87 ++- .../rocprofiler-sdk/counters/dimensions.hpp | 61 +- .../rocprofiler-sdk/counters/evaluate_ast.cpp | 85 +-- .../rocprofiler-sdk/counters/evaluate_ast.hpp | 20 +- .../rocprofiler-sdk/counters/id_decode.cpp | 59 +- .../rocprofiler-sdk/counters/id_decode.hpp | 9 +- .../counters/tests/CMakeLists.txt | 5 +- .../rocprofiler-sdk/counters/tests/core.cpp | 648 ++++++++++++++++++ .../counters/tests/dimension.cpp | 270 +++++++- .../counters/tests/init_order.cpp | 3 + .../counters/tests/metrics_test.cpp | 16 + .../source/lib/rocprofiler-sdk/hsa/queue.hpp | 12 +- .../rocprofiler-sdk/hsa/queue_controller.cpp | 19 + .../rocprofiler-sdk/hsa/queue_controller.hpp | 17 +- .../lib/rocprofiler-sdk/rocprofiler.cpp | 13 +- 29 files changed, 1670 insertions(+), 293 deletions(-) create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/core.cpp diff --git a/projects/rocprofiler-sdk/samples/counter_collection/main.cpp b/projects/rocprofiler-sdk/samples/counter_collection/main.cpp index bb35ae126c..205346bc9d 100644 --- a/projects/rocprofiler-sdk/samples/counter_collection/main.cpp +++ b/projects/rocprofiler-sdk/samples/counter_collection/main.cpp @@ -62,7 +62,7 @@ kernelC(T* C_d, const T* A_d, size_t N) void launchKernels() { - const int NUM_LAUNCH = 50000; + const int NUM_LAUNCH = 20000; // Normal HIP Calls int* gpuMem; [[maybe_unused]] hipDeviceProp_t devProp; diff --git a/projects/rocprofiler-sdk/samples/counter_collection/print_functional_counters.cpp b/projects/rocprofiler-sdk/samples/counter_collection/print_functional_counters.cpp index 297169977a..7e1190ba2b 100644 --- a/projects/rocprofiler-sdk/samples/counter_collection/print_functional_counters.cpp +++ b/projects/rocprofiler-sdk/samples/counter_collection/print_functional_counters.cpp @@ -13,6 +13,8 @@ #include #include +#define PRINT_ONLY_FAILING true + /** * Tests the collection of all counters on the agent the test is run on. */ @@ -55,13 +57,92 @@ get_buffer() return buf; } +// Struct to validate that all dimension values are present. Does +// so by creating a tree of dimension values expected. If all are marked as +// having values, then all values are present in the output. +struct validate_dim_presence +{ + validate_dim_presence() {} + + void maybe_forward(const rocprofiler_record_dimension_info_t& dim) + { + if(sub_vectors.empty()) + { + for(size_t i = 0; i < dim.instance_size; i++) + { + sub_vectors.emplace_back(std::make_unique()); + sub_vectors.back()->vector_pos = std::make_pair(dim, i); + } + } + else + { + for(auto& vec : sub_vectors) + { + vec->maybe_forward(dim); + } + } + } + + void mark_seen(const rocprofiler_counter_instance_id_t& id) + { + if(sub_vectors.empty()) + { + has_value = true; + return; + } + size_t pos = 0; + ROCPROFILER_CALL(rocprofiler_query_record_dimension_position( + id, sub_vectors.at(0)->vector_pos.first.id, &pos), + "Could not query position"); + sub_vectors.at(pos)->mark_seen(id); + } + + bool check_seen(std::stringstream& out, + std::vector>& pos_stack) + { + bool ret = true; + if(sub_vectors.empty()) + { + if(!has_value) + { + ret = false; + out << "\tMissing Value at ["; + } + else + { + out << "\tHas Value at ["; + } + for(const auto& [dim, pos] : pos_stack) + { + out << dim.name << ":" << pos << ","; + } + out << "]\n"; + return ret; + } + + for(size_t i = 0; i < sub_vectors.size(); i++) + { + pos_stack.push_back(sub_vectors[i]->vector_pos); + if(!sub_vectors[i]->check_seen(out, pos_stack)) ret = false; + pos_stack.pop_back(); + } + return ret; + } + + std::pair vector_pos; + std::vector> sub_vectors; + bool has_value{false}; +}; + struct CaptureRecords { std::shared_mutex m_mutex{}; // - std::map expected{}; - std::map expected_counter_names{}; - std::vector remaining{}; + std::map expected{}; + // expected dims that we should see data for + std::map expected_data_dims{}; + std::map expected_counter_names{}; + std::vector remaining{}; // std::map captured{}; }; @@ -82,8 +163,9 @@ buffered_callback(rocprofiler_context_id_t, void*, uint64_t) { - auto& cap = *get_capture(); - auto wlock = std::unique_lock{cap.m_mutex}; + auto& cap = *get_capture(); + auto wlock = std::unique_lock{cap.m_mutex}; + std::map seen_counters; for(size_t i = 0; i < num_headers; ++i) { @@ -95,6 +177,7 @@ buffered_callback(rocprofiler_context_id_t, rocprofiler_counter_id_t counter; auto* record = static_cast(header->payload); rocprofiler_query_record_counter_id(record->id, &counter); + cap.expected_data_dims.at(counter.handle).mark_seen(record->id); seen_counters.emplace(counter.handle, 0).first->second++; } } @@ -146,15 +229,38 @@ dispatch_callback(rocprofiler_queue_id_t /*queue_id*/, for(auto& found_counter : counters_needed) { - size_t expected = 0; - rocprofiler_query_counter_instance_count(agent->id, found_counter, &expected); - cap.remaining.push_back(found_counter); - cap.expected.emplace(found_counter.handle, expected); const char* name; size_t name_size; ROCPROFILER_CALL(rocprofiler_query_counter_name(found_counter, &name, &name_size), "Could not query name"); cap.expected_counter_names.emplace(found_counter.handle, std::string(name)); + size_t expected = 0; + ROCPROFILER_CALL( + rocprofiler_query_counter_instance_count(agent->id, found_counter, &expected), + "COULD NOT QUERY INSTANCES"); + cap.remaining.push_back(found_counter); + cap.expected.emplace(found_counter.handle, expected); + + auto& info_vector = + cap.expected_data_dims.emplace(found_counter.handle, validate_dim_presence{}) + .first->second; + + ROCPROFILER_CALL(rocprofiler_iterate_counter_dimensions( + found_counter, + [](rocprofiler_counter_id_t, + const rocprofiler_record_dimension_info_t* dim_info, + size_t num_dims, + void* user_data) { + validate_dim_presence* dim_presence = + static_cast(user_data); + for(size_t i = 0; i < num_dims; i++) + { + dim_presence->maybe_forward(dim_info[i]); + } + return ROCPROFILER_STATUS_SUCCESS; + }, + static_cast(&info_vector)), + "Could not fetch dimension info"); } if(cap.expected.empty()) { @@ -247,11 +353,25 @@ tool_fini(void*) << " (" << name << ")" << " expected " << expected << " instances and got " << *actual_size << "\n"; } - else + else if(!actual_size) { std::clog << "[ERROR] Counter ID: " << counter_id << " (" << name << ") is missing from output\n"; } + else + { + // Counter collected OK + std::stringstream ss; + std::vector> stack; + bool passed = cap.expected_data_dims.at(counter_id).check_seen(ss, stack); + if(!PRINT_ONLY_FAILING || !passed) + { + std::clog << (passed ? "[OK] " : "[ERROR] ") << "Counter ID: " << counter_id << " (" + << name << ")" + << " Expected: " << expected << " Got: " << *actual_size << "\n"; + std::clog << ss.str(); + } + } } } diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/counters.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/counters.h index 5ea937de87..47871cdb49 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/counters.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/counters.h @@ -63,22 +63,38 @@ rocprofiler_query_record_dimension_position(rocprofiler_counter_instance_id_t i size_t* pos) ROCPROFILER_NONNULL(3); /** - * @brief Return information about the dimension for a specified counter. This call - * is primary for future use not related to this alpha since the only dimension - * supported is 0 (flat array without dimension). + * @brief Callback that gives a list of available dimensions for a counter + * + * @param [in] id Counter id the dimension data is for + * @param [in] dim_info An array of dimensions for the counter + * @ref rocprofiler_iterate_counter_dimensions was called on. + * @param [in] num_dims Number of dimensions + * @param [in] user_data User data supplied by + * @ref rocprofiler_iterate_agent_supported_counters + */ +typedef rocprofiler_status_t (*rocprofiler_available_dimensions_cb_t)( + rocprofiler_counter_id_t id, + const rocprofiler_record_dimension_info_t* dim_info, + size_t num_dims, + void* user_data); + +/** + * @brief Return information about the dimensions that exists for a specific counter + * and the extent of each dimension. * * @param [in] id counter id to query dimension info for. - * @param [in] dim dimension (currently only 0 is allowed) - * @param [out] info info on the dimension (name, instance_size) + * @param [in] info_cb Callback to return dimension information for counter + * @param [in] user_data data to pass into the callback * @return ::rocprofiler_status_t * @retval ROCPROFILER_STATUS_SUCCESS if dimension exists - * @retval ROCPROFILER_STATUS_ERROR if the dimension does not + * @retval ROCPROFILER_STATUS_ERROR_HSA_NOT_LOADED if HSA is not loaded when this is called + * @retval ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND if counter is not found + * @retval ROCPROFILER_STATUS_ERROR_DIM_NOT_FOUND if counter does not have this dimension */ rocprofiler_status_t ROCPROFILER_API -rocprofiler_query_record_dimension_info(rocprofiler_counter_id_t id, - rocprofiler_counter_dimension_id_t dim, - rocprofiler_record_dimension_info_t* info) - ROCPROFILER_NONNULL(3); +rocprofiler_iterate_counter_dimensions(rocprofiler_counter_id_t id, + rocprofiler_available_dimensions_cb_t info_cb, + void* user_data); /** * @brief Query Counter name. Name is a pointer controlled by rocprofiler and diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/dispatch_profile.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/dispatch_profile.h index e6f741bd1a..dcd68c66c4 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/dispatch_profile.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/dispatch_profile.h @@ -68,6 +68,13 @@ typedef void (*rocprofiler_profile_counting_dispatch_callback_t)( * one dispatch (denoted by correlation id). Will trigger the * callback based on the parameters setup in buffer_id_t. * + * NOTE: Interface is up for comment as to whether restrictions + * on agent should be made here (limiting the CB based on agent) + * or if the restriction should be performed by the tool in + * @ref rocprofiler_profile_counting_dispatch_callback_t (i.e. + * tool code checking the agent param to see if they want to profile + * it). + * * Interface is up for comment as to whether restrictions * on agent should be made here (limiting the CB based on agent) * or if the restriction should be performed by the tool in diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h index 0848a98418..d5637feddf 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h @@ -79,7 +79,14 @@ typedef enum // NOLINT(performance-enum-size) ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT, ///< Function invoked with one or more invalid ///< arguments ROCPROFILER_STATUS_ERROR_METRIC_NOT_VALID_FOR_AGENT, ///< Invalid metric supplied to agent. - ROCPROFILER_STATUS_ERROR_FINALIZED, ///< invalid because rocprofiler has been finalized + ROCPROFILER_STATUS_ERROR_FINALIZED, ///< invalid because rocprofiler has been finalized + ROCPROFILER_STATUS_ERROR_HSA_NOT_LOADED, ///< Call requires HSA to be loaded before performed + ROCPROFILER_STATUS_ERROR_DIM_NOT_FOUND, ///< Dimension is not found for counter + ROCPROFILER_STATUS_ERROR_PROFILE_COUNTER_NOT_FOUND, ///< Profile could not find counter for GPU + ///< agent + ROCPROFILER_STATUS_ERROR_AST_GENERATION_FAILED, ///< AST could not be generated correctly + ROCPROFILER_STATUS_ERROR_AST_NOT_FOUND, ///< AST was not found + ROCPROFILER_STATUS_ERROR_AQL_NO_EVENT_COORD, ///< Event coordinate was not found by AQL profile ROCPROFILER_STATUS_LAST, } rocprofiler_status_t; @@ -479,6 +486,8 @@ typedef struct { const char* name; size_t instance_size; + rocprofiler_counter_dimension_id_t + id; //<< Id for this dimension used by @ref rocprofiler_query_record_dimension_position } rocprofiler_record_dimension_info_t; /** diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/CMakeLists.txt index 890aeec5b4..01f7f0683d 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/CMakeLists.txt @@ -1,5 +1,5 @@ set(ROCPROFILER_LIB_AQL_SOURCES helpers.cpp packet_construct.cpp) -set(ROCPROFILER_LIB_AQL_HEADERS helpers.hpp packet_construct.hpp) +set(ROCPROFILER_LIB_AQL_HEADERS helpers.hpp packet_construct.hpp aql_profile_v2.h) target_sources(rocprofiler-object-library PRIVATE ${ROCPROFILER_LIB_AQL_SOURCES} ${ROCPROFILER_LIB_AQL_HEADERS}) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h new file mode 100644 index 0000000000..88d4ce925a --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h @@ -0,0 +1,33 @@ +// This file should be removed when it appears in AQL Profile +#pragma ONCE + +#include + +#define PUBLIC_API __attribute__((visibility("default"))) + +extern "C" { +/** + * @brief Callback for iteration of all possible event coordinate IDs and coordinate names. + * @param [in] id Integer identifying the dimension. + * @param [in] name Name of the dimension + * @param [in] data User data supplied to @ref aqlprofile_iterate_event_ids + * @return hsa_status_t + * @retval HSA_STATUS_SUCCESS Continues iteration + * @retval OTHERS Any other HSA return values stops iteration, passing back this value through + * @ref aqlprofile_iterate_event_ids + */ +typedef hsa_status_t (*aqlprofile_eventname_callback_t)(int id, const char* name, void* data); + +/** + * @brief Iterate over all possible event coordinate IDs and their names. + * @param [in] callback Callback to use for iteration of dimensions + * @param [in] user_data Data to supply to callback @ref aqlprofile_eventname_callback_t + * @return hsa_status_t + * @retval HSA_STATUS_SUCCESS if successful + * @retval HSA_STATUS_ERROR if error on interation + * @retval OTHERS If @ref aqlprofile_eventname_callback_t returns non-HSA_STATUS_SUCCESS, + * that value is returned. + */ +PUBLIC_API hsa_status_t +aqlprofile_iterate_event_ids(aqlprofile_eventname_callback_t callback, void* user_data); +} diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/helpers.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/helpers.cpp index f5ddbe8d72..bf62d610e8 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/helpers.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/helpers.cpp @@ -25,6 +25,12 @@ #include #include +#include + +#include "lib/common/synchronized.hpp" +#include "lib/common/utility.hpp" +#include "lib/rocprofiler-sdk/counters/id_decode.hpp" + namespace rocprofiler { namespace aql @@ -60,5 +66,52 @@ get_block_counters(hsa_agent_t agent, const hsa_ven_amd_aqlprofile_event_t& even } return max_block_counters; } + +rocprofiler_status_t +set_dim_id_from_sample(rocprofiler_counter_instance_id_t& id, + hsa_agent_t agent, + hsa_ven_amd_aqlprofile_event_t event, + uint32_t sample_id) +{ + auto callback = + [](int, int sid, int, int coordinate, const char*, void* userdata) -> hsa_status_t { + if(const auto* rocprof_id = + rocprofiler::common::get_val(counters::aqlprofile_id_to_rocprof_instance(), sid)) + { + counters::set_dim_in_rec(*static_cast(userdata), + *rocprof_id, + static_cast(coordinate)); + } + return HSA_STATUS_SUCCESS; + }; + + if(hsa_ven_amd_aqlprofile_iterate_event_coord( + agent, event, sample_id, callback, static_cast(&id)) != HSA_STATUS_SUCCESS) + { + return ROCPROFILER_STATUS_ERROR_AQL_NO_EVENT_COORD; + } + return ROCPROFILER_STATUS_SUCCESS; +} + +rocprofiler_status_t +get_dim_info(hsa_agent_t agent, + hsa_ven_amd_aqlprofile_event_t event, + uint32_t sample_id, + std::map& dims) +{ + auto callback = [](int, int id, int extent, int, const char*, void* userdata) -> hsa_status_t { + auto& map = *static_cast*>(userdata); + map.emplace(id, extent); + return HSA_STATUS_SUCCESS; + }; + + if(hsa_ven_amd_aqlprofile_iterate_event_coord( + agent, event, sample_id, callback, static_cast(&dims)) != HSA_STATUS_SUCCESS) + { + return ROCPROFILER_STATUS_ERROR_AQL_NO_EVENT_COORD; + } + + return ROCPROFILER_STATUS_SUCCESS; +} } // namespace aql -} // namespace rocprofiler +} // namespace rocprofiler \ No newline at end of file diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/helpers.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/helpers.hpp index bb3de4ddb0..44df378086 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/helpers.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/helpers.hpp @@ -23,9 +23,13 @@ #pragma once #include +#include +#include #include +#include + #include "lib/rocprofiler-sdk/counters/metrics.hpp" namespace rocprofiler @@ -39,5 +43,19 @@ get_query_info(hsa_agent_t agent, const counters::Metric& metric); // Query HSA_VEN_AMD_AQLPROFILE_INFO_BLOCK_COUNTERS from aqlprofiler uint32_t get_block_counters(hsa_agent_t agent, const hsa_ven_amd_aqlprofile_event_t& event); + +// Query dimimension ids for counter event. Returns AQLProfiler ID -> extent +rocprofiler_status_t +get_dim_info(hsa_agent_t agent, + hsa_ven_amd_aqlprofile_event_t event, + uint32_t sample_id, + std::map& dims); + +// Set dimension ids into id for sample +rocprofiler_status_t +set_dim_id_from_sample(rocprofiler_counter_instance_id_t& id, + hsa_agent_t agent, + hsa_ven_amd_aqlprofile_event_t event, + uint32_t sample_id); } // namespace aql } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp index a7075d58d9..4c1ac85251 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp @@ -190,6 +190,19 @@ AQLPacketConstruct::event_to_metric(const hsa_ven_amd_aqlprofile_event_t& event) return nullptr; } +const std::vector& +AQLPacketConstruct::get_counter_events(const counters::Metric& metric) const +{ + for(const auto& prof_metric : _metrics) + { + if(prof_metric.metric.id() == metric.id()) + { + return prof_metric.instances; + } + } + throw std::runtime_error(fmt::format("Cannot Find Events for {}", metric)); +} + void AQLPacketConstruct::can_collect() { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.hpp index 2fe247d2fd..c38b5e8326 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.hpp @@ -56,6 +56,11 @@ public: const counters::Metric* event_to_metric(const hsa_ven_amd_aqlprofile_event_t& event) const; std::vector get_all_events() const; + const std::vector& get_counter_events( + const counters::Metric&) const; + + hsa_agent_t hsa_agent() const { return _agent.get_hsa_agent(); } + private: struct AQLProfileMetric { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters.cpp index adda2d8176..b0e0fbc602 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters.cpp @@ -25,9 +25,12 @@ #include +#include "lib/common/container/small_vector.hpp" +#include "lib/common/static_object.hpp" #include "lib/common/synchronized.hpp" #include "lib/rocprofiler-sdk/agent.hpp" #include "lib/rocprofiler-sdk/aql/helpers.hpp" +#include "lib/rocprofiler-sdk/counters/dimensions.hpp" #include "lib/rocprofiler-sdk/counters/evaluate_ast.hpp" #include "lib/rocprofiler-sdk/counters/id_decode.hpp" #include "lib/rocprofiler-sdk/counters/metrics.hpp" @@ -73,55 +76,27 @@ rocprofiler_query_counter_name(rocprofiler_counter_id_t counter_id, const char** * @return rocprofiler_status_t */ rocprofiler_status_t -rocprofiler_query_counter_instance_count(rocprofiler_agent_id_t agent_id, +rocprofiler_query_counter_instance_count(rocprofiler_agent_id_t, rocprofiler_counter_id_t counter_id, size_t* instance_count) { - const rocprofiler_agent_t* agent = rocprofiler::agent::get_agent(agent_id); - - if(!agent) return ROCPROFILER_STATUS_ERROR_AGENT_NOT_FOUND; - if(agent->type != ROCPROFILER_AGENT_TYPE_GPU) return ROCPROFILER_STATUS_ERROR; - - const auto& id_map = *CHECK_NOTNULL(rocprofiler::counters::getMetricIdMap()); - const auto* metric_ptr = rocprofiler::common::get_val(id_map, counter_id.handle); - if(!metric_ptr) return ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND; - *instance_count = 0; - // Special counters do not have hardware metrics and will always have an instance - // count of 1 (i.e. MAX_WAVE_SIZE) - if(!metric_ptr->special().empty()) + + if(rocprofiler::counters::get_dimension_cache().empty()) { - *instance_count = 1; - return ROCPROFILER_STATUS_SUCCESS; + return ROCPROFILER_STATUS_ERROR_HSA_NOT_LOADED; } - // Returns the set of hardware counters needed to evaluate the metric. - // For derived metrics, this can be more than one counter. In that case, - // we return the maximum instance count among all underlying counters. - auto req_counters = rocprofiler::counters::get_required_hardware_counters( - rocprofiler::counters::get_ast_map(), std::string(agent->name), *metric_ptr); - if(!req_counters) return ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND; + const auto* dims = rocprofiler::common::get_val(rocprofiler::counters::get_dimension_cache(), + counter_id.handle); + if(!dims) return ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND; - for(const auto& counter : *req_counters) + for(const auto& metric_dim : *dims) { - if(!counter.special().empty()) - { - *instance_count = std::max(size_t(1), *instance_count); - continue; - } - - try - { - auto dims = rocprofiler::counters::getBlockDimensions(agent->name, counter); - for(const auto& dim : dims) - { - *instance_count = std::max(static_cast(dim.size()), *instance_count); - } - } catch(std::runtime_error& err) - { - LOG(ERROR) << fmt::format("Could not lookup instance count for counter {}", counter); - return ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND; - } + if(*instance_count == 0) + *instance_count = metric_dim.size(); + else if(metric_dim.size() > 0) + *instance_count = metric_dim.size() * *instance_count; } return ROCPROFILER_STATUS_SUCCESS; @@ -180,20 +155,36 @@ rocprofiler_query_record_dimension_position(rocprofiler_counter_instance_id_t i } rocprofiler_status_t -rocprofiler_query_record_dimension_info(rocprofiler_counter_id_t, - rocprofiler_counter_dimension_id_t dim, - rocprofiler_record_dimension_info_t* info) +rocprofiler_iterate_counter_dimensions(rocprofiler_counter_id_t id, + rocprofiler_available_dimensions_cb_t info_cb, + void* user_data) { - if(const auto* ptr = rocprofiler::common::get_val( - rocprofiler::counters::dimension_map(), - static_cast(dim))) + if(rocprofiler::counters::get_dimension_cache().empty()) { - info->name = ptr->c_str(); - // TODO: Needs info on the instance size per block to fill in. - // counter_id will be used to lookup this information. - info->instance_size = 0; - return ROCPROFILER_STATUS_SUCCESS; + return ROCPROFILER_STATUS_ERROR_HSA_NOT_LOADED; } - return ROCPROFILER_STATUS_ERROR; + + const auto* dims = + rocprofiler::common::get_val(rocprofiler::counters::get_dimension_cache(), id.handle); + if(!dims) return ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND; + + // This is likely faster than a map lookup given the limited number of dims. + rocprofiler::common::container::small_vector user_dims; + for(const auto& internal_dim : *dims) + { + auto& dim = user_dims.emplace_back(); + dim.name = internal_dim.name().c_str(); + dim.instance_size = internal_dim.size(); + dim.id = static_cast(internal_dim.type()); + } + + if(user_dims.empty()) + { + return ROCPROFILER_STATUS_ERROR_DIM_NOT_FOUND; + } + + info_cb(id, user_dims.data(), user_dims.size(), user_data); + + return ROCPROFILER_STATUS_SUCCESS; } } diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.cpp index 32a506b165..aea9db5383 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.cpp @@ -37,9 +37,6 @@ namespace rocprofiler { namespace counters { -using ClientID = int64_t; -using inst_pkt_t = common::container:: - small_vector, ClientID>, 4>; class CounterController { public: @@ -133,6 +130,113 @@ 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) +{ + if(profile->pkt_generator || !profile->reqired_hw_counters.empty()) + { + return ROCPROFILER_STATUS_SUCCESS; + } + + // Sets up the packet generator for the profile. This must be delayed until after HSA is loaded. + // This call needs to be thread protected in that only one thread must be setting up profile at + // the same time. + + auto& config = *profile; + auto agent_name = std::string(config.agent->name); + for(const auto& metric : config.metrics) + { + auto req_counters = get_required_hardware_counters(get_ast_map(), agent_name, metric); + + if(!req_counters) + { + LOG(ERROR) << fmt::format("Could not find counter {}", metric.name()); + return ROCPROFILER_STATUS_ERROR_PROFILE_COUNTER_NOT_FOUND; + } + + // Special metrics are those that are not hw counters but other + // constants like MAX_WAVE_SIZE + for(const auto& req_metric : *req_counters) + { + if(req_metric.special().empty()) + { + config.reqired_hw_counters.insert(req_metric); + } + else + { + config.required_special_counters.insert(req_metric); + } + } + + const auto& asts = get_ast_map(); + const auto* agent_map = rocprofiler::common::get_val(asts, agent_name); + if(!agent_map) + { + LOG(ERROR) << fmt::format("Coult not build AST for {}", agent_name); + return ROCPROFILER_STATUS_ERROR_AST_GENERATION_FAILED; + } + + const auto* counter_ast = rocprofiler::common::get_val(*agent_map, metric.name()); + if(!counter_ast) + { + LOG(ERROR) << fmt::format("Coult not find AST for {}", metric.name()); + return ROCPROFILER_STATUS_ERROR_AST_NOT_FOUND; + } + config.asts.push_back(*counter_ast); + config.asts.back().set_dimensions(); + } + + profile->pkt_generator = std::make_unique( + agent, + std::vector{profile->reqired_hw_counters.begin(), + profile->reqired_hw_counters.end()}); + return ROCPROFILER_STATUS_SUCCESS; +} + +rocprofiler_status_t +counter_callback_info::get_packet(std::unique_ptr& ret_pkt, + const hsa::AgentCache& agent, + std::shared_ptr& profile) +{ + rocprofiler_status_t status; + // Check packet cache + profile->packets.wlock([&](auto& pkt_vector) { + status = counter_callback_info::setup_profile_config(agent, profile); + if(!pkt_vector.empty() && status == ROCPROFILER_STATUS_SUCCESS) + { + ret_pkt = std::move(pkt_vector.back()); + pkt_vector.pop_back(); + } + }); + + if(status != ROCPROFILER_STATUS_SUCCESS) return status; + if(!ret_pkt) + { + // If we do not have a packet in the cache, create one. + ret_pkt = + profile->pkt_generator->construct_packet(hsa::get_queue_controller().get_ext_table()); + } + + ret_pkt->before_krn_pkt.clear(); + ret_pkt->after_krn_pkt.clear(); + packet_return_map.wlock([&](auto& data) { data.emplace(ret_pkt.get(), profile); }); + + return ROCPROFILER_STATUS_SUCCESS; +} + /** * Callback we get from HSA interceptor when a kernel packet is being enqueued. * @@ -174,85 +278,11 @@ queue_cb(const std::shared_ptr& info, 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); - // Check packet cache - prof_config->packets.wlock([&](auto& pkt_vector) { - // Delay packet generator construction until first HSA packet is processed - // This ensures that HSA exists - if(!prof_config->pkt_generator) - { - // One time setup of profile config - if(prof_config->reqired_hw_counters.empty()) - { - auto& config = *prof_config; - auto agent_name = std::string(config.agent->name); - for(const auto& metric : config.metrics) - { - auto req_counters = - get_required_hardware_counters(get_ast_map(), agent_name, metric); - - if(!req_counters) - { - throw std::runtime_error( - fmt::format("Could not find counter {}", metric.name())); - } - - // Special metrics are those that are not hw counters but other - // constants like MAX_WAVE_SIZE - for(const auto& req_metric : *req_counters) - { - if(req_metric.special().empty()) - { - config.reqired_hw_counters.insert(req_metric); - } - else - { - config.required_special_counters.insert(req_metric); - } - } - - const auto& asts = get_ast_map(); - const auto* agent_map = rocprofiler::common::get_val(asts, agent_name); - if(!agent_map) - throw std::runtime_error( - fmt::format("Coult not build AST for {}", agent_name)); - const auto* counter_ast = - rocprofiler::common::get_val(*agent_map, metric.name()); - if(!counter_ast) - { - throw std::runtime_error( - fmt::format("Coult not find AST for {}", metric.name())); - } - config.asts.push_back(*counter_ast); - config.asts.back().set_dimensions(); - } - } - - prof_config->pkt_generator = std::make_unique( - queue.get_agent(), - std::vector{prof_config->reqired_hw_counters.begin(), - prof_config->reqired_hw_counters.end()}); - } - - if(!pkt_vector.empty()) - { - ret_pkt = std::move(pkt_vector.back()); - pkt_vector.pop_back(); - } - }); - - if(!ret_pkt) - { - // If we do not have a packet in the cache, create one. - ret_pkt = prof_config->pkt_generator->construct_packet( - hsa::get_queue_controller().get_ext_table()); - } - ret_pkt->before_krn_pkt.clear(); - ret_pkt->after_krn_pkt.clear(); if(ret_pkt->empty) return ret_pkt; - info->packet_return_map.wlock([&](auto& data) { data.emplace(ret_pkt.get(), prof_config); }); - auto&& CreateBarrierPacket = [](hsa_signal_t* dependency_signal, hsa_signal_t* completion_signal, @@ -312,7 +342,11 @@ completed_cb(const std::shared_ptr& info, }); if(!pkt) return; - hsa::profiler_serializer_kernel_completion_signal(session.queue.block_signal); + + if(!pkt->empty) + { + hsa::profiler_serializer_kernel_completion_signal(session.queue.block_signal); + } auto decoded_pkt = EvaluateAST::read_pkt(prof_config->pkt_generator.get(), *pkt); EvaluateAST::read_special_counters( diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.hpp index 2ed3638dc2..6897ed14bf 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.hpp @@ -88,6 +88,13 @@ struct counter_callback_info rocprofiler::common::Synchronized< std::unordered_map>> packet_return_map{}; + + static rocprofiler_status_t setup_profile_config(const hsa::AgentCache&, + std::shared_ptr&); + + rocprofiler_status_t get_packet(std::unique_ptr&, + const hsa::AgentCache&, + std::shared_ptr&); }; uint64_t @@ -107,5 +114,26 @@ start_context(const context::context*); void stop_context(const context::context*); + +std::unique_ptr +queue_cb(const std::shared_ptr& info, + const hsa::Queue& queue, + const hsa::rocprofiler_packet& pkt, + uint64_t kernel_id, + 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 std::shared_ptr&, + const hsa::Queue&, + hsa::rocprofiler_packet, + const hsa::Queue::queue_info_session_t&, + inst_pkt_t& pkts); + +std::shared_ptr get_profile_config(rocprofiler_profile_config_id_t); } // namespace counters } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/dimensions.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/dimensions.cpp index f67f305151..8d97863041 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/dimensions.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/dimensions.cpp @@ -29,10 +29,12 @@ #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 @@ -45,35 +47,82 @@ getBlockDimensions(std::string_view agent, const Metric& metric) if(!metric.special().empty()) { // Special non-hardware counters without dimension data - return std::vector{ - {dimension_map().at(ROCPROFILER_DIMENSION_NONE), 1, ROCPROFILER_DIMENSION_NONE}}; + return std::vector{{dimension_map().at(ROCPROFILER_DIMENSION_INSTANCE), + 1, + ROCPROFILER_DIMENSION_INSTANCE}}; } + std::unordered_map count; + + std::vector ret; + for(const auto& [_, maybe_agent] : hsa::get_queue_controller().get_supported_agents()) { if(maybe_agent.name() == agent) { - // To be returned when instance counting is functional with AQL profiler - // return std::vector{ - // {dimension_map().at(ROCPROFILER_DIMENSION_SHADER_ENGINE), - // maybe_agent.get_rocp_agent()->num_shader_banks, - // ROCPROFILER_DIMENSION_SHADER_ENGINE}, - // {dimension_map().at(ROCPROFILER_DIMENSION_XCC), - // maybe_agent.get_rocp_agent()->num_xcc, - // ROCPROFILER_DIMENSION_XCC}, - // {dimension_map().at(ROCPROFILER_DIMENSION_CU), - // maybe_agent.get_rocp_agent()->cu_count, - // ROCPROFILER_DIMENSION_CU}, - // {dimension_map().at(ROCPROFILER_DIMENSION_AGENT), - // maybe_agent.get_rocp_agent()->id.handle, - // ROCPROFILER_DIMENSION_AGENT}}; aql::AQLPacketConstruct pkt_gen(maybe_agent, {metric}); - return std::vector{ - {metric.block(), pkt_gen.get_all_events().size(), ROCPROFILER_DIMENSION_NONE}}; + const auto& events = pkt_gen.get_counter_events(metric); + + for(const auto& event : events) + { + std::map dims; + auto status = aql::get_dim_info(maybe_agent.get_hsa_agent(), event, 0, dims); + CHECK_EQ(status, ROCPROFILER_STATUS_SUCCESS) + << rocprofiler_get_status_string(status); + + for(const auto& [id, extent] : dims) + { + if(const auto* inst_type = + rocprofiler::common::get_val(aqlprofile_id_to_rocprof_instance(), id)) + { + count.emplace(*inst_type, 0).first->second = extent; + } + else + { + LOG(ERROR) << "Unknown AQL Profiler Dimension " << id << " " << extent; + } + } + } } } - return {}; + ret.reserve(count.size()); + for(const auto& [dim, size] : count) + { + ret.emplace_back(dimension_map().at(dim), size, dim); + } + + return ret; +} + +const std::unordered_map>& +get_dimension_cache() +{ + static auto*& cache = + common::static_object>>:: + construct([]() -> std::unordered_map> { + std::unordered_map> dims; + /** + * Fails if HSA is not loaded by retruning nothing. This should not remain after + * AQL is transistioned away from HSA. + */ + if(rocprofiler::hsa::get_queue_controller().get_supported_agents().empty()) + { + return {}; + } + + const auto& asts = counters::get_ast_map(); + for(const auto& [gfx, metrics] : asts) + { + for(const auto& [_, ast] : metrics) + { + auto ast_copy = ast; + dims.emplace(ast.out_id().handle, ast_copy.set_dimensions()); + } + } + return dims; + }()); + return *cache; } } // namespace counters diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/dimensions.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/dimensions.hpp index 982c9dc795..c7b437d0d8 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/dimensions.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/dimensions.hpp @@ -39,10 +39,10 @@ namespace counters class MetricDimension { public: - MetricDimension(std::string name, + MetricDimension(std::string_view name, uint64_t dim_size, rocprofiler_profile_counter_instance_types type) - : name_(std::move(name)) + : name_(name) , size_(dim_size) , type_(type){}; @@ -60,44 +60,31 @@ private: rocprofiler_profile_counter_instance_types type_; }; -/* -{ - AgentBlockDimensionsMap = { - "gfx906":{}, - "gfx908":{ - "TCC": [dim_1, dim_2 ... dim_n] - "SQ": [dim_1, dim_2 ... dim_n] - "TCP": [dim_1, dim_2 ... dim_n] - } - } -} -*/ - -// // map block_name -> MetricDimension -// using BlockDimensionMap = std::unordered_map>; -// // map agent_name -> BlockDimensionMap -// using AgentBlockDimensionsMap = std::unordered_map; - -// // map dimension_id -> MetricDimension -// using DimensionIdMap = std::unordered_map; - -// // get the complete AgentBlockDimensionsMap -// const AgentBlockDimensionsMap& -// getAgentBlockDimensionsMap(); - -// // get specific dimensions for an agent, block_name -// const std::vector& -// getBlockDimension(const std::string& agent, -// std::string block_name, -// rocprofiler_profile_counter_instance_types dim); - // get all dimensions for an agent, block_name std::vector getBlockDimensions(std::string_view agent, const counters::Metric&); -// // get a specific dimension by id -// const MetricDimension& -// getDimensionById(uint64_t id); - +const std::unordered_map>& +get_dimension_cache(); } // namespace counters } // namespace rocprofiler + +namespace fmt +{ +// fmt::format support for metric +template <> +struct formatter +{ + template + constexpr auto parse(ParseContext& ctx) + { + return ctx.begin(); + } + + template + auto format(rocprofiler::counters::MetricDimension const& dims, Ctx& ctx) const + { + return fmt::format_to(ctx.out(), "[{}, {}]", dims.name(), dims.size()); + } +}; +} // namespace fmt \ No newline at end of file diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp index 7624201c24..c2e0ebbbea 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp @@ -21,13 +21,16 @@ // SOFTWARE. #include "lib/rocprofiler-sdk/counters/evaluate_ast.hpp" -#include #include -#include - #include +#include #include + +#include +#include +#include + #include "lib/common/synchronized.hpp" #include "lib/common/utility.hpp" #include "lib/rocprofiler-sdk/counters/dimensions.hpp" @@ -221,41 +224,47 @@ EvaluateAST::EvaluateAST(rocprofiler_counter_id_t out_id, } } -DimensionTypes +std::vector EvaluateAST::set_dimensions() { - if(_dimension_types != DIMENSION_LAST) + if(!_dimension_types.empty()) { return _dimension_types; } - auto get_dim_types = [&](auto& metric) { - int dim_types = 0; - for(const auto& dim : getBlockDimensions(_agent, metric)) - { - dim_types |= (dim.type() != ROCPROFILER_DIMENSION_NONE) ? (0x1 << dim.type()) : 0; - } - return static_cast(dim_types); - }; + auto get_dim_types = [&](auto& metric) { return getBlockDimensions(_agent, metric); }; switch(_type) { case NONE: case RANGE_NODE: case CONSTANT_NODE: - case NUMBER_NODE: break; + case NUMBER_NODE: + { + _dimension_types = + std::vector{{dimension_map().at(ROCPROFILER_DIMENSION_INSTANCE), + 1, + ROCPROFILER_DIMENSION_INSTANCE}}; + } + break; case ADDITION_NODE: case SUBTRACTION_NODE: case MULTIPLY_NODE: case DIVIDE_NODE: { - if(_children[0].set_dimensions() != _children[1].set_dimensions() && - _children[0].type() != NUMBER_NODE && _children[1].type() != NUMBER_NODE) - throw std::runtime_error(fmt::format("Dimension mis-mismatch: {} and {}", - _children[0].metric(), - _children[1].metric())); - _dimension_types = (_children[0].type() != NUMBER_NODE) ? _children[0].set_dimensions() - : _children[1].set_dimensions(); + auto first = _children[0].set_dimensions(); + auto second = _children[1].set_dimensions(); + // - first.size() > 1 && second.size() > 1 + // This is an explicit compatibility change to allow existing integer * COUNTER + // derived counters to function + if(first != second && first.size() > 1 && second.size() > 1) + throw std::runtime_error( + fmt::format("Dimension mis-mismatch: {} (dims: {}) and {} (dims: {})", + _children[0].metric(), + fmt::join(_children[0].set_dimensions(), ","), + _children[1].metric(), + fmt::join(_children[1].set_dimensions(), ","))); + _dimension_types = first.size() > second.size() ? first : second; } break; case REFERENCE_NODE: @@ -265,19 +274,11 @@ EvaluateAST::set_dimensions() break; case REDUCE_NODE: { - // There is only one child node in case of REDUCE_NODE and that - // child node denotes the expression on which the reduce is applied. - // The resulting dimension of REDUCE_NODE will be the child's dimension - // minus the dimensions specified in the reduce_dimension_set. - int original_dim = static_cast(_children[0].set_dimensions()); - int turn_off_dims = 0; - for(auto dim : _reduce_dimension_set) - { - turn_off_dims |= (dim != ROCPROFILER_DIMENSION_NONE) ? (0x1 << dim) : 1; - } - int final_dims = _reduce_dimension_set.empty() ? ROCPROFILER_DIMENSION_NONE - : (original_dim & ~turn_off_dims); - _dimension_types = static_cast(final_dims); + // Reduction down to a single instance supported for now. + _dimension_types = + std::vector{{dimension_map().at(ROCPROFILER_DIMENSION_INSTANCE), + 1, + ROCPROFILER_DIMENSION_INSTANCE}}; } break; case SELECT_NODE: @@ -387,11 +388,12 @@ using property_function_t = int64_t (*)(const rocprofiler_agent_t&); return static_cast(value); \ }) \ } +} // namespace int64_t -get_agent_property(const std::string& property, const rocprofiler_agent_t& agent) +get_agent_property(std::string_view property, const rocprofiler_agent_t& agent) { - static std::unordered_map props = { + static std::unordered_map props = { GEN_MAP_ENTRY("cpu_cores_count", agent_info.cpu_cores_count), GEN_MAP_ENTRY("simd_count", agent_info.simd_count), GEN_MAP_ENTRY("mem_banks_count", agent_info.mem_banks_count), @@ -427,10 +429,8 @@ get_agent_property(const std::string& property, const rocprofiler_agent_t& agent return (*func)(agent); } - LOG(ERROR) << fmt::format("Unsupported special property {}", property); return 0.0; } -} // namespace void EvaluateAST::read_special_counters( @@ -477,7 +477,14 @@ EvaluateAST::read_pkt(const aql::AQLPacketConstruct* pkt_gen, hsa::AQLPacket& pk auto& next_rec = vec.emplace_back(); set_counter_in_rec(next_rec.id, {.handle = metric->id()}); // Actual dimension info needs to be used here in the future - set_dim_in_rec(next_rec.id, ROCPROFILER_DIMENSION_NONE, vec.size() - 1); + auto aql_status = aql::set_dim_id_from_sample(next_rec.id, + it.pkt_gen->hsa_agent(), + info_data->pmc_data.event, + info_data->sample_id); + CHECK_EQ(aql_status, ROCPROFILER_STATUS_SUCCESS) + << rocprofiler_get_status_string(aql_status); + + // set_dim_in_rec(next_rec.id, ROCPROFILER_DIMENSION_NONE, vec.size() - 1); // Note: in the near future we need to use hw_counter here instead next_rec.counter_value = info_data->pmc_data.result; return HSA_STATUS_SUCCESS; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/evaluate_ast.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/evaluate_ast.hpp index 37387b6f65..c610c7ec1c 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/evaluate_ast.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/evaluate_ast.hpp @@ -104,9 +104,9 @@ public: * of this AST. Can throw if the AST is invalid (i.e. dimension mismatch in * child nodes of this AST). This is done in a recursive fashion. * - * @return DimensionTypes dimension of the output of this AST. + * @return std::vector dimension of the output of this AST. */ - DimensionTypes set_dimensions(); + std::vector set_dimensions(); bool validate_raw_ast(const std::unordered_map& metrics); @@ -152,11 +152,11 @@ public: const std::set& required_special_counters, std::unordered_map>& out_map); - NodeType type() const { return _type; } - ReduceOperation reduce_op() const { return _reduce_op; } - const std::vector& children() const { return _children; } - const Metric& metric() const { return _metric; } - DimensionTypes dimension_types() const { return _dimension_types; } + NodeType type() const { return _type; } + ReduceOperation reduce_op() const { return _reduce_op; } + const std::vector& children() const { return _children; } + const Metric& metric() const { return _metric; } + const std::vector& dimension_types() const { return _dimension_types; } /** * @brief When an evaluation is complete, set the output id of the results. This is called @@ -166,6 +166,8 @@ public: */ void set_out_id(std::vector& results) const; + const rocprofiler_counter_id_t& out_id() const { return _out_id; } + private: NodeType _type{NONE}; ReduceOperation _reduce_op{REDUCE_NONE}; @@ -173,7 +175,7 @@ private: double _raw_value{0}; std::vector _children; std::string _agent; - DimensionTypes _dimension_types{DIMENSION_LAST}; + std::vector _dimension_types{}; std::vector _static_value; std::unordered_set _reduce_dimension_set; bool _expanded{false}; @@ -197,6 +199,8 @@ std::optional> get_required_hardware_counters(const std::unordered_map& asts, const std::string& agent, const Metric& metric); +int64_t +get_agent_property(std::string_view property, const rocprofiler_agent_t& agent); } // namespace counters } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/id_decode.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/id_decode.cpp index c724d448df..a1bb2c22e7 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/id_decode.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/id_decode.cpp @@ -22,28 +22,71 @@ #include "lib/rocprofiler-sdk/counters/id_decode.hpp" +#include #include #include +#include "lib/common/static_object.hpp" #include "lib/common/utility.hpp" +#include "lib/rocprofiler-sdk/aql/aql_profile_v2.h" namespace rocprofiler { namespace counters { -const std::unordered_map& +const std::unordered_map& dimension_map() { - static std::unordered_map map = { - {ROCPROFILER_DIMENSION_NONE, "DIMENSION_NONE"}, - {ROCPROFILER_DIMENSION_XCC, "DIMENSION_XCC"}, - {ROCPROFILER_DIMENSION_SHADER_ENGINE, "DIMENSION_SHADER_ENGINE"}, - {ROCPROFILER_DIMENSION_AGENT, "DIMENSION_AGENT"}, - {ROCPROFILER_DIMENSION_PMC_CHANNEL, "DIMENSION_PMC_CHANNEL"}, - {ROCPROFILER_DIMENSION_CU, "DIMENSION_CU"}, + static std::unordered_map map = { + {ROCPROFILER_DIMENSION_NONE, std::string_view("DIMENSION_NONE")}, + {ROCPROFILER_DIMENSION_XCC, std::string_view("DIMENSION_XCC")}, + {ROCPROFILER_DIMENSION_SHADER_ENGINE, std::string_view("DIMENSION_SHADER_ENGINE")}, + {ROCPROFILER_DIMENSION_AGENT, std::string_view("DIMENSION_AGENT")}, + {ROCPROFILER_DIMENSION_SHADER_ARRAY, std::string_view("DIMENSION_SHADER_ARRAY")}, + {ROCPROFILER_DIMENSION_CU, std::string_view("DIMENSION_CU")}, + {ROCPROFILER_DIMENSION_INSTANCE, std::string_view("DIMENSION_INSTANCE")}, }; return map; } +const std::unordered_map& +aqlprofile_id_to_rocprof_instance() +{ + using dims_map_t = std::unordered_map; + + static auto*& aql_to_rocprof_dims = + common::static_object::construct([]() -> dims_map_t { + dims_map_t data; + + aqlprofile_iterate_event_ids( + [](int id, const char* name, void* userdata) -> hsa_status_t { + const std::unordered_map + aql_string_to_dim = { + {"XCD", ROCPROFILER_DIMENSION_XCC}, + {"SE", ROCPROFILER_DIMENSION_SHADER_ENGINE}, + {"SA", ROCPROFILER_DIMENSION_SHADER_ARRAY}, + {"CU", ROCPROFILER_DIMENSION_CU}, + {"INSTANCE", ROCPROFILER_DIMENSION_INSTANCE}, + }; + + if(const auto* inst_type = + rocprofiler::common::get_val(aql_string_to_dim, name)) + { + // Supported instance type + auto& map = *static_cast< + std::unordered_map*>( + userdata); + map.emplace(id, *inst_type); + } + return HSA_STATUS_SUCCESS; + }, + static_cast(&data)); + return data; + }()); + + return *aql_to_rocprof_dims; +} + } // namespace counters } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/id_decode.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/id_decode.hpp index 82a6f8da86..e15f21e9bb 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/id_decode.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/id_decode.hpp @@ -22,6 +22,7 @@ #pragma once +#include #include #include @@ -41,12 +42,13 @@ enum rocprofiler_profile_counter_instance_types ROCPROFILER_DIMENSION_XCC, ///< XCC dimension of result ROCPROFILER_DIMENSION_SHADER_ENGINE, ///< SE dimension of result ROCPROFILER_DIMENSION_AGENT, ///< Agent dimension - ROCPROFILER_DIMENSION_PMC_CHANNEL, ///< PMC Channel dimension of result + ROCPROFILER_DIMENSION_SHADER_ARRAY, ///< Number of shader arrays ROCPROFILER_DIMENSION_CU, ///< Number of compute units + ROCPROFILER_DIMENSION_INSTANCE, ///< Number of instances ROCPROFILER_DIMENSION_LAST }; -const std::unordered_map& +const std::unordered_map& dimension_map(); inline rocprofiler_counter_id_t @@ -62,6 +64,9 @@ inline size_t rec_to_dim_pos(rocprofiler_counter_instance_id_t id, rocprofiler_profile_counter_instance_types dim); +const std::unordered_map& +aqlprofile_id_to_rocprof_instance(); + } // namespace counters } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/CMakeLists.txt index 131ea05ff4..3ce3ef6440 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/CMakeLists.txt @@ -3,7 +3,7 @@ rocprofiler_deactivate_clang_tidy() include(GoogleTest) set(ROCPROFILER_LIB_COUNTER_TEST_SOURCES metrics_test.cpp evaluate_ast_test.cpp - dimension.cpp init_order.cpp) + dimension.cpp init_order.cpp core.cpp) add_executable(counter-test) @@ -11,7 +11,8 @@ target_sources(counter-test PRIVATE ${ROCPROFILER_LIB_COUNTER_TEST_SOURCES}) target_link_libraries( counter-test - PRIVATE rocprofiler::rocprofiler-hip rocprofiler::rocprofiler-common-library + PRIVATE rocprofiler::rocprofiler-hsa-runtime rocprofiler::rocprofiler-hip + rocprofiler::rocprofiler-common-library rocprofiler::rocprofiler-static-library GTest::gtest GTest::gtest_main) gtest_add_tests( diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/core.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/core.cpp new file mode 100644 index 0000000000..fadd4ef320 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/core.cpp @@ -0,0 +1,648 @@ +// 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 +#include +#include +#include + +#include +#include +#include +#include +#include + +#include + +#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/core.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" +#include "rocprofiler-sdk/registration.h" + +using namespace rocprofiler::counters; +using namespace rocprofiler; + +#define ROCPROFILER_CALL(result, msg) \ + { \ + rocprofiler_status_t CHECKSTATUS = result; \ + if(CHECKSTATUS != ROCPROFILER_STATUS_SUCCESS) \ + { \ + std::string status_msg = rocprofiler_get_status_string(CHECKSTATUS); \ + std::cerr << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " << msg \ + << " failed with error code " << CHECKSTATUS << ": " << status_msg \ + << std::endl; \ + std::stringstream errmsg{}; \ + errmsg << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " << msg " failure (" \ + << status_msg << ")"; \ + ASSERT_EQ(CHECKSTATUS, ROCPROFILER_STATUS_SUCCESS) << errmsg.str(); \ + } \ + } + +namespace +{ +AmdExtTable& +get_ext_table() +{ + static auto _v = []() { + auto val = AmdExtTable{}; + val.hsa_amd_memory_pool_get_info_fn = hsa_amd_memory_pool_get_info; + val.hsa_amd_agent_iterate_memory_pools_fn = hsa_amd_agent_iterate_memory_pools; + val.hsa_amd_memory_pool_allocate_fn = hsa_amd_memory_pool_allocate; + val.hsa_amd_memory_pool_free_fn = hsa_amd_memory_pool_free; + val.hsa_amd_agent_memory_pool_get_info_fn = hsa_amd_agent_memory_pool_get_info; + val.hsa_amd_agents_allow_access_fn = hsa_amd_agents_allow_access; + return val; + }(); + return _v; +} + +CoreApiTable& +get_api_table() +{ + static auto _v = []() { + auto val = CoreApiTable{}; + val.hsa_iterate_agents_fn = hsa_iterate_agents; + val.hsa_agent_get_info_fn = hsa_agent_get_info; + val.hsa_queue_create_fn = hsa_queue_create; + val.hsa_queue_destroy_fn = hsa_queue_destroy; + return val; + }(); + return _v; +} + +auto +findDeviceMetrics(const hsa::AgentCache& agent, const std::unordered_set& metrics) +{ + std::vector ret; + auto all_counters = counters::getMetricMap(); + + LOG(ERROR) << "Looking up counters for " << std::string(agent.name()); + auto gfx_metrics = common::get_val(*all_counters, std::string(agent.name())); + if(!gfx_metrics) + { + LOG(ERROR) << "No counters found for " << std::string(agent.name()); + return ret; + } + + for(auto& counter : *gfx_metrics) + { + if(metrics.count(counter.name()) > 0 || metrics.empty()) + { + ret.push_back(counter); + } + } + return ret; +} + +void +test_init() +{ + HsaApiTable table; + table.amd_ext_ = &get_ext_table(); + table.core_ = &get_api_table(); + agent::construct_agent_cache(&table); + hsa::get_queue_controller().init(get_api_table(), get_ext_table()); +} + +} // namespace + +namespace +{ +rocprofiler_context_id_t& +get_client_ctx() +{ + static rocprofiler_context_id_t ctx; + return ctx; +} + +struct buf_check +{ + size_t expected_size{0}; + bool is_special{false}; + double special_val{0.0}; +}; + +void +buffered_callback(rocprofiler_context_id_t, + rocprofiler_buffer_id_t, + rocprofiler_record_header_t** headers, + size_t num_headers, + void* user_data, + uint64_t) +{ + buf_check& expected = *static_cast(user_data); + if(expected.is_special) + { + // Special values are single value constants (from agent_t) + expected.expected_size = 1; + } + + std::set seen_data; + std::set seen_dims; + for(size_t i = 0; i < num_headers; ++i) + { + auto* header = headers[i]; + if(header->category == ROCPROFILER_BUFFER_CATEGORY_COUNTERS && header->kind == 0) + { + // Print the returned counter data. + auto* record = static_cast(header->payload); + seen_dims.insert(record->id); + seen_data.insert(record->counter_value); + } + } + + /** + * Specific counters default to a size of 128 even if they have less data (primarily + * TCP). This is a known quirk on AQL profile's end where it will allocate for 128 entries + * but return less (and the data may be duplicated across entries). Skip these entires for + * testing purposes since we cannot determine what mock data will be in the return (and its + * arch dependent). + */ + if(expected.expected_size == 128) return; + + EXPECT_EQ(seen_dims.size(), expected.expected_size); + EXPECT_EQ(seen_data.size(), expected.expected_size); + + ASSERT_FALSE(seen_data.empty()); + if(expected.is_special) + { + EXPECT_EQ(*seen_data.begin(), expected.special_val); + } + else + { + EXPECT_EQ(*seen_data.begin(), 1.0); + EXPECT_EQ(*seen_data.rbegin(), double(seen_data.size())); + } +} + +void +null_dispatch_callback(rocprofiler_queue_id_t, + const rocprofiler_agent_t*, + rocprofiler_correlation_id_t, + const hsa_kernel_dispatch_packet_t*, + uint64_t, + void*, + rocprofiler_profile_config_id_t*) +{} + +void +null_buffered_callback(rocprofiler_context_id_t, + rocprofiler_buffer_id_t, + rocprofiler_record_header_t**, + size_t, + void*, + uint64_t) +{} + +} // namespace + +TEST(core, check_packet_generation) +{ + ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); + test_init(); + + auto agents = hsa::get_queue_controller().get_supported_agents(); + ASSERT_GT(agents.size(), 0); + for(const auto& [_, agent] : agents) + { + auto metrics = findDeviceMetrics(agent, {}); + ASSERT_FALSE(metrics.empty()); + ASSERT_TRUE(agent.get_rocp_agent()); + for(auto& metric : metrics) + { + /** + * Check profile construction + */ + rocprofiler_profile_config_id_t cfg_id = {}; + rocprofiler_counter_id_t id = {.handle = metric.id()}; + ROCPROFILER_CALL( + rocprofiler_create_profile_config(agent.get_rocp_agent()->id, &id, 1, &cfg_id), + "Unable to create profile"); + auto profile = counters::get_profile_config(cfg_id); + ASSERT_TRUE(profile); + EXPECT_EQ(counters::counter_callback_info::setup_profile_config(agent, profile), + ROCPROFILER_STATUS_SUCCESS) + << fmt::format("Could not build profile for {}", metric.name()); + + /** + * Check that a packet generator was created and there is an AST with constructed + * dimensions + */ + EXPECT_TRUE(profile->pkt_generator) << "No packet generator created"; + EXPECT_EQ(profile->asts.size(), 1); + EXPECT_FALSE(profile->asts.at(0).dimension_types().empty()); + + /** + * Check packet generation + */ + counters::counter_callback_info cb_info; + std::unique_ptr pkt; + EXPECT_EQ(cb_info.get_packet(pkt, agent, profile), ROCPROFILER_STATUS_SUCCESS) + << "Unable to generate packet"; + EXPECT_TRUE(pkt) << "Expected a packet to be generated"; + cb_info.packet_return_map.wlock([&](const auto& data) { + EXPECT_EQ(data.size(), 1) << "Incorrect packet size"; + const auto* ptr = common::get_val(data, pkt.get()); + EXPECT_TRUE(ptr) << "Could not find pkt"; + }); + + /** + * Check that required hardware counters match + */ + ASSERT_TRUE(profile->agent); + auto name_str = std::string(profile->agent->name); + auto req_counters = + counters::get_required_hardware_counters(counters::get_ast_map(), name_str, metric); + for(const auto& req_metric : *req_counters) + { + if(req_metric.special().empty()) + { + EXPECT_GT(profile->reqired_hw_counters.count(req_metric), 0) + << "Could not find metric - " << req_metric.name(); + profile->reqired_hw_counters.erase(req_metric); + } + else + { + EXPECT_GT(profile->required_special_counters.count(req_metric), 0) + << "Could not find metric - " << req_metric.name(); + profile->required_special_counters.erase(req_metric); + } + } + EXPECT_TRUE(profile->required_special_counters.empty()) + << "Special counters list larger than expected"; + EXPECT_TRUE(profile->reqired_hw_counters.empty()) + << "HW Counter list larger than expected"; + } + } +} + +namespace rocprofiler +{ +namespace hsa +{ +class FakeQueue : public Queue +{ +public: + FakeQueue(const AgentCache& a, rocprofiler_queue_id_t id) + : Queue(a) + , _agent(a) + , _id(id) + {} + virtual const AgentCache& get_agent() const override final { return _agent; }; + virtual rocprofiler_queue_id_t get_id() const override final { return _id; }; + + ~FakeQueue() {} + +private: + const AgentCache& _agent; + rocprofiler_queue_id_t _id = {}; +}; + +} // namespace hsa +} // namespace rocprofiler + +namespace +{ +struct expected_dispatch +{ + // To pass back + rocprofiler_profile_config_id_t id; + + rocprofiler_queue_id_t queue_id; + const rocprofiler_agent_t* agent; + rocprofiler_correlation_id_t correlation_id; + hsa_kernel_dispatch_packet_t* dispatch_packet; + uint64_t kernel_id; + rocprofiler_profile_config_id_t* config; +}; + +void +user_dispatch_cb(rocprofiler_queue_id_t queue_id, + const rocprofiler_agent_t* agent, + rocprofiler_correlation_id_t correlation_id, + const hsa_kernel_dispatch_packet_t* dispatch_packet, + uint64_t kernel_id, + void* callback_data_args, + rocprofiler_profile_config_id_t* config) +{ + expected_dispatch& expected = *static_cast(callback_data_args); + ASSERT_EQ(expected.agent, agent); + ASSERT_EQ(expected.queue_id.handle, queue_id.handle); + ASSERT_EQ(expected.correlation_id.internal, correlation_id.internal); + ASSERT_EQ(expected.correlation_id.external.ptr, correlation_id.external.ptr); + ASSERT_EQ(expected.correlation_id.external.value, correlation_id.external.value); + ASSERT_EQ(expected.dispatch_packet, dispatch_packet); + ASSERT_EQ(expected.kernel_id, kernel_id); + config->handle = expected.id.handle; +} + +} // namespace + +namespace rocprofiler +{ +namespace buffer +{ +uint64_t +get_buffer_offset(); +} +} // namespace rocprofiler + +TEST(core, check_callbacks) +{ + int64_t count = 0; + ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); + test_init(); + + registration::init_logging(); + registration::set_init_status(-1); + context::push_client(1); + ROCPROFILER_CALL(rocprofiler_create_context(&get_client_ctx()), "context creation failed"); + + auto agents = hsa::get_queue_controller().get_supported_agents(); + ASSERT_GT(agents.size(), 0); + hsa::get_queue_controller().disable_serialization(); + + for(const auto& [_, agent] : agents) + { + /** + * Setup + */ + rocprofiler_queue_id_t qid = {.handle = static_cast(count++)}; + hsa::FakeQueue fq(agent, qid); + auto metrics = findDeviceMetrics(agent, {}); + ASSERT_FALSE(metrics.empty()); + ASSERT_TRUE(agent.get_rocp_agent()); + for(auto& metric : metrics) + { + /** + * Do not check expression evaluation here. This is checked as part of evaluate_ast + * tests in a more controlled manner (aka, not requiring construction of the AST here). + */ + if(!metric.expression().empty()) continue; + + /** + * Setup + */ + expected_dispatch expected = {}; + rocprofiler_counter_id_t id = {.handle = metric.id()}; + ROCPROFILER_CALL( + rocprofiler_create_profile_config(agent.get_rocp_agent()->id, &id, 1, &expected.id), + "Unable to create profile"); + auto profile = counters::get_profile_config(expected.id); + ASSERT_TRUE(profile); + + std::shared_ptr cb_info = + std::make_shared(); + cb_info->user_cb = user_dispatch_cb; + cb_info->callback_args = static_cast(&expected); + + context::correlation_id corr_id; + corr_id.internal = count++; + + hsa::rocprofiler_packet pkt; + pkt.ext_amd_aql_pm4.header = count++; + + expected.correlation_id = {.internal = corr_id.internal, + .external = context::null_user_data}; + expected.dispatch_packet = &pkt.kernel_dispatch; + expected.kernel_id = count++; + expected.queue_id = qid; + expected.agent = fq.get_agent().get_rocp_agent(); + + hsa::Queue::queue_info_session_t::external_corr_id_map_t extern_ids = {}; + + auto ret_pkt = + counters::queue_cb(cb_info, fq, pkt, expected.kernel_id, extern_ids, &corr_id); + + ASSERT_TRUE(ret_pkt) << fmt::format("Expected a packet to be generated for - {}", + metric.name()); + + /** + * Fake some data for the counter + */ + size_t* fake_data = static_cast(ret_pkt->profile.output_buffer.ptr); + for(size_t i = 0; i < (ret_pkt->profile.output_buffer.size / sizeof(size_t)); i++) + { + fake_data[i] = i + 1; + } + + /** + * Create the buffer and run test + */ + rocprofiler_buffer_id_t opt_buff_id = {.handle = 0}; + buf_check check = { + .expected_size = ret_pkt->profile.output_buffer.size / sizeof(size_t), + .is_special = !metric.special().empty(), + .special_val = (metric.special().empty() ? 0.0 + : double(counters::get_agent_property( + std::string_view(metric.name()), + *agent.get_rocp_agent())))}; + + ROCPROFILER_CALL(rocprofiler_create_buffer(get_client_ctx(), + 500 * sizeof(size_t), + 500 * sizeof(size_t), + ROCPROFILER_BUFFER_POLICY_LOSSLESS, + buffered_callback, + &check, + &opt_buff_id), + "Could not create buffer"); + cb_info->buffer = opt_buff_id; + hsa::Queue::queue_info_session_t sess = {.queue = fq, .correlation_id = &corr_id}; + + counters::inst_pkt_t pkts; + pkts.emplace_back( + std::make_pair(std::move(ret_pkt), static_cast(0))); + completed_cb(cb_info, fq, pkt, sess, pkts); + rocprofiler_flush_buffer(opt_buff_id); + rocprofiler_destroy_buffer(opt_buff_id); + } + } + registration::set_init_status(1); + + registration::finalize(); +} + +TEST(core, destroy_counter_profile) +{ + ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); + test_init(); + + registration::init_logging(); + registration::set_init_status(-1); + context::push_client(1); + ROCPROFILER_CALL(rocprofiler_create_context(&get_client_ctx()), "context creation failed"); + + auto agents = hsa::get_queue_controller().get_supported_agents(); + ASSERT_GT(agents.size(), 0); + for(const auto& [_, agent] : agents) + { + auto metrics = findDeviceMetrics(agent, {}); + ASSERT_FALSE(metrics.empty()); + ASSERT_TRUE(agent.get_rocp_agent()); + for(auto& metric : metrics) + { + expected_dispatch expected = {}; + rocprofiler_counter_id_t id = {.handle = metric.id()}; + ROCPROFILER_CALL( + rocprofiler_create_profile_config(agent.get_rocp_agent()->id, &id, 1, &expected.id), + "Unable to create profile"); + ROCPROFILER_CALL(rocprofiler_destroy_profile_config(expected.id), + "Could not delete profile id"); + /** + * Check the profile was actually destroyed + */ + auto profile = counters::get_profile_config(expected.id); + EXPECT_FALSE(profile); + } + } + registration::set_init_status(1); + + registration::finalize(); +} + +TEST(core, start_stop_ctx) +{ + ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); + test_init(); + + registration::init_logging(); + registration::set_init_status(-1); + context::push_client(1); + ROCPROFILER_CALL(rocprofiler_create_context(&get_client_ctx()), "context creation failed"); + + rocprofiler_buffer_id_t opt_buff_id = {.handle = 0}; + ROCPROFILER_CALL(rocprofiler_create_buffer(get_client_ctx(), + 500 * sizeof(size_t), + 500 * sizeof(size_t), + ROCPROFILER_BUFFER_POLICY_LOSSLESS, + null_buffered_callback, + nullptr, + &opt_buff_id), + "Could not create buffer"); + + ROCPROFILER_CALL(rocprofiler_configure_buffered_dispatch_profile_counting_service( + get_client_ctx(), opt_buff_id, null_dispatch_callback, (void*) 0x12345), + "Could not setup buffered service"); + ROCPROFILER_CALL(rocprofiler_start_context(get_client_ctx()), "start context"); + + /** + * Check that the context was actually started + */ + auto* ctx_p = context::get_mutable_registered_context(get_client_ctx()); + ASSERT_TRUE(ctx_p); + auto& ctx = *ctx_p; + + ASSERT_TRUE(ctx.counter_collection); + ASSERT_EQ(ctx.counter_collection->callbacks.size(), 1); + EXPECT_EQ(ctx.counter_collection->callbacks.at(0)->user_cb, null_dispatch_callback); + EXPECT_EQ(ctx.counter_collection->callbacks.at(0)->callback_args, (void*) 0x12345); + EXPECT_EQ(ctx.counter_collection->callbacks.at(0)->context.handle, get_client_ctx().handle); + ASSERT_TRUE(ctx.counter_collection->callbacks.at(0)->buffer); + EXPECT_EQ(ctx.counter_collection->callbacks.at(0)->buffer->handle, opt_buff_id.handle); + + bool found = false; + ctx.counter_collection->enabled.rlock([&](const auto& data) { found = data; }); + EXPECT_TRUE(found); + + found = false; + hsa::get_queue_controller().iterate_callbacks([&](auto cid, const auto&) { + if(cid == ctx.counter_collection->callbacks.at(0)->queue_id) + { + found = true; + } + }); + EXPECT_TRUE(found); + + /** + * Check if context can be disabled correctly + */ + ROCPROFILER_CALL(rocprofiler_stop_context(get_client_ctx()), "stop context"); + + found = false; + hsa::get_queue_controller().iterate_callbacks([&](auto cid, const auto&) { + if(cid == ctx.counter_collection->callbacks.at(0)->queue_id) + { + found = true; + } + }); + EXPECT_FALSE(found); + found = false; + ctx.counter_collection->enabled.rlock([&](const auto& data) { found = data; }); + EXPECT_FALSE(found); + + rocprofiler_flush_buffer(opt_buff_id); + rocprofiler_destroy_buffer(opt_buff_id); + + registration::set_init_status(1); + + registration::finalize(); +} + +TEST(core, public_api_iterate_agents) +{ + ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); + test_init(); + + registration::init_logging(); + registration::set_init_status(-1); + context::push_client(1); + auto agents = hsa::get_queue_controller().get_supported_agents(); + for(const auto& [_, agent] : agents) + { + std::set from_api; + + // Iterate through the agents and get the counters available on that agent + ROCPROFILER_CALL(rocprofiler_iterate_agent_supported_counters( + agent.get_rocp_agent()->id, + [](rocprofiler_agent_id_t, + rocprofiler_counter_id_t* counters, + size_t num_counters, + void* user_data) { + std::set* vec = + static_cast*>(user_data); + for(size_t i = 0; i < num_counters; i++) + { + vec->insert(counters[i].handle); + } + return ROCPROFILER_STATUS_SUCCESS; + }, + static_cast(&from_api)), + "Could not fetch supported counters"); + + auto expected = findDeviceMetrics(agent, {}); + for(const auto& x : expected) + { + ASSERT_GT(from_api.count(x.id()), 0); + from_api.erase(x.id()); + } + EXPECT_TRUE(from_api.empty()); + } +} \ No newline at end of file diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp index 3e0d46ffe6..a8290c8b81 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp @@ -22,8 +22,51 @@ #include +#include +#include +#include +#include +#include + +#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" +#include "rocprofiler-sdk/registration.h" + +namespace +{ +void +check_dim_pos(rocprofiler_counter_instance_id_t test_id, + rocprofiler::counters::rocprofiler_profile_counter_instance_types dim, + size_t expected) +{ + EXPECT_EQ(rec_to_dim_pos(test_id, dim), expected); + size_t pos = 0; + rocprofiler_query_record_dimension_position( + test_id, static_cast(dim), &pos); + EXPECT_EQ(pos, expected); +} + +void +check_counter_id(rocprofiler_counter_instance_id_t id, uint64_t expected_handle) +{ + rocprofiler_counter_id_t api_id = {.handle = 0}; + rocprofiler_query_record_counter_id(id, &api_id); + EXPECT_EQ(rocprofiler::counters::rec_to_counter_id(id).handle, expected_handle); + EXPECT_EQ(rocprofiler::counters::rec_to_counter_id(id).handle, api_id.handle); +} +} // namespace TEST(dimension, set_get) { @@ -41,7 +84,7 @@ TEST(dimension, set_get) set_counter_in_rec(test_id, test_counter); // 0x0141000000000000 = decimal counter id 321 << DIM_BIT_LENGTH EXPECT_EQ(test_id, 0x0141000000000000); - EXPECT_EQ(rec_to_counter_id(test_id).handle, 321); + check_counter_id(test_id, 321); // Test multiples of i, setting/getting those values across all // dimensions @@ -51,7 +94,7 @@ TEST(dimension, set_get) { auto dim = static_cast(i); set_dim_in_rec(test_id, dim, i); - EXPECT_EQ(rec_to_dim_pos(test_id, dim), i); + check_dim_pos(test_id, dim, i); set_dim_in_rec(test_id, dim, i * multi_factor); for(size_t j = 1; j < static_cast(ROCPROFILER_DIMENSION_LAST); j++) { @@ -59,10 +102,10 @@ TEST(dimension, set_get) set_dim_in_rec(test_id, static_cast(j), max_counter_val); - EXPECT_EQ(rec_to_dim_pos( - test_id, static_cast(j)), - max_counter_val); - EXPECT_EQ(rec_to_dim_pos(test_id, dim), i * multi_factor); + check_dim_pos(test_id, + static_cast(j), + max_counter_val); + check_dim_pos(test_id, dim, i * multi_factor); } for(size_t j = static_cast(ROCPROFILER_DIMENSION_LAST - 1); j > 0; j--) @@ -71,9 +114,9 @@ TEST(dimension, set_get) set_dim_in_rec(test_id, static_cast(j), max_counter_val); - EXPECT_EQ(rec_to_dim_pos(test_id, (rocprofiler_profile_counter_instance_types) j), - max_counter_val); - EXPECT_EQ(rec_to_dim_pos(test_id, dim), i * multi_factor); + check_dim_pos( + test_id, (rocprofiler_profile_counter_instance_types) j, max_counter_val); + check_dim_pos(test_id, dim, i * multi_factor); } // Check that name exists @@ -87,18 +130,217 @@ TEST(dimension, set_get) { auto dim = static_cast(i); set_dim_in_rec(test_id, dim, i * 5); - EXPECT_EQ(rec_to_dim_pos(test_id, dim), i * 5); + check_dim_pos(test_id, dim, i * 5); set_dim_in_rec(test_id, dim, i * 3); - EXPECT_EQ(rec_to_dim_pos(test_id, dim), i * 3); + check_dim_pos(test_id, dim, i * 3); } test_counter.handle = 123; set_counter_in_rec(test_id, test_counter); - EXPECT_EQ(rec_to_counter_id(test_id).handle, 123); + check_counter_id(test_id, 123); // Test that all bits can be set/fetched for dims, 0xFAFBFCFDFEFF is a random // collection of 48 bits. set_dim_in_rec(test_id, ROCPROFILER_DIMENSION_NONE, 0xFAFBFCFDFEFF); - EXPECT_EQ(rec_to_dim_pos(test_id, ROCPROFILER_DIMENSION_NONE), 0xFAFBFCFDFEFF); - EXPECT_EQ(rec_to_counter_id(test_id).handle, 123); + check_dim_pos(test_id, ROCPROFILER_DIMENSION_NONE, 0xFAFBFCFDFEFF); + check_counter_id(test_id, 123); +} + +using namespace rocprofiler; + +namespace +{ +AmdExtTable& +get_ext_table() +{ + static auto _v = []() { + auto val = AmdExtTable{}; + val.hsa_amd_memory_pool_get_info_fn = hsa_amd_memory_pool_get_info; + val.hsa_amd_agent_iterate_memory_pools_fn = hsa_amd_agent_iterate_memory_pools; + val.hsa_amd_memory_pool_allocate_fn = hsa_amd_memory_pool_allocate; + val.hsa_amd_memory_pool_free_fn = hsa_amd_memory_pool_free; + val.hsa_amd_agent_memory_pool_get_info_fn = hsa_amd_agent_memory_pool_get_info; + val.hsa_amd_agents_allow_access_fn = hsa_amd_agents_allow_access; + return val; + }(); + return _v; +} + +CoreApiTable& +get_api_table() +{ + static auto _v = []() { + auto val = CoreApiTable{}; + val.hsa_iterate_agents_fn = hsa_iterate_agents; + val.hsa_agent_get_info_fn = hsa_agent_get_info; + val.hsa_queue_create_fn = hsa_queue_create; + val.hsa_queue_destroy_fn = hsa_queue_destroy; + return val; + }(); + return _v; +} + +auto +findDeviceMetrics(const hsa::AgentCache& agent, const std::unordered_set& metrics) +{ + std::vector ret; + auto all_counters = counters::getMetricMap(); + + LOG(ERROR) << "Looking up counters for " << std::string(agent.name()); + auto gfx_metrics = common::get_val(*all_counters, std::string(agent.name())); + if(!gfx_metrics) + { + LOG(ERROR) << "No counters found for " << std::string(agent.name()); + return ret; + } + + for(auto& counter : *gfx_metrics) + { + if(metrics.count(counter.name()) > 0 || metrics.empty()) + { + ret.push_back(counter); + } + } + return ret; +} + +void +test_init() +{ + HsaApiTable table; + table.amd_ext_ = &get_ext_table(); + table.core_ = &get_api_table(); + agent::construct_agent_cache(&table); + hsa::get_queue_controller().init(get_api_table(), get_ext_table()); +} + +} // namespace + +TEST(dimension, block_dim_test) +{ + ASSERT_EQ(hsa_init(), HSA_STATUS_SUCCESS); + test_init(); + + auto agents = hsa::get_queue_controller().get_supported_agents(); + ASSERT_GT(agents.size(), 0); + for(const auto& [_, agent] : agents) + { + auto metrics = findDeviceMetrics(agent, {}); + ASSERT_FALSE(metrics.empty()); + ASSERT_TRUE(agent.get_rocp_agent()); + // aql::AQLPacketConstruct pkt(agent, metrics); + // auto test_pkt = pkt.construct_packet(get_ext_table()); + for(const auto& metric : metrics) + { + /** + * Calculate expected dimensions from AQL Profiler + */ + std::unordered_map + rocp_dims; + LOG(ERROR) << metric.name() << " " << metric.special(); + if(!metric.special().empty()) + { + rocp_dims[counters::rocprofiler_profile_counter_instance_types:: + ROCPROFILER_DIMENSION_INSTANCE] = 1; + } + else if(!metric.expression().empty()) + { + continue; + } + else + { + aql::AQLPacketConstruct pkt_gen(agent, {metric}); + const auto& events = pkt_gen.get_counter_events(metric); + for(const auto& event : events) + { + std::map dims; + auto status = aql::get_dim_info(agent.get_hsa_agent(), event, 0, dims); + CHECK_EQ(status, ROCPROFILER_STATUS_SUCCESS) + << rocprofiler_get_status_string(status); + for(const auto& [id, extent] : dims) + { + if(const auto* inst_type = rocprofiler::common::get_val( + counters::aqlprofile_id_to_rocprof_instance(), id)) + { + rocp_dims.emplace(*inst_type, 0).first->second = extent; + } + } + } + } + + /** + * Compare with actual + */ + auto dims = getBlockDimensions(agent.name(), metric); + EXPECT_FALSE(dims.empty()); + EXPECT_EQ(dims.size(), rocp_dims.size()); + for(const auto& dim : dims) + { + const auto* ptr = rocprofiler::common::get_val(rocp_dims, dim.type()); + ASSERT_TRUE(ptr) << fmt::format("{}", dim); + EXPECT_EQ(*ptr, dim.size()) << fmt::format("{}", dim); + EXPECT_EQ(std::string(counters::dimension_map().at(dim.type())), dim.name()) + << fmt::format("{}", dim); + } + + /** + * Check this value exists in the dimension cache + */ + const auto* dim_cache = + rocprofiler::common::get_val(counters::get_dimension_cache(), metric.id()); + ASSERT_TRUE(dim_cache); + EXPECT_EQ(fmt::format("{}", fmt::join(dims, "|")), + fmt::format("{}", fmt::join(*dim_cache, "|"))); + + /** + * Check counter instance count public API + */ + size_t instance_count = 0; + size_t calculated_instance_count = 0; + rocprofiler_query_counter_instance_count( + agent.get_rocp_agent()->id, {.handle = metric.id()}, &instance_count); + for(const auto& dim : dims) + { + if(calculated_instance_count == 0) + calculated_instance_count = dim.size(); + else if(dim.size() > 0) + calculated_instance_count = dim.size() * calculated_instance_count; + } + EXPECT_EQ(instance_count, calculated_instance_count); + + /** + * Check the public API returns this value + */ + rocprofiler_iterate_counter_dimensions( + {.handle = metric.id()}, + [](rocprofiler_counter_id_t, + const rocprofiler_record_dimension_info_t* dim_info, + size_t num_dims, + void* user_data) -> rocprofiler_status_t { + auto expected_dims = *static_cast< + std::unordered_map*>(user_data); + EXPECT_EQ(num_dims, expected_dims.size()); + for(size_t i = 0; i < num_dims; i++) + { + const auto* lookup_ptr = rocprofiler::common::get_val( + expected_dims, + static_cast( + dim_info[i].id)); + EXPECT_TRUE(lookup_ptr); + if(!lookup_ptr) return ROCPROFILER_STATUS_ERROR; + EXPECT_EQ(*lookup_ptr, dim_info[i].instance_size); + EXPECT_EQ( + counters::dimension_map().at( + static_cast( + dim_info[i].id)), + std::string(dim_info[i].name)); + } + return ROCPROFILER_STATUS_SUCCESS; + }, + static_cast(&rocp_dims)); + } + } + + hsa_shut_down(); } diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/init_order.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/init_order.cpp index 2c78e48aac..7e618616da 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/init_order.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/init_order.cpp @@ -106,6 +106,8 @@ private: MetricIdMap copy_ = *CHECK_NOTNULL(rocprofiler::counters::getMetricIdMap()); }; +namespace +{ metric_map_order& get_metric_map() { @@ -145,6 +147,7 @@ get_buffer() static rocprofiler_buffer_id_t buf = {}; return buf; } +} // namespace // Test that metrics map remains in scope at exit TEST(counters_init_order, metric_map_order) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/metrics_test.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/metrics_test.cpp index 5b0238680b..d6259492a6 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/metrics_test.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/metrics_test.cpp @@ -27,6 +27,8 @@ #include +#include + #include "lib/rocprofiler-sdk/agent.hpp" #include "lib/rocprofiler-sdk/counters/metrics.hpp" @@ -175,3 +177,17 @@ TEST(metrics, check_agent_valid) } } } + +TEST(metrics, check_public_api_query) +{ + const auto* id_map = counters::getMetricIdMap(); + for(const auto& [id, metric] : *id_map) + { + const char* name = nullptr; + size_t size = 0; + ASSERT_EQ(rocprofiler_query_counter_name({.handle = id}, &name, &size), + ROCPROFILER_STATUS_SUCCESS); + EXPECT_EQ(std::string(name), metric.name()); + EXPECT_EQ(size, metric.name().size()); + } +} \ No newline at end of file diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp index d2b05b9337..95f9c05c93 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp @@ -156,10 +156,14 @@ public: AmdExtTable ext_api, hsa_queue_t** queue); - ~Queue(); + Queue(const AgentCache& agent) + : _agent(agent) + {} - const hsa_queue_t* intercept_queue() const { return _intercept_queue; }; - const AgentCache& get_agent() const { return _agent; } + virtual ~Queue(); + + const hsa_queue_t* intercept_queue() const { return _intercept_queue; }; + virtual const AgentCache& get_agent() const { return _agent; } void create_signal(uint32_t attribute, hsa_signal_t* signal) const; void signal_async_handler(const hsa_signal_t& signal, Queue::queue_info_session_t* data) const; @@ -167,7 +171,7 @@ public: template void signal_callback(FuncT&& func) const; - rocprofiler_queue_id_t get_id() const; + virtual rocprofiler_queue_id_t get_id() const; // Fast check to see if we have any callbacks we need to notify int get_notifiers() const { return _notifiers; } diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp index 106bdacd37..179025d4a5 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp @@ -269,6 +269,12 @@ QueueController::profiler_serializer(FuncT&& lambda) _profiler_serializer.wlock(std::forward(lambda)); } +void +QueueController::disable_serialization() +{ + profiler_serializer([](auto& serializer) { serializer.enabled = false; }); +} + namespace { /* @@ -291,6 +297,7 @@ profiler_serializer_ready_signal_handler(hsa_signal_value_t /* signal_value */, auto* hsa_queue = static_cast(data); const auto* queue = get_queue_controller().get_queue(*hsa_queue); get_queue_controller().profiler_serializer([&](auto& serializer) { + if(!serializer.enabled) return; { std::lock_guard cv_lock(queue->cv_mutex); if(queue->get_state() == queue_state::to_destroy) @@ -322,6 +329,7 @@ void profiler_serializer_kernel_completion_signal(hsa_signal_t queue_block_signal) { get_queue_controller().profiler_serializer([queue_block_signal](auto& serializer) { + if(!serializer.enabled) return; assert(serializer.dispatch_queue != nullptr); serializer.dispatch_queue = nullptr; get_queue_controller().get_core_table().hsa_signal_store_screlease_fn(queue_block_signal, @@ -370,6 +378,17 @@ QueueController::iterate_queues(const queue_iterator_cb_t& cb) const }); } +void +QueueController::iterate_callbacks(const callback_iterator_cb_t& cb) const +{ + _callback_cache.rlock([&cb](const auto& map) { + for(const auto& [cid, tuple] : map) + { + cb(cid, tuple); + } + }); +} + QueueController& get_queue_controller() { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp index 9d4832447e..aaba4ca309 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp @@ -54,13 +54,17 @@ struct profiler_serializer_t { const Queue* dispatch_queue{nullptr}; std::deque dispatch_ready; + bool enabled{true}; }; // Tracks and manages HSA queues class QueueController { public: - using queue_iterator_cb_t = std::function; + using agent_callback_tuple_t = + std::tuple; + using queue_iterator_cb_t = std::function; + using callback_iterator_cb_t = std::function; QueueController() = default; // Initializes the QueueInterceptor. This must be delayed until @@ -96,9 +100,16 @@ public: template void profiler_serializer(FuncT&& lambda); + void iterate_callbacks(const callback_iterator_cb_t&) const; + + /** + * Disable serialization for QueueController, has no effect if counter collection + * is not in use (which defaults to no serialization mechanism). Should only be used for + * testing. + */ + void disable_serialization(); + private: - using agent_callback_tuple_t = - std::tuple; using queue_map_t = std::unordered_map>; using client_id_map_t = std::unordered_map; using agent_cache_map_t = std::unordered_map; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocprofiler.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocprofiler.cpp index 585408403a..69bc37487d 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocprofiler.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocprofiler.cpp @@ -75,7 +75,18 @@ ROCPROFILER_STATUS_STRING(ROCPROFILER_STATUS_ERROR_METRIC_NOT_VALID_FOR_AGENT, "Metric is not valid for the agent") ROCPROFILER_STATUS_STRING(ROCPROFILER_STATUS_ERROR_FINALIZED, "Invalid request because rocprofiler has finalized") - +ROCPROFILER_STATUS_STRING(ROCPROFILER_STATUS_ERROR_HSA_NOT_LOADED, + "Function call requires that HSA is loaded") +ROCPROFILER_STATUS_STRING(ROCPROFILER_STATUS_ERROR_DIM_NOT_FOUND, + "Dimension is not found for counter") +ROCPROFILER_STATUS_STRING(ROCPROFILER_STATUS_ERROR_PROFILE_COUNTER_NOT_FOUND, + "Profile could not find counter for GPU") +ROCPROFILER_STATUS_STRING(ROCPROFILER_STATUS_ERROR_AST_GENERATION_FAILED, + "AST could not be generated correctly") +ROCPROFILER_STATUS_STRING(ROCPROFILER_STATUS_ERROR_AST_NOT_FOUND, "AST was not found") +ROCPROFILER_STATUS_STRING( + ROCPROFILER_STATUS_ERROR_AQL_NO_EVENT_COORD, + "AQL Profiler was not able to find event coordinates for defined counters") template const char* get_status_name(rocprofiler_status_t status, std::index_sequence)