multiple changes update

src/core/roctracer.cpp

@@ -35,12 +35,18 @@ THE SOFTWARE.
 #include <sys/syscall.h>
 
 #include "core/loader.h"
+#include "core/trace_buffer.h"
 #include "proxy/tracker.h"
 #include "ext/hsa_rt_utils.hpp"
 #include "util/exception.h"
 #include "util/hsa_rsrc_factory.h"
 #include "util/logger.h"
 
+#include "proxy/hsa_queue.h"
+#include "proxy/intercept_queue.h"
+#include "proxy/proxy_queue.h"
+#include "proxy/simple_proxy_queue.h"
+
 #define PUBLIC_API __attribute__((visibility("default")))
 #define CONSTRUCTOR_API __attribute__((constructor))
 #define DESTRUCTOR_API __attribute__((destructor))
@@ -85,17 +91,86 @@ THE SOFTWARE.
 ///////////////////////////////////////////////////////////////////////////////////////////////////
 // Internal library methods
 //
+namespace rocprofiler {
+decltype(hsa_queue_create)* hsa_queue_create_fn;
+decltype(hsa_queue_destroy)* hsa_queue_destroy_fn;
+
+decltype(hsa_signal_store_relaxed)* hsa_signal_store_relaxed_fn;
+decltype(hsa_signal_store_relaxed)* hsa_signal_store_screlease_fn;
+
+decltype(hsa_queue_load_write_index_relaxed)* hsa_queue_load_write_index_relaxed_fn;
+decltype(hsa_queue_store_write_index_relaxed)* hsa_queue_store_write_index_relaxed_fn;
+decltype(hsa_queue_load_read_index_relaxed)* hsa_queue_load_read_index_relaxed_fn;
+
+decltype(hsa_queue_load_write_index_scacquire)* hsa_queue_load_write_index_scacquire_fn;
+decltype(hsa_queue_store_write_index_screlease)* hsa_queue_store_write_index_screlease_fn;
+decltype(hsa_queue_load_read_index_scacquire)* hsa_queue_load_read_index_scacquire_fn;
+
+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) {
+  util::HsaRsrcFactory::InitHsaApiTable(table);
+
+  kHsaApiTable = table;
+  hsa_queue_create_fn = table->core_->hsa_queue_create_fn;
+  hsa_queue_destroy_fn = table->core_->hsa_queue_destroy_fn;
+
+  hsa_signal_store_relaxed_fn = table->core_->hsa_signal_store_relaxed_fn;
+  hsa_signal_store_screlease_fn = table->core_->hsa_signal_store_screlease_fn;
+
+  hsa_queue_load_write_index_relaxed_fn = table->core_->hsa_queue_load_write_index_relaxed_fn;
+  hsa_queue_store_write_index_relaxed_fn = table->core_->hsa_queue_store_write_index_relaxed_fn;
+  hsa_queue_load_read_index_relaxed_fn = table->core_->hsa_queue_load_read_index_relaxed_fn;
+
+  hsa_queue_load_write_index_scacquire_fn = table->core_->hsa_queue_load_write_index_scacquire_fn;
+  hsa_queue_store_write_index_screlease_fn = table->core_->hsa_queue_store_write_index_screlease_fn;
+  hsa_queue_load_read_index_scacquire_fn = table->core_->hsa_queue_load_read_index_scacquire_fn;
+
+  hsa_amd_queue_intercept_create_fn = table->amd_ext_->hsa_amd_queue_intercept_create_fn;
+  hsa_amd_queue_intercept_register_fn = table->amd_ext_->hsa_amd_queue_intercept_register_fn;
+}
+
+void RestoreHsaApi() {
+  ::HsaApiTable* table = kHsaApiTable;
+  table->core_->hsa_queue_create_fn = hsa_queue_create_fn;
+  table->core_->hsa_queue_destroy_fn = hsa_queue_destroy_fn;
+
+  table->core_->hsa_signal_store_relaxed_fn = hsa_signal_store_relaxed_fn;
+  table->core_->hsa_signal_store_screlease_fn = hsa_signal_store_screlease_fn;
+
+  table->core_->hsa_queue_load_write_index_relaxed_fn = hsa_queue_load_write_index_relaxed_fn;
+  table->core_->hsa_queue_store_write_index_relaxed_fn = hsa_queue_store_write_index_relaxed_fn;
+  table->core_->hsa_queue_load_read_index_relaxed_fn = hsa_queue_load_read_index_relaxed_fn;
+
+  table->core_->hsa_queue_load_write_index_scacquire_fn = hsa_queue_load_write_index_scacquire_fn;
+  table->core_->hsa_queue_store_write_index_screlease_fn = hsa_queue_store_write_index_screlease_fn;
+  table->core_->hsa_queue_load_read_index_scacquire_fn = hsa_queue_load_read_index_scacquire_fn;
+
+  table->amd_ext_->hsa_amd_queue_intercept_create_fn = hsa_amd_queue_intercept_create_fn;
+  table->amd_ext_->hsa_amd_queue_intercept_register_fn = hsa_amd_queue_intercept_register_fn;
+}
+}
+
 namespace roctracer {
 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;
 
+TraceBuffer<trace_entry_t> trace_buffer(0x200000);
+
 namespace hsa_support {
 // callbacks table
 cb_table_t cb_table;
 // asyc copy activity callback
+bool async_copy_callback_enabled = false;
 activity_async_callback_t async_copy_callback_fun = NULL;
 void* async_copy_callback_arg = NULL;
-bool async_copy_callback_enabled = false;
+const char* output_prefix = NULL;
 // Table of function pointers to HSA Core Runtime
 CoreApiTable CoreApiTable_saved{};
 // Table of function pointers to AMD extensions
@@ -288,15 +363,6 @@ class MemoryPool {
   pthread_cond_t read_cond_;
 };
 
-CONSTRUCTOR_API void constructor() {
-  util::Logger::Create();
-}
-
-DESTRUCTOR_API void destructor() {
-  ::util::HsaRsrcFactory::Destroy();
-  util::Logger::Destroy();
-}
-
 // Correlation id storage
 static thread_local activity_correlation_id_t correlation_id_tls = 0;
 typedef std::map<activity_correlation_id_t, activity_correlation_id_t> correlation_id_map_t;
@@ -367,18 +433,73 @@ void HCC_AsyncActivityCallback(uint32_t op_id, void* record, void* arg) {
   pool->Write(*record_ptr);
 }
 
-bool hsa_async_copy_handler(hsa_signal_value_t value, void* arg) {
-  ::proxy::Tracker::entry_t* entry = reinterpret_cast<::proxy::Tracker::entry_t*>(arg);
+// Open output file
+FILE* open_output_file(const char* prefix, const char* name) {
+  FILE* file_handle = NULL;
+  if (prefix != NULL) {
+    std::ostringstream oss;
+    oss << prefix << "/" << name;
+    file_handle = fopen(oss.str().c_str(), "w");
+    if (file_handle == NULL) {
+      std::ostringstream errmsg;
+      errmsg << "ROCTracer: fopen error, file '" << oss.str().c_str() << "'";
+      perror(errmsg.str().c_str());
+      abort();
+    }
+  } else file_handle = stdout;
+  return file_handle;
+}
 
+FILE* kernel_file_handle = NULL;
+void hsa_kernel_handler(::proxy::Tracker::entry_t* entry) {
+  static uint64_t index = 0;
+  if (index == 0) {
+    kernel_file_handle = open_output_file(hsa_support::output_prefix, "results.txt");
+  }
+  fprintf(kernel_file_handle, "dispatch[%lu], gpu-id(%u), tid(%u), kernel-name(\"%s\"), time(%lu,%lu,%lu,%lu)\n",
+    index,
+    ::util::HsaRsrcFactory::Instance().GetAgentInfo(entry->agent)->dev_index,
+    entry->kernel.tid,
+    entry->kernel.name,
+    entry->dispatch,
+    entry->begin,
+    entry->end,
+    entry->complete);
+#if 0
+  fprintf(file_handle, "dispatch[%u], gpu-id(%u), queue-id(%u), queue-index(%lu), tid(%lu), grd(%u), wgr(%u), lds(%u), scr(%u), vgpr(%u), sgpr(%u), fbar(%u), sig(0x%lx), kernel-name(\"%s\")",
+    index,
+    HsaRsrcFactory::Instance().GetAgentInfo(entry->agent)->dev_index,
+    entry->data.queue_id,
+    entry->data.queue_index,
+    entry->data.thread_id,
+    entry->kernel_properties.grid_size,
+    entry->kernel_properties.workgroup_size,
+    entry->kernel_properties.lds_size,
+    entry->kernel_properties.scratch_size,
+    entry->kernel_properties.vgpr_count,
+    entry->kernel_properties.sgpr_count,
+    entry->kernel_properties.fbarrier_count,
+    entry->kernel_properties.signal.handle,
+    nik_name.c_str());
+  if (record) fprintf(file_handle, ", time(%lu,%lu,%lu,%lu)",
+    record->dispatch,
+    record->begin,
+    record->end,
+    record->complete);
+  fprintf(file_handle, "\n");
+  fflush(file_handle);
+#endif
+  index++;
+}
+
+void hsa_async_copy_handler(::proxy::Tracker::entry_t* entry) {
   activity_record_t record{};
   record.domain = ACTIVITY_DOMAIN_HSA_OPS;   // activity domain id
-  record.correlation_id = entry->index;      // activity ID
-  record.begin_ns = entry->record->begin;    // host begin timestamp
-  record.end_ns = entry->record->end;        // host end timestamp
+  record.begin_ns = entry->begin;    // host begin timestamp
+  record.end_ns = entry->end;        // host end timestamp
   record.device_id = 0;                      // device id
 
   hsa_support::async_copy_callback_fun(hsa_support::HSA_OP_ID_async_copy, &record, hsa_support::async_copy_callback_arg);
-  return false;
 }
 
 hsa_status_t hsa_amd_memory_async_copy_interceptor(
@@ -388,17 +509,15 @@ hsa_status_t hsa_amd_memory_async_copy_interceptor(
 {
   hsa_status_t status = HSA_STATUS_SUCCESS;
   if (hsa_support::async_copy_callback_enabled) {
-    ::proxy::Tracker* tracker = &::proxy::Tracker::Instance();
-    ::proxy::Tracker::entry_t* tracker_entry = tracker->Alloc(hsa_agent_t{}, completion_signal);
+    trace_entry_t* entry = trace_buffer.GetEntry();
+    ::proxy::Tracker::Enable(COPY_ENTRY_TYPE, hsa_agent_t{}, completion_signal, entry);
     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, hsa_async_copy_handler, reinterpret_cast<void*>(tracker_entry));
-    } else {
-      tracker->Delete(tracker_entry);
-    }
-  } else {
+                                          dep_signals, entry->signal);
+    if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, "hsa_amd_memory_async_copy interceptor");
+  }
+  else
+  {
     status = hsa_amd_memory_async_copy_fn(dst, dst_agent, src,
                                           src_agent, size, num_dep_signals,
                                           dep_signals, completion_signal);
@@ -414,18 +533,16 @@ hsa_status_t hsa_amd_memory_async_copy_rect_interceptor(
 {
   hsa_status_t status = HSA_STATUS_SUCCESS;
   if (hsa_support::async_copy_callback_enabled) {
-    ::proxy::Tracker* tracker = &::proxy::Tracker::Instance();
-    ::proxy::Tracker::entry_t* tracker_entry = tracker->Alloc(hsa_agent_t{}, completion_signal);
+    trace_entry_t* entry = trace_buffer.GetEntry();
+    ::proxy::Tracker::Enable(COPY_ENTRY_TYPE, hsa_agent_t{}, completion_signal, entry);
     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, hsa_async_copy_handler, reinterpret_cast<void*>(tracker_entry));
-    } else {
-      tracker->Delete(tracker_entry);
-    }
-  } else {
+                                               entry->signal);
+    if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, "hsa_amd_memory_async_copy interceptor");
+  }
+  else
+  {
     status = hsa_amd_memory_async_copy_rect_fn(dst, dst_offset, src,
                                                src_offset, range, copy_agent,
                                                dir, num_dep_signals, dep_signals,
@@ -443,10 +560,6 @@ memory_pool_mutex_t memory_pool_mutex;
 
 LOADER_INSTANTIATE();
 
-std::atomic<proxy::Tracker*> proxy::Tracker::instance_{};
-proxy::Tracker::mutex_t proxy::Tracker::glob_mutex_;
-proxy::Tracker::counter_t proxy::Tracker::counter_ = 0;
-
 ///////////////////////////////////////////////////////////////////////////////////////////////////
 // Public library methods
 //
@@ -685,6 +798,7 @@ static void roctracer_enable_activity_impl(
   switch (domain) {
     case ACTIVITY_DOMAIN_HSA_OPS: {
       roctracer::hsa_support::async_copy_callback_enabled = true;
+      rocprofiler::InterceptQueue::Enable(true);
       break;
     }
     case ACTIVITY_DOMAIN_HSA_API: break;
@@ -747,6 +861,7 @@ static void roctracer_disable_activity_impl(
   switch (domain) {
     case ACTIVITY_DOMAIN_HSA_OPS: {
       roctracer::hsa_support::async_copy_callback_enabled = false;
+      rocprofiler::InterceptQueue::Enable(false);
       break;
     }
     case ACTIVITY_DOMAIN_HSA_API: break;
@@ -812,8 +927,25 @@ PUBLIC_API roctracer_status_t roctracer_set_properties(
     case ACTIVITY_DOMAIN_HSA_OPS: {
       // HSA OPS properties
       roctracer::hsa_ops_properties_t* ops_properties = reinterpret_cast<roctracer::hsa_ops_properties_t*>(properties);
+      HsaApiTable* table = reinterpret_cast<HsaApiTable*>(ops_properties->table);
       roctracer::hsa_support::async_copy_callback_fun = ops_properties->async_copy_callback_fun;
       roctracer::hsa_support::async_copy_callback_arg = ops_properties->async_copy_callback_arg;
+      roctracer::hsa_support::output_prefix = ops_properties->output_prefix;
+
+      // HSA dispatches intercepting
+      rocprofiler::SaveHsaApi(table);
+      rocprofiler::ProxyQueue::InitFactory();
+      rocprofiler::ProxyQueue::HsaIntercept(table);
+      rocprofiler::InterceptQueue::HsaIntercept(table);
+
+      // HSA async-copy tracing
+      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");
+      roctracer::hsa_amd_memory_async_copy_fn = table->amd_ext_->hsa_amd_memory_async_copy_fn;
+      roctracer::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 = roctracer::hsa_amd_memory_async_copy_interceptor;
+      table->amd_ext_->hsa_amd_memory_async_copy_rect_fn = roctracer::hsa_amd_memory_async_copy_rect_interceptor;
+
       break;
     }
     case ACTIVITY_DOMAIN_HSA_API: {
@@ -834,17 +966,39 @@ PUBLIC_API roctracer_status_t roctracer_set_properties(
 }
 
 // HSA-runtime tool on-load method
-PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
+PUBLIC_API bool roctracer_load(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
                        const char* const* failed_tool_names) {
-  // enabled HSA async-copy tracing
-  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");
-  roctracer::hsa_amd_memory_async_copy_fn = table->amd_ext_->hsa_amd_memory_async_copy_fn;
-  roctracer::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 = roctracer::hsa_amd_memory_async_copy_interceptor;
-  table->amd_ext_->hsa_amd_memory_async_copy_rect_fn = roctracer::hsa_amd_memory_async_copy_rect_interceptor;
+  static bool is_loaded = false;
+  if (is_loaded) return true;
+  is_loaded = true;
 
   return true;
 }
 
+PUBLIC_API void roctracer_unload() {
+  static bool is_unloaded = false;
+  if (is_unloaded) return;
+  is_unloaded = true;
+
+  roctracer::trace_buffer.Flush(roctracer::COPY_ENTRY_TYPE, roctracer::hsa_async_copy_handler);
+  roctracer::trace_buffer.Flush(roctracer::KERNEL_ENTRY_TYPE, roctracer::hsa_kernel_handler);
+  if ((roctracer::hsa_support::output_prefix != NULL) && (roctracer::kernel_file_handle != NULL)) fclose(roctracer::kernel_file_handle);
+}
+
+PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
+                       const char* const* failed_tool_names) {
+  return roctracer_load(table, runtime_version, failed_tool_count, failed_tool_names);
+}
+PUBLIC_API void OnUnload() { roctracer_unload(); }
+
+CONSTRUCTOR_API void constructor() {
+  roctracer::util::Logger::Create();
+}
+
+DESTRUCTOR_API void destructor() {
+  roctracer_unload();
+  util::HsaRsrcFactory::Destroy();
+  roctracer::util::Logger::Destroy();
+}
+
 }  // extern "C"

src/core/trace_buffer.h

@@ -0,0 +1,79 @@
+#ifndef SRC_CORE_TRACE_BUFFER_H_
+#define SRC_CORE_TRACE_BUFFER_H_
+
+namespace roctracer {
+enum {
+  TRACE_ENTRY_INV = 0,
+  TRACE_ENTRY_INIT = 1,
+  TRACE_ENTRY_COMPL = 2
+};
+enum {
+  API_ENTRY_TYPE,
+  COPY_ENTRY_TYPE,
+  KERNEL_ENTRY_TYPE
+};
+
+struct trace_entry_t {
+  std::atomic<uint32_t> valid;
+  uint32_t type;
+  uint64_t dispatch;
+  uint64_t begin;                                      // kernel begin timestamp, ns
+  uint64_t end;                                        // kernel end timestamp, ns
+  uint64_t complete;
+  hsa_agent_t agent;
+  hsa_signal_t orig;
+  hsa_signal_t signal;
+  union {
+    struct {
+    } copy;
+    struct {
+      const char* name;
+      hsa_agent_t agent;
+      uint32_t tid;
+    } kernel;
+  };
+};
+
+template <typename Entry>
+class TraceBuffer {
+  public:
+  typedef void (*callabck_t)(Entry*);
+
+  TraceBuffer(uint32_t size) {
+    size_ = size;
+    data_ = new Entry[size_];;
+    memset(data_, 0, size_ * sizeof(Entry));
+    read_pointer_ = data_;
+  }
+
+  Entry* GetEntry() {
+    Entry* ptr = read_pointer_.fetch_add(1);
+    if (ptr >= (data_ + size_))  {
+      fprintf(stderr, "GetEntry: trace buffer is out of range\n");
+      abort();
+    }
+    return ptr;
+  }
+
+  void Flush(uint32_t type, callabck_t fun) {
+    Entry* ptr = data_;
+    for (; (ptr < read_pointer_) && (ptr < (data_ + size_)); ptr++) {
+      if (ptr->type == type) {
+        if (ptr->valid == TRACE_ENTRY_COMPL) {
+          fun(ptr);
+        }
+      }
+    }
+    if (ptr >= (data_ + size_))  {
+      fprintf(stderr, "Flush: trace buffer is out of range\n");
+    }
+  }
+
+  private:
+  Entry* data_;
+  uint32_t size_;
+  std::atomic<Entry*> read_pointer_;
+};
+}  // namespace roctracer
+
+#endif  // SRC_CORE_TRACE_BUFFER_H_