From 7b097599c4191fe13866d30237b5f49ea9989492 Mon Sep 17 00:00:00 2001 From: David Yat Sin <77975354+dayatsin-amd@users.noreply.github.com> Date: Wed, 12 Nov 2025 13:54:26 -0500 Subject: [PATCH] rocr: Fix race condition in SetAsyncSignalHandler (#1642) Fix race condition when SetAsyncSignalHandler for the first time because async_events_thread_ could be null and launched twice. Refactored async-events to use lazy_pointer. --- .../runtime/hsa-runtime/core/inc/runtime.h | 30 +++---- .../hsa-runtime/core/runtime/runtime.cpp | 83 ++++++++++--------- 2 files changed, 58 insertions(+), 55 deletions(-) diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/runtime.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/runtime.h index 251298857a..ec2e20a622 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/runtime.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/runtime.h @@ -69,6 +69,7 @@ #include "core/inc/interrupt_signal.h" #include "core/inc/memory_region.h" #include "core/inc/signal.h" +#include "core/util/lazy_ptr.h" #include "core/inc/svm_profiler.h" #include "core/inc/thunk_loader.h" #include "core/util/flag.h" @@ -565,13 +566,13 @@ class Runtime { amdgpu_bo_handle ldrm_bo; }; + struct AsyncEventsInfo; struct AsyncEventsControl { - AsyncEventsControl() : async_events_thread_(NULL) {} + AsyncEventsControl(AsyncEventsInfo *asyncInfo); void Shutdown(); hsa_signal_t wake; - os::Thread async_events_thread_; - HybridMutex lock; + os::Thread thread_; bool exit; }; @@ -680,6 +681,16 @@ class Runtime { AsyncEventsPool asyncEventPool_; }; + struct AsyncEventsInfo { + AsyncEventsControl control; + AsyncEvents events; + ConcurrentAsyncEvents new_events; + bool monitor_exceptions; + + AsyncEventsInfo(bool exceptions); + ~AsyncEventsInfo(); + }; + struct PrefetchRange; typedef std::map prefetch_map_t; @@ -829,17 +840,8 @@ class Runtime { // Deprecated HSA Region API GPU (for legacy APU support only) Agent* region_gpu_; - struct AsyncEventsInfo { - AsyncEventsControl control; - AsyncEvents events; - ConcurrentAsyncEvents new_events; - bool monitor_exceptions; - - AsyncEventsInfo() : control(), events(), new_events(), monitor_exceptions(false) {} - }; - - struct AsyncEventsInfo asyncSignals_; - struct AsyncEventsInfo asyncExceptions_; + lazy_ptr asyncSignals_; + lazy_ptr asyncExceptions_; // System clock frequency. uint64_t sys_clock_freq_; diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp index 16736780b6..f875f01300 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp @@ -845,41 +845,18 @@ hsa_status_t Runtime::SetAsyncSignalHandler(hsa_signal_t signal, hsa_signal_value_t value, hsa_amd_signal_handler handler, void* arg) { - struct AsyncEventsInfo* asyncInfo = &asyncSignals_; - int priority = runtime_singleton_->flag().async_events_thread_priority(); + bool exception = false; - if (signal.handle != 0) { + if (signal.handle) { // Indicate that this signal is in use. hsa_signal_handle(signal)->Retain(); core::Signal* coreSignal = core::Signal::Convert(signal); - if (coreSignal->EopEvent() && coreSignal->EopEvent()->EventData.EventType != HSA_EVENTTYPE_SIGNAL) { - priority = os::OS_THREAD_PRIORITY_DEFAULT; - asyncInfo = &asyncExceptions_; - } + exception = !!(coreSignal->EopEvent() && coreSignal->EopEvent()->EventData.EventType != HSA_EVENTTYPE_SIGNAL); } - - // Lazy initializer - if (asyncInfo->control.async_events_thread_ == NULL) { - // Create monitoring thread control signal - auto err = HSA::hsa_signal_create(0, 0, NULL, &asyncInfo->control.wake); - if (err != HSA_STATUS_SUCCESS) { - assert(false && "Asyncronous events control signal creation error."); - return HSA_STATUS_ERROR_OUT_OF_RESOURCES; - } - asyncInfo->events.PushBack(asyncInfo->control.wake, HSA_SIGNAL_CONDITION_NE, - 0, NULL, NULL); - - // Start event monitoring thread - asyncInfo->control.exit = false; - asyncInfo->control.async_events_thread_ = - os::CreateThread(AsyncEventsLoop, asyncInfo, 0, priority); - if (asyncInfo->control.async_events_thread_ == NULL) { - assert(false && "Asyncronous events thread creation error."); - return HSA_STATUS_ERROR_OUT_OF_RESOURCES; - } - } + // Lazy initializer asyncExceptions_ and asyncSignals_ will be constructed on first dereference + struct AsyncEventsInfo* asyncInfo = exception ? (*asyncExceptions_).get() : (*asyncSignals_).get(); asyncInfo->new_events.PushBack(signal, cond, value, handler, arg); @@ -2329,6 +2306,32 @@ void Runtime::PrintMemoryMapNear(void* ptr) { } } +Runtime::AsyncEventsInfo::AsyncEventsInfo(bool exceptions_) + : control(this), events(), new_events(), monitor_exceptions(exceptions_) { + + events.PushBack(control.wake, HSA_SIGNAL_CONDITION_NE, 0, NULL, NULL); +} + +Runtime::AsyncEventsInfo::~AsyncEventsInfo() { + control.Shutdown(); +} + +Runtime::AsyncEventsControl::AsyncEventsControl(AsyncEventsInfo *asyncInfo) + : exit(false) { + + int priority = asyncInfo->monitor_exceptions ? os::OS_THREAD_PRIORITY_DEFAULT : + runtime_singleton_->flag().async_events_thread_priority(); + + auto err = HSA::hsa_signal_create(0, 0, NULL, &wake); + if (err != HSA_STATUS_SUCCESS) + throw AMD::hsa_exception(HSA_STATUS_ERROR, "Failed to allocate async handler signal"); + + thread_ = os::CreateThread(AsyncEventsLoop, asyncInfo, 0, priority); + if (!asyncInfo->control.thread_) + throw AMD::hsa_exception(HSA_STATUS_ERROR, "Failed to initialize async handler thread"); + +} + Runtime::Runtime() : loader_(nullptr), region_gpu_(nullptr), @@ -2346,8 +2349,6 @@ Runtime::Runtime() virtual_mem_api_supported_ = false; ipc_dmabuf_supported_ = false; xnack_enabled_ = false; - asyncSignals_.monitor_exceptions = false; - asyncExceptions_.monitor_exceptions = true; g_use_interrupt_wait = true; g_use_mwaitx = true; ::_amdgpu_r_debug = {11, @@ -2356,7 +2357,6 @@ Runtime::Runtime() &_loader_debug_state), r_debug::RT_CONSISTENT, 0}; - log_file = stderr; } @@ -2388,6 +2388,9 @@ hsa_status_t Runtime::Load() { return HSA_STATUS_ERROR_OUT_OF_RESOURCES; } + asyncSignals_.reset(new AsyncEventsInfo(false)); + asyncExceptions_.reset(new AsyncEventsInfo(true)); + // Setup system clock frequency for the first time. if (sys_clock_freq_ == 0) { sys_clock_freq_ = os::SystemClockFrequency(); @@ -2445,8 +2448,8 @@ void Runtime::Unload() { agent->ReleaseResources(); } - asyncSignals_.control.Shutdown(); - asyncExceptions_.control.Shutdown(); + asyncSignals_.reset(); + asyncExceptions_.reset(); if (vm_fault_signal_ != nullptr) { vm_fault_signal_->DestroySignal(); @@ -2831,14 +2834,12 @@ void Runtime::CloseTools() { } void Runtime::AsyncEventsControl::Shutdown() { - if (async_events_thread_ != NULL) { - exit = true; - hsa_signal_handle(wake)->StoreRelaxed(1); - os::WaitForThread(async_events_thread_); - os::CloseThread(async_events_thread_); - async_events_thread_ = NULL; - HSA::hsa_signal_destroy(wake); - } + exit = true; + hsa_signal_handle(wake)->StoreRelaxed(1); + os::WaitForThread(thread_); + os::CloseThread(thread_); + thread_ = NULL; + HSA::hsa_signal_destroy(wake); } void Runtime::AsyncEvents::PushBack(hsa_signal_t signal,