diff --git a/projects/rocprofiler-sdk/.cmake-format.yaml b/projects/rocprofiler-sdk/.cmake-format.yaml
index 0e41b4f54a2..5cb250f302a 100644
--- a/projects/rocprofiler-sdk/.cmake-format.yaml
+++ b/projects/rocprofiler-sdk/.cmake-format.yaml
@@ -367,6 +367,19 @@ parse:
         PASS_REGULAR_EXPRESSION: '*'
         FAIL_REGULAR_EXPRESSION: '*'
         SKIP_REGULAR_EXPRESSION: '*'
+    rocprofiler_add_integration_cleanup_test:
+      kwargs:
+        DIRECTORIES: '*'
+        FILES: '*'
+        LABELS: '*'
+        FIXTURES_SETUP: '*'
+    add_signal_handler_test:
+      kwargs:
+        PROCESS_TYPE: '*'
+        VALIDATE_EXPECT: '*'
+        PASS_REGEX: '*'
+        APP_EXTRA_FLAGS: '*'
+        ROCPROFV3_EXTRA_FLAGS: '*'
   override_spec: {}
   vartags: []
   proptags: []
diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp
index 8a7af7a254e..4f03cec7464 100644
--- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp
+++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp
@@ -104,6 +104,7 @@
 #include <iomanip>
 #include <limits>
 #include <mutex>
+#include <new>
 #include <optional>
 #include <shared_mutex>
 #include <string>
@@ -115,7 +116,10 @@
 #include <vector>
 
 #include <dlfcn.h>
+#include <linux/futex.h>
+#include <sys/eventfd.h>
 #include <sys/mman.h>
+#include <sys/syscall.h>
 #include <sys/types.h>
 #include <sys/wait.h>
 
@@ -204,6 +208,55 @@ is_handled_signal(int signum)
     return false;
 }
 
+struct signal_worker_state
+{
+    int                   eventfd       = -1;   // handler -> worker wakeup fd
+    std::atomic<uint32_t> finalize_done = {};   // worker sets when flush done (futex word)
+    std::atomic<uint32_t> handling      = {0};  // 1 once our handler owns the path
+    std::atomic_flag      finalized     = ATOMIC_FLAG_INIT;  // runs finalize_rocprofv3 once
+    int                   signo         = {0};               // signal handled (0 == normal exit)
+    std::thread           thread        = {};                // the finalization worker
+
+    ~signal_worker_state() { join(); }
+
+    // Join the finalization worker and close the eventfd. Idempotent and self-safe (a call from the
+    // worker itself skips the join). Called from finalize_rocprofv3 (atexit/main teardown) and from
+    // the destructor, so the worker is never left joinable at static destruction
+    void join()
+    {
+        if(thread.joinable() && thread.get_id() != std::this_thread::get_id())
+        {
+            // Wake the worker's blocking read() so it can exit
+            // closing the fd while read() blocks is UB.
+            if(eventfd >= 0)
+            {
+                uint64_t val = 1;
+                if(write(eventfd, &val, sizeof(val)) < 0)
+                    ROCP_WARNING << "signal worker: eventfd wake before join failed: "
+                                 << strerror(errno);
+            }
+            thread.join();
+            if(eventfd >= 0)
+            {
+                if(close(eventfd) < 0)
+                    ROCP_WARNING << "signal worker: close(eventfd) failed: " << strerror(errno);
+                eventfd = -1;
+            }
+        }
+    }
+};
+
+static_assert(std::atomic<uint32_t>::is_always_lock_free,
+              "rocprofv3 signal path requires lock-free finalize_done atomic support");
+
+auto&
+get_signal_worker()
+{
+    static auto*& _v = common::static_object<signal_worker_state>::construct();
+    auto&         sw = *CHECK_NOTNULL(_v);
+    return sw;
+}
+
 struct buffer_ids
 {
     rocprofiler_buffer_id_t hsa_api_trace           = {};
@@ -2085,7 +2138,13 @@ initialize_signal_handler(sigaction_func_t sigaction_func)
 
     struct sigaction sig_act = {};
     sigemptyset(&sig_act.sa_mask);
-    sig_act.sa_flags     = (SA_SIGINFO | SA_RESETHAND | SA_NOCLDSTOP);
+    // No SA_RESETHAND: a one-shot handler resets the disposition to SIG_DFL the instant it fires,
+    // so a *second* delivery of the same signal (e.g. a chained app handler like LLVM/comgr
+    // re-raising) would land on SIG_DFL and terminate the process mid-flush -> data loss. Keeping
+    // our handler installed routes every re-delivery back through the re-entry guard, which
+    // swallows until the flush completes (finalize_done) and only then escalates. Termination is
+    // guaranteed by the worker restoring the real disposition and re-raising once the flush done.
+    sig_act.sa_flags     = (SA_SIGINFO | SA_NOCLDSTOP);
     sig_act.sa_sigaction = &rocprofv3_error_signal_handler;
     for(auto signal_v : rocprofv3_handled_signals)
     {
@@ -2186,6 +2245,13 @@ wait_peer_finished(const pid_t& pid, const pid_t& ppid)
 void
 finalize_rocprofv3(std::string_view context)
 {
+    auto& sw = get_signal_worker();
+    if(sw.finalized.test_and_set(std::memory_order_acq_rel))
+    {
+        ROCP_INFO << "finalize_rocprofv3('" << context << "') ignored: already finalized";
+        return;
+    }
+
     ROCP_INFO << "invoked: finalize_rocprofv3";
     if(client_finalizer && client_identifier)
     {
@@ -2194,10 +2260,9 @@ finalize_rocprofv3(std::string_view context)
         client_finalizer  = nullptr;
         client_identifier = nullptr;
     }
-    else
-    {
-        ROCP_INFO << "finalize_rocprofv3('" << context << "') ignored: already finalized";
-    }
+
+    // Join the worker thread (from atexit / rocprofv3_main; a call from the worker itself no-ops).
+    sw.join();
 }
 
 bool
@@ -3951,6 +4016,11 @@ get_sigaction_function()
 
 bool signal_handler_exit =
     rocprofiler::tool::get_env("ROCPROF_INTERNAL_TEST_SIGNAL_HANDLER_VIA_EXIT", false);
+
+// Read once here because getenv() is not async-signal-safe; <= 0 waits indefinitely.
+int signal_abort_flush_timeout_sec =
+    rocprofiler::tool::get_env("ROCPROF_ABORT_FLUSH_TIMEOUT_SECONDS", 10);
+
 }  // namespace
 
 #define ROCPROFV3_INTERNAL_API __attribute__((visibility("internal")));
@@ -4106,43 +4176,95 @@ diagnose_status(pid_t _pid, int _status)
 }
 
 void
-rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext)
+wait_for_children(pid_t this_pid, pid_t this_ppid, uint64_t this_tid, std::string_view context)
 {
-    // Only the first fatal signal in the process runs the body below. A later one, whether it
-    // is a re-entrant abort from finalization on this thread or a signal on another thread,
-    // must not wait on the std::call_once: the first is a recursive lock, the second blocks
-    // until the finalization it is waiting on completes, and either way the process can no
-    // longer be terminated. This runs before any logging so that a stuck logger or allocator
-    // cannot get in the way.
-    static auto _handling = std::atomic_flag{};
-    if(_handling.test_and_set())
-    {
-        constexpr auto _msg =
-            std::string_view{"[rocprofv3] fatal signal while already handling one... "
-                             "terminating\n"};
-        // write() rather than the logger: async-signal-safe and takes no locks
-        auto _written = ::write(STDERR_FILENO, _msg.data(), _msg.size());
-        (void) _written;
-
-        // rocprofv3 interposes signal()/sigaction() to keep this handler installed, so the
-        // default disposition has to be restored through the real symbol before re-raising.
-        auto _default_action       = sigaction_t{};
-        _default_action.sa_handler = SIG_DFL;
-        sigemptyset(&_default_action.sa_mask);
-        if(auto* _real_sigaction = get_sigaction_function();
-           _real_sigaction != nullptr && _real_sigaction(signo, &_default_action, nullptr) == 0)
+    auto get_children = [&this_pid]() {
+        auto fname    = fmt::format("/proc/{}/task/{}/children", this_pid, this_pid);
+        auto ifs      = std::ifstream{fname};
+        auto children = std::vector<pid_t>{};
+        if(!ifs.is_open())
         {
-            // the handler was entered with signo blocked, so unblock it or the raise below
-            // only marks it pending and the process exits rather than dying from the signal
-            auto _blocked = sigset_t{};
-            sigemptyset(&_blocked);
-            sigaddset(&_blocked, signo);
-            ::pthread_sigmask(SIG_UNBLOCK, &_blocked, nullptr);
-            ::raise(signo);
+            ROCP_WARNING << "signal worker: failed to open " << fname << ": " << strerror(errno);
+            return children;
         }
+        while(ifs)
+        {
+            pid_t child_pid = 0;
+            ifs >> child_pid;
+            if(ifs && !ifs.eof() && child_pid > 0) children.emplace_back(child_pid);
+        }
+        return children;
+    };
 
-        // only reached if the signal is not fatal by default or could not be restored
-        ::_exit(128 + signo);
+    auto _children = get_children();
+    ROCP_WARNING << fmt::format(
+        "[PPID={}][PID={}][TID={}][{}] rocprofv3 waiting for {} children to exit",
+        this_ppid,
+        this_pid,
+        this_tid,
+        context,
+        _children.size());
+
+    for(auto itr : _children)
+    {
+        auto status = wait_pid(itr, WUNTRACED | WNOHANG);
+        if(status) diagnose_status(itr, status.value());
+    }
+}
+
+// Reinstall the disposition we wrapped for `signo`: the saved chained handler if any, else
+// SIG_DFL. Uses the real sigaction (bypasses our interceptor).
+void
+restore_signal_disposition(int signo)
+{
+    if(auto& chained = get_chained_signals().at(signo); chained)
+    {
+        if(chained->action)
+        {
+            get_sigaction_function()(signo, &(*chained->action), nullptr);
+        }
+        else
+        {
+            struct sigaction sa = {};
+            sa.sa_handler       = chained->handler ? chained->handler : SIG_DFL;
+            sigemptyset(&sa.sa_mask);
+            get_sigaction_function()(signo, &sa, nullptr);
+        }
+    }
+    else
+    {
+        struct sigaction sa = {};
+        sa.sa_handler       = SIG_DFL;
+        sigemptyset(&sa.sa_mask);
+        get_sigaction_function()(signo, &sa, nullptr);
+    }
+}
+
+void
+signal_finalization_worker()
+{
+    auto& sw = get_signal_worker();
+
+    sigset_t mask;
+    sigemptyset(&mask);
+    for(auto sig : rocprofv3_handled_signals)
+        sigaddset(&mask, sig);
+    pthread_sigmask(SIG_BLOCK, &mask, nullptr);
+
+    // Retry across EINTR: the worker blocks only the handled signals, so any other signal delivered
+    // here (SIGCHLD, SIGALRM, ...) would otherwise abort the read and kill the worker, after which
+    // nothing flushes and every repeat signal is swallowed forever.
+    uint64_t val   = 0;
+    ssize_t  nread = 0;
+    do
+    {
+        nread = read(sw.eventfd, &val, sizeof(val));
+    } while(nread < 0 && errno == EINTR);
+
+    if(nread < 0)
+    {
+        ROCP_WARNING << "signal worker: eventfd read failed: " << strerror(errno);
+        return;
     }
 
     auto this_pid  = getpid();
@@ -4150,132 +4272,159 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext)
     auto this_tid  = common::get_tid();
     auto this_func = std::string_view{__FUNCTION__};
 
-    ROCP_WARNING << fmt::format("[PPID={}][PID={}][TID={}][{}] rocprofv3 caught signal {}...",
-                                this_ppid,
-                                this_pid,
-                                this_tid,
-                                this_func,
-                                signo);
+    ROCP_WARNING << fmt::format(
+        "[PPID={}][PID={}][TID={}][{}] rocprofv3 finalizing after signal {}...",
+        this_ppid,
+        this_pid,
+        this_tid,
+        this_func,
+        sw.signo);
 
-    static auto _once = std::once_flag{};
-    std::call_once(_once, [&]() {
-        auto get_children = [&this_pid]() {
-            auto fname    = fmt::format("/proc/{}/task/{}/children", this_pid, this_pid);
-            auto ifs      = std::ifstream{fname};
-            auto children = std::vector<pid_t>{};
-            while(ifs)
-            {
-                pid_t val = 0;
-                ifs >> val;
-                if(ifs && !ifs.eof() && val > 0) children.emplace_back(val);
-            }
-            return children;
-        };
+    finalize_rocprofv3(this_func);
 
-        auto _children = get_children();
-        ROCP_WARNING << fmt::format(
-            "[PPID={}][PID={}][TID={}][{}] rocprofv3 will wait for {} children to exit",
-            this_ppid,
-            this_pid,
-            this_tid,
-            this_func,
-            _children.size());
+    if(tool::get_config().enable_process_sync) wait_peer_finished(this_pid, this_ppid);
 
-        // wait for children
-        for(auto itr : _children)
+    // Signal completion (only the SIGABRT handler path waits on this).
+    sw.finalize_done.store(1u, std::memory_order_release);
+    syscall(SYS_futex, &sw.finalize_done, FUTEX_WAKE, 1, nullptr, nullptr, 0);
+
+    // Re-raise BEFORE reaping children. The app's own handler needs the signal to run its
+    // coordinated shutdown, and that shutdown is what makes the children exit. Reaping first
+    // deadlocks multi-process apps (e.g. tensor-parallel servers) whose worker processes only
+    // exit once the main tells them to. signo == 0 is normal-exit finalization; SIGABRT
+    // terminates itself once its (synchronously waiting) handler returns into abort().
+    if(sw.signo != 0 && sw.signo != SIGABRT)
+    {
+        const bool have_chained = static_cast<bool>(get_chained_signals().at(sw.signo));
+        restore_signal_disposition(sw.signo);
+
+        // Re-raise process-directed (kill, not raise) so it lands on an app thread, not this
+        // worker (which blocks these signals): the chained handler runs, or SIG_DFL exits.
+        if(!have_chained)
         {
-            auto status = wait_pid(itr, WUNTRACED | WNOHANG);
-            if(status) diagnose_status(itr, status.value());
+            // No chained handler: also unblock here so SIG_DFL is guaranteed to terminate.
+            sigset_t unblock{};
+            sigemptyset(&unblock);
+            sigaddset(&unblock, sw.signo);
+            pthread_sigmask(SIG_UNBLOCK, &unblock, nullptr);
         }
+        kill(getpid(), sw.signo);
+    }
 
-        ROCP_WARNING << fmt::format(
-            "[PPID={}][PID={}][TID={}][{}] rocprofv3 finalizing after signal {}...",
-            this_ppid,
-            this_pid,
-            this_tid,
-            this_func,
-            signo);
-
-        finalize_rocprofv3(this_func);
-        if(tool::get_config().enable_process_sync) wait_peer_finished(this_pid, this_ppid);
+    // Best-effort reap to avoid leaving zombies if the app keeps running (e.g. a chained handler
+    // that returns). We do NOT drive the signal into children -- delivering it to a separate PID
+    // is the app's/OS's job; a child that received the signal finalizes via its own worker.
+    wait_for_children(this_pid, this_ppid, this_tid, this_func);
 
-        ROCP_INFO << fmt::format(
-            "[PPID={}][PID={}][TID={}][{}] rocprofv3 finalizing after signal {}... complete",
-            this_ppid,
-            this_pid,
-            this_tid,
-            this_func,
-            signo);
+    ROCP_INFO << fmt::format(
+        "[PPID={}][PID={}][TID={}][{}] rocprofv3 finalizing after signal... complete",
+        this_ppid,
+        this_pid,
+        this_tid,
+        this_func);
+}
 
-        if(get_chained_signals().at(signo))
+void
+rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext)
+{
+    (void) info;
+    (void) ucontext;
+
+    auto& sw = get_signal_worker();
+
+    // Our handler being installed means the app doesn't coordinate shutdown.
+    // Well-behaved apps use --disable-signal-handlers (atexit handles everything).
+    //
+    // Re-entry: a second signal arrived while we're still handling the first
+    // (a chained handler re-raising, or a double Ctrl+C).
+    // While the flush is running, swallow it. Delivering it now would cut the
+    // flush short and truncate the profile.
+    // After the flush completes (finalize_done), force SIG_DFL and re-raise.
+    // The worker re-raises on its own once done, so the process still exits.
+    uint32_t expected = 0;
+    if(sw.handling.compare_exchange_strong(expected, 1u, std::memory_order_acquire) == false)
+    {
+        if(sw.finalize_done.load(std::memory_order_acquire) != 0)
         {
-            ROCP_INFO << fmt::format(
-                "[PPID={}][PID={}][TID={}][{}] rocprofv3 found chained signal handler for {}",
-                this_ppid,
-                this_pid,
-                this_tid,
-                this_func,
-                signo);
+            struct sigaction _dfl = {};
+            _dfl.sa_handler       = SIG_DFL;
+            sigemptyset(&_dfl.sa_mask);
+            get_sigaction_function()(signo, &_dfl, nullptr);
+
+            sigset_t _unblock{};
+            sigemptyset(&_unblock);
+            sigaddset(&_unblock, signo);
+            pthread_sigmask(SIG_UNBLOCK, &_unblock, nullptr);
+            raise(signo);
+        }
+        return;
+    }
 
-            if(auto& _chained = *get_chained_signals().at(signo); _chained.action)
+    // For testing: allow quick_exit path
+    if(signal_handler_exit) ::quick_exit(signo);
+
+    // Hand the signal number to the worker (ordered by the eventfd write/read below).
+    sw.signo = signo;
+
+    // Wake the finalization worker: the flush is NOT async-signal-safe, so it runs there.
+    bool worker_notified = false;
+    if(sw.eventfd >= 0)
+    {
+        uint64_t val    = 1;
+        worker_notified = (write(sw.eventfd, &val, sizeof(val)) == sizeof(val));
+    }
+
+    // SIGABRT can't defer: returning lets abort() re-raise SIG_DFL before the worker flushes.
+    // So wait for the flush here, then return into abort(). Unlike the async signals, this can
+    // block on a lock the aborting thread holds as abort() often fires from lock-holding runtime
+    // paths, e.g. heap-corruption detection. Bound the wait with
+    // ROCPROF_ABORT_FLUSH_TIMEOUT_SECONDS and fall through into abort() on expiry: a core dump
+    // beats a hung process. <= 0 waits forever.
+    if(signo == SIGABRT)
+    {
+        if(worker_notified)
+        {
+            // FUTEX_WAIT_BITSET takes an absolute CLOCK_MONOTONIC deadline, so compute it once and
+            // let the kernel track the time remaining across wakeups. <= 0 waits indefinitely.
+            const auto bounded  = (signal_abort_flush_timeout_sec > 0);
+            auto       deadline = timespec{};
+            if(bounded)
             {
-                ROCP_TRACE << fmt::format("[PPID={}][PID={}][TID={}][{}] rocprofv3 found chained "
-                                          "signal handler for {}... executing chained sigaction",
-                                          this_ppid,
-                                          this_pid,
-                                          this_tid,
-                                          this_func,
-                                          signo);
-                if((_chained.action->sa_flags & SA_SIGINFO) == SA_SIGINFO &&
-                   _chained.action->sa_sigaction &&
-                   _chained.action->sa_sigaction != &rocprofv3_error_signal_handler)
-                {
-                    ROCP_WARNING << fmt::format(
-                        "[PPID={}][PID={}][TID={}][{}] rocprofv3 found chained signal handler for "
-                        "{}... executing chained sigaction (SIGINFO)",
-                        this_ppid,
-                        this_pid,
-                        this_tid,
-                        this_func,
-                        signo);
-                    _chained.action->sa_sigaction(signo, info, ucontext);
-                }
-                else if((_chained.action->sa_flags & SA_SIGINFO) != SA_SIGINFO &&
-                        _chained.action->sa_handler &&
-                        _chained.action->sa_sigaction != &rocprofv3_error_signal_handler)
-                {
-                    ROCP_WARNING << fmt::format(
-                        "[PPID={}][PID={}][TID={}][{}] rocprofv3 found chained signal handler for "
-                        "{}... executing chained sigaction (HANDLER)",
-                        this_ppid,
-                        this_pid,
-                        this_tid,
-                        this_func,
-                        signo);
-                    _chained.action->sa_handler(signo);
-                }
+                clock_gettime(CLOCK_MONOTONIC, &deadline);
+                deadline.tv_sec += signal_abort_flush_timeout_sec;
             }
-            else
+
+            while(sw.finalize_done.load(std::memory_order_acquire) == 0)
             {
-                if(_chained.handler)
+                if(syscall(SYS_futex,
+                           &sw.finalize_done,
+                           FUTEX_WAIT_BITSET,
+                           0,
+                           bounded ? &deadline : nullptr,
+                           nullptr,
+                           FUTEX_BITSET_MATCH_ANY) == -1 &&
+                   errno == ETIMEDOUT)
                 {
-                    ROCP_WARNING << fmt::format(
-                        "[PPID={}][PID={}][TID={}][{}] rocprofv3 found chained signal handler for "
-                        "{}... executing chained handler",
-                        this_ppid,
-                        this_pid,
-                        this_tid,
-                        this_func,
-                        signo);
-                    _chained.handler(signo);
+                    break;
                 }
             }
         }
-    });
+        return;
+    }
 
