From 150869be6aad4ea6f39e4de5708ce46217560bc9 Mon Sep 17 00:00:00 2001 From: Saurabh Verma Date: Mon, 13 May 2024 17:24:22 -0500 Subject: [PATCH] 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: 34606db8c3b5cc5bb812fe4ff0c982363b060799] --- projects/rocprofiler/src/core/profile.h | 19 ++++++++++++++++++- .../rocprofiler/src/util/hsa_rsrc_factory.cpp | 18 ++++++++++++++++-- .../rocprofiler/src/util/hsa_rsrc_factory.h | 1 + 3 files changed, 35 insertions(+), 3 deletions(-) diff --git a/projects/rocprofiler/src/core/profile.h b/projects/rocprofiler/src/core/profile.h index 8c9c7561eb..a60adf210d 100644 --- a/projects/rocprofiler/src/core/profile.h +++ b/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_; diff --git a/projects/rocprofiler/src/util/hsa_rsrc_factory.cpp b/projects/rocprofiler/src/util/hsa_rsrc_factory.cpp index a419f49a47..2c4718643e 100644 --- a/projects/rocprofiler/src/util/hsa_rsrc_factory.cpp +++ b/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_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 diff --git a/projects/rocprofiler/src/util/hsa_rsrc_factory.h b/projects/rocprofiler/src/util/hsa_rsrc_factory.h index 6426cf953d..f23b31fd9a 100644 --- a/projects/rocprofiler/src/util/hsa_rsrc_factory.h +++ b/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;