unload and destruct normalizing

2019-08-10 00:14:04 -05:00
@@ -526,7 +526,7 @@ hsa_status_t hsa_amd_memory_async_copy_interceptor(
    status = hsa_amd_memory_async_copy_fn(dst, dst_agent, src,
                                          src_agent, size, num_dep_signals,
                                          dep_signals, entry->signal);
-    if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, "hsa_amd_memory_async_copy interceptor");
+    if (status != HSA_STATUS_SUCCESS) ::proxy::Tracker::Disable(entry);
  }
  else
  {
@@ -551,7 +551,7 @@ hsa_status_t hsa_amd_memory_async_copy_rect_interceptor(
                                               src_offset, range, copy_agent,
                                               dir, num_dep_signals, dep_signals,
                                               entry->signal);
-    if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, "hsa_amd_memory_async_copy interceptor");
+    if (status != HSA_STATUS_SUCCESS) ::proxy::Tracker::Disable(entry);
  }
  else
  {
@@ -985,36 +985,51 @@ PUBLIC_API roctracer_status_t roctracer_set_properties(
 // HSA-runtime tool on-load method
 PUBLIC_API bool roctracer_load(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
                       const char* const* failed_tool_names) {
+//  printf("LIB roctracer_load\n"); fflush(stdout);
  static bool is_loaded = false;
  if (is_loaded) return true;
  is_loaded = true;

+//  printf("LIB roctracer_load end\n"); fflush(stdout);
  return true;
 }

 PUBLIC_API void roctracer_unload(bool destruct) {
+//  printf("LIB roctracer_unload\n"); fflush(stdout);
  static bool is_unloaded = false;
  if (is_unloaded) return;
  is_unloaded = true;

  if (destruct == false) roctracer::trace_buffer.Flush();
  if ((roctracer::hsa_support::output_prefix != NULL) && (roctracer::kernel_file_handle != NULL)) fclose(roctracer::kernel_file_handle);
+//  printf("LIB roctracer_unload end\n"); fflush(stdout);
 }

 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);
+//  printf("LIB OnLoad\n"); fflush(stdout);
+  const bool ret = roctracer_load(table, runtime_version, failed_tool_count, failed_tool_names);
+//  printf("LIB OnLoad end\n"); fflush(stdout);
+  return ret;
+}
+PUBLIC_API void OnUnload() {
+//  printf("LIB OnUnload\n"); fflush(stdout);
+  roctracer_unload(false);
+//  printf("LIB OnUnload end\n"); fflush(stdout);
 }
-PUBLIC_API void OnUnload() { roctracer_unload(false); }

 CONSTRUCTOR_API void constructor() {
+//  printf("LIB constructor\n"); fflush(stdout);
  roctracer::util::Logger::Create();
+//  printf("LIB constructor end\n"); fflush(stdout);
 }

 DESTRUCTOR_API void destructor() {
+//  printf("LIB destructor\n"); fflush(stdout);
  roctracer_unload(true);
  util::HsaRsrcFactory::Destroy();
  roctracer::util::Logger::Destroy();
+//  printf("LIB destructor end\n"); fflush(stdout);
 }

 }  // extern "C"
@@ -4,6 +4,7 @@
 #include <list>
 #include <mutex>
 #include <pthread.h>
