hsa memcopy tracking

Change-Id: I61bc4ea01a62b587d5941f9d0e9e358a6ed71b5a


[ROCm/rocprofiler commit: 2b2a279ce9]

projects/rocprofiler/inc/rocprofiler.h

@@ -65,6 +65,7 @@ uint32_t rocprofiler_version_minor();
 
 typedef struct {
   uint32_t intercept_mode;
+  uint32_t memcopy_tracking;
   uint32_t sqtt_size;
   uint32_t sqtt_local;
   uint64_t timeout;

projects/rocprofiler/src/core/intercept_queue.h

@@ -71,7 +71,7 @@ class InterceptQueue {
     if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, "ProxyQueue::Create()");
 
     if (tracker_on || tracker_on_) {
-      if (tracker_ == NULL) tracker_ = new Tracker;
+      if (tracker_ == NULL) tracker_ = &Tracker::Instance();
       status = hsa_amd_profiling_set_profiler_enabled(*queue, true);
       if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, "hsa_amd_profiling_set_profiler_enabled()");
     }
@@ -179,7 +179,7 @@ class InterceptQueue {
             if (tracker_entry != NULL) {
               Group* context_group = context->GetGroup(group.index);
               context_group->IncrRefsCount();
-              tracker_->Enable(tracker_entry, Context::Handler, reinterpret_cast<void*>(context_group));
+              tracker_->EnableContext(tracker_entry, Context::Handler, reinterpret_cast<void*>(context_group));
             }
 
             const pkt_vector_t& start_vector = context->StartPackets(group.index);
@@ -197,7 +197,7 @@ class InterceptQueue {
             if (tracker_entry != NULL) {
               void* context_handler_arg = NULL;
               rocprofiler_handler_t context_handler_fun = context->GetHandler(&context_handler_arg);
-              tracker_->Enable(tracker_entry, context_handler_fun, context_handler_arg);
+              tracker_->EnableDispatch(tracker_entry, context_handler_fun, context_handler_arg);
             }
           }
         }

projects/rocprofiler/src/core/rocprofiler.cpp

@@ -83,6 +83,9 @@ decltype(hsa_queue_load_read_index_scacquire)* hsa_queue_load_read_index_scacqui
 decltype(hsa_amd_queue_intercept_create)* hsa_amd_queue_intercept_create_fn;
 decltype(hsa_amd_queue_intercept_register)* hsa_amd_queue_intercept_register_fn;
 
