SWDEV-297569: Support for xgmi counters

Change-Id: If14c0f63821b721fe719d3b76915afc2f2333d5c

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{

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<std::mutex> 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<rocprofiler_record_att_tracer_t*>(
-                  reinterpret_cast<const rocprofiler_record_att_tracer_t*>(begin));
+          rocprofiler_record_att_tracer_t* att_record = const_cast<rocprofiler_record_att_tracer_t*>(
+              reinterpret_cast<const rocprofiler_record_att_tracer_t*>(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;

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)
+
 # ################################################################################################################
 
 # ############################################################################################################################################

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<std::string> 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<char*>(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;
+}

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::DFPerfmonMMIO*>(mmio::MMIOManager::CreateMMIO(mmio::DF_PERFMON, info));
+}
+
+DFPerfMonMI200::~DFPerfMonMI200() {
+  mmio::MMIOManager::DestroyMMIOInstance(dynamic_cast<mmio::MMIO*>(mmio_));
+}
+
+void DFPerfMonMI200::SetCounterNames(std::vector<std::string>& 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

src/core/counters/mmio/df_counters_mi200.h

@@ -0,0 +1,44 @@
+#ifndef DF_COUNTERS_MI200_H
+#define DF_COUNTERS_MI200_H
+
+#include <cstdint>
+#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<rocprofiler_counters_sampler_counter_output_t>& values) {};
+  void SetCounterNames(std::vector<std::string>& 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

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 <stdint.h>
+#include <map>
+
+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<std::string, uint64_t> 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

src/core/counters/mmio/mmio.cpp

@@ -22,34 +22,37 @@
 #include <cstdint>
 #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<std::mutex> 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<std::mutex> 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<std::mutex> 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<std::mutex> 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<MMIO*>(new PciePerfmonMMIO(info));
         AddInstance(mmio);
       }
       break;
     }
     case DF_PERFMON: {
+      mmio = GetMMIOInstance(type, info);
+      if (mmio == nullptr) {
+        mmio = dynamic_cast<MMIO*>(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<decltype(hsa_agent_t::handle), std::vector<MMIO*>> MMIOManager::mmio_instances_; 
+std::map<decltype(hsa_agent_t::handle), std::vector<MMIO*>> MMIOManager::mmio_instances_;
 
 
-} // namespace mmio
+}  // namespace mmio
 
-} // namespace rocmtools
-
- 
+}  // namespace rocmtools

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 <hsa/hsa.h>
 #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<decltype(hsa_agent_t::handle), std::vector<MMIO*>> mmio_instances_;
 };

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<hsa_agent_t> agents;
   rocmtools::hsa_support::GetCoreApiTable().hsa_iterate_agents_fn(
@@ -57,25 +57,39 @@ CountersSampler::CountersSampler(
   // PCIE counters
   std::vector<std::string> 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<std::string> 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<std::mutex> 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<rocprofiler_counters_sampler_counter_output_t> values;
-    for (auto& perfmon : perfmon_instances_){
+    for (auto& perfmon : perfmon_instances_) {
       perfmon->Read(values);
     }
-    record.counters =
-        static_cast<rocprofiler_counters_sampler_counter_output_t*>(
-            malloc(values.size() * sizeof(rocprofiler_counters_sampler_counter_output_t)));
+    record.counters = static_cast<rocprofiler_counters_sampler_counter_output_t*>(
+        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
+}  // namespace rocmtools