Merge "setting code-obj tracking by default" into amd-master

bin/rpl_run.sh

@@ -433,8 +433,8 @@ while [ 1 ] ; do
     errck "Option '$ARG_IN', rate value"
     export ROCP_FLUSH_RATE="$period_rate"
   elif [ "$1" = "--obj-tracking" ] ; then
-    if [ "$2" = "on" ] ; then
-      export ROCP_OBJ_TRACKING=1
+    if [ "$2" = "off" ] ; then
+      export ROCP_OBJ_TRACKING=0
     fi
   elif [ "$1" = "--verbose" ] ; then
     ARG_VAL=0

src/core/rocprofiler.cpp

@@ -203,6 +203,7 @@ uint32_t LoadTool() {
     settings.trace_local = TraceProfile::IsLocal() ? 1: 0;
     settings.timeout = util::HsaRsrcFactory::GetTimeoutNs();
     settings.timestamp_on = InterceptQueue::IsTrackerOn() ? 1 : 0;
+    settings.code_obj_tracking = 1;
 
     if (handler) handler();
     else if (handler_prop) handler_prop(&settings);

src/util/hsa_rsrc_factory.cpp

@@ -155,6 +155,9 @@ HsaRsrcFactory::HsaRsrcFactory(bool initialize_hsa) : initialize_hsa_(initialize
 
   // System timeout
   timeout_ = (timeout_ns_ == HsaTimer::TIMESTAMP_MAX) ? timeout_ns_ : timer_->ns_to_sysclock(timeout_ns_);
+
+  // To dump code objects
+  to_dump_code_obj_ = getenv("ROCP_DUMP_CODEOBJ");
 }
 
 // Destructor of the class
@@ -195,6 +198,7 @@ void HsaRsrcFactory::InitHsaApiTable(HsaApiTable* table) {
       hsa_api_.hsa_executable_create_alt = table->core_->hsa_executable_create_alt_fn;
       hsa_api_.hsa_executable_load_agent_code_object = table->core_->hsa_executable_load_agent_code_object_fn;
       hsa_api_.hsa_executable_freeze = table->core_->hsa_executable_freeze_fn;
+      hsa_api_.hsa_executable_destroy = table->core_->hsa_executable_destroy_fn;
       hsa_api_.hsa_executable_get_symbol = table->core_->hsa_executable_get_symbol_fn;
       hsa_api_.hsa_executable_symbol_get_info = table->core_->hsa_executable_symbol_get_info_fn;
       hsa_api_.hsa_executable_iterate_symbols = table->core_->hsa_executable_iterate_symbols_fn;
@@ -235,6 +239,7 @@ void HsaRsrcFactory::InitHsaApiTable(HsaApiTable* table) {
       hsa_api_.hsa_executable_create_alt = hsa_executable_create_alt;
       hsa_api_.hsa_executable_load_agent_code_object = hsa_executable_load_agent_code_object;
       hsa_api_.hsa_executable_freeze = hsa_executable_freeze;
+      hsa_api_.hsa_executable_destroy = hsa_executable_destroy;
       hsa_api_.hsa_executable_get_symbol = hsa_executable_get_symbol;
       hsa_api_.hsa_executable_symbol_get_info = hsa_executable_symbol_get_info;
       hsa_api_.hsa_executable_iterate_symbols = hsa_executable_iterate_symbols;
@@ -710,6 +715,7 @@ void HsaRsrcFactory::EnableExecutableTracking(HsaApiTable* table) {
   std::lock_guard<mutex_t> lck(mutex_);
   executable_tracking_on_ = true;
   table->core_->hsa_executable_freeze_fn = hsa_executable_freeze_interceptor;
+  table->core_->hsa_executable_destroy_fn = hsa_executable_destroy_interceptor;
 }
 
 hsa_status_t HsaRsrcFactory::executable_symbols_cb(hsa_executable_t exec, hsa_executable_symbol_t symbol, void *data) {
@@ -727,10 +733,14 @@ hsa_status_t HsaRsrcFactory::executable_symbols_cb(hsa_executable_t exec, hsa_ex
     status = hsa_api_.hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME, name);
     CHECK_STATUS("Error in getting kernel name", status);
     name[len] = 0;
-    auto ret = symbols_map_->insert({addr, name});
-    if (ret.second == false) {
-      delete[] ret.first->second;
-      ret.first->second = name;
+    if (data == NULL) {
+      auto ret = symbols_map_->insert({addr, name});
+      if (ret.second == false) {
+        delete[] ret.first->second;
+        ret.first->second = name;
+      }
+    } else {
+      symbols_map_->erase(addr);
     }
   }
   return HSA_STATUS_SUCCESS;
@@ -741,7 +751,16 @@ hsa_status_t HsaRsrcFactory::hsa_executable_freeze_interceptor(hsa_executable_t
   if (symbols_map_ == NULL) symbols_map_ = new symbols_map_t;
   hsa_status_t status = hsa_api_.hsa_executable_iterate_symbols(executable, executable_symbols_cb, NULL);
   CHECK_STATUS("Error in iterating executable symbols", status);
-  return hsa_api_.hsa_executable_freeze(executable, options);;
+  return hsa_api_.hsa_executable_freeze(executable, options);
+}
+
+hsa_status_t HsaRsrcFactory::hsa_executable_destroy_interceptor(hsa_executable_t executable) {
+  std::lock_guard<mutex_t> lck(mutex_);
+  if (symbols_map_ != NULL) {
+    hsa_status_t status = hsa_api_.hsa_executable_iterate_symbols(executable, executable_symbols_cb, (void*)1);
+    CHECK_STATUS("Error in iterating executable symbols", status);
+  }
+  return hsa_api_.hsa_executable_destroy(executable);
 }
 
 std::atomic<HsaRsrcFactory*> HsaRsrcFactory::instance_{};
@@ -750,6 +769,7 @@ HsaRsrcFactory::timestamp_t HsaRsrcFactory::timeout_ns_ = HsaTimer::TIMESTAMP_MA
 hsa_pfn_t HsaRsrcFactory::hsa_api_{};
 bool HsaRsrcFactory::executable_tracking_on_ = false;
 HsaRsrcFactory::symbols_map_t* HsaRsrcFactory::symbols_map_ = NULL;
+void* HsaRsrcFactory::to_dump_code_obj_ = NULL;
 
 }  // namespace util
 }  // namespace rocprofiler

src/util/hsa_rsrc_factory.h

@@ -97,6 +97,7 @@ struct hsa_pfn_t {
   decltype(hsa_executable_create_alt)* hsa_executable_create_alt;
   decltype(hsa_executable_load_agent_code_object)* hsa_executable_load_agent_code_object;
   decltype(hsa_executable_freeze)* hsa_executable_freeze;
+  decltype(hsa_executable_destroy)* hsa_executable_destroy;
   decltype(hsa_executable_get_symbol)* hsa_executable_get_symbol;
   decltype(hsa_executable_symbol_get_info)* hsa_executable_symbol_get_info;
   decltype(hsa_executable_iterate_symbols)* hsa_executable_iterate_symbols;
@@ -498,7 +499,9 @@ class HsaRsrcFactory {
   typedef std::map<uint64_t, const char*> symbols_map_t;
   static symbols_map_t* symbols_map_;
   static bool executable_tracking_on_;
+  static void* to_dump_code_obj_;
   static hsa_status_t hsa_executable_freeze_interceptor(hsa_executable_t executable, const char *options);
+  static hsa_status_t hsa_executable_destroy_interceptor(hsa_executable_t executable);
   static hsa_status_t executable_symbols_cb(hsa_executable_t exec, hsa_executable_symbol_t symbol, void *data);
 
   // HSA runtime API table