SWDEV-297574: Support for pcie counters

Change-Id: I4a662c43a9d0cf883f336574baa09fc33b78b9af

samples/CMakeLists.txt

@@ -84,6 +84,16 @@ target_link_options(profiler_device_profiling PRIVATE "-Wl,--build-id=md5")
 add_dependencies(samples profiler_device_profiling)
 install(TARGETS profiler_device_profiling RUNTIME DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/${PROJECT_NAME}/samples COMPONENT samples)
 
+## Build Counters Sampling example
+set_source_files_properties(counters_sampler/pcie_counters_example.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1)
+hip_add_executable(pcie_counters_sampler counters_sampler/pcie_counters_example.cpp ${ROCPROFILER_UTIL_SRC_FILES})
+target_include_directories(pcie_counters_sampler PRIVATE ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/inc ${CMAKE_CURRENT_SOURCE_DIR}/common)
+target_link_libraries(pcie_counters_sampler PRIVATE ${ROCPROFILER_TARGET} systemd amd_comgr)
+add_dependencies(samples pcie_counters_sampler)
+install(TARGETS pcie_counters_sampler RUNTIME DESTINATION ${CMAKE_INSTALL_DATAROOTDIR}/${PROJECT_NAME}/samples COMPONENT samples)
+
+# ################################################################################################################
+
 # ############################################################################################################################################
 # Tracer Samples
 # ############################################################################################################################################

samples/common/common.h

@@ -285,6 +285,16 @@ void FlushPCSamplingRecord(
               << std::endl;
 }
 
+void FlushCountersSamplerRecord(
+  const rocprofiler_record_counters_sampler_t *counters_sampler_record) {
+  for (uint32_t i = 0; i < counters_sampler_record->num_counters; i++) {
+    output_file << ",Counter_" << i << "("
+                << std::to_string(counters_sampler_record->counters[i].value.value) << ")"
+                << std::endl;
+  }
+  output_file << std::endl;
+}
+
 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) {
   while (begin < end) {
@@ -308,9 +318,16 @@ int WriteBufferRecords(const rocprofiler_record_header_t* begin, const rocprofil
         FlushPCSamplingRecord(pc_sampling_record);
         break;
       }
-      default: {
+      case ROCPROFILER_COUNTERS_SAMPLER_RECORD: {
+        const rocprofiler_record_counters_sampler_t *counters_sampler_record =
+         reinterpret_cast<const rocprofiler_record_counters_sampler_t *>(begin);
+        FlushCountersSamplerRecord(counters_sampler_record);
         break;
       }
+      default: {
+        std::cout <<"unknown record\n";
+          break;
+      }
     }
     rocprofiler_next_record(begin, &begin, session_id, buffer_id);
   }

samples/counters_sampler/pcie_counters_example.cpp

@@ -0,0 +1,103 @@
+#include "../common/common.h"
+
+int main(int argc, char** argv) {
+  int* gpuMem;
+  int counter_option = 0;
+  
+  std::vector<std::string> pcie_counters = {
+      "CI_PERF_slv_MemRd_Bandwidth0", "CI_PERF_slv_MemWr_Bandwidth0", "CI_PERF_slv_totalMemRdTx",
+      "CI_PERF_slv_totalMemWrTx", "CI_PERF_slv_totalTx"};
+
+  if(argc > 1) {
+    counter_option = atoi(argv[1]);
+  }
+  else{
+    std::cout<< "Please provide one of the counter index options as argument:\n";
+    for(int i = 0; i < pcie_counters.size(); i++){
+      std::cout<< "[" << i << "]: " << pcie_counters[i] << std::endl;
+    }
+    std::cout<< "Example:\n ./pcie_counters_sampler 1\n";
+    exit(0);
+  }
+
+  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 = {};
+
+
+  rocprofiler_counters_sampler_counter_input_t counters_input[2] = {
+      {.name = const_cast<char*>(pcie_counters[counter_option].c_str()),
+       .type = ROCPROFILER_COUNTERS_SAMPLER_PCIE_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: " << pcie_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;
+}