From 34e62464ee4a6e860a7585784c2128eedced4c81 Mon Sep 17 00:00:00 2001 From: gobhardw Date: Mon, 13 Feb 2023 20:03:51 +0530 Subject: [PATCH] SWDEV-297569: Support for xgmi counters Change-Id: If14c0f63821b721fe719d3b76915afc2f2333d5c --- inc/rocprofiler.h | 3 +- plugin/att/att.cpp | 25 ++-- samples/CMakeLists.txt | 8 ++ .../xgmi_counters_sampler_example.cpp | 92 +++++++++++++ src/core/counters/mmio/df_counters_mi200.cpp | 76 +++++++++++ src/core/counters/mmio/df_counters_mi200.h | 44 ++++++ .../mmio/df_perfmon_registers_mi200.h | 32 +++++ src/core/counters/mmio/mmio.cpp | 128 +++++++++--------- src/core/counters/mmio/mmio.h | 94 +++++++------ src/core/session/counters_sampler.cpp | 76 ++++++----- 10 files changed, 430 insertions(+), 148 deletions(-) create mode 100644 samples/counters_sampler/xgmi_counters_sampler_example.cpp create mode 100644 src/core/counters/mmio/df_counters_mi200.cpp create mode 100644 src/core/counters/mmio/df_counters_mi200.h create mode 100644 src/core/counters/mmio/df_perfmon_registers_mi200.h diff --git a/inc/rocprofiler.h b/inc/rocprofiler.h index 82bc4252b3..4784a4894f 100644 --- a/inc/rocprofiler.h +++ b/inc/rocprofiler.h @@ -1904,7 +1904,8 @@ typedef struct { } rocprofiler_spm_parameter_t; typedef enum{ - ROCPROFILER_COUNTERS_SAMPLER_PCIE_COUNTERS = 0 + ROCPROFILER_COUNTERS_SAMPLER_PCIE_COUNTERS = 0, + ROCPROFILER_COUNTERS_SAMPLER_XGMI_COUNTERS = 1 } rocprofiler_counters_sampler_counter_type_t; typedef struct{ diff --git a/plugin/att/att.cpp b/plugin/att/att.cpp index aa646867c6..00227ca8d7 100644 --- a/plugin/att/att.cpp +++ b/plugin/att/att.cpp @@ -55,14 +55,14 @@ class att_plugin_t { bool is_valid_{true}; inline bool att_file_exists(const std::string& name) { - struct stat buffer; + struct stat buffer; return stat(name.c_str(), &buffer) == 0; } bool IsValid() const { return is_valid_; } void FlushATTRecord(const rocprofiler_record_att_tracer_t* att_tracer_record, - rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id) { + rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id) { std::lock_guard lock(writing_lock); if (!att_tracer_record) { @@ -124,8 +124,8 @@ class att_plugin_t { } int WriteBufferRecords(const rocprofiler_record_header_t* begin, - const rocprofiler_record_header_t* end, - rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id) { + const rocprofiler_record_header_t* end, rocprofiler_session_id_t session_id, + rocprofiler_buffer_id_t buffer_id) { while (begin < end) { if (!begin) return 0; switch (begin->kind) { @@ -133,13 +133,13 @@ class att_plugin_t { case ROCPROFILER_TRACER_RECORD: case ROCPROFILER_PC_SAMPLING_RECORD: case ROCPROFILER_SPM_RECORD: + case ROCPROFILER_COUNTERS_SAMPLER_RECORD: printf("Invalid record Kind: %d", begin->kind); break; case ROCPROFILER_ATT_TRACER_RECORD: { - rocprofiler_record_att_tracer_t* att_record = - const_cast( - reinterpret_cast(begin)); + rocprofiler_record_att_tracer_t* att_record = const_cast( + reinterpret_cast(begin)); FlushATTRecord(att_record, session_id, buffer_id); break; } @@ -158,7 +158,7 @@ att_plugin_t* att_plugin = nullptr; } // namespace ROCPROFILER_EXPORT int rocprofiler_plugin_initialize(uint32_t rocprofiler_major_version, - uint32_t rocprofiler_minor_version) { + uint32_t rocprofiler_minor_version) { if (rocprofiler_major_version != ROCPROFILER_VERSION_MAJOR || rocprofiler_minor_version < ROCPROFILER_VERSION_MINOR) return -1; @@ -180,15 +180,16 @@ ROCPROFILER_EXPORT void rocprofiler_plugin_finalize() { att_plugin = nullptr; } -ROCPROFILER_EXPORT int rocprofiler_plugin_write_buffer_records( - const rocprofiler_record_header_t* begin, const rocprofiler_record_header_t* end, - rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id) { +ROCPROFILER_EXPORT int rocprofiler_plugin_write_buffer_records(const rocprofiler_record_header_t* begin, + const rocprofiler_record_header_t* end, + rocprofiler_session_id_t session_id, + rocprofiler_buffer_id_t buffer_id) { if (!att_plugin || !att_plugin->IsValid()) return -1; return att_plugin->WriteBufferRecords(begin, end, session_id, buffer_id); } ROCPROFILER_EXPORT int rocprofiler_plugin_write_record(rocprofiler_record_tracer_t record, - rocprofiler_session_id_t session_id) { + rocprofiler_session_id_t session_id) { if (!att_plugin || !att_plugin->IsValid()) return -1; if (record.header.id.handle == 0) return 0; return 0; diff --git a/samples/CMakeLists.txt b/samples/CMakeLists.txt index 3e939f6e4d..544bc2bd9d 100644 --- a/samples/CMakeLists.txt +++ b/samples/CMakeLists.txt @@ -92,6 +92,14 @@ target_link_libraries(pcie_counters_sampler PRIVATE ${ROCPROFILER_TARGET} system add_dependencies(samples pcie_counters_sampler) install(TARGETS pcie_counters_sampler RUNTIME DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/${PROJECT_NAME}/samples COMPONENT samples) +## Build XGMI Counters Sampling example +set_source_files_properties(counters_sampler/xgmi_counters_sampler_example.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) +hip_add_executable(xgmi_counters_sampler counters_sampler/xgmi_counters_sampler_example.cpp ${ROCPROFILER_UTIL_SRC_FILES}) +target_include_directories(xgmi_counters_sampler PRIVATE ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/inc ${CMAKE_CURRENT_SOURCE_DIR}/common) +target_link_libraries(xgmi_counters_sampler PRIVATE ${ROCPROFILER_TARGET} systemd amd_comgr) +add_dependencies(samples xgmi_counters_sampler) +install(TARGETS xgmi_counters_sampler RUNTIME DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/${PROJECT_NAME}/samples COMPONENT samples) + # ################################################################################################################ # ############################################################################################################################################ diff --git a/samples/counters_sampler/xgmi_counters_sampler_example.cpp b/samples/counters_sampler/xgmi_counters_sampler_example.cpp new file mode 100644 index 0000000000..e360fd92a9 --- /dev/null +++ b/samples/counters_sampler/xgmi_counters_sampler_example.cpp @@ -0,0 +1,92 @@ +#include "../common/common.h" + +int main(int argc, char** argv) { + int* gpuMem; + int counter_option = 0; + if (argc > 1) counter_option = atoi(argv[1]); + + prepare(); + // Initialize the tools + CHECK_ROCPROFILER(rocprofiler_initialize()); + + // Creating the session with given replay mode + rocprofiler_session_id_t session_id; + CHECK_ROCPROFILER(rocprofiler_create_session(ROCPROFILER_KERNEL_REPLAY_MODE, &session_id)); + + // Creating Output Buffer for the data + rocprofiler_buffer_id_t buffer_id; + CHECK_ROCPROFILER(rocprofiler_create_buffer( + session_id, + [](const rocprofiler_record_header_t* record, const rocprofiler_record_header_t* end_record, + rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id) { + WriteBufferRecords(record, end_record, session_id, buffer_id); + }, + 0x999999, &buffer_id)); + + // Counters Sampler Filter + rocprofiler_filter_id_t filter_id; + [[maybe_unused]] rocprofiler_filter_property_t property = {}; + + std::vector xgmi_counters = {"xgmi_link0_data_outbound", "xgmi_link1_data_outbound", + "xgmi_link2_data_outbound", "xgmi_link3_data_outbound", + "xgmi_link4_data_outbound", "xgmi_link5_data_outbound", + "xgmi_link5_data_outbound", "xgmi_link6_data_outbound"}; + + rocprofiler_counters_sampler_counter_input_t counters_input[2] = { + {.name = const_cast(xgmi_counters[counter_option].c_str()), + .type = ROCPROFILER_COUNTERS_SAMPLER_XGMI_COUNTERS}}; + + uint32_t rate = 1000; + uint32_t duration = 5000; + + rocprofiler_counters_sampler_parameters_t cs_parameters = {.counters = counters_input, + .counters_num = 1, + .sampling_rate = rate, + .sampling_duration = duration, + .gpu_agent_index = 0}; + CHECK_ROCPROFILER( + rocprofiler_create_filter(session_id, ROCPROFILER_COUNTERS_SAMPLER, + rocprofiler_filter_data_t{.counters_sampler_parameters = cs_parameters}, + 0, &filter_id, property)); + CHECK_ROCPROFILER(rocprofiler_set_filter_buffer(session_id, filter_id, buffer_id)); + + // Normal HIP Calls + hipDeviceProp_t devProp; + HIP_CALL(hipGetDeviceProperties(&devProp, 0)); + HIP_CALL(hipMalloc((void**)&gpuMem, 1 * sizeof(int))); + + // KernelA and KernelB won't be profiled + kernelCalls('A'); + kernelCalls('B'); + + std::cout << "Collecting samples for: " << xgmi_counters[counter_option] + << " ; sampling rate: " << rate << " ms; duration: " << duration << " ms" << std::endl; + // Activating the session + CHECK_ROCPROFILER(rocprofiler_start_session(session_id)); + + // KernelC, KernelD, KernelE and KernelF to be profiled as part of the session + kernelCalls('C'); + kernelCalls('D'); + kernelCalls('E'); + kernelCalls('F'); + // Normal HIP Calls + HIP_CALL(hipFree(gpuMem)); + + // allow sampler to run for 10 secs + sleep(6); + + // Deactivating session + CHECK_ROCPROFILER(rocprofiler_terminate_session(session_id)); + + // Manual Flush user buffer request + CHECK_ROCPROFILER(rocprofiler_flush_data(session_id, buffer_id)); + + // Destroy sessions + CHECK_ROCPROFILER(rocprofiler_destroy_session(session_id)); + + // Destroy all profiling related objects(User buffer, sessions, filters, + // etc..) + CHECK_ROCPROFILER(rocprofiler_finalize()); + + return 0; +} \ No newline at end of file diff --git a/src/core/counters/mmio/df_counters_mi200.cpp b/src/core/counters/mmio/df_counters_mi200.cpp new file mode 100644 index 0000000000..4a0cd6631f --- /dev/null +++ b/src/core/counters/mmio/df_counters_mi200.cpp @@ -0,0 +1,76 @@ +#include "df_counters_mi200.h" +#include "df_perfmon_registers_mi200.h" +#include "mmio.h" + +namespace rocmtools { + +#define DF_V3_6_MAX_COUNTERS 4 + +/* get flags from df perfmon config */ +#define DF_V3_6_GET_EVENT(x) (x & 0xFFUL) +#define DF_V3_6_GET_INSTANCE(x) ((x >> 8) & 0xFFUL) +#define DF_V3_6_GET_UNITMASK(x) ((x >> 16) & 0xFFUL) +#define DF_V3_6_PERFMON_OVERFLOW 0xFFFFFFFFFFFFULL + +/* get ficaa value for accessing CakeDlwmActiveTransferCount */ +#define AMDGPU_PMU_SET_FICAA(o) ((o << 16) | 0x1AF5) + +DFPerfMonMI200::DFPerfMonMI200(const Agent::AgentInfo& info) : PerfMon(), mmio_(nullptr) { + mmio_ = dynamic_cast(mmio::MMIOManager::CreateMMIO(mmio::DF_PERFMON, info)); +} + +DFPerfMonMI200::~DFPerfMonMI200() { + mmio::MMIOManager::DestroyMMIOInstance(dynamic_cast(mmio_)); +} + +void DFPerfMonMI200::SetCounterNames(std::vector& counter_names) { + counter_names_ = counter_names; + // TODO: only one event at a time is supported at the moment + auto it = DF_MI200::xgmi_events_table.find(counter_names[0]); + if (it != DF_MI200::xgmi_events_table.end()) { + instance_id_ = it->second; + } +} + +void DFPerfMonMI200::writeRegister(uint32_t reg_offset, uint32_t value) { + mmio_->RegisterWriteAPI(reg_offset, value); +} + +void DFPerfMonMI200::readRegister(uint32_t reg_offset, uint32_t& value) { + mmio_->RegisterReadAPI(reg_offset, value); +} + +uint64_t DFPerfMonMI200::GetFicaNodeOutboundBw(uint32_t ficaa_val) { + uint32_t ficadl_val, ficadh_val; + + // setting up FICAA for address + writeRegister(DF_MI200::smnDF_PIE_AON_FabricIndirectConfigAccessAddress3, ficaa_val); + + // setting up FICADL for data + readRegister(DF_MI200::smnDF_PIE_AON_FabricIndirectConfigAccessDataLo3, ficadl_val); + + // setting up FICADH for data + readRegister(DF_MI200::smnDF_PIE_AON_FabricIndirectConfigAccessDataHi3, ficadh_val); + + return (((ficadh_val & 0xFFFFFFFFFFFFFFFF) << 32) | ficadl_val); +} + +void DFPerfMonMI200::SetFicaNodeOutboundBw(uint64_t node_instance, uint32_t& ficaa_in_val) { + uint32_t instance = node_instance; + ficaa_in_val = AMDGPU_PMU_SET_FICAA(instance); +} + + +void DFPerfMonMI200::Start() { + uint32_t ficaa_in_val; + uint32_t ficaa_out_val; + + SetFicaNodeOutboundBw(instance_id_, ficaa_in_val); + + ficaa_out_val = GetFicaNodeOutboundBw(ficaa_in_val); + printf("CakeDlwmActiveTransferCount_1=%u\n", ficaa_out_val); + ficaa_out_val = ficaa_out_val & 0xFFFFFFFF; + printf("CakeDlwmActiveTransferCount_2=%u\n", ficaa_out_val); +} + +} // namespace rocmtools diff --git a/src/core/counters/mmio/df_counters_mi200.h b/src/core/counters/mmio/df_counters_mi200.h new file mode 100644 index 0000000000..0b0962e2a1 --- /dev/null +++ b/src/core/counters/mmio/df_counters_mi200.h @@ -0,0 +1,44 @@ +#ifndef DF_COUNTERS_MI200_H +#define DF_COUNTERS_MI200_H + +#include +#include "mmio.h" +#include "perfmon.h" + +namespace rocmtools { + +/* + One perfmon per GPU. + Only one instance per GPU, per process +*/ + +class DFPerfMonMI200 : public PerfMon { + public: + DFPerfMonMI200(const Agent::AgentInfo& info); + ~DFPerfMonMI200(); + void Start() override; + void Stop() {}; + void Read(std::vector& values) {}; + void SetCounterNames(std::vector& counter_names); + mmio::mmap_type_t Type() override { return mmio::mmap_type_t::DF_PERFMON; } + + private: + void writeRegister(uint32_t reg_offset, uint32_t value); + void readRegister(uint32_t reg_offset, uint32_t& value); + + // outboud bandwidth for xgmi nodes + void SetFicaNodeOutboundBw(uint64_t node_instance, uint32_t& ficaa_in_val); + uint64_t GetFicaNodeOutboundBw(uint32_t ficaa_val); + + + + private: + mmio::DFPerfmonMMIO* mmio_; + static std::mutex mutex_; // should be an MMIO member + static DFPerfMonMI200* instance_; + uint64_t instance_id_; +}; + +} // namespace rocmtools + +#endif // DF_COUNTERS_MI200_H \ No newline at end of file diff --git a/src/core/counters/mmio/df_perfmon_registers_mi200.h b/src/core/counters/mmio/df_perfmon_registers_mi200.h new file mode 100644 index 0000000000..51ea9f9842 --- /dev/null +++ b/src/core/counters/mmio/df_perfmon_registers_mi200.h @@ -0,0 +1,32 @@ +#ifndef DF_PERFMON_REGISTERS_MI200_H +#define DF_PERFMON_REGISTERS_MI200_H + +#include +#include + +namespace DF_MI200 { + +const static uint32_t smnDF_PIE_AON_FabricIndirectConfigAccessAddress3 = 0x1d05cUL; +const static uint32_t smnDF_PIE_AON_FabricIndirectConfigAccessDataLo3 = 0x1d098UL; +const static uint32_t smnDF_PIE_AON_FabricIndirectConfigAccessDataHi3 = 0x1d09cUL; + +#define NUM_EVENT_TYPES_ALDEBARAN 1 +#define NUM_EVENTS_ALDEBARAN_XGMI 8 +#define NUM_EVENTS_ALDEBARAN_MAX NUM_EVENTS_ALDEBARAN_XGMI + +#define mmPCIE_INDEX2 0x000e +#define mmPCIE_INDEX2_BASE_IDX 0 +#define mmPCIE_DATA2 0x000f +#define mmPCIE_DATA2_BASE_IDX 0 + +/* MI200 events */ +const static std::map xgmi_events_table = { + {"xgmi_link0_data_outbound", 0x4b}, {"xgmi_link1_data_outbound", 0x4c}, + {"xgmi_link2_data_outbound", 0x4d}, {"xgmi_link3_data_outbound", 0x4e}, + {"xgmi_link4_data_outbound", 0x4f}, {"xgmi_link5_data_outbound", 0x50}, + {"xgmi_link6_data_outbound", 0x51}, {"xgmi_link7_data_outbound", 0x52}}; + +} // namespace DF_MI200 + + +#endif // DF_PERFMON_REGISTERS_MI200_H \ No newline at end of file diff --git a/src/core/counters/mmio/mmio.cpp b/src/core/counters/mmio/mmio.cpp index c336630568..c9af41312c 100644 --- a/src/core/counters/mmio/mmio.cpp +++ b/src/core/counters/mmio/mmio.cpp @@ -22,34 +22,37 @@ #include #include "../../../utils/helper.h" #include "pcie_perfmon_registers_mi200.h" - +#include "df_perfmon_registers_mi200.h" namespace rocmtools { namespace mmio { -void PrintFunctionPhase(const char* function_name, int phase){ +void PrintFunctionPhase(const char* function_name, int phase) { #if defined(DEBUG_TRACE) - if (phase == 0) - std::clog << function_name << "() START" << std::endl; - else - std::clog << function_name << "() END" << std::endl; + if (phase == 0) + std::clog << function_name << "() START" << std::endl; + else + std::clog << function_name << "() END" << std::endl; #endif } -void PrintRegisterData(uint32_t& index_value, uint32_t& data_value, const char* function_name, int phase){ +void PrintRegisterData(uint32_t& index_value, uint32_t& data_value, const char* function_name, + int phase) { #if defined(DEBUG_TRACE) - if (phase == 0) { - PrintFunctionPhase(function_name, phase); - std::clog << "Old (index, data) : "<< std::hex << index_value << " " << data_value << std::endl; - } - else { - std::clog << "New (index, data) : "<< std::hex << index_value << " " << data_value << std::endl; - PrintFunctionPhase(function_name, phase); - } + if (phase == 0) { + PrintFunctionPhase(function_name, phase); + std::clog << "Old (index, data) : " << std::hex << index_value << " " << data_value + << std::endl; + } else { + std::clog << "New (index, data) : " << std::hex << index_value << " " << data_value + << std::endl; + PrintFunctionPhase(function_name, phase); + } #endif } -MMIO::MMIO(const Agent::AgentInfo& info) : agent_info_(&info), pci_memory_(nullptr), type_(DEFAULT_MMAP) { +MMIO::MMIO(const Agent::AgentInfo& info) + : agent_info_(&info), pci_memory_(nullptr), type_(DEFAULT_MMAP) { const auto pci_domain = agent_info_->getPCIDomain(); const auto pci_location_id = agent_info_->getPCILocationID(); @@ -67,10 +70,9 @@ MMIO::MMIO(const Agent::AgentInfo& info) : agent_info_(&info), pci_memory_(nullp pci_device_->regions[region].size <= (4096UL * 1024)) { pci_memory_size_ = pci_device_->regions[region].size; int err = pci_device_map_range(pci_device_, pci_device_->regions[region].base_addr, - pci_device_->regions[region].size, PCI_DEV_MAP_FLAG_WRITABLE, - (void**)&pci_memory_); - if(err) - fatal("failed to map the registers. Error code: %d\n", err); + pci_device_->regions[region].size, PCI_DEV_MAP_FLAG_WRITABLE, + (void**)&pci_memory_); + if (err) fatal("failed to map the registers. Error code: %d\n", err); } if (pci_memory_ == nullptr) fatal("could not find the pci memory address\n"); @@ -79,46 +81,45 @@ MMIO::MMIO(const Agent::AgentInfo& info) : agent_info_(&info), pci_memory_(nullp } MMIO::~MMIO() { - if (pci_memory_) - { + if (pci_memory_) { int err = pci_device_unmap_range(pci_device_, pci_memory_, pci_memory_size_); - if(err) warning("failed to unmap the pci memory. Error code: %d\n", err); + if (err) warning("failed to unmap the pci memory. Error code: %d\n", err); } } -bool MMIO::RegisterWriteAPI(uint32_t reg_offset, uint32_t value){ - // access the mmap - // write register offset to index register 0x38 of index/data pair (indirect addressing) - // write register bits to data register 0x3c of index/data pair (indirect addressing) - - // std::lock_guard lock(mutex_); - PrintRegisterData(*index_reg_addr, *data_reg_addr, __FUNCTION__, 0); +bool MMIO::RegisterWriteAPI(uint32_t reg_offset, uint32_t value) { + // access the mmap + // write register offset to index register 0x38 of index/data pair (indirect addressing) + // write register bits to data register 0x3c of index/data pair (indirect addressing) - // TODO: should work only if map is created + // std::lock_guard lock(mutex_); + PrintRegisterData(*index_reg_addr, *data_reg_addr, __FUNCTION__, 0); - *index_reg_addr = reg_offset; - *data_reg_addr = value; + // TODO: should work only if map is created - PrintRegisterData(*index_reg_addr, *data_reg_addr, __FUNCTION__, 1); - return true; + *index_reg_addr = reg_offset; + *data_reg_addr = value; + + PrintRegisterData(*index_reg_addr, *data_reg_addr, __FUNCTION__, 1); + return true; } -bool MMIO::RegisterReadAPI(uint32_t reg_offset, uint32_t& value){ - // access the mmap - // write register offset to index register 0x38 of index/data pair (indirect addressing) - // read register bits to data register 0x3c of index/data pair (indirect addressing) - - // std::lock_guard lock(mutex_); - PrintRegisterData(*index_reg_addr, *data_reg_addr, __FUNCTION__, 0); +bool MMIO::RegisterReadAPI(uint32_t reg_offset, uint32_t& value) { + // access the mmap + // write register offset to index register 0x38 of index/data pair (indirect addressing) + // read register bits to data register 0x3c of index/data pair (indirect addressing) - // TODO: should work only if map is created - - *index_reg_addr = reg_offset; - // TODO: add delay here?? - value = *data_reg_addr; + // std::lock_guard lock(mutex_); + PrintRegisterData(*index_reg_addr, *data_reg_addr, __FUNCTION__, 0); - PrintRegisterData(*index_reg_addr, *data_reg_addr, __FUNCTION__, 1); - return true; + // TODO: should work only if map is created + + *index_reg_addr = reg_offset; + // TODO: add delay here?? + value = *data_reg_addr; + + PrintRegisterData(*index_reg_addr, *data_reg_addr, __FUNCTION__, 1); + return true; } @@ -127,13 +128,18 @@ MMIO* MMIOManager::CreateMMIO(mmap_type_t type, const Agent::AgentInfo& info) { switch (type) { case PCIE_PERFMON: { mmio = GetMMIOInstance(type, info); - if(mmio == nullptr){ + if (mmio == nullptr) { mmio = dynamic_cast(new PciePerfmonMMIO(info)); AddInstance(mmio); } break; } case DF_PERFMON: { + mmio = GetMMIOInstance(type, info); + if (mmio == nullptr) { + mmio = dynamic_cast(new DFPerfmonMMIO(info)); + AddInstance(mmio); + } break; } case UMC_PERFMON: { @@ -149,9 +155,9 @@ MMIO* MMIOManager::CreateMMIO(mmap_type_t type, const Agent::AgentInfo& info) { MMIO* MMIOManager::GetMMIOInstance(mmap_type_t type, const Agent::AgentInfo& info) { MMIO* mmio = nullptr; auto it = mmio_instances_.find(info.getHandle()); - if(it != mmio_instances_.end()){ - for(auto& mmio_instance: it->second){ - if(mmio_instance->Type() == type){ + if (it != mmio_instances_.end()) { + for (auto& mmio_instance : it->second) { + if (mmio_instance->Type() == type) { mmio = mmio_instance; } } @@ -165,25 +171,23 @@ void MMIOManager::AddInstance(MMIO* in_mmio_instance) { } void MMIOManager::DestroyMMIOInstance(MMIO* in_mmio_instance) { - if(in_mmio_instance == nullptr) - return; + if (in_mmio_instance == nullptr) return; uint64_t handle = in_mmio_instance->GetAgentInfo().getHandle(); auto it = mmio_instances_.find(handle); - if(it != mmio_instances_.end()){ + if (it != mmio_instances_.end()) { auto& mmio_array = it->second; // find instance in the array and remove it from the array - mmio_array.erase(std::remove(mmio_array.begin(), mmio_array.end(), in_mmio_instance), mmio_array.end()); + mmio_array.erase(std::remove(mmio_array.begin(), mmio_array.end(), in_mmio_instance), + mmio_array.end()); } delete in_mmio_instance; } -std::map> MMIOManager::mmio_instances_; +std::map> MMIOManager::mmio_instances_; -} // namespace mmio +} // namespace mmio -} // namespace rocmtools - - \ No newline at end of file +} // namespace rocmtools diff --git a/src/core/counters/mmio/mmio.h b/src/core/counters/mmio/mmio.h index eb1e6dd69c..a0c1bb7629 100644 --- a/src/core/counters/mmio/mmio.h +++ b/src/core/counters/mmio/mmio.h @@ -18,8 +18,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ - #ifndef SRC_CORE_NON_GFXIP_COUNTERS_MMIO_H - #define SRC_CORE_NON_GFXIP_COUNTERS_MMIO_H +#ifndef SRC_CORE_NON_GFXIP_COUNTERS_MMIO_H +#define SRC_CORE_NON_GFXIP_COUNTERS_MMIO_H #include #include "src/core/hardware/hsa_info.h" @@ -36,13 +36,14 @@ namespace rocmtools { namespace mmio { #define FUNCTION_START() PrintFunctionPhase(__FUNCTION__, 0) -#define FUNCTION_END() PrintFunctionPhase(__FUNCTION__, 1) +#define FUNCTION_END() PrintFunctionPhase(__FUNCTION__, 1) // uncomment below to see register write sequences // #define DEBUG_TRACE = 1 void PrintFunctionPhase(const char* function_name, int phase); -void PrintRegisterData(uint32_t& index_value, uint32_t& data_value, const char* function_name, int phase); +void PrintRegisterData(uint32_t& index_value, uint32_t& data_value, const char* function_name, + int phase); // Index/Data registers const static uint32_t INDIRECT_REG_INDEX = 0x38; @@ -53,50 +54,58 @@ typedef enum { DEFAULT_MMAP, DF_PERFMON, UMC_PERFMON, PCIE_PERFMON } mmap_type_t class MMIOManager; class MMIO { - public: - virtual bool RegisterWriteAPI(uint32_t reg_offset, uint32_t value); - virtual bool RegisterReadAPI(uint32_t reg_offset, uint32_t& value); - virtual void SetIndexDataRegisters(const uint32_t index_reg, const uint32_t data_reg) { - index_reg_addr = (uint32_t*)((char*)pci_memory_ + index_reg); - data_reg_addr = (uint32_t*)((char*)pci_memory_ + data_reg); - } + public: + virtual bool RegisterWriteAPI(uint32_t reg_offset, uint32_t value); + virtual bool RegisterReadAPI(uint32_t reg_offset, uint32_t& value); + virtual void SetIndexDataRegisters(const uint32_t index_reg, const uint32_t data_reg) { + index_reg_addr = (uint32_t*)((char*)pci_memory_ + index_reg); + data_reg_addr = (uint32_t*)((char*)pci_memory_ + data_reg); + } - MMIO(MMIO& other) = delete; - void operator=(const MMIO&) = delete; - virtual ~MMIO(); - friend class MMIOManager; + MMIO(MMIO& other) = delete; + void operator=(const MMIO&) = delete; + virtual ~MMIO(); + friend class MMIOManager; - const Agent::AgentInfo& GetAgentInfo() { return *agent_info_; } - mmap_type_t Type() { return type_; } + const Agent::AgentInfo& GetAgentInfo() { return *agent_info_; } + mmap_type_t Type() { return type_; } - protected: - MMIO(const Agent::AgentInfo& info); + protected: + MMIO(const Agent::AgentInfo& info); - // default constructor; helpful for derived classes - // which want to setup mmio construction differently - MMIO(){ type_ = DEFAULT_MMAP; }; + // default constructor; helpful for derived classes + // which want to setup mmio construction differently + MMIO() { type_ = DEFAULT_MMAP; }; - const Agent::AgentInfo* agent_info_; - struct pci_device* pci_device_; - size_t pci_memory_size_; - uint32_t* pci_memory_; - mmap_type_t type_; + const Agent::AgentInfo* agent_info_; + struct pci_device* pci_device_; + size_t pci_memory_size_; + uint32_t* pci_memory_; + mmap_type_t type_; - uint32_t* index_reg_addr; - uint32_t* data_reg_addr; + uint32_t* index_reg_addr; + uint32_t* data_reg_addr; }; -// PciePerfmonMMIO has same mmio setup approach as +// PciePerfmonMMIO has same mmio setup approach as // done in MMIO class -class PciePerfmonMMIO : public MMIO{ - public: - friend class MMIOManager; - protected: - PciePerfmonMMIO(const Agent::AgentInfo& info): MMIO(info) { - type_ = PCIE_PERFMON; - }; +class PciePerfmonMMIO : public MMIO { + public: + friend class MMIOManager; + + protected: + PciePerfmonMMIO(const Agent::AgentInfo& info) : MMIO(info) { type_ = PCIE_PERFMON; }; }; +// DFPerfmonMMIO has same mmio setup approach as +// done in MMIO class +class DFPerfmonMMIO : public MMIO { + public: + friend class MMIOManager; + + protected: + DFPerfmonMMIO(const Agent::AgentInfo& info) : MMIO(info) { type_ = DF_PERFMON; }; +}; /* Class to manage mmio for UMC/DF/PCIe etc. The mmio approach for the different IPs may @@ -104,11 +113,12 @@ class PciePerfmonMMIO : public MMIO{ the same mmio and index/data registers */ class MMIOManager { - public: - static MMIO* CreateMMIO(mmap_type_t type, const Agent::AgentInfo& info); - static MMIO* GetMMIOInstance(mmap_type_t type, const Agent::AgentInfo& info); - static void DestroyMMIOInstance(MMIO* instance); - private: + public: + static MMIO* CreateMMIO(mmap_type_t type, const Agent::AgentInfo& info); + static MMIO* GetMMIOInstance(mmap_type_t type, const Agent::AgentInfo& info); + static void DestroyMMIOInstance(MMIO* instance); + + private: static void AddInstance(MMIO* instance); static std::map> mmio_instances_; }; diff --git a/src/core/session/counters_sampler.cpp b/src/core/session/counters_sampler.cpp index 18adb675fd..7ab43ee5d7 100644 --- a/src/core/session/counters_sampler.cpp +++ b/src/core/session/counters_sampler.cpp @@ -22,23 +22,23 @@ #include "src/core/hsa/hsa_support.h" #include "src/api/rocmtool.h" #include "src/core/counters/mmio/pcie_counters_mi200.h" +#include "src/core/counters/mmio/df_counters_mi200.h" namespace rocmtools { -CountersSampler::CountersSampler( - rocprofiler_buffer_id_t buffer_id, - rocprofiler_filter_id_t filter_id, - rocprofiler_session_id_t session_id) -: buffer_id_(buffer_id) -, filter_id_(filter_id) -, session_id_(session_id) -, pci_system_initialized_(pci_system_init() == 0) +CountersSampler::CountersSampler(rocprofiler_buffer_id_t buffer_id, + rocprofiler_filter_id_t filter_id, + rocprofiler_session_id_t session_id) + : buffer_id_(buffer_id), + filter_id_(filter_id), + session_id_(session_id), + pci_system_initialized_(pci_system_init() == 0) { params_ = rocmtools::GetROCMToolObj() - ->GetSession(session_id_) - ->GetFilter(filter_id_) - ->GetCountersSamplerParameterData(); + ->GetSession(session_id_) + ->GetFilter(filter_id_) + ->GetCountersSamplerParameterData(); std::vector agents; rocmtools::hsa_support::GetCoreApiTable().hsa_iterate_agents_fn( @@ -57,25 +57,39 @@ CountersSampler::CountersSampler( // PCIE counters std::vector pcie_counter_names; for (int i = 0; i < params_.counters_num; i++) { - if(params_.counters[i].type == ROCPROFILER_COUNTERS_SAMPLER_PCIE_COUNTERS) + if (params_.counters[i].type == ROCPROFILER_COUNTERS_SAMPLER_PCIE_COUNTERS) pcie_counter_names.push_back(params_.counters[i].name); } if (pcie_counter_names.size() > 0) { auto agentInfo = rocmtools::hsa_support::GetAgentInfo(agents[params_.gpu_agent_index].handle); - if(agentInfo.getName() == "gfx90a"){ + if (agentInfo.getName() == "gfx90a") { PciePerfMonMI200* perfmon = new PciePerfMonMI200(agentInfo); perfmon->SetCounterNames(pcie_counter_names); perfmon_instances_.push_back(perfmon); } } + // XGMI counters + std::vector xgmi_counter_names; + for (int i = 0; i < params_.counters_num; i++) { + if (params_.counters[i].type == ROCPROFILER_COUNTERS_SAMPLER_XGMI_COUNTERS) + xgmi_counter_names.push_back(params_.counters[i].name); + } + + if (xgmi_counter_names.size() > 0) { + auto agentInfo = rocmtools::hsa_support::GetAgentInfo(agents[params_.gpu_agent_index].handle); + if (agentInfo.getName() == "gfx90a") { + DFPerfMonMI200* perfmon = new DFPerfMonMI200(agentInfo); + perfmon->SetCounterNames(xgmi_counter_names); + perfmon_instances_.push_back(perfmon); + } + } } CountersSampler::~CountersSampler() { // cleanup perfmon instancess - for (auto &perfmon : perfmon_instances_) { - if(perfmon != nullptr) - delete perfmon; + for (auto& perfmon : perfmon_instances_) { + if (perfmon != nullptr) delete perfmon; } // clean up libpcieaccess resources // TODO: should be part of mmio class in future @@ -86,7 +100,9 @@ CountersSampler::~CountersSampler() { } void CountersSampler::Start() { - if (sampler_thread_.joinable()) { return; } + if (sampler_thread_.joinable()) { + return; + } std::cout << "Sampler Start\n"; // Start all Perfmons @@ -100,7 +116,9 @@ void CountersSampler::Start() { } void CountersSampler::Stop() { - if (!sampler_thread_.joinable()) { return; } + if (!sampler_thread_.joinable()) { + return; + } std::cout << "Sampler Stop\n"; // Stop all Perfmons @@ -113,20 +131,17 @@ void CountersSampler::Stop() { sampler_thread_.join(); } -void CountersSampler::AddRecord(rocprofiler_record_counters_sampler_t &record) { +void CountersSampler::AddRecord(rocprofiler_record_counters_sampler_t& record) { const auto tool = rocmtools::GetROCMToolObj(); const auto session = tool->GetSession(session_id_); const auto buffer = session->GetBuffer(buffer_id_); std::lock_guard lk(session->GetSessionLock()); - record.header = { - ROCPROFILER_COUNTERS_SAMPLER_RECORD, - { tool->GetUniqueRecordId() } - }; + record.header = {ROCPROFILER_COUNTERS_SAMPLER_RECORD, {tool->GetUniqueRecordId()}}; - // Add the record to the buffer(a deep-copy operation) along with - // a lambda function to deep-copy the record.counters member to + // Add the record to the buffer(a deep-copy operation) along with + // a lambda function to deep-copy the record.counters member to // the newly created buffer record buffer->AddRecord( record, record.counters, @@ -139,7 +154,7 @@ void CountersSampler::AddRecord(rocprofiler_record_counters_sampler_t &record) { void CountersSampler::SamplerLoop() { std::this_thread::sleep_until(std::chrono::steady_clock::now() + - std::chrono::milliseconds(params_.initial_delay)); + std::chrono::milliseconds(params_.initial_delay)); uint32_t elapsed = 0; while (keep_running_ && (elapsed <= params_.sampling_duration)) { auto next_tick = @@ -147,12 +162,11 @@ void CountersSampler::SamplerLoop() { rocprofiler_record_counters_sampler_t record; std::vector values; - for (auto& perfmon : perfmon_instances_){ + for (auto& perfmon : perfmon_instances_) { perfmon->Read(values); } - record.counters = - static_cast( - malloc(values.size() * sizeof(rocprofiler_counters_sampler_counter_output_t))); + record.counters = static_cast( + malloc(values.size() * sizeof(rocprofiler_counters_sampler_counter_output_t))); ::memcpy(record.counters, &(values)[0], values.size() * sizeof(rocprofiler_counters_sampler_counter_output_t)); record.num_counters = values.size(); @@ -165,4 +179,4 @@ void CountersSampler::SamplerLoop() { } } -} // namespace rocmtools \ No newline at end of file +} // namespace rocmtools \ No newline at end of file