diff --git a/projects/clr/rocclr/device/devhostcall.cpp b/projects/clr/rocclr/device/devhostcall.cpp index 99da349697..d59e4f799b 100644 --- a/projects/clr/rocclr/device/devhostcall.cpp +++ b/projects/clr/rocclr/device/devhostcall.cpp @@ -273,7 +273,7 @@ class HostcallListener { }; HostcallListener* hostcallListener = nullptr; -amd::Monitor listenerLock("Hostcall listener lock"); +extern amd::Monitor listenerLock; constexpr static uint64_t kTimeoutFloor = K * K * 4; constexpr static uint64_t kTimeoutCeil = K * K * 16; static struct Init { diff --git a/projects/clr/rocclr/device/rocm/rocdevice.cpp b/projects/clr/rocclr/device/rocm/rocdevice.cpp index 1cc442795c..e9bafafa3b 100644 --- a/projects/clr/rocclr/device/rocm/rocdevice.cpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.cpp @@ -246,6 +246,11 @@ void Device::checkAtomicSupport() { } Device::~Device() { + if (coopHostcallBuffer_) { + amd::disableHostcalls(coopHostcallBuffer_); + context().svmFree(coopHostcallBuffer_); + coopHostcallBuffer_ = nullptr; + } // Release cached map targets for (uint i = 0; mapCache_ != nullptr && i < mapCache_->size(); ++i) { if ((*mapCache_)[i] != nullptr) { @@ -301,12 +306,6 @@ Device::~Device() { delete[] p2p_agents_list_; - if (coopHostcallBuffer_) { - amd::disableHostcalls(coopHostcallBuffer_); - context().svmFree(coopHostcallBuffer_); - coopHostcallBuffer_ = nullptr; - } - if (0 != prefetch_signal_.handle) { hsa_signal_destroy(prefetch_signal_); } diff --git a/projects/clr/rocclr/platform/runtime.cpp b/projects/clr/rocclr/platform/runtime.cpp index 444f736468..1b710ca207 100644 --- a/projects/clr/rocclr/platform/runtime.cpp +++ b/projects/clr/rocclr/platform/runtime.cpp @@ -102,6 +102,10 @@ void Runtime::tearDown() { initialized_ = false; } +// ~RuntimeTearDown() will reference listenerLock. +// listenerLock will be constructed ealier and destructed later than +// runtime_tear_down. +amd::Monitor listenerLock("Hostcall listener lock"); std::vector RuntimeTearDown::external_; RuntimeTearDown::~RuntimeTearDown() {