-    // below is for testing purposes. re-raising the signal causes CTest to ignore WILL_FAIL ON
-    if(signal_handler_exit) ::quick_exit(signo);
-    ::raise(signo);
+    // Async signals (SIGINT/SIGQUIT/SIGTERM): don't block or re-raise here. Returning lets this
+    // thread resume and release any lock it holds, so the worker's flush can't deadlock against
+    // it; the worker terminates via kill() once the flush is done.
+    if(!worker_notified)
+    {
+        // No worker to hand off to — best-effort terminate on this thread.
+        restore_signal_disposition(signo);
+        sigset_t unblock{};
+        sigemptyset(&unblock);
+        sigaddset(&unblock, signo);
+        pthread_sigmask(SIG_UNBLOCK, &unblock, nullptr);
+        raise(signo);
+    }
 }
 
 int
@@ -4393,14 +4542,18 @@ rocprofv3_set_main(main_func_t main_func)
 sighandler_t
 rocprofv3_signal(int signum, sighandler_t handler)
 {
-    static auto _once = std::once_flag{};
-    std::call_once(_once,
-                   []() { get_signal_function() = (signal_func_t) dlsym(RTLD_NEXT, "signal"); });
+    if(!get_signal_function()) get_signal_function() = (signal_func_t) dlsym(RTLD_NEXT, "signal");
+
+    if(get_signal_worker().handling.load(std::memory_order_relaxed) != 0)
+        return get_signal_function()(signum, handler);
 
     if(!is_handled_signal(signum) || !tool::get_config().enable_signal_handlers)
         return CHECK_NOTNULL(get_signal_function())(signum, handler);
 
-    get_chained_signals().at(signum) = chained_siginfo{signum, handler, std::nullopt};
+    // Save the app's disposition (first install only; keep SIG_IGN, skip SIG_DFL). See the detailed
+    // rationale in rocprofv3_sigaction.
+    if(!get_chained_signals().at(signum) && handler != SIG_DFL)
+        get_chained_signals().at(signum) = chained_siginfo{signum, handler, std::nullopt};
 
     return get_signal_function()(
         signum, [](int signum_v) { rocprofv3_error_signal_handler(signum_v, nullptr, nullptr); });
@@ -4411,21 +4564,41 @@ rocprofv3_sigaction(int signum,
                     const struct sigaction* __restrict__ act,
                     struct sigaction* __restrict__ oldact)
 {
-    static auto _once = std::once_flag{};
-    std::call_once(_once, []() {
+    if(!get_sigaction_function())
         get_sigaction_function() = (sigaction_func_t) dlsym(RTLD_NEXT, "sigaction");
-    });
+
+    if(get_signal_worker().handling.load(std::memory_order_relaxed) != 0)
+        return get_sigaction_function()(signum, act, oldact);
 
     if(!is_handled_signal(signum) || !act || !tool::get_config().enable_signal_handlers)
         return CHECK_NOTNULL(get_sigaction_function())(signum, act, oldact);
 
-    // make sure rocprofv3_error_signal_handler doesn't call itself
-    if((act->sa_flags & SA_SIGINFO) == SA_SIGINFO &&
-       act->sa_sigaction != &rocprofv3_error_signal_handler)
-        get_chained_signals().at(signum) = chained_siginfo{signum, nullptr, *act};
+    // Save the app's disposition so we can restore it on signal delivery.
+    // Only save the first one registered per signal. Later installs (e.g., LLVM comgr) often
+    // don't chain properly — they clobber the disposition on re-raise. Preserving the app's
+    // disposition ensures the application sees its signal as-if the profiler wasn't there.
+    if(!get_chained_signals().at(signum))
+    {
+        if((act->sa_flags & SA_SIGINFO) == SA_SIGINFO)
+        {
+            if(act->sa_sigaction != &rocprofv3_error_signal_handler)
+                get_chained_signals().at(signum) = chained_siginfo{signum, nullptr, *act};
+        }
+        else
+        {
+            // Save SIG_IGN too (matching signal()): an app that ignores the signal — e.g. a
+            // deliberately un-interruptible process — must still see it ignored after we flush.
+            // SIG_DFL is skipped because an empty entry already restores as SIG_DFL.
+            if(act->sa_handler != SIG_DFL)
+                get_chained_signals().at(signum) = chained_siginfo{signum, nullptr, *act};
+        }
+    }
 
     struct sigaction _upd_act = *act;
-    _upd_act.sa_flags |= (SA_SIGINFO | SA_RESETHAND | SA_NOCLDSTOP);
+    // See initialize_signal_handler: no SA_RESETHAND so re-deliveries route back through our
+    // re-entry guard instead of hitting SIG_DFL mid-flush.
+    _upd_act.sa_flags &= ~SA_RESETHAND;
+    _upd_act.sa_flags |= (SA_SIGINFO | SA_NOCLDSTOP);
     _upd_act.sa_sigaction = &rocprofv3_error_signal_handler;
 
     return get_sigaction_function()(signum, &_upd_act, oldact);
@@ -4456,6 +4629,42 @@ rocprofv3_main(int argc, char** argv, char** envp)
 
     initialize_rocprofv3();
 
+    // Resolve dlsym'd function pointers at init time (before interceptors run).
+    if(!get_signal_function()) get_signal_function() = (signal_func_t) dlsym(RTLD_NEXT, "signal");
+    if(!get_sigaction_function())
+        get_sigaction_function() = (sigaction_func_t) dlsym(RTLD_NEXT, "sigaction");
+
+    // fork+exec and spawn children re-init through the normal init path. Plain fork children
+    // re-init worker state via the pthread_atfork handler below. GPU use in a forked child is
+    // illegal (runtime not fork-safe) so it usually dies first, but host-side data like roctx
+    // markers still flushes if reached (covered by a test).
+    {
+        auto& sw   = get_signal_worker();
+        sw.eventfd = eventfd(0, EFD_CLOEXEC);
+        if(sw.eventfd >= 0) sw.thread = std::thread{signal_finalization_worker};
+
+        // Register atfork handler so fork() children get a fresh worker thread.
+        // Without this, fork children inherit the parent's stale eventfd/thread.
+        static bool atfork_registered = false;
+        if(!atfork_registered)
+        {
+            atfork_registered = true;
+            pthread_atfork(nullptr, nullptr, []() {
+                // Child handler: reset stale state and spawn a fresh worker.
+                auto& child_sw = get_signal_worker();
+                if(child_sw.eventfd >= 0 && ::close(child_sw.eventfd) != 0)
+                {
+                    ROCP_WARNING << "signal worker: close(eventfd) after fork failed: "
+                                 << strerror(errno);
+                }
+
+                ::new(static_cast<void*>(&child_sw)) signal_worker_state{};
+                child_sw.eventfd = eventfd(0, EFD_CLOEXEC);
+                if(child_sw.eventfd >= 0) child_sw.thread = std::thread{signal_finalization_worker};
+            });
+        }
+    }
+
     initialize_signal_handler(get_sigaction_function());
 
     ROCP_INFO << "rocprofv3: main function wrapper will be invoked...";
diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/correlation_id.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/correlation_id.cpp
index 0cb9633fc8a..aad337957f1 100644
--- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/correlation_id.cpp
+++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/correlation_id.cpp
@@ -76,7 +76,9 @@ correlation_id::sub_ref_count()
 {
     if(m_ref_count == 0)
     {
-        ROCP_CI_LOG(WARNING) << fmt::format(
+        // Below-zero is a real bug pre-finalization, but once finalize starts the retirement sweep
+        // races in-flight decrements (benign teardown), so only treat it as CI-fatal before then.
+        ROCP_CI_LOG_IF(WARNING, registration::get_fini_status() == 0) << fmt::format(
             "attempt to decrement correlation id {} reference count but reference count is zero",
             internal);
         return 0;
@@ -86,7 +88,10 @@ correlation_id::sub_ref_count()
 
     if(registration::get_fini_status() > 0) return 0;
 
-    ROCP_CI_LOG_IF(WARNING, _ret == 0) << fmt::format("correlation id underflow on {}", internal);
+    // Same finalization race as the zero-check above: a concurrent decrement can slip between that
+    // check and this fetch_sub during the retirement sweep, so only flag underflow before finalize.
+    ROCP_CI_LOG_IF(WARNING, _ret == 0 && registration::get_fini_status() == 0)
+        << fmt::format("correlation id underflow on {}", internal);
 
     if(_ret == 1)
     {
@@ -253,7 +258,27 @@ correlation_id_finalize()
                 {}
             }
         }
-        ROCP_CI_LOG_IF(INFO, ndangling > 0) << "retired dangling correlation IDs: " << ndangling;
+
+        if(ndangling > 0)
+        {
+            auto retirement_ctxs = get_active_contexts([](const context* ctx) {
+                return (ctx->buffered_tracer &&
+                        (ctx->buffered_tracer->domains(
+                            ROCPROFILER_BUFFER_TRACING_CORRELATION_ID_RETIREMENT)));
+            });
+
+            if(!retirement_ctxs.empty())
+            {
+                ROCP_CI_LOG_IF(INFO, ndangling > 0)
+                    << fmt::format("retired dangling correlation IDs: {}", ndangling);
+            }
+            else
+            {
+                ROCP_INFO << fmt::format(
+                    "retired dangling correlation IDs: {} (no retirement consumer active)",
+                    ndangling);
+            }
+        }
     });
 }
 }  // namespace context
diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.cpp
index ae668b0f33e..ff79b185a83 100644
--- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.cpp
+++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.cpp
@@ -231,12 +231,14 @@ hip_api_impl<TableIdx, OpIdx>::functor(Args... args)
             return;
     }
 
-    auto  buffer_record    = common::init_public_api_struct(buffered_api_data_t{});
-    auto  extended_record  = common::init_public_api_struct(buffered_ext_data_t{});
-    auto  tracer_data      = common::init_public_api_struct(callback_api_data_t{});
-    auto* corr_id          = tracing::correlation_service::construct(ref_count);
-    auto  internal_corr_id = corr_id->internal;
-    auto  ancestor_corr_id = corr_id->ancestor;
+    auto  buffer_record   = common::init_public_api_struct(buffered_api_data_t{});
+    auto  extended_record = common::init_public_api_struct(buffered_ext_data_t{});
+    auto  tracer_data     = common::init_public_api_struct(callback_api_data_t{});
+    auto* corr_id         = tracing::correlation_service::construct(ref_count);
+    RETURN_UNTRACED_ON_NULL_CORRELATION_ID(
+        corr_id, RetT, exec(info_type::get_table_func(), std::forward<Args>(args)...));
+    auto internal_corr_id = corr_id->internal;
+    auto ancestor_corr_id = corr_id->ancestor;
 
     tracing::populate_external_correlation_ids(external_corr_ids,
                                                thr_id,
diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hipfile/hipfile.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hipfile/hipfile.cpp
index 56314ee9573..46dc76d5c1a 100644
--- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hipfile/hipfile.cpp
+++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hipfile/hipfile.cpp
@@ -242,12 +242,14 @@ hipfile_api_impl<TableIdx, OpIdx>::functor(Args... args)
             return;
     }
 
-    auto  buffer_record    = common::init_public_api_struct(buffered_api_data_t{});
-    auto  extended_record  = common::init_public_api_struct(buffered_ext_data_t{});
-    auto  tracer_data      = common::init_public_api_struct(callback_api_data_t{});
-    auto* corr_id          = tracing::correlation_service::construct(ref_count);
-    auto  internal_corr_id = corr_id->internal;
-    auto  ancestor_corr_id = corr_id->ancestor;
+    auto  buffer_record   = common::init_public_api_struct(buffered_api_data_t{});
+    auto  extended_record = common::init_public_api_struct(buffered_ext_data_t{});
+    auto  tracer_data     = common::init_public_api_struct(callback_api_data_t{});
+    auto* corr_id         = tracing::correlation_service::construct(ref_count);
+    RETURN_UNTRACED_ON_NULL_CORRELATION_ID(
+        corr_id, RetT, exec(info_type::get_table_func(), std::forward<Args>(args)...));
+    auto internal_corr_id = corr_id->internal;
+    auto ancestor_corr_id = corr_id->ancestor;
 
     tracing::populate_external_correlation_ids(external_corr_ids,
                                                thr_id,
diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/marker.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/marker.cpp
index dd6b7c424df..19f8251973d 100644
--- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/marker.cpp
+++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/marker.cpp
@@ -167,12 +167,14 @@ roctx_api_impl<TableIdx, OpIdx>::functor(Args... args)
             return;
     }
 
-    auto  ref_count        = 2;
-    auto  buffer_record    = common::init_public_api_struct(buffered_api_data_t{});
-    auto  tracer_data      = common::init_public_api_struct(callback_api_data_t{});
-    auto* corr_id          = tracing::correlation_service::construct(ref_count);
-    auto  internal_corr_id = corr_id->internal;
-    auto  ancestor_corr_id = corr_id->ancestor;
+    auto  ref_count     = 2;
+    auto  buffer_record = common::init_public_api_struct(buffered_api_data_t{});
+    auto  tracer_data   = common::init_public_api_struct(callback_api_data_t{});
+    auto* corr_id       = tracing::correlation_service::construct(ref_count);
+    RETURN_UNTRACED_ON_NULL_CORRELATION_ID(
+        corr_id, RetT, exec(info_type::get_table_func(), std::forward<Args>(args)...));
+    auto internal_corr_id = corr_id->internal;
+    auto ancestor_corr_id = corr_id->ancestor;
 
     tracing::populate_external_correlation_ids(external_corr_ids,
                                                thr_id,
diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/range_marker.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/range_marker.cpp
index ab3ac78832e..8bf7bf5e050 100644
--- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/range_marker.cpp
+++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/range_marker.cpp
@@ -193,12 +193,14 @@ roctx_api_impl<TableIdx, OpIdx>::functor(Args... args)
             return;
     }
 
-    auto  ref_count        = 2;
-    auto  buffer_record    = common::init_public_api_struct(buffered_api_data_t{});
-    auto  callback_data    = common::init_public_api_struct(callback_api_data_t{});
-    auto* corr_id          = tracing::correlation_service::construct(ref_count);
-    auto  internal_corr_id = CHECK_NOTNULL(corr_id)->internal;
-    auto  ancestor_corr_id = CHECK_NOTNULL(corr_id)->ancestor;
+    auto  ref_count     = 2;
+    auto  buffer_record = common::init_public_api_struct(buffered_api_data_t{});
+    auto  callback_data = common::init_public_api_struct(callback_api_data_t{});
+    auto* corr_id       = tracing::correlation_service::construct(ref_count);
+    RETURN_UNTRACED_ON_NULL_CORRELATION_ID(
+        corr_id, RetT, exec(info_type::get_table_func(), std::forward<Args>(args)...));
+    auto internal_corr_id = corr_id->internal;
+    auto ancestor_corr_id = corr_id->ancestor;
 
     tracing::populate_external_correlation_ids(external_corr_ids,
                                                thr_id,
@@ -310,7 +312,9 @@ roctx_api_impl<TableIdx, OpIdx>::push_functor(Args... args)
     auto&  callback_data = range_data.callback_data;
     auto*& corr_id       = range_data.corr_id;
 
-    corr_id               = tracing::correlation_service::construct(ref_count);
+    corr_id = tracing::correlation_service::construct(ref_count);
+    RETURN_UNTRACED_ON_NULL_CORRELATION_ID(
+        corr_id, RetT, exec(info_type::get_push_table_func(), std::forward<Args>(args)...));
     auto internal_corr_id = corr_id->internal;
     auto ancestor_corr_id = corr_id->ancestor;
 
diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/ompt/ompt.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/ompt/ompt.cpp
index f83d3260028..3b422cc83f5 100644
--- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/ompt/ompt.cpp
+++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/ompt/ompt.cpp
@@ -766,9 +766,15 @@ ompt_impl<OpIdx>::begin(ompt_data_t* data, Args... args)
                                buffered_contexts,
                                external_corr_ids);
 
-    auto* corr_id          = tracing::correlation_service::construct(ref_count);
-    auto  internal_corr_id = corr_id->internal;
-    auto  ancestor_corr_id = corr_id->ancestor;
+    auto* corr_id = tracing::correlation_service::construct(ref_count);
+    if(!corr_id)
+    {
+        // finalization began mid-call: construct() returns null. Skip this OMPT region; no state
+        // is stashed, so the paired end() skips too (see the null-state guard there).
+        return;
+    }
+    auto internal_corr_id = corr_id->internal;
+    auto ancestor_corr_id = corr_id->ancestor;
 
     tracing::populate_external_correlation_ids(external_corr_ids,
                                                thr_id,
@@ -830,9 +836,12 @@ ompt_impl<OpIdx>::end(ompt_data_t* data, Args... args)
     ompt_save_state* state = nullptr;
     if(data != nullptr)
         state = static_cast<ompt_save_state*>(data->ptr);
-    else
-        state = get_ompt_state_stack().pop_back_val();
-    assert(state != nullptr);
+    else if(auto& _state_stack = get_ompt_state_stack(); !_state_stack.empty())
+        state = _state_stack.pop_back_val();
+
+    // begin() does not stash state when it cannot construct a correlation id (finalization in
+    // progress), so there is nothing to pair this end() with -- skip instead of dereferencing null.
+    if(!state) return;
 
     ROCP_FATAL_IF(state->operation_idx != info_type::operation_idx)
         << "Mismatch of OMPT operation: begin=" << state->operation_idx
diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rccl/rccl.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rccl/rccl.cpp
index dce911a4aff..56e385129eb 100644
--- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rccl/rccl.cpp
+++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rccl/rccl.cpp
@@ -186,11 +186,13 @@ rccl_api_impl<TableIdx, OpIdx>::functor(Args... args)
             return;
     }
 
-    auto  buffer_record    = common::init_public_api_struct(buffered_api_data_t{});
-    auto  tracer_data      = common::init_public_api_struct(callback_api_data_t{});
-    auto* corr_id          = tracing::correlation_service::construct(ref_count);
-    auto  internal_corr_id = corr_id->internal;
-    auto  ancestor_corr_id = corr_id->internal;
+    auto  buffer_record = common::init_public_api_struct(buffered_api_data_t{});
+    auto  tracer_data   = common::init_public_api_struct(callback_api_data_t{});
+    auto* corr_id       = tracing::correlation_service::construct(ref_count);
+    RETURN_UNTRACED_ON_NULL_CORRELATION_ID(
+        corr_id, RetT, exec(info_type::get_table_func(), std::forward<Args>(args)...));
+    auto internal_corr_id = corr_id->internal;
+    auto ancestor_corr_id = corr_id->internal;
 
     tracing::populate_external_correlation_ids(external_corr_ids,
                                                thr_id,
diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocdecode/rocdecode.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocdecode/rocdecode.cpp
index 597ad9d6a92..fe9e0f79a2d 100644
--- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocdecode/rocdecode.cpp
+++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocdecode/rocdecode.cpp
@@ -203,12 +203,14 @@ rocdecode_api_impl<TableIdx, OpIdx>::functor(Args... args)
             return;
     }
 