+#include <string.h>

 #define PTHREAD_CALL(call)                                                                         \
  do {                                                                                             \
@@ -61,8 +62,10 @@ class TraceBuffer {
    callback_t fun;
  };

-  TraceBuffer(const char* name, uint32_t size, flush_prm_t* flush_prm_arr, uint32_t flush_prm_count) {
-    (void) name;
+  TraceBuffer(const char* name, uint32_t size, flush_prm_t* flush_prm_arr, uint32_t flush_prm_count) :
+    is_flushed_(ATOMIC_FLAG_INIT)
+  {
+    name_ = strdup(name);
    size_ = size;
    data_ = allocate_fun();
    next_ = NULL;
@@ -72,7 +75,6 @@ class TraceBuffer {

    flush_prm_arr_ = flush_prm_arr;
    flush_prm_count_ = flush_prm_count;
-    is_flushed_ = false;

    PTHREAD_CALL(pthread_mutex_init(&work_mutex_, NULL));
    PTHREAD_CALL(pthread_cond_init(&work_cond_, NULL));
@@ -85,7 +87,7 @@ class TraceBuffer {
    PTHREAD_CALL(pthread_join(work_thread_, &res));
    if (res != PTHREAD_CANCELED) abort_run("~TraceBuffer: consumer thread wasn't stopped correctly");

-    if (is_flushed_ == false) flush_buf();
+    Flush();
  }


@@ -96,28 +98,28 @@ class TraceBuffer {
  }

  void Flush() {
-    PTHREAD_CALL(pthread_mutex_lock(&work_mutex_));
    flush_buf();
-    PTHREAD_CALL(pthread_mutex_unlock(&work_mutex_));
  }

  private:
  void flush_buf() {
-    is_flushed_ = true;
-    for (flush_prm_t* prm = flush_prm_arr_; prm < flush_prm_arr_ + flush_prm_count_; prm++) {
-      uint32_t type = prm->type;
-      callback_t fun = prm->fun;
-      pointer_t pointer = 0;
-      for (Entry* ptr : buf_list_) {
-        Entry* end = ptr + size_;
-        while ((ptr < end) && (pointer < read_pointer_)) {
-          if (ptr->type == type) {
-            if (ptr->valid == TRACE_ENTRY_COMPL) {
-              fun(ptr);
+    const bool is_flushed = atomic_flag_test_and_set_explicit(&is_flushed_, std::memory_order_acquire);
+    if (is_flushed == false) {
+      for (flush_prm_t* prm = flush_prm_arr_; prm < flush_prm_arr_ + flush_prm_count_; prm++) {
+        uint32_t type = prm->type;
+        callback_t fun = prm->fun;
+        pointer_t pointer = 0;
+        for (Entry* ptr : buf_list_) {
+          Entry* end = ptr + size_;
+          while ((ptr < end) && (pointer < read_pointer_)) {
+            if (ptr->type == type) {
+              if (ptr->valid == TRACE_ENTRY_COMPL) {
+                fun(ptr);
+              }
            }
+            ptr++;
+            pointer++;
          }
-          ptr++;
-          pointer++;
        }
      }
    }
@@ -164,6 +166,7 @@ class TraceBuffer {
    abort();
  }

+  const char* name_;
  uint32_t size_;
  Entry* data_;
  Entry* next_;
@@ -173,7 +176,7 @@ class TraceBuffer {

  flush_prm_t* flush_prm_arr_;
  uint32_t flush_prm_count_;
-  bool is_flushed_;
+  volatile std::atomic_flag is_flushed_;

  pthread_t work_thread_;
  pthread_mutex_t work_mutex_;
@@ -60,6 +60,12 @@ class Tracker {
    if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "hsa_amd_signal_async_handler");
  }

+  // Delete tracker entry
+  inline static void Disable(entry_t* entry) {
+    hsa_signal_destroy(entry->signal);
+    entry->valid.store(roctracer::TRACE_ENTRY_INV, std::memory_order_release);
+  }
+
  private:
  // Entry completion
  inline static void Complete(hsa_signal_value_t signal_value, entry_t* entry) {
@@ -367,6 +367,7 @@ FILE* open_output_file(const char* prefix, const char* name) {
 // HSA-runtime tool on-load method
 extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
                       const char* const* failed_tool_names) {
+//  printf("TOOL OnLoad\n"); fflush(stdout);
  timer = new hsa_rt_utils::Timer(table->core_->hsa_system_get_info_fn);

  // API traces switches
@@ -497,11 +498,13 @@ extern "C" PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version,
    roctracer_set_properties(ACTIVITY_DOMAIN_HIP_API, (void*)mark_api_callback);
  }

+//  printf("TOOL OnLoad end\n"); fflush(stdout);
  return roctracer_load(table, runtime_version, failed_tool_count, failed_tool_names);
 }

 // tool unload method
 void tool_unload(bool destruct) {
+//  printf("TOOL tool_unload\n"); fflush(stdout);
  static bool is_unloaded = false;
  if (is_unloaded) {
    return;
@@ -530,13 +533,24 @@ void tool_unload(bool destruct) {

    if (destruct == false) hip_api_trace_buffer.Flush();

-    fclose(hip_api_file_handle);
-    fclose(hcc_activity_file_handle);
+    if (hip_api_file_handle != stdout) fclose(hip_api_file_handle);
+    if (hcc_activity_file_handle != stdout) fclose(hcc_activity_file_handle);
  }
+//  printf("TOOL tool_unload end\n"); fflush(stdout);
 }

 // HSA-runtime on-unload method
-extern "C" PUBLIC_API void OnUnload() { tool_unload(false); }
+extern "C" PUBLIC_API void OnUnload() {
+//  printf("TOOL OnUnload\n"); fflush(stdout);
+  tool_unload(false);
+//  printf("TOOL OnUnload end\n"); fflush(stdout);
+}

-extern "C" CONSTRUCTOR_API void constructor() {}
-extern "C" DESTRUCTOR_API void destructor() { tool_unload(true); }
+extern "C" CONSTRUCTOR_API void constructor() {
+//  printf("TOOL constructor ...end\n"); fflush(stdout);
+}
+extern "C" DESTRUCTOR_API void destructor() {
+//  printf("TOOL destructor\n"); fflush(stdout);
+  tool_unload(true);
+//  printf("TOOL destructor end\n"); fflush(stdout);
+}