Fix for crash Error: HsaRsrcFactory::SignalWait: signal_value(1), ret_value(-1)

Context::GetData() would crash when it tries to wait on the completion signal of the stop packet. The expected signal value after waiting in HsaRsrcFactory::SignalWait() is supposed to be 0 but sometimes it comes up as -1. The signal being checked has already decremented before. Profile::Finalize() was assigning the same completion signal to the read and stop packet. So those two packets have been assigned different completion signals. This fix only affects standalone profiling mode.

Change-Id: I4b16825019e58d95d70188a72b2cc5871e09dd81


[ROCm/rocprofiler commit: 34606db8c3]

projects/rocprofiler/src/core/profile.h

@@ -104,7 +104,9 @@ class Profile {
   Profile(const util::AgentInfo* agent_info) : agent_info_(agent_info) {
     profile_ = {};
     profile_.agent = agent_info->dev_id;
+    is_standalone_mode = util::check_standalone_mode();
     completion_signal_ = {};
+    stop_completion_signal_ = {};
     dispatch_signal_ = {};
     barrier_signal_ = {};
     read_signal_ = {};
@@ -121,6 +123,10 @@ class Profile {
       hsa_status_t status = hsa_signal_destroy(completion_signal_);
       if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, "signal_destroy " << std::hex << status);
     }
+    if (is_standalone_mode == true && stop_completion_signal_.handle) {
+      hsa_status_t status = hsa_signal_destroy(stop_completion_signal_);
+      if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, "signal_destroy " << std::hex << status);
+    }
     if (dispatch_signal_.handle) {
       hsa_status_t status = hsa_signal_destroy(dispatch_signal_);
       if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, "signal_destroy " << std::hex << status);
@@ -214,9 +220,17 @@ class Profile {
       dummy_signal.handle = 0;
       start.completion_signal = dummy_signal;
 
-      // Set completion signal of read/stop
+      // Set completion signal of read
       status = hsa_signal_create(1, 0, NULL, &completion_signal_);
       if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "signal_create " << std::hex << status);
+
+      if (is_standalone_mode) {
+        // Set completion signal of stop
+        status = hsa_signal_create(1, 0, NULL, &stop_completion_signal_);
+        if (status != HSA_STATUS_SUCCESS)
+          EXC_RAISING(status, "signal_create " << std::hex << status);
+      }
+
       if (is_concurrent) {
         status = hsa_signal_create(1, 0, NULL, &read_signal_);
         if (status != HSA_STATUS_SUCCESS)
@@ -227,6 +241,7 @@ class Profile {
         read.completion_signal = completion_signal_;
       }
       stop.completion_signal = completion_signal_;
+      if (is_standalone_mode) stop.completion_signal = stop_completion_signal_;
 
       status = hsa_signal_create(1, 0, NULL, &dispatch_signal_);
       if (status != HSA_STATUS_SUCCESS) EXC_RAISING(status, "signal_create " << std::hex << status);
@@ -312,7 +327,9 @@ class Profile {
   bool is_legacy_;
   profile_t profile_;
   info_vector_t info_vector_;
+  bool is_standalone_mode;
   hsa_signal_t completion_signal_;
+  hsa_signal_t stop_completion_signal_;
   hsa_signal_t dispatch_signal_;
   hsa_signal_t barrier_signal_;
   hsa_signal_t read_signal_;

projects/rocprofiler/src/util/hsa_rsrc_factory.cpp

@@ -49,6 +49,20 @@ POSSIBILITY OF SUCH DAMAGE.
 namespace rocprofiler {
 namespace util {
 
+bool check_standalone_mode() {
+  static bool is_standalone_mode = [] {
+    // Checking environment variable to see if interception is enabled.
+    // value of zero indicates standalone mode
+    const char* intercept_env = getenv("ROCP_HSA_INTERCEPT");
+    int intercept_env_value = 0;
+    if (intercept_env != NULL) {
+      intercept_env_value = atoi(intercept_env);
+    }
+    return intercept_env_value == 0;
+  }();
+  return is_standalone_mode;
+}
+
 // Demangle C++ symbol name
 static const char* cpp_demangle(const char* symname) {
   size_t size = 0;
@@ -392,8 +406,8 @@ const AgentInfo* HsaRsrcFactory::AddAgentInfo(const hsa_agent_t agent) {
     agent_info->vgpr_block_size = 4;
 
     if (hsa_api_.hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_XCC),
-                        &agent_info->xcc_num) != HSA_STATUS_SUCCESS) {
-        agent_info->xcc_num = 1;
+                                    &agent_info->xcc_num) != HSA_STATUS_SUCCESS) {
+      agent_info->xcc_num = 1;
     };
 
     // Set GPU index

projects/rocprofiler/src/util/hsa_rsrc_factory.h

@@ -73,6 +73,7 @@ namespace util {
 static const size_t MEM_PAGE_BYTES = 0x1000;
 static const size_t MEM_PAGE_MASK = MEM_PAGE_BYTES - 1;
 typedef decltype(hsa_agent_t::handle) hsa_agent_handle_t;
+bool check_standalone_mode();
 
 struct hsa_pfn_t {
   decltype(::hsa_init)* hsa_init;