-    auto  buffer_record    = common::init_public_api_struct(buffered_api_data_t{});
-    auto  extended_record  = common::init_public_api_struct(buffered_ext_data_t{});
-    auto  tracer_data      = common::init_public_api_struct(callback_api_data_t{});
-    auto* corr_id          = tracing::correlation_service::construct(ref_count);
-    auto  internal_corr_id = corr_id->internal;
-    auto  ancestor_corr_id = corr_id->ancestor;
+    auto  buffer_record   = common::init_public_api_struct(buffered_api_data_t{});
+    auto  extended_record = common::init_public_api_struct(buffered_ext_data_t{});
+    auto  tracer_data     = common::init_public_api_struct(callback_api_data_t{});
+    auto* corr_id         = tracing::correlation_service::construct(ref_count);
+    RETURN_UNTRACED_ON_NULL_CORRELATION_ID(
+        corr_id, RetT, exec(info_type::get_table_func(), std::forward<Args>(args)...));
+    auto internal_corr_id = corr_id->internal;
+    auto ancestor_corr_id = corr_id->ancestor;
 
     tracing::populate_external_correlation_ids(external_corr_ids,
                                                thr_id,
diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocjpeg/rocjpeg.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocjpeg/rocjpeg.cpp
index fcc1ee58645..aa34f91b458 100644
--- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocjpeg/rocjpeg.cpp
+++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocjpeg/rocjpeg.cpp
@@ -185,11 +185,13 @@ rocjpeg_api_impl<TableIdx, OpIdx>::functor(Args... args)
             return;
     }
 
-    auto  buffer_record    = common::init_public_api_struct(buffered_api_data_t{});
-    auto  tracer_data      = common::init_public_api_struct(callback_api_data_t{});
-    auto* corr_id          = tracing::correlation_service::construct(ref_count);
-    auto  internal_corr_id = corr_id->internal;
-    auto  ancestor_corr_id = corr_id->ancestor;
+    auto  buffer_record = common::init_public_api_struct(buffered_api_data_t{});
+    auto  tracer_data   = common::init_public_api_struct(callback_api_data_t{});
+    auto* corr_id       = tracing::correlation_service::construct(ref_count);
+    RETURN_UNTRACED_ON_NULL_CORRELATION_ID(
+        corr_id, RetT, exec(info_type::get_table_func(), std::forward<Args>(args)...));
+    auto internal_corr_id = corr_id->internal;
+    auto ancestor_corr_id = corr_id->ancestor;
 
     tracing::populate_external_correlation_ids(external_corr_ids,
                                                thr_id,
diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocshmem/rocshmem.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocshmem/rocshmem.cpp
index 2a85a53b2b2..deea24d672d 100644
--- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocshmem/rocshmem.cpp
+++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocshmem/rocshmem.cpp
@@ -177,12 +177,14 @@ rocshmem_api_impl<TableIdx, OpIdx>::functor(Args... args)
             return;
     }
 
