From 5ebd50c0b4bef28814b416507e1b088c0bb8260a Mon Sep 17 00:00:00 2001
From: David Yat Sin <77975354+dayatsin-amd@users.noreply.github.com>
Date: Wed, 17 Dec 2025 23:52:20 -0500
Subject: [PATCH] rocr: Fix asyncHandler segfault (#2261)

Fix initialization order for the async events handler. The polling
thread would launch before the wake signal is initialized.
---
 .../runtime/hsa-runtime/core/inc/runtime.h     | 18 +++++++++++-------
 .../hsa-runtime/core/runtime/runtime.cpp       | 18 ++++++++++--------
 2 files changed, 21 insertions(+), 15 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 bfffcb6c84..800bc94ca5 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/runtime.h
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/runtime.h
@@ -569,12 +569,16 @@ class Runtime {
   struct AsyncEventsInfo;
   struct AsyncEventsControl {
     AsyncEventsControl(AsyncEventsInfo *asyncInfo);
+    void Start();
     void Shutdown();
 
     hsa_signal_t wake;
-    os::Thread thread_;
     bool exit;
-  };
+
+    private:
+    AsyncEventsInfo* info_;
+    os::Thread thread_;
+ };
 
   struct AsyncEvents {
     void PushBack(hsa_signal_t signal, hsa_signal_condition_t cond,
@@ -682,13 +686,13 @@ class Runtime {
   };
 
   struct AsyncEventsInfo {
-    bool monitor_exceptions;
-    AsyncEventsControl control;
-    AsyncEvents events;
-    ConcurrentAsyncEvents new_events;
-
     AsyncEventsInfo(bool exceptions);
     ~AsyncEventsInfo();
+
+    bool monitor_exceptions;
+    AsyncEvents events;
+    ConcurrentAsyncEvents new_events;
+    AsyncEventsControl control;
   };
 
   struct PrefetchRange;
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 a0edafdfe2..3d34133c7f 100644
--- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp
+++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/runtime.cpp
@@ -2314,9 +2314,10 @@ void Runtime::PrintMemoryMapNear(void* ptr) {
 }
 
 Runtime::AsyncEventsInfo::AsyncEventsInfo(bool exceptions_)
-  : control(this), events(), new_events(), monitor_exceptions(exceptions_) {
+  : monitor_exceptions(exceptions_), events(), new_events(), control(this) {
 
   events.PushBack(control.wake, HSA_SIGNAL_CONDITION_NE, 0, NULL, NULL);
+  control.Start();
 }
 
 Runtime::AsyncEventsInfo::~AsyncEventsInfo() {
@@ -2324,19 +2325,20 @@ Runtime::AsyncEventsInfo::~AsyncEventsInfo() {
 }
 
 Runtime::AsyncEventsControl::AsyncEventsControl(AsyncEventsInfo *asyncInfo)
-  : exit(false) {
-
-  int priority = asyncInfo->monitor_exceptions ? os::OS_THREAD_PRIORITY_DEFAULT :
-                  runtime_singleton_->flag().async_events_thread_priority();
+  : info_(asyncInfo), exit(false) {
 
   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");
+}
 
-  asyncInfo->control.thread_ = os::CreateThread(AsyncEventsLoop, asyncInfo, 0, priority);
-  if (!asyncInfo->control.thread_)
+void Runtime::AsyncEventsControl::Start() {
+  int priority = info_->monitor_exceptions ? os::OS_THREAD_PRIORITY_DEFAULT :
+                  runtime_singleton_->flag().async_events_thread_priority();
+
+  thread_ = os::CreateThread(AsyncEventsLoop, info_, 0, priority);
+  if (!thread_)
     throw AMD::hsa_exception(HSA_STATUS_ERROR, "Failed to initialize async handler thread");
-
 }
 
 Runtime::Runtime()