AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

Add support for AQL dimensions (#262)

Telusuri Sumber 
* 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 <bwelton@users.noreply.github.com>

* Minor fix

* Test change

---------

Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
Co-authored-by: bwelton <bwelton@users.noreply.github.com>

[ROCm/rocprofiler-sdk commit: 3eb6a27bc6]
This commit is contained in:
Benjamin Welton
2024-02-07 20:03:21 -08:00
committed by GitHub
orang tua 4daf20d431
melakukan 8e620cc11d
29 mengubah file dengan 1670 tambahan dan 293 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocprofiler-sdk/samples/counter_collection/main.cpp
+1 -1
Melihat File
@@ -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;
projects/rocprofiler-sdk/samples/counter_collection/print_functional_counters.cpp
+130 -10
Melihat File
@@ -13,6 +13,8 @@
#include <rocprofiler-sdk/registration.h>
#include <rocprofiler-sdk/rocprofiler.h>
#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<validate_dim_presence>());
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<std::pair<rocprofiler_record_dimension_info_t, size_t>>& 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<rocprofiler_record_dimension_info_t, size_t> vector_pos;
std::vector<std::unique_ptr<validate_dim_presence>> sub_vectors;
bool has_value{false};
};
struct CaptureRecords
{
std::shared_mutex m_mutex{};
// <counter id handle, expected instances>
std::map<uint64_t, size_t> expected{};
std::map<uint64_t, std::string> expected_counter_names{};
std::vector<rocprofiler_counter_id_t> remaining{};
std::map<uint64_t, size_t> expected{};
// expected dims that we should see data for
std::map<uint64_t, validate_dim_presence> expected_data_dims{};
std::map<uint64_t, std::string> expected_counter_names{};
std::vector<rocprofiler_counter_id_t> remaining{};
// <counter_id handle, instances seen>
std::map<uint64_t, size_t> 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<uint64_t, size_t> 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<rocprofiler_record_counter_t*>(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<validate_dim_presence*>(user_data);
for(size_t i = 0; i < num_dims; i++)
{
dim_presence->maybe_forward(dim_info[i]);
}
return ROCPROFILER_STATUS_SUCCESS;
},
static_cast<void*>(&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<std::pair<rocprofiler_record_dimension_info_t, size_t>> 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();
}
}
}
}
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/counters.h
+26 -10
Melihat File
@@ -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
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/dispatch_profile.h
+7
Melihat File
@@ -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
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h
+10 -1
Melihat File
@@ -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;
/**
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/CMakeLists.txt
+1 -1
Melihat File
@@ -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})
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/aql_profile_v2.h
+33
Melihat File
@@ -0,0 +1,33 @@
// This file should be removed when it appears in AQL Profile
#pragma ONCE
#include <hsa/hsa.h>
#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);
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/helpers.cpp
+54 -1
Melihat File
@@ -25,6 +25,12 @@
#include <fmt/core.h>
#include <glog/logging.h>
#include <rocprofiler-sdk/fwd.h>
#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<rocprofiler_counter_instance_id_t*>(userdata),
*rocprof_id,
static_cast<size_t>(coordinate));
}
return HSA_STATUS_SUCCESS;
};
if(hsa_ven_amd_aqlprofile_iterate_event_coord(
agent, event, sample_id, callback, static_cast<void*>(&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<int, uint64_t>& dims)
{
auto callback = [](int, int id, int extent, int, const char*, void* userdata) -> hsa_status_t {
auto& map = *static_cast<std::map<int, uint64_t>*>(userdata);
map.emplace(id, extent);
return HSA_STATUS_SUCCESS;
};
if(hsa_ven_amd_aqlprofile_iterate_event_coord(
agent, event, sample_id, callback, static_cast<void*>(&dims)) != HSA_STATUS_SUCCESS)
{
return ROCPROFILER_STATUS_ERROR_AQL_NO_EVENT_COORD;
}
return ROCPROFILER_STATUS_SUCCESS;
}
} // namespace aql
} // namespace rocprofiler
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/helpers.hpp
+18
Melihat File
@@ -23,9 +23,13 @@
#pragma once
#include <functional>
#include <map>
#include <string>
#include <hsa/hsa_ven_amd_aqlprofile.h>
#include <rocprofiler-sdk/fwd.h>
#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<int, uint64_t>& 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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp
+13
Melihat File
@@ -190,6 +190,19 @@ AQLPacketConstruct::event_to_metric(const hsa_ven_amd_aqlprofile_event_t& event)
return nullptr;
}
const std::vector<hsa_ven_amd_aqlprofile_event_t>&
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()
{
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.hpp
+5
Melihat File
@@ -56,6 +56,11 @@ public:
const counters::Metric* event_to_metric(const hsa_ven_amd_aqlprofile_event_t& event) const;
std::vector<hsa_ven_amd_aqlprofile_event_t> get_all_events() const;
const std::vector<hsa_ven_amd_aqlprofile_event_t>& get_counter_events(
const counters::Metric&) const;
hsa_agent_t hsa_agent() const { return _agent.get_hsa_agent(); }
private:
struct AQLProfileMetric
{
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters.cpp
+43 -52
Melihat File
@@ -25,9 +25,12 @@
#include <fmt/core.h>
#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<size_t>(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<rocprofiler::counters::rocprofiler_profile_counter_instance_types>(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<rocprofiler_record_dimension_info_t, 6> 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<rocprofiler_counter_dimension_id_t>(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;
}
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.cpp
+114 -80
Melihat File
@@ -37,9 +37,6 @@ namespace rocprofiler
{
namespace counters
{
using ClientID = int64_t;
using inst_pkt_t = common::container::
small_vector<std::pair<std::unique_ptr<rocprofiler::hsa::AQLPacket>, ClientID>, 4>;
class CounterController
{
public:
@@ -133,6 +130,113 @@ destroy_counter_profile(uint64_t id)
get_controller().destroy_profile(id);
}
std::shared_ptr<profile_config>
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_config>& 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<rocprofiler::aql::AQLPacketConstruct>(
agent,
std::vector<counters::Metric>{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<rocprofiler::hsa::AQLPacket>& ret_pkt,
const hsa::AgentCache& agent,
std::shared_ptr<profile_config>& 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<counter_callback_info>& info,
CHECK(prof_config);
std::unique_ptr<rocprofiler::hsa::AQLPacket> 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<rocprofiler::aql::AQLPacketConstruct>(
queue.get_agent(),
std::vector<counters::Metric>{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<counter_callback_info>& 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(
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.hpp
+28
Melihat File
@@ -88,6 +88,13 @@ struct counter_callback_info
rocprofiler::common::Synchronized<
std::unordered_map<rocprofiler::hsa::AQLPacket*, std::shared_ptr<profile_config>>>
packet_return_map{};
static rocprofiler_status_t setup_profile_config(const hsa::AgentCache&,
std::shared_ptr<profile_config>&);
rocprofiler_status_t get_packet(std::unique_ptr<rocprofiler::hsa::AQLPacket>&,
const hsa::AgentCache&,
std::shared_ptr<profile_config>&);
};
uint64_t
@@ -107,5 +114,26 @@ start_context(const context::context*);
void
stop_context(const context::context*);
std::unique_ptr<rocprofiler::hsa::AQLPacket>
queue_cb(const std::shared_ptr<counter_callback_info>& 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<std::pair<std::unique_ptr<rocprofiler::hsa::AQLPacket>, ClientID>, 4>;
void
completed_cb(const std::shared_ptr<counter_callback_info>&,
const hsa::Queue&,
hsa::rocprofiler_packet,
const hsa::Queue::queue_info_session_t&,
inst_pkt_t& pkts);
std::shared_ptr<profile_config> get_profile_config(rocprofiler_profile_config_id_t);
} // namespace counters
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/dimensions.cpp
+68 -19
Melihat File
@@ -29,10 +29,12 @@
#include <fmt/core.h>
#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<MetricDimension>{
{dimension_map().at(ROCPROFILER_DIMENSION_NONE), 1, ROCPROFILER_DIMENSION_NONE}};
return std::vector<MetricDimension>{{dimension_map().at(ROCPROFILER_DIMENSION_INSTANCE),
1,
ROCPROFILER_DIMENSION_INSTANCE}};
}
std::unordered_map<rocprofiler_profile_counter_instance_types, uint64_t> count;
std::vector<MetricDimension> 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<MetricDimension>{
// {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<MetricDimension>{
{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<int, uint64_t> 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<uint64_t, std::vector<MetricDimension>>&
get_dimension_cache()
{
static auto*& cache =
common::static_object<std::unordered_map<uint64_t, std::vector<MetricDimension>>>::
construct([]() -> std::unordered_map<uint64_t, std::vector<MetricDimension>> {
std::unordered_map<uint64_t, std::vector<MetricDimension>> 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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/dimensions.hpp
+24 -37
Melihat File
@@ -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<std::string, std::vector<MetricDimension>>;
// // map agent_name -> BlockDimensionMap
// using AgentBlockDimensionsMap = std::unordered_map<std::string, BlockDimensionMap&>;
// // map dimension_id -> MetricDimension
// using DimensionIdMap = std::unordered_map<uint64_t, MetricDimension>;
// // get the complete AgentBlockDimensionsMap
// const AgentBlockDimensionsMap&
// getAgentBlockDimensionsMap();
// // get specific dimensions for an agent, block_name
// const std::vector<MetricDimension>&
// 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<MetricDimension>
getBlockDimensions(std::string_view agent, const counters::Metric&);
// // get a specific dimension by id
// const MetricDimension&
// getDimensionById(uint64_t id);
const std::unordered_map<uint64_t, std::vector<MetricDimension>>&
get_dimension_cache();
} // namespace counters
} // namespace rocprofiler
namespace fmt
{
// fmt::format support for metric
template <>
struct formatter<rocprofiler::counters::MetricDimension>
{
template <typename ParseContext>
constexpr auto parse(ParseContext& ctx)
{
return ctx.begin();
}
template <typename Ctx>
auto format(rocprofiler::counters::MetricDimension const& dims, Ctx& ctx) const
{
return fmt::format_to(ctx.out(), "[{}, {}]", dims.name(), dims.size());
}
};
} // namespace fmt
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp
+46 -39
Melihat File
@@ -21,13 +21,16 @@
// SOFTWARE.
#include "lib/rocprofiler-sdk/counters/evaluate_ast.hpp"
#include <fmt/core.h>
#include <exception>
#include <optional>
#include <numeric>
#include <optional>
#include <stdexcept>
#include <fmt/core.h>
#include <fmt/ranges.h>
#include <rocprofiler-sdk/rocprofiler.h>
#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<MetricDimension>
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<DimensionTypes>(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<MetricDimension>{{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<int>(_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<DimensionTypes>(final_dims);
// Reduction down to a single instance supported for now.
_dimension_types =
std::vector<MetricDimension>{{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<int64_t>(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<std::string, property_function_t> props = {
static std::unordered_map<std::string_view, property_function_t> 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;
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/evaluate_ast.hpp
+12 -8
Melihat File
@@ -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<MetricDimension> dimension of the output of this AST.
*/
DimensionTypes set_dimensions();
std::vector<MetricDimension> set_dimensions();
bool validate_raw_ast(const std::unordered_map<std::string, Metric>& metrics);
@@ -152,11 +152,11 @@ public:
const std::set<counters::Metric>& required_special_counters,
std::unordered_map<uint64_t, std::vector<rocprofiler_record_counter_t>>& out_map);
NodeType type() const { return _type; }
ReduceOperation reduce_op() const { return _reduce_op; }
const std::vector<EvaluateAST>& 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<EvaluateAST>& children() const { return _children; }
const Metric& metric() const { return _metric; }
const std::vector<MetricDimension>& 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<rocprofiler_record_counter_t>& 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<EvaluateAST> _children;
std::string _agent;
DimensionTypes _dimension_types{DIMENSION_LAST};
std::vector<MetricDimension> _dimension_types{};
std::vector<rocprofiler_record_counter_t> _static_value;
std::unordered_set<rocprofiler_profile_counter_instance_types> _reduce_dimension_set;
bool _expanded{false};
@@ -197,6 +199,8 @@ std::optional<std::set<Metric>>
get_required_hardware_counters(const std::unordered_map<std::string, EvaluateASTMap>& 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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/id_decode.cpp
+51 -8
Melihat File
@@ -22,28 +22,71 @@
#include "lib/rocprofiler-sdk/counters/id_decode.hpp"
#include <hsa/hsa_ven_amd_aqlprofile.h>
#include <string>
#include <unordered_map>
#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<rocprofiler_profile_counter_instance_types, std::string>&
const std::unordered_map<rocprofiler_profile_counter_instance_types, std::string_view>&
dimension_map()
{
static std::unordered_map<rocprofiler_profile_counter_instance_types, std::string> 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<rocprofiler_profile_counter_instance_types, std::string_view> 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<int, rocprofiler_profile_counter_instance_types>&
aqlprofile_id_to_rocprof_instance()
{
using dims_map_t = std::unordered_map<int, rocprofiler_profile_counter_instance_types>;
static auto*& aql_to_rocprof_dims =
common::static_object<dims_map_t>::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<std::string_view,
rocprofiler_profile_counter_instance_types>
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<int, rocprofiler_profile_counter_instance_types>*>(
userdata);
map.emplace(id, *inst_type);
}
return HSA_STATUS_SUCCESS;
},
static_cast<void*>(&data));
return data;
}());
return *aql_to_rocprof_dims;
}
} // namespace counters
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/id_decode.hpp
+7 -2
Melihat File
@@ -22,6 +22,7 @@
#pragma once
#include <map>
#include <unordered_map>
#include <rocprofiler-sdk/fwd.h>
@@ -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<rocprofiler_profile_counter_instance_types, std::string>&
const std::unordered_map<rocprofiler_profile_counter_instance_types, std::string_view>&
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<int, rocprofiler_profile_counter_instance_types>&
aqlprofile_id_to_rocprof_instance();
} // namespace counters
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/CMakeLists.txt
+3 -2
Melihat File
@@ -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(
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/core.cpp
+648
Melihat File
@@ -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 <algorithm>
#include <cstdint>
#include <sstream>
#include <tuple>
#include <fmt/core.h>
#include <gtest/gtest.h>
#include <hsa/hsa.h>
#include <hsa/hsa_api_trace.h>
#include <hsa/hsa_ext_amd.h>
#include <rocprofiler-sdk/rocprofiler.h>
#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<std::string>& metrics)
{
std::vector<counters::Metric> 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<buf_check*>(user_data);
if(expected.is_special)
{
// Special values are single value constants (from agent_t)
expected.expected_size = 1;
}
std::set<double> seen_data;
std::set<uint64_t> 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<rocprofiler_record_counter_t*>(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<hsa::AQLPacket> 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<expected_dispatch*>(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<uint64_t>(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<counters::counter_callback_info> cb_info =
std::make_shared<counters::counter_callback_info>();
cb_info->user_cb = user_dispatch_cb;
cb_info->callback_args = static_cast<void*>(&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<size_t*>(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<counters::ClientID>(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<uint64_t> 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<uint64_t>* vec =
static_cast<std::set<uint64_t>*>(user_data);
for(size_t i = 0; i < num_counters; i++)
{
vec->insert(counters[i].handle);
}
return ROCPROFILER_STATUS_SUCCESS;
},
static_cast<void*>(&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());
}
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp
+256 -14
Melihat File
@@ -22,8 +22,51 @@
#include <gtest/gtest.h>
#include <fmt/core.h>
#include <hsa/hsa.h>
#include <hsa/hsa_api_trace.h>
#include <hsa/hsa_ext_amd.h>
#include <rocprofiler-sdk/rocprofiler.h>
#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<rocprofiler_counter_dimension_id_t>(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<rocprofiler_profile_counter_instance_types>(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<size_t>(ROCPROFILER_DIMENSION_LAST); j++)
{
@@ -59,10 +102,10 @@ TEST(dimension, set_get)
set_dim_in_rec(test_id,
static_cast<rocprofiler_profile_counter_instance_types>(j),
max_counter_val);
EXPECT_EQ(rec_to_dim_pos(
test_id, static_cast<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,
static_cast<rocprofiler_profile_counter_instance_types>(j),
max_counter_val);
check_dim_pos(test_id, dim, i * multi_factor);
}
for(size_t j = static_cast<size_t>(ROCPROFILER_DIMENSION_LAST - 1); j > 0; j--)
@@ -71,9 +114,9 @@ TEST(dimension, set_get)
set_dim_in_rec(test_id,
static_cast<rocprofiler_profile_counter_instance_types>(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<rocprofiler_profile_counter_instance_types>(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<std::string>& metrics)
{
std::vector<counters::Metric> 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<counters::rocprofiler_profile_counter_instance_types, uint64_t>
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<int, uint64_t> 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<counters::rocprofiler_profile_counter_instance_types,
uint64_t>*>(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<counters::rocprofiler_profile_counter_instance_types>(
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<counters::rocprofiler_profile_counter_instance_types>(
dim_info[i].id)),
std::string(dim_info[i].name));
}
return ROCPROFILER_STATUS_SUCCESS;
},
static_cast<void*>(&rocp_dims));
}
}
hsa_shut_down();
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/init_order.cpp
+3
Melihat File
@@ -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)
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/metrics_test.cpp
+16
Melihat File
@@ -27,6 +27,8 @@
#include <algorithm>
#include <rocprofiler-sdk/rocprofiler.h>
#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());
}
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp
+8 -4
Melihat File
@@ -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 <typename FuncT>
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; }
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp
+19
Melihat File
@@ -269,6 +269,12 @@ QueueController::profiler_serializer(FuncT&& lambda)
_profiler_serializer.wlock(std::forward<FuncT>(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<hsa_queue_t*>(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<std::mutex> 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()
{
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.hpp
+14 -3
Melihat File
@@ -54,13 +54,17 @@ struct profiler_serializer_t
{
const Queue* dispatch_queue{nullptr};
std::deque<const Queue*> dispatch_ready;
bool enabled{true};
};
// Tracks and manages HSA queues
class QueueController
{
public:
using queue_iterator_cb_t = std::function<void(const Queue*)>;
using agent_callback_tuple_t =
std::tuple<rocprofiler_agent_t, Queue::queue_cb_t, Queue::completed_cb_t>;
using queue_iterator_cb_t = std::function<void(const Queue*)>;
using callback_iterator_cb_t = std::function<void(ClientID, const agent_callback_tuple_t&)>;
QueueController() = default;
// Initializes the QueueInterceptor. This must be delayed until
@@ -96,9 +100,16 @@ public:
template <typename FuncT>
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<rocprofiler_agent_t, Queue::queue_cb_t, Queue::completed_cb_t>;
using queue_map_t = std::unordered_map<hsa_queue_t*, std::unique_ptr<Queue>>;
using client_id_map_t = std::unordered_map<ClientID, agent_callback_tuple_t>;
using agent_cache_map_t = std::unordered_map<uint32_t, AgentCache>;
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocprofiler.cpp
+12 -1
Melihat File
@@ -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 <size_t Idx, size_t... Tail>
const char*
get_status_name(rocprofiler_status_t status, std::index_sequence<Idx, Tail...>)