-    auto  buffer_record    = common::init_public_api_struct(buffered_api_data_t{});
-    auto  extended_record  = common::init_public_api_struct(buffered_ext_data_t{});
-    auto  tracer_data      = common::init_public_api_struct(callback_api_data_t{});
-    auto* corr_id          = tracing::correlation_service::construct(ref_count);
-    auto  internal_corr_id = corr_id->internal;
-    auto  ancestor_corr_id = corr_id->ancestor;
+    auto  buffer_record   = common::init_public_api_struct(buffered_api_data_t{});
+    auto  extended_record = common::init_public_api_struct(buffered_ext_data_t{});
+    auto  tracer_data     = common::init_public_api_struct(callback_api_data_t{});
+    auto* corr_id         = tracing::correlation_service::construct(ref_count);
+    RETURN_UNTRACED_ON_NULL_CORRELATION_ID(
+        corr_id, RetT, exec(info_type::get_table_func(), std::forward<Args>(args)...));
+    auto internal_corr_id = corr_id->internal;
+    auto ancestor_corr_id = corr_id->ancestor;
 
     tracing::populate_external_correlation_ids(external_corr_ids,
                                                thr_id,
diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tracing/tracing.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tracing/tracing.hpp
index ffa451f266f..4b365cc5ca2 100644
--- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tracing/tracing.hpp
+++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tracing/tracing.hpp
@@ -33,8 +33,21 @@
 
 #include <functional>
 #include <unordered_map>
+#include <utility>
 #include <vector>
 
+// construct() returns nullptr once finalization has begun. Rather than dereference it, run the
+// underlying API untraced and return
+#define RETURN_UNTRACED_ON_NULL_CORRELATION_ID(CORR_ID, RETURN_T, UNTRACED_CALL)                   \
+    if(!(CORR_ID))                                                                                 \
+    {                                                                                              \
+        [[maybe_unused]] auto _untraced_ret = (UNTRACED_CALL);                                     \
+        if constexpr(!std::is_void<RETURN_T>::value)                                               \
+            return _untraced_ret;                                                                  \
+        else                                                                                       \
+            return;                                                                                \
+    }
+
 namespace rocprofiler
 {
 namespace tracing
diff --git a/projects/rocprofiler-sdk/tests/bin/CMakeLists.txt b/projects/rocprofiler-sdk/tests/bin/CMakeLists.txt
index 4cb2e2fe802..8a77dff15ac 100644
--- a/projects/rocprofiler-sdk/tests/bin/CMakeLists.txt
+++ b/projects/rocprofiler-sdk/tests/bin/CMakeLists.txt
@@ -34,6 +34,7 @@ add_subdirectory(reproducible-runtime)
 add_subdirectory(reproducible-dispatch-count)
 add_subdirectory(transpose)
 add_subdirectory(openmp)
+add_subdirectory(signal-handler)
 
 set(CMAKE_BUILD_RPATH "\$ORIGIN:\$ORIGIN/../lib")
 
diff --git a/projects/rocprofiler-sdk/tests/bin/signal-handler/CMakeLists.txt b/projects/rocprofiler-sdk/tests/bin/signal-handler/CMakeLists.txt
new file mode 100644
index 00000000000..092011603e9
--- /dev/null
+++ b/projects/rocprofiler-sdk/tests/bin/signal-handler/CMakeLists.txt
@@ -0,0 +1,45 @@
+#
+#
+#
+cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
+
+if(NOT CMAKE_HIP_COMPILER)
+    find_program(
+        amdclangpp_EXECUTABLE
+        NAMES amdclang++
+        HINTS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
+        PATHS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm
+        PATH_SUFFIXES bin llvm/bin NO_CACHE)
+    mark_as_advanced(amdclangpp_EXECUTABLE)
+
+    if(amdclangpp_EXECUTABLE)
+        set(CMAKE_HIP_COMPILER "${amdclangpp_EXECUTABLE}")
+    endif()
+endif()
+
+project(rocprofiler-sdk-tests-bin-signal-handler LANGUAGES CXX HIP)
+
+foreach(_TYPE DEBUG MINSIZEREL RELEASE RELWITHDEBINFO)
+    if("${CMAKE_HIP_FLAGS_${_TYPE}}" STREQUAL "")
+        set(CMAKE_HIP_FLAGS_${_TYPE} "${CMAKE_CXX_FLAGS_${_TYPE}}")
+    endif()
+endforeach()
+
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_EXTENSIONS OFF)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_HIP_STANDARD 17)
+set(CMAKE_HIP_EXTENSIONS OFF)
+set(CMAKE_HIP_STANDARD_REQUIRED ON)
+
+set_source_files_properties(signal-handler.cpp PROPERTIES LANGUAGE HIP)
+add_executable(signal-handler-test)
+target_sources(signal-handler-test PRIVATE signal-handler.cpp)
+target_compile_options(signal-handler-test PRIVATE -W -Wall -Wextra -Wpedantic -Wshadow)
+
+find_package(Threads REQUIRED)
+target_link_libraries(signal-handler-test PRIVATE Threads::Threads)
+
+find_package(rocprofiler-sdk-roctx REQUIRED)
+target_link_libraries(signal-handler-test
+                      PRIVATE rocprofiler-sdk-roctx::rocprofiler-sdk-roctx)
diff --git a/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler.cpp b/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler.cpp
new file mode 100644
index 00000000000..e877ae17335
--- /dev/null
+++ b/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler.cpp
@@ -0,0 +1,482 @@
+// MIT License
+//
+// Copyright (c) 2023-2026 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in all
+// copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+
+// Signal handler integration test.
+//
+// Tests signal handling behavior with rocprofv3 in various process configurations.
+// The test runs indefinitely until killed by an external SIGINT (use `timeout`).
+//
+// Usage:
+//   timeout --signal=INT 3s signal-handler-test [--app-signal-handler|--no-app-signal-handler]
+//                                               [--single-process|--fork|--fork-exec|--spawn]
+//
+// Good case (--app-signal-handler): app installs handler, coordinates shutdown.
+//   Used with rocprofv3 --disable-signal-handlers. Profiler flushes via atexit.
+//
+// Bad case (--no-app-signal-handler): app has SIG_DFL everywhere.
+//   Profiler's signal handler is the only thing that can flush data before death.
+//
+// TODO: --raw-fork (_Fork without exec) is not supported. _Fork skips
+// pthread_atfork handlers, leaving stale profiler state. Extremely rare in practice.
+
+#include <hip/hip_runtime.h>
+#include <rocprofiler-sdk-roctx/roctx.h>
+
+#include <fcntl.h>
+#include <signal.h>
+#include <sys/wait.h>
+#include <unistd.h>
+#include <cerrno>
+#include <cstdarg>
+#include <cstring>
+#include <sstream>
+#include <string>
+
+#include <atomic>
+#include <stdexcept>
+
+namespace
+{
+std::atomic<bool> g_shutdown{false};
+
+template <typename... Args>
+std::string
+str_join(Args&&... _args)
+{
+    auto _ss = std::stringstream{};
+    ((_ss << std::forward<Args>(_args)), ...);
+    return _ss.str();
+}
+
+template <typename... Args>
+void
+emit_roctx_marker(const char* fmt, ...)
+{
+    char    buf[128];
+    va_list args;
+    va_start(args, fmt);
+    vsnprintf(buf, sizeof(buf), fmt, args);
+    va_end(args);
+    roctxMark(buf);
+}
+
+#define HIP_CHECK(call)                                                                            \
+    do                                                                                             \
+    {                                                                                              \
+        hipError_t err = (call);                                                                   \
+        if(err != hipSuccess)                                                                      \
+        {                                                                                          \
+            char msg[256];                                                                         \
+            snprintf(msg,                                                                          \
+                     sizeof(msg),                                                                  \
+                     "HIP error %d at %s:%d: %s",                                                  \
+                     err,                                                                          \
+                     __FILE__,                                                                     \
+                     __LINE__,                                                                     \
+                     hipGetErrorString(err));                                                      \
+            fprintf(stderr, "%s\n", msg);                                                          \
+            throw std::runtime_error(msg);                                                         \
+        }                                                                                          \
+    } while(0)
+
+__global__ void
+test_kernel(float* out, int n)
+{
+    int idx = threadIdx.x + blockIdx.x * blockDim.x;
+    if(idx < n)
+    {
+        float v = static_cast<float>(idx);
+        for(int i = 0; i < 50; i++)
+            v = v * 0.999f + 0.001f;
+        out[idx] = v;
+    }
+}
+
+void
+sigint_handler(int)
+{
+    g_shutdown.store(true, std::memory_order_relaxed);
+}
+
+// Runs HIP kernels in a loop until g_shutdown is set.
+void
+run_kernels(const char* label)
+{
+    roctxRangePush(str_join(label, "_pid_", getpid()).c_str());
+
+    float* d_buf = nullptr;
+    HIP_CHECK(hipMalloc(&d_buf, 1024 * sizeof(float)));
+
+    int iter = 0;
+    while(!g_shutdown.load(std::memory_order_relaxed))
+    {
+        roctxRangePush(str_join(label, "_iter_", iter).c_str());
+        test_kernel<<<4, 256>>>(d_buf, 1024);
+        roctxRangePop();
+
+        iter++;
+        usleep(10000);
+    }
+
+    HIP_CHECK(hipDeviceSynchronize());
+    HIP_CHECK(hipFree(d_buf));
+    roctxRangePop();
+
+    fprintf(stderr, "  %s PID=%d: completed %d iterations\n", label, getpid(), iter);
+}
+
+// Fork child worker for the good case (app handles signals).
+// Child ignores SIGINT and waits for 'q' on pipe from parent.
+void
+fork_child_worker(int id, int pipe_rd)
+{
+    signal(SIGINT, SIG_IGN);
+    fcntl(pipe_rd, F_SETFL, O_NONBLOCK);
+
+    float* d_buf = nullptr;
+    HIP_CHECK(hipMalloc(&d_buf, 1024 * sizeof(float)));
+
+    roctxRangePush(str_join("child_", id, "_pid_", getpid()).c_str());
+
+    int iter = 0;
+    while(true)
+    {
+        roctxRangePush(str_join("child_", id, "_iter_", iter).c_str());
+
+        test_kernel<<<4, 256>>>(d_buf, 1024);
+
+        roctxRangePop();
+        iter++;
+
+        char cmd = 0;
+        if(read(pipe_rd, &cmd, 1) == 1 && cmd == 'q') break;
+        usleep(10000);
+    }
+
+    HIP_CHECK(hipDeviceSynchronize());
+    HIP_CHECK(hipFree(d_buf));
+    roctxRangePop();
+
+    emit_roctx_marker("exit_marker child fork ppid:%d pid:%d", getppid(), getpid());
+    fprintf(stderr, "  child_%d PID=%d: exiting after %d iters\n", id, getpid(), iter);
+    close(pipe_rd);
+}
+
+// ============================================================================
+// GOOD CASE modes: app handles signals, coordinates shutdown
+// ============================================================================
+
+int
+mode_good_single_process()
+{
+    fprintf(stderr, "Mode: good/single-process, PID=%d\n", getpid());
+
+    struct sigaction sa = {};
+    sa.sa_handler       = sigint_handler;
+    sigemptyset(&sa.sa_mask);
+    sigaction(SIGINT, &sa, nullptr);
+
+    run_kernels("parent");
+
+    emit_roctx_marker("exit_marker parent single-process ppid:%d pid:%d", getppid(), getpid());
+    fprintf(stderr, "Parent PID=%d: clean exit\n", getpid());
+    return 0;
+}
+
+int
+mode_good_fork()
+{
+    fprintf(stderr, "Mode: good/fork, PID=%d\n", getpid());
+
+    constexpr int NUM_CHILDREN = 2;
+    int           pipes[NUM_CHILDREN][2];
+    pid_t         children[NUM_CHILDREN];
+
+    for(int i = 0; i < NUM_CHILDREN; i++)
+    {
+        if(pipe(pipes[i]) != 0)
+        {
+            for(int j = 0; j < i; j++)
+            {
+                close(pipes[j][0]);
+                close(pipes[j][1]);
+            }
+
+            throw std::runtime_error(
+                std::string("signal-handler-test pipe() failed with error code ") +
+                std::to_string(errno));
+        }
+        pid_t pid = fork();
+        if(pid == 0)
+        {
+            close(pipes[i][1]);
+            fork_child_worker(i, pipes[i][0]);
+            exit(0);
+        }
+        children[i] = pid;
+        close(pipes[i][0]);
+    }
+
+    struct sigaction sa = {};
+    sa.sa_handler       = sigint_handler;
+    sigemptyset(&sa.sa_mask);
+    sigaction(SIGINT, &sa, nullptr);
+
+    run_kernels("parent");
+
+    fprintf(stderr, "Parent: sending shutdown to children\n");
+    for(int i = 0; i < NUM_CHILDREN; i++)
+    {
+        write(pipes[i][1], "q", 1);
+        close(pipes[i][1]);
+    }
+    for(int i = 0; i < NUM_CHILDREN; i++)
+        waitpid(children[i], nullptr, 0);
+
+    emit_roctx_marker("exit_marker parent fork ppid:%d pid:%d", getppid(), getpid());
+    fprintf(stderr, "Parent PID=%d: clean exit\n", getpid());
+    return 0;
+}
+
+int
+mode_good_fork_exec(const char* self_path)
+{
+    fprintf(stderr, "Mode: good/fork-exec, PID=%d\n", getpid());
+
+    constexpr int NUM_CHILDREN = 2;
+    pid_t         children[NUM_CHILDREN];
+
+    for(int i = 0; i < NUM_CHILDREN; i++)
+    {
+        pid_t pid = fork();
+        if(pid == 0)
+        {
+            execl(self_path, self_path, "--single-process", "--app-signal-handler", nullptr);
+            _exit(127);
+        }
+        children[i] = pid;
+    }
+
+    struct sigaction sa = {};
+    sa.sa_handler       = sigint_handler;
+    sigemptyset(&sa.sa_mask);
+    sigaction(SIGINT, &sa, nullptr);
+
+    run_kernels("parent");
+
+    for(int i = 0; i < NUM_CHILDREN; i++)
+        waitpid(children[i], nullptr, 0);
+
+    emit_roctx_marker("exit_marker parent fork-exec ppid:%d pid:%d", getppid(), getpid());
+    fprintf(stderr, "Parent PID=%d: clean exit\n", getpid());
+    return 0;
+}
+
+int
+mode_good_spawn(const char* self_path)
+{
+    fprintf(stderr, "Mode: good/spawn, PID=%d\n", getpid());
+
+    constexpr int NUM_CHILDREN = 2;
+    pid_t         children[NUM_CHILDREN];
+
+    for(int i = 0; i < NUM_CHILDREN; i++)
+    {
+        pid_t pid = vfork();
+        if(pid == 0)
+        {
+            execl(self_path, self_path, "--single-process", "--app-signal-handler", nullptr);
+            _exit(127);
+        }
+        children[i] = pid;
+    }
+
+    struct sigaction sa = {};
+    sa.sa_handler       = sigint_handler;
+    sigemptyset(&sa.sa_mask);
+    sigaction(SIGINT, &sa, nullptr);
+
+    run_kernels("parent");
+
+    for(int i = 0; i < NUM_CHILDREN; i++)
+        waitpid(children[i], nullptr, 0);
+
+    emit_roctx_marker("exit_marker parent spawn ppid:%d pid:%d", getppid(), getpid());
+    fprintf(stderr, "Parent PID=%d: clean exit\n", getpid());
+    return 0;
+}
+
+// ============================================================================
+// BAD CASE modes: app does NOT handle signals (SIG_DFL everywhere)
+// Profiler's signal handler is the only thing that can flush data.
+// ============================================================================
+
+int
+mode_bad_single_process()
+{
+    fprintf(stderr, "Mode: bad/single-process, PID=%d (no signal handler)\n", getpid());
+    run_kernels("parent");
+    emit_roctx_marker("exit_marker parent single-process ppid:%d pid:%d", getppid(), getpid());
+    fprintf(stderr, "Parent PID=%d: exit\n", getpid());
+    return 0;
+}
+
+int
+mode_bad_fork()
+{
+    fprintf(stderr, "Mode: bad/fork, PID=%d (no signal handler)\n", getpid());
+
+    constexpr int NUM_CHILDREN = 2;
+    pid_t         children[NUM_CHILDREN];
+
+    for(int i = 0; i < NUM_CHILDREN; i++)
+    {
+        pid_t pid = fork();
+        if(pid == 0)
+        {
+            run_kernels(str_join("child_", i).c_str());
+            emit_roctx_marker("exit_marker child fork ppid:%d pid:%d", getppid(), getpid());
+            exit(0);
+        }
+        children[i] = pid;
+    }
+
+    run_kernels("parent");
+
+    for(int i = 0; i < NUM_CHILDREN; i++)
+        waitpid(children[i], nullptr, 0);
+
+    emit_roctx_marker("exit_marker parent fork ppid:%d pid:%d", getppid(), getpid());
+    fprintf(stderr, "Parent PID=%d: exit\n", getpid());
+    return 0;
+}
+
+int
+mode_bad_fork_exec(const char* self_path)
+{
+    fprintf(stderr, "Mode: bad/fork-exec, PID=%d (no signal handler)\n", getpid());
+
+    constexpr int NUM_CHILDREN = 2;
+    pid_t         children[NUM_CHILDREN];
+
+    for(int i = 0; i < NUM_CHILDREN; i++)
+    {
+        pid_t pid = fork();
+        if(pid == 0)
+        {
+            execl(self_path, self_path, "--single-process", "--no-app-signal-handler", nullptr);
+            _exit(127);
+        }
+        children[i] = pid;
+    }
+
+    run_kernels("parent");
+
+    for(int i = 0; i < NUM_CHILDREN; i++)
+        waitpid(children[i], nullptr, 0);
+
+    emit_roctx_marker("exit_marker parent fork-exec ppid:%d pid:%d", getppid(), getpid());
+    fprintf(stderr, "Parent PID=%d: exit\n", getpid());
+    return 0;
+}
+
+int
+mode_bad_spawn(const char* self_path)
+{
+    fprintf(stderr, "Mode: bad/spawn, PID=%d (no signal handler)\n", getpid());
+
+    constexpr int NUM_CHILDREN = 2;
+    pid_t         children[NUM_CHILDREN];
+
+    for(int i = 0; i < NUM_CHILDREN; i++)
+    {
+        pid_t pid = vfork();
+        if(pid == 0)
+        {
+            execl(self_path, self_path, "--single-process", "--no-app-signal-handler", nullptr);
+            _exit(127);
+        }
+        children[i] = pid;
+    }
+
+    run_kernels("parent");
+
+    for(int i = 0; i < NUM_CHILDREN; i++)
+        waitpid(children[i], nullptr, 0);
+
+    emit_roctx_marker("exit_marker parent spawn ppid:%d pid:%d", getppid(), getpid());
+    fprintf(stderr, "Parent PID=%d: exit\n", getpid());
+    return 0;
+}
+
+}  // namespace
+// ============================================================================
+// Main
+// ============================================================================
+
+int
+main(int argc, char** argv)
+{
+    const char* mode                = "--single-process";
+    bool        app_handles_signals = true;
+
+    for(int i = 1; i < argc; i++)
+    {
+        if(std::strcmp(argv[i], "--single-process") == 0 || std::strcmp(argv[i], "--fork") == 0 ||
+           std::strcmp(argv[i], "--fork-exec") == 0 || std::strcmp(argv[i], "--spawn") == 0)
+        {
+            mode = argv[i];
+        }
+        else if(std::strcmp(argv[i], "--app-signal-handler") == 0)
+        {
+            app_handles_signals = true;
+        }
+        else if(std::strcmp(argv[i], "--no-app-signal-handler") == 0)
+        {
+            app_handles_signals = false;
+        }
+    }
+
+    fprintf(stderr,
+            "signal-handler-test: mode=%s app_handles_signals=%d PID=%d\n",
+            mode,
+            static_cast<int>(app_handles_signals),
+            getpid());
+
+    if(app_handles_signals)
+    {
+        if(std::strcmp(mode, "--single-process") == 0) return mode_good_single_process();
+        if(std::strcmp(mode, "--fork") == 0) return mode_good_fork();
+        if(std::strcmp(mode, "--fork-exec") == 0) return mode_good_fork_exec(argv[0]);
+        if(std::strcmp(mode, "--spawn") == 0) return mode_good_spawn(argv[0]);
+    }
+    else
+    {
+        if(std::strcmp(mode, "--single-process") == 0) return mode_bad_single_process();
+        if(std::strcmp(mode, "--fork") == 0) return mode_bad_fork();
+        if(std::strcmp(mode, "--fork-exec") == 0) return mode_bad_fork_exec(argv[0]);
+        if(std::strcmp(mode, "--spawn") == 0) return mode_bad_spawn(argv[0]);
+    }
+
+    fprintf(stderr, "Unknown mode: %s\n", mode);
+    return 1;
+}
diff --git a/projects/rocprofiler-sdk/tests/common/CMakeLists.txt b/projects/rocprofiler-sdk/tests/common/CMakeLists.txt
index bf2959484f7..573d385fdbb 100644
--- a/projects/rocprofiler-sdk/tests/common/CMakeLists.txt
+++ b/projects/rocprofiler-sdk/tests/common/CMakeLists.txt
@@ -520,3 +520,44 @@ function(rocprofiler_add_integration_validate_test NAME)
         PROPERTIES ${_props})
 
 endfunction()