+decltype(hsa_amd_memory_async_copy)* hsa_amd_memory_async_copy_fn;
+decltype(hsa_amd_memory_async_copy_rect)* hsa_amd_memory_async_copy_rect_fn;
+
 ::HsaApiTable* kHsaApiTable;
 
 void SaveHsaApi(::HsaApiTable* table) {
@@ -136,12 +139,16 @@ void * tool_handle = NULL;
 
 // Load profiling tool library
 // Return true if intercepting mode is enabled
-bool LoadTool() {
-  bool intercept_mode = false;
+enum {
+  DISPATCH_INTERCEPT_MODE = 0x1,
+  MEMCOPY_INTERCEPT_MODE = 0x2
+};
+uint32_t LoadTool() {
+  uint32_t intercept_mode = 0;
   const char* tool_lib = getenv("ROCP_TOOL_LIB");
 
   if (tool_lib) {
-    intercept_mode = true;
+    intercept_mode = DISPATCH_INTERCEPT_MODE;
 
     tool_handle = dlopen(tool_lib, RTLD_NOW);
     if (tool_handle == NULL) {
@@ -164,7 +171,7 @@ bool LoadTool() {
     }
 
     rocprofiler_settings_t settings{};
-    settings.intercept_mode = (intercept_mode) ? 1 : 0;
+    settings.intercept_mode = (intercept_mode != 0) ? 1 : 0;
     settings.sqtt_size = SqttProfile::GetSize();
     settings.sqtt_local = SqttProfile::IsLocal() ? 1: 0;
     settings.timeout = util::HsaRsrcFactory::GetTimeoutNs();
@@ -173,11 +180,12 @@ bool LoadTool() {
     if (handler) handler();
     else if (handler_prop) handler_prop(&settings);
 
-    intercept_mode = (settings.intercept_mode != 0);
     SqttProfile::SetSize(settings.sqtt_size);
     SqttProfile::SetLocal(settings.sqtt_local != 0);
     util::HsaRsrcFactory::SetTimeoutNs(settings.timeout);
     InterceptQueue::TrackerOn(settings.timestamp_on != 0);
+    if (settings.intercept_mode != 0) intercept_mode = DISPATCH_INTERCEPT_MODE;
+    if (settings.memcopy_tracking) intercept_mode |= MEMCOPY_INTERCEPT_MODE;
   }
 
   return intercept_mode;
@@ -310,9 +318,56 @@ hsa_status_t CreateQueuePro(
   return HSA_STATUS_SUCCESS;
 }
 
+bool async_copy_handler(hsa_signal_value_t value, void* arg) {
+  Tracker::entry_t* entry = reinterpret_cast<Tracker::entry_t*>(arg);
+  printf("%lu: async-copy time(%lu,%lu)\n", entry->index, entry->record->begin, entry->record->end);
+  return false;
+}
+
+hsa_status_t hsa_amd_memory_async_copy_interceptor(
+    void* dst, hsa_agent_t dst_agent, const void* src,
+    hsa_agent_t src_agent, size_t size, uint32_t num_dep_signals,
+    const hsa_signal_t* dep_signals, hsa_signal_t completion_signal)
+{
+  Tracker* tracker = &Tracker::Instance();
+  Tracker::entry_t* tracker_entry = tracker->Alloc(hsa_agent_t{}, completion_signal);
+  hsa_status_t status = hsa_amd_memory_async_copy_fn(dst, dst_agent, src,
+                                                     src_agent, size, num_dep_signals,
+                                                     dep_signals, tracker_entry->signal);
+  if (status == HSA_STATUS_SUCCESS) {
+    tracker->EnableMemcopy(tracker_entry, async_copy_handler, reinterpret_cast<void*>(tracker_entry));
+  } else {
+    tracker->Delete(tracker_entry);
+  }
+  return status;
+}
+
+hsa_status_t hsa_amd_memory_async_copy_rect_interceptor(
+    const hsa_pitched_ptr_t* dst, const hsa_dim3_t* dst_offset, const hsa_pitched_ptr_t* src,
+    const hsa_dim3_t* src_offset, const hsa_dim3_t* range, hsa_agent_t copy_agent,
+    hsa_amd_copy_direction_t dir, uint32_t num_dep_signals, const hsa_signal_t* dep_signals,
+    hsa_signal_t completion_signal)
+{
+  Tracker* tracker = &Tracker::Instance();
+  Tracker::entry_t* tracker_entry = tracker->Alloc(hsa_agent_t{}, completion_signal);
+  hsa_status_t status = hsa_amd_memory_async_copy_rect_fn(dst, dst_offset, src,
+                                                          src_offset, range, copy_agent,
+                                                          dir, num_dep_signals, dep_signals,
+                                                          tracker_entry->signal);
+  if (status == HSA_STATUS_SUCCESS) {
+    tracker->EnableMemcopy(tracker_entry, async_copy_handler, reinterpret_cast<void*>(tracker_entry));
+  } else {
+    tracker->Delete(tracker_entry);
+  }
+  return status;
+}
+
 rocprofiler_properties_t rocprofiler_properties;
 uint32_t SqttProfile::output_buffer_size_ = 0x2000000;  // 32M
 bool SqttProfile::output_buffer_local_ = true;
+Tracker* Tracker::instance_ = NULL;
+Tracker::mutex_t Tracker::glob_mutex_;
+Tracker::counter_t Tracker::counter_ = 0;
 util::Logger::mutex_t util::Logger::mutex_;
 util::Logger* util::Logger::instance_ = NULL;
 }
@@ -355,8 +410,16 @@ PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t fa
   }
 
   // Loading a tool lib and setting of intercept mode
-  const bool intercept_mode_on = rocprofiler::LoadTool();
-  if (intercept_mode_on) intercept_mode = true;
+  const uint32_t intercept_mode_mask = rocprofiler::LoadTool();
+  if (intercept_mode_mask & rocprofiler::DISPATCH_INTERCEPT_MODE) intercept_mode = true;
+  if (intercept_mode_mask & rocprofiler::MEMCOPY_INTERCEPT_MODE) {
+    hsa_status_t status = hsa_amd_profiling_async_copy_enable(true);
+    if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, "hsa_amd_profiling_async_copy_enable");
+    rocprofiler::hsa_amd_memory_async_copy_fn = table->amd_ext_->hsa_amd_memory_async_copy_fn;
+    rocprofiler::hsa_amd_memory_async_copy_rect_fn = table->amd_ext_->hsa_amd_memory_async_copy_rect_fn;
+    table->amd_ext_->hsa_amd_memory_async_copy_fn = rocprofiler::hsa_amd_memory_async_copy_interceptor;
+    table->amd_ext_->hsa_amd_memory_async_copy_rect_fn = rocprofiler::hsa_amd_memory_async_copy_rect_interceptor;
+  }
 
   // HSA intercepting
   if (intercept_mode) {
@@ -371,6 +434,7 @@ PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t fa
 
 // HSA-runtime tool on-unload method
 PUBLIC_API void OnUnload() {
+  rocprofiler::Tracker::Destroy();
   rocprofiler::UnloadTool();
   rocprofiler::RestoreHsaApi();
 }

projects/rocprofiler/src/core/tracker.h

@@ -47,8 +47,10 @@ class Tracker {
   struct entry_t;
   typedef std::list<entry_t*> sig_list_t;
   typedef sig_list_t::iterator sig_list_it_t;
+  typedef uint64_t counter_t;
 
   struct entry_t {
+    counter_t index;
     std::atomic<bool> valid;
     Tracker* tracker;
     sig_list_t::iterator it;
@@ -58,22 +60,25 @@ class Tracker {
     record_t* record;
     std::atomic<void*> handler;
     void* arg;
-    bool context_active;
+    bool is_context;
+    bool is_memcopy;
   };
 
-  Tracker() :
-    outstanding_(0),
-    hsa_rsrc_(&(util::HsaRsrcFactory::Instance()))
-  {}
+  static Tracker* Create() {
+    std::lock_guard<mutex_t> lck(glob_mutex_);
+    if (instance_ == NULL) instance_ = new Tracker;
+    return instance_;
+  }
 
-  ~Tracker() {
-    auto it = sig_list_.begin();
-    auto end = sig_list_.end();
-    while (it != end) {
-      auto cur = it++;
-      hsa_rsrc_->SignalWait((*cur)->signal);
-      Erase(cur);
-    }
+  static Tracker& Instance() {
+    if (instance_ == NULL) instance_ = Create();
+    return *instance_;
+  }
+
+  static void Destroy() {
+    std::lock_guard<mutex_t> lck(glob_mutex_);
+    if (instance_ != NULL) delete instance_;
+    instance_ = NULL;
   }
 
   // Add tracker entry
@@ -102,6 +107,7 @@ class Tracker {
     // Adding antry to the list
     mutex_.lock();
     entry->it = sig_list_.insert(sig_list_.end(), entry);
+    entry->index = counter_++;
     mutex_.unlock();
 
     return entry;
@@ -130,20 +136,39 @@ class Tracker {
     }
   }
 
-  void Enable(entry_t* entry, hsa_amd_signal_handler handler, void* arg) {
-    entry->context_active = true;
+  void EnableContext(entry_t* entry, hsa_amd_signal_handler handler, void* arg) {
+    entry->is_context = true;
     Enable(entry, reinterpret_cast<void*>(handler), arg);
   }
-  void Enable(entry_t* entry, rocprofiler_handler_t handler, void* arg) {
+  void EnableDispatch(entry_t* entry, rocprofiler_handler_t handler, void* arg) {
+    Enable(entry, reinterpret_cast<void*>(handler), arg);
+  }
+  void EnableMemcopy(entry_t* entry, hsa_amd_signal_handler handler, void* arg) {
+    entry->is_memcopy = true;
     Enable(entry, reinterpret_cast<void*>(handler), arg);
   }
 
   private:
+  Tracker() :
+    outstanding_(0),
+    hsa_rsrc_(&(util::HsaRsrcFactory::Instance()))
+  {}
+
+  ~Tracker() {
+    auto it = sig_list_.begin();
+    auto end = sig_list_.end();
+    while (it != end) {
+      auto cur = it++;
+      hsa_rsrc_->SignalWait((*cur)->signal);
+      Erase(cur);
+    }
+  }
+
   // Delete an entry by iterator
   void Erase(const sig_list_it_t& it) { Delete(*it); }
 
   // Entry completion
-  inline void Complete(entry_t* entry) {
+  inline void Complete(hsa_signal_value_t signal_value, entry_t* entry) {
     record_t* record = entry->record;
 
     // Debug trace
@@ -154,12 +179,20 @@ class Tracker {
     }
 
     // Query begin/end and complete timestamps
-    hsa_amd_profiling_dispatch_time_t dispatch_time{};
-    hsa_status_t status = hsa_amd_profiling_get_dispatch_time(entry->agent, entry->signal, &dispatch_time);
-    if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_amd_profiling_get_dispatch_time");
+    if (entry->is_memcopy) {
+      hsa_amd_profiling_async_copy_time_t async_copy_time{};
+      hsa_status_t status = hsa_amd_profiling_get_async_copy_time(entry->signal, &async_copy_time);
+      if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_amd_profiling_get_async_copy_time");
+      record->begin = hsa_rsrc_->SysclockToNs(async_copy_time.start);
+      record->end = hsa_rsrc_->SysclockToNs(async_copy_time.end);
+    } else {
+      hsa_amd_profiling_dispatch_time_t dispatch_time{};
+      hsa_status_t status = hsa_amd_profiling_get_dispatch_time(entry->agent, entry->signal, &dispatch_time);
+      if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_amd_profiling_get_dispatch_time");
+      record->begin = hsa_rsrc_->SysclockToNs(dispatch_time.start);
+      record->end = hsa_rsrc_->SysclockToNs(dispatch_time.end);
+    }
 
-    record->begin = hsa_rsrc_->SysclockToNs(dispatch_time.start);
-    record->end = hsa_rsrc_->SysclockToNs(dispatch_time.end);
     record->complete = hsa_rsrc_->TimestampNs();
     entry->valid.store(true, std::memory_order_release);
 
@@ -171,16 +204,17 @@ class Tracker {
       orig_signal_ptr->start_ts = prof_signal_ptr->start_ts;
       orig_signal_ptr->end_ts = prof_signal_ptr->end_ts;
 
-      const hsa_signal_value_t value = hsa_signal_load_relaxed(orig);
-      hsa_signal_store_screlease(orig, value - 1);
+      const hsa_signal_value_t new_value = hsa_signal_load_relaxed(orig) - 1;
+      if (signal_value != new_value) EXC_ABORT(HSA_STATUS_ERROR, "Tracker::Complete bad signal value");
+      hsa_signal_store_screlease(orig, signal_value);
     }
   }
 
-  inline static void HandleEntry(entry_t* entry) {
+  inline static void HandleEntry(hsa_signal_value_t signal_value, entry_t* entry) {
     // Call entry handler
     void* handler = static_cast<void*>(entry->handler);
-    if (entry->context_active) {
-      reinterpret_cast<hsa_amd_signal_handler>(handler)(0, entry->arg);
+    if (entry->is_context || entry->is_memcopy) {
+      reinterpret_cast<hsa_amd_signal_handler>(handler)(signal_value, entry->arg);
     } else {
       rocprofiler_group_t group{};
       reinterpret_cast<rocprofiler_handler_t>(handler)(group, entry->arg);
@@ -190,7 +224,7 @@ class Tracker {
   }
 
   // Handler for packet completion
-  static bool Handler(hsa_signal_value_t, void* arg) {
+  static bool Handler(hsa_signal_value_t signal_value, void* arg) {
     // Acquire entry
     entry_t* entry = reinterpret_cast<entry_t*>(arg);
     volatile std::atomic<void*>* ptr = &entry->handler;
@@ -198,10 +232,10 @@ class Tracker {
 
     // Complete entry
     Tracker* tracker = entry->tracker;
-    tracker->Complete(entry);
+    tracker->Complete(signal_value, entry);
 
     if (ordering_enabled_ == false) {
-      HandleEntry(entry);
+      HandleEntry(signal_value, entry);
     } else {
       // Acquire last entry
       entry_t* back = tracker->sig_list_.back();
@@ -214,7 +248,7 @@ class Tracker {
       while (it != end) {
         entry = *(it++);
         if (entry->valid.load(std::memory_order_acquire)) {
-          HandleEntry(entry);
+          HandleEntry(signal_value, entry);
         } else {
           break;
         }
@@ -225,6 +259,11 @@ class Tracker {
     return false;
   }
 
+  // instance
+  static Tracker* instance_;
+  static mutex_t glob_mutex_;
+  static counter_t counter_;
+
   // Tracked signals list
   sig_list_t sig_list_;
   // Inter-thread synchronization
@@ -235,7 +274,7 @@ class Tracker {
   // HSA resources factory
   util::HsaRsrcFactory* hsa_rsrc_;
   // Handling ordering enabled
-  static const bool ordering_enabled_ = true;
+  static const bool ordering_enabled_ = false;
   // Enable tracing
   static const bool trace_on_ = false;
 };

projects/rocprofiler/test/run.sh

@@ -56,11 +56,6 @@ if [ ! -e $ROCP_TOOL_LIB ] ; then
   export ROCP_TOOL_LIB=test/libtool.so
 fi
 
-export ROCP_KITER=1
-export ROCP_DITER=4
-export ROCP_INPUT=input1.xml
-eval ./test/ctrl
-
 export ROCP_KITER=50
 export ROCP_DITER=50
 export ROCP_AGENTS=1
@@ -68,6 +63,14 @@ export ROCP_THRS=1
 export ROCP_INPUT=input.xml
 eval ./test/ctrl
 
+# Memcopies tracking
+export ROCP_MCOPY_TRACKING=1
+
+export ROCP_KITER=1
+export ROCP_DITER=4
+export ROCP_INPUT=input1.xml
+eval ./test/ctrl
+
 #valgrind --leak-check=full $tbin
 #valgrind --tool=massif $tbin
 #ms_print massif.out.<N>

projects/rocprofiler/test/tool/tool.cpp

@@ -818,6 +818,8 @@ extern "C" PUBLIC_API void OnLoadToolProp(rocprofiler_settings_t* settings)
       }
       it = opts.find("sqtt-local");
       if (it != opts.end()) { settings->sqtt_local = (it->second == "on"); }
+      it = opts.find("memcopies");
+      if (it != opts.end()) { settings->memcopy_tracking = (it->second == "on"); }
     }
   }
   // Enable verbose mode
@@ -835,6 +837,8 @@ extern "C" PUBLIC_API void OnLoadToolProp(rocprofiler_settings_t* settings)
   check_env_var("ROCP_SQTT_SIZE", settings->sqtt_size);
   // Set SQTT local buffer
   check_env_var("ROCP_SQTT_LOCAL", settings->sqtt_local);
+  // Set memcopies tracking
+  check_env_var("ROCP_MCOPY_TRACKING", settings->memcopy_tracking);
 
   is_sqtt_local = settings->sqtt_local;