+
+# Registers a ctest that removes the given DIRECTORIES and/or FILES via `cmake -E rm -rf`.
+# Intended to be used as a fixture setup step so that stale output from a prior run is
+# cleared before an execute test writes fresh artifacts, e.g.:
+#
+# rocprofiler_add_integration_cleanup_test(foo DIRECTORIES ${out} FIXTURES_SETUP
+# foo-clean) rocprofiler_add_integration_execute_test(foo ... FIXTURES_REQUIRED foo-clean)
+#
+# TODO: `rm -rf` is intentionally blunt for now; replace with a more targeted cleanup once
+# the set of artifacts each test produces is nailed down.
+function(rocprofiler_add_integration_cleanup_test NAME)
+    cmake_parse_arguments(arg "" "" "DIRECTORIES;FILES;LABELS;FIXTURES_SETUP" ${ARGN})
+
+    rocprofiler_set_integration_test_name(arg_NAME "${NAME}" "cleanup")
+
+    if(NOT arg_DIRECTORIES AND NOT arg_FILES)
+        message(
+            SEND_ERROR
+                "rocprofiler_add_integration_cleanup_test(${NAME}): one of DIRECTORIES or FILES must be specified"
+            )
+        return()
+    endif()
+
+    if(NOT arg_FIXTURES_SETUP)
+        set(arg_FIXTURES_SETUP ${NAME})
+    endif()
+
+    if(NOT "integration-tests" IN_LIST arg_LABELS)
+        list(APPEND arg_LABELS "integration-tests")
+    endif()
+
+    if(NOT "cleanup" IN_LIST arg_LABELS)
+        list(APPEND arg_LABELS "cleanup")
+    endif()
+
+    add_test(NAME "${arg_NAME}" COMMAND ${CMAKE_COMMAND} -E rm -rf ${arg_DIRECTORIES}
+                                        ${arg_FILES})
+
+    set_tests_properties(${arg_NAME} PROPERTIES LABELS "${arg_LABELS}" FIXTURES_SETUP
+                                                "${arg_FIXTURES_SETUP}")
+endfunction()
diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt
index af8cdd5cbe7..90b97155f17 100644
--- a/projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt
+++ b/projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt
@@ -63,3 +63,4 @@ add_subdirectory(roctx-pause-resume)
 add_subdirectory(mpi-ranks)
 add_subdirectory(hip-graph-bubbles-test)
 add_subdirectory(hip-graph-attribution-test)
+add_subdirectory(signal-handler)
diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/CMakeLists.txt
new file mode 100644
index 00000000000..0af5beaead5
--- /dev/null
+++ b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/CMakeLists.txt
@@ -0,0 +1,166 @@
+#
+# rocprofv3 signal handler tests
+#
+cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
+
+project(
+    rocprofiler-sdk-tests-rocprofv3-signal-handler
+    LANGUAGES CXX
+    VERSION 0.0.0)
+
+find_package(rocprofiler-sdk REQUIRED)
+
+string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV
+               "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}")
+
+# Signal delay: how long the app runs before timeout sends SIGINT
+set(SIGNAL_DELAY 3)
+set(CTEST_TIMEOUT 30)
+
+# timeout sends SIGINT after SIGNAL_DELAY seconds (simulates Ctrl+C). --preserve-status
+# returns the child's actual exit code. If our signal handling hangs, the test fails via
+# ctest TIMEOUT ${CTEST_TIMEOUT}
+set(TIMEOUT_CMD timeout --preserve-status --signal=INT ${SIGNAL_DELAY})
+
+# Output directory base
+set(BASE_OUT_DIR ${CMAKE_CURRENT_BINARY_DIR})
+
+# cmake-format: off
+# add_signal_handler_test(<name> [keyword args])
+#
+#   <name>                            : scenario name; ctest is registered as "app-<name>-<type>"
+#   PROCESS_TYPE <type>               : single-process | fork | fork-exec | spawn
+#   VALIDATE_EXPECT <outcome>         : what validate.py asserts:
+#                                          clean-exit: app exits cleanly via atexit, exit_marker present
+#                                          flushed-markers: profiler flushed >10 markers before signal death
+#   PASS_REGEX <regex>                : execute PASS_REGULAR_EXPRESSION
+#   APP_EXTRA_FLAGS <flags...>        : appended to the signal-handler-test command
+#                                       (e.g. --app-signal-handler / --no-app-signal-handler)
+#   ROCPROFV3_EXTRA_FLAGS <flags...>  : inserted into the rocprofv3 command
+#                                       (e.g. --disable-signal-handlers)
+#
+# Registers the {cleanup, execute, validate} triad for one scenario:
+#   cleanup wipes the output dir (FIXTURES_SETUP <name>-clean)
+#   before execute runs (FIXTURES_REQUIRED <name>-clean), and
+#   validate runs after execute (FIXTURES_REQUIRED <name>).
+# cmake-format: on
+function(add_signal_handler_test NAME)
+    cmake_parse_arguments(arg "" "PROCESS_TYPE;VALIDATE_EXPECT;PASS_REGEX"
+                          "APP_EXTRA_FLAGS;ROCPROFV3_EXTRA_FLAGS" ${ARGN})
+
+    set(_name "app-${NAME}-${arg_PROCESS_TYPE}")
+    set(_subdir "${NAME}-${arg_PROCESS_TYPE}")
+    set(_outdir ${BASE_OUT_DIR}/${_subdir})
+
+    rocprofiler_add_integration_cleanup_test(
+        ${_name}
+        DIRECTORIES ${_outdir}
+        LABELS "integration-tests;signal-handler"
+        FIXTURES_SETUP ${_name}-clean)
+
+    rocprofiler_add_integration_execute_test(
+        ${_name}
+        COMMAND
+            ${TIMEOUT_CMD} $<TARGET_FILE:rocprofiler-sdk::rocprofv3> --hip-trace
+            --marker-trace --output-format json ${arg_ROCPROFV3_EXTRA_FLAGS} -d ${_outdir}
+            -o out_%pid% -- $<TARGET_FILE:signal-handler-test> --${arg_PROCESS_TYPE}
+            ${arg_APP_EXTRA_FLAGS}
+        DEPENDS signal-handler-test
+        TIMEOUT ${CTEST_TIMEOUT}
+        LABELS "integration-tests;signal-handler"
+        PRELOAD "${PRELOAD_ENV}"
+        PASS_REGULAR_EXPRESSION "${arg_PASS_REGEX}"
+        FIXTURES_REQUIRED ${_name}-clean
+        FIXTURES_SETUP ${_name})
+
+    rocprofiler_add_integration_validate_test(
+        ${_name}
+        TEST_PATHS validate.py
+        COPY conftest.py
+        CONFIG pytest.ini
+        TIMEOUT ${CTEST_TIMEOUT}
+        LABELS "integration-tests;signal-handler"
+        FIXTURES_REQUIRED ${_name}
+        ARGS --output-dir ${_outdir} --expect ${arg_VALIDATE_EXPECT} --process-type
+             ${arg_PROCESS_TYPE})
+endfunction()
+
+# good case: app installs its own signal handler, rocprofv3 runs with
+# --disable-signal-handlers, so the app exits cleanly through atexit.
+add_signal_handler_test(
+    with-signal-handling
+    PROCESS_TYPE single-process
+    VALIDATE_EXPECT clean-exit
+    PASS_REGEX "clean exit"
+    APP_EXTRA_FLAGS --app-signal-handler
+    ROCPROFV3_EXTRA_FLAGS --disable-signal-handlers)
+
+add_signal_handler_test(
+    with-signal-handling
+    PROCESS_TYPE fork
+    VALIDATE_EXPECT clean-exit
+    PASS_REGEX "clean exit"
+    APP_EXTRA_FLAGS --app-signal-handler
+    ROCPROFV3_EXTRA_FLAGS --disable-signal-handlers)
+
+add_signal_handler_test(
+    with-signal-handling
+    PROCESS_TYPE fork-exec
+    VALIDATE_EXPECT clean-exit
+    PASS_REGEX "clean exit"
+    APP_EXTRA_FLAGS --app-signal-handler
+    ROCPROFV3_EXTRA_FLAGS --disable-signal-handlers)
+
+add_signal_handler_test(
+    with-signal-handling
+    PROCESS_TYPE spawn
+    VALIDATE_EXPECT clean-exit
+    PASS_REGEX "clean exit"
+    APP_EXTRA_FLAGS --app-signal-handler
+    ROCPROFV3_EXTRA_FLAGS --disable-signal-handlers)
+
+# bad case: app does NOT handle signals, so rocprofv3's handler catches the signal and
+# flushes profiling data before the process dies.
+add_signal_handler_test(
+    without-signal-handling
+    PROCESS_TYPE single-process
+    VALIDATE_EXPECT flushed-markers
+    PASS_REGEX "tool finalization"
+    APP_EXTRA_FLAGS --no-app-signal-handler)
+
+add_signal_handler_test(
+    without-signal-handling
+    PROCESS_TYPE fork
+    VALIDATE_EXPECT flushed-markers
+    PASS_REGEX "tool finalization"
+    APP_EXTRA_FLAGS --no-app-signal-handler)
+
+add_signal_handler_test(
+    without-signal-handling
+    PROCESS_TYPE fork-exec
+    VALIDATE_EXPECT flushed-markers
+    PASS_REGEX "tool finalization"
+    APP_EXTRA_FLAGS --no-app-signal-handler)
+
+add_signal_handler_test(
+    without-signal-handling
+    PROCESS_TYPE spawn
+    VALIDATE_EXPECT flushed-markers
+    PASS_REGEX "tool finalization"
+    APP_EXTRA_FLAGS --no-app-signal-handler)
+
+# Coordinated-shutdown case: profiler signal handlers are ON while the app runs its own
+# coordinated multi-process shutdown (the sglang/vLLM tensor-parallel pattern). Worker
+# children ignore SIGINT and only exit once the parent tells them to. The profiler must
+# re-raise to the app's handler before waiting for children to exit, else it deadlocks
+# because parent waits for children that are waiting for the parent and the test hangs,
+# which triggers a ctest TIMEOUT.
+#
+# NOTE: unlike the good case, this does NOT pass --disable-signal-handlers, so the
+# profiler's handlers are exercised alongside the app's own --app-signal-handler.
+add_signal_handler_test(
+    coordinated-shutdown
+    PROCESS_TYPE fork
+    VALIDATE_EXPECT flushed-markers
+    PASS_REGEX "clean exit"
+    APP_EXTRA_FLAGS --app-signal-handler)
diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/conftest.py b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/conftest.py
new file mode 100644
index 00000000000..e085ef5cb50
--- /dev/null
+++ b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/conftest.py
@@ -0,0 +1,67 @@
+#!/usr/bin/env python3
+
+# MIT License
+#
+# Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in all
+# copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+# SOFTWARE.
+
+import pytest
+
+# Recognized --expect outcomes. An unknown value must fail loudly rather than silently skip
+# every check (which would let a scenario pass with zero assertions run).
+VALID_EXPECT = ("clean-exit", "flushed-markers")
+
+
+def pytest_addoption(parser):
+    parser.addoption(
+        "--output-dir", action="store", help="Output directory from rocprofv3"
+    )
+    parser.addoption(
+        "--expect",
+        action="store",
+        help="Expected outcome to assert: clean-exit or flushed-markers",
+    )
+    parser.addoption("--process-type", action="store", help="Process type tested")
+
+
+@pytest.fixture
+def output_dir(request):
+    val = request.config.getoption("--output-dir")
+    if val is None:
+        pytest.fail("--output-dir not provided")
+    return val
+
+
+@pytest.fixture
+def expect(request):
+    val = request.config.getoption("--expect")
+    if val is None:
+        pytest.fail("--expect not provided")
+    if val not in VALID_EXPECT:
+        pytest.fail(f"invalid --expect '{val}'; expected one of {list(VALID_EXPECT)}")
+    return val
+
+
+@pytest.fixture
+def process_type(request):
+    val = request.config.getoption("--process-type")
+    if val is None:
+        pytest.fail("--process-type not provided")
+    return val
diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/pytest.ini b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/pytest.ini
new file mode 100644
index 00000000000..8bf72b0989b
--- /dev/null
+++ b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/pytest.ini
@@ -0,0 +1,5 @@
+
+[pytest]
+addopts = --durations=20 -rA -s
+testpaths = validate.py
+pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages
diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/validate.py
new file mode 100644
index 00000000000..fe355480ea4
--- /dev/null
+++ b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/validate.py
@@ -0,0 +1,139 @@
+#!/usr/bin/env python3
+
+# MIT License
+#
+# Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in all
+# copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+# SOFTWARE.
+
+"""
+Validation for signal handler integration tests.
+
+Every scenario must produce valid and complete JSON output. The profiler intercepts
+signals to flush cleanly, so a truncated or invalid JSON is a bug in signal handler.
+
+The per-scenario expectation is passed explicitly via --expect:
+  clean-exit      : exit_marker present in JSON (proves clean atexit finalization). Used
+                    when the app handles signals and rocprofv3 doesn't (--disable-signal-handlers)
+  flushed-markers : valid JSON with >10 marker entries (proves the profiler flushed before
+                    the process died on the signal). Covers both the "app doesn't handle
+                    signals" case and the coordinated-shutdown case. For the latter,
+                    simply reaching this validator proves there was no deadlock (a
+                    regression trips the execute-step ctest TIMEOUT instead).
+"""
+
+import json
+import os
+import glob
+
+import pytest
+
+
+def find_json_files(output_dir):
+    """Find all JSON result files in the output directory."""
+    pattern = os.path.join(output_dir, "*.json")
+    return sorted(glob.glob(pattern))
+
+
+def load_json(path):
+    """Load and validate a JSON file. Raises on truncated/corrupt JSON."""
+    with open(path, "r") as f:
+        return json.load(f)
+
+
+def count_markers_in_json(data):
+    """Count marker events in the rocprofv3 JSON structure."""
+    count = 0
+    try:
+        for tool_entry in data.get("rocprofiler-sdk-tool", []):
+            buffer_records = tool_entry.get("buffer_records", {})
+            count += len(buffer_records.get("marker_api", []))
+            callback_records = tool_entry.get("callback_records", {})
+            count += len(callback_records.get("marker_api", []))
+    except (AttributeError, TypeError):
+        pass
+    return count
+
+
+def test_output_files_exist(output_dir):
+    """JSON output files must exist."""
+    files = find_json_files(output_dir)
+    assert len(files) > 0, f"No JSON output files found in {output_dir}"
+
+
+def test_json_is_valid(output_dir):
+    """All JSON output files must be valid (not truncated)."""
+    files = find_json_files(output_dir)
+    assert len(files) > 0, f"No JSON files in {output_dir}"
+
+    for path in files:
+        try:
+            load_json(path)
+        except json.JSONDecodeError as e:
+            assert False, (
+                f"JSON file is truncated/corrupt: {path}\n"
+                f"Error: {e}\n"
+                f"This means the profiler did not flush cleanly before process death."
+            )
+
+
+def test_clean_exit(output_dir, expect, process_type):
+    """clean-exit: the app-specific exit_marker must be present (proves the app reached
+    atexit and finalized cleanly)."""
+    if expect != "clean-exit":
+        pytest.skip(f"expectation is '{expect}', not clean-exit")
+
+    files = find_json_files(output_dir)
+    assert len(files) > 0
+
+    expected_marker = f"exit_marker parent {process_type}"
+
+    all_content = ""
+    for path in files:
+        with open(path, "r") as f:
+            all_content += f.read()
+
+    assert expected_marker in all_content, (
+        f"Expected marker '{expected_marker}' not found in output files in {output_dir}. "
+        f"App did not exit cleanly (atexit finalization failed)."
+    )
+
+
+def test_flushed_markers(output_dir, expect):
+    """flushed-markers: JSON must contain >10 marker entries, proving the profiler flushed
+    before the process died on the signal. Covers the "app doesn't handle signals" case and
+    the coordinated-shutdown case (reaching this validator at all already proves the
+    coordinated case did not deadlock -- a regression trips the execute-step ctest
+    TIMEOUT). exit_marker is intentionally NOT required here."""
+    if expect != "flushed-markers":
+        pytest.skip(f"expectation is '{expect}', not flushed-markers")
+
+    files = find_json_files(output_dir)
+    assert len(files) > 0
+
+    total_markers = 0
+    for path in files:
+        data = load_json(path)
+        total_markers += count_markers_in_json(data)
+
+    assert total_markers > 10, (
+        f"Expected >10 marker events in JSON output, got {total_markers}. "
+        f"Profiler may not have flushed marker data before signal death. "
+        f"Files: {files}"
+    )