From d6322f4deeedd265f31ddf5f996221a7be70212d Mon Sep 17 00:00:00 2001 From: "Mythreya, Kuricheti" Date: Tue, 2 Jun 2026 21:03:28 +0000 Subject: [PATCH 01/18] [rocprofiler-sdk] rework signal handling --- .../source/lib/rocprofiler-sdk-tool/tool.cpp | 376 ++++++++++-------- 1 file changed, 212 insertions(+), 164 deletions(-) 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..e76e0eb1b02 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -115,7 +115,10 @@ #include #include +#include +#include #include +#include #include #include @@ -204,6 +207,24 @@ is_handled_signal(int signum) return false; } +struct signal_worker_state +{ + int eventfd = -1; + std::atomic finalize_done = {0}; + std::atomic handling = {false}; + std::atomic finalized = {false}; + int signo = 0; + std::thread thread = {}; + int timeout_sec = 10; +}; + +auto& +get_signal_worker() +{ + static auto*& _v = common::static_object::construct(); + return *CHECK_NOTNULL(_v); +} + struct buffer_ids { rocprofiler_buffer_id_t hsa_api_trace = {}; @@ -2186,6 +2207,9 @@ 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.exchange(true)) return; + ROCP_INFO << "invoked: finalize_rocprofv3"; if(client_finalizer && client_identifier) { @@ -2194,9 +2218,29 @@ finalize_rocprofv3(std::string_view context) client_finalizer = nullptr; client_identifier = nullptr; } - else + + // Join the worker thread if called from a non-worker context (atexit / rocprofv3_main). + // The worker thread calls this function too, so avoid self-join. + if(sw.thread.joinable() && sw.thread.get_id() != std::this_thread::get_id()) { - ROCP_INFO << "finalize_rocprofv3('" << context << "') ignored: already finalized"; + if(sw.eventfd >= 0) + { + // Write to unblock the worker's read() — closing the fd is UB while read() blocks + uint64_t val = 1; + if(write(sw.eventfd, &val, sizeof(val)) < 0) + { + ROCP_WARNING << "failed to write to signal worker eventfd: " << strerror(errno); + } + } + sw.thread.join(); + if(sw.eventfd >= 0) + { + if(close(sw.eventfd) < 0) + { + ROCP_WARNING << "failed to close signal worker eventfd: " << strerror(errno); + } + sw.eventfd = -1; + } } } @@ -3951,6 +3995,8 @@ get_sigaction_function() bool signal_handler_exit = rocprofiler::tool::get_env("ROCPROF_INTERNAL_TEST_SIGNAL_HANDLER_VIA_EXIT", false); + +int signal_handler_timeout = rocprofiler::tool::get_env("ROCPROF_SIGNAL_HANDLER_TIMEOUT", 10); } // namespace #define ROCPROFV3_INTERNAL_API __attribute__((visibility("internal"))); @@ -4106,176 +4152,144 @@ diagnose_status(pid_t _pid, int _status) } void -rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext) +signal_finalization_worker() { - // 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) - { - // 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); - } + auto& sw = get_signal_worker(); - // only reached if the signal is not fatal by default or could not be restored - ::_exit(128 + signo); - } + sigset_t mask; + sigemptyset(&mask); + for(auto sig : rocprofv3_handled_signals) + sigaddset(&mask, sig); + pthread_sigmask(SIG_BLOCK, &mask, nullptr); + + uint64_t val = 0; + if(read(sw.eventfd, &val, sizeof(val)) < 0) return; auto this_pid = getpid(); auto this_ppid = getppid(); 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); + 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{}; + while(ifs) + { + pid_t child_pid = 0; + ifs >> child_pid; + if(ifs && !ifs.eof() && child_pid > 0) children.emplace_back(child_pid); + } + return children; + }; - 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{}; - while(ifs) - { - pid_t val = 0; - ifs >> val; - if(ifs && !ifs.eof() && val > 0) children.emplace_back(val); - } - return children; - }; + 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()); - 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()); + for(auto itr : _children) + { + auto status = wait_pid(itr, WUNTRACED | WNOHANG); + if(status) diagnose_status(itr, status.value()); + } - // wait for children - for(auto itr : _children) - { - auto status = wait_pid(itr, WUNTRACED | WNOHANG); - if(status) diagnose_status(itr, status.value()); - } + ROCP_WARNING << fmt::format( + "[PPID={}][PID={}][TID={}][{}] rocprofv3 finalizing after signal {}...", + this_ppid, + this_pid, + this_tid, + this_func, + 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); - finalize_rocprofv3(this_func); - if(tool::get_config().enable_process_sync) wait_peer_finished(this_pid, this_ppid); + if(tool::get_config().enable_process_sync) wait_peer_finished(this_pid, this_ppid); - 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); + + sw.finalize_done.store(1, std::memory_order_release); + syscall(SYS_futex, &sw.finalize_done, FUTEX_WAKE, 1, nullptr, nullptr, 0); +} + +void +rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext) +{ + (void) info; + (void) ucontext; - if(get_chained_signals().at(signo)) + auto& sw = get_signal_worker(); + + // Re-entry guard (lock-free atomic, async-signal-safe) + bool expected = false; + if(!sw.handling.compare_exchange_strong(expected, true, std::memory_order_acquire)) return; + + // For testing: allow quick_exit path + if(signal_handler_exit) ::quick_exit(signo); + + // Store signal number for the worker thread to log (serialized by eventfd write/read) + sw.signo = signo; + + // Wake the worker thread (write is async-signal-safe) + bool worker_notified = false; + if(sw.eventfd >= 0) + { + uint64_t val = 1; + worker_notified = (write(sw.eventfd, &val, sizeof(val)) == sizeof(val)); + } + + // Wait for worker with bounded timeout (futex syscall is async-signal-safe) + if(worker_notified) + { + struct timespec timeout = {}; + timeout.tv_sec = sw.timeout_sec; + while(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); + syscall(SYS_futex, &sw.finalize_done, FUTEX_WAIT, 0, &timeout, nullptr, 0); + break; + } + } - if(auto& _chained = *get_chained_signals().at(signo); _chained.action) - { - 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); - } - } - else - { - if(_chained.handler) - { - 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); - } - } + // Restore original handler. sigaction() is async-signal-safe per POSIX. + // sw.handling is set, so our interceptor passes through to the real sigaction. + 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); + } - // 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); + // Unblock the signal (it's auto-blocked during handler execution) and re-raise. + // This ensures immediate delivery to the just-restored handler. + sigset_t unblock; + sigemptyset(&unblock); + sigaddset(&unblock, signo); + sigprocmask(SIG_UNBLOCK, &unblock, nullptr); + raise(signo); } int @@ -4317,7 +4331,9 @@ rocprofiler_configure(uint32_t version, add_destructor(execution_profile); // in case main wrapper is not used - ::atexit([]() { finalize_rocprofv3("atexit"); }); + ::atexit([]() { + if(!get_signal_worker().finalized.load()) finalize_rocprofv3("atexit"); + }); tool::get_tmp_file_name_callback() = [](domain_type type) -> std::string { return compose_tmp_file_name(tool::get_config(), type); @@ -4393,9 +4409,10 @@ 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) == true) + 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); @@ -4411,18 +4428,33 @@ 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) == true) + 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 handler so we can restore it on signal delivery. + // Only save the first non-trivial handler registered for each signal. + // Later installs (e.g., LLVM comgr) often don't chain properly — they clobber + // the disposition on re-raise. Preserving the app's handler 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 + { + if(act->sa_handler != SIG_DFL && act->sa_handler != SIG_IGN) + 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); @@ -4456,6 +4488,22 @@ rocprofv3_main(int argc, char** argv, char** envp) initialize_rocprofv3(); + // Resolve dlsym'd function pointers at init time (not in call_once later) + 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"); + + // Spawn signal finalization worker thread before installing signal handlers + { + auto& sw = get_signal_worker(); + sw.eventfd = eventfd(0, EFD_CLOEXEC); + sw.timeout_sec = signal_handler_timeout; + // TODO: evaluate if tools need pre/post thread create callbacks + // if(auto* cb = get_thread_create_callback()) cb->pre("signal_finalization_worker"); + if(sw.eventfd >= 0) sw.thread = std::thread{signal_finalization_worker}; + // if(auto* cb = get_thread_create_callback()) cb->post("signal_finalization_worker"); + } + initialize_signal_handler(get_sigaction_function()); ROCP_INFO << "rocprofv3: main function wrapper will be invoked..."; From 730667ad45feb61b1ed22f7defb2f8ce42c725d1 Mon Sep 17 00:00:00 2001 From: "Mythreya, Kuricheti" Date: Fri, 5 Jun 2026 02:13:12 +0000 Subject: [PATCH 02/18] [rocprofiler-sdk] rework signal handling to use worker thread Replace async-signal-unsafe rocprofv3_error_signal_handler with a worker thread architecture. Signal handler is now fully async-signal-safe: writes eventfd, waits on futex with bounded timeout, restores original handler, re-raises. Worker thread performs finalization in normal context where all APIs are safe to call. Co-authored-by: Cursor --- .../source/lib/rocprofiler-sdk-tool/tool.cpp | 107 +++++++++++++----- 1 file changed, 76 insertions(+), 31 deletions(-) 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 e76e0eb1b02..537bd11be9c 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -140,6 +140,10 @@ rocprofv3_error_signal_handler(int signo, siginfo_t*, void*); namespace { +// Initialized once in rocprofv3_main. Used by the signal handler and finalization +// to detect child processes (fork/fork+exec). +pid_t ROCPROF_ROOT_PID = 0; + // Thread for safe cleanup output generation auto output_generation_thread = common::Synchronized>{}; @@ -219,9 +223,15 @@ struct signal_worker_state }; auto& -get_signal_worker() +get_signal_worker(bool reset = false) { static auto*& _v = common::static_object::construct(); + if(reset) + { + // After fork, parent's state (eventfd, thread, atomics) is stale. + // Zero everything and reconstruct in place. + std::memset(static_cast(_v), 0, sizeof(signal_worker_state)); + } return *CHECK_NOTNULL(_v); } @@ -2220,8 +2230,9 @@ finalize_rocprofv3(std::string_view context) } // Join the worker thread if called from a non-worker context (atexit / rocprofv3_main). - // The worker thread calls this function too, so avoid self-join. - if(sw.thread.joinable() && sw.thread.get_id() != std::this_thread::get_id()) + // Only in the root process — forked children have a stale thread object. + if(ROCPROF_ROOT_PID == getpid() && sw.thread.joinable() && + sw.thread.get_id() != std::this_thread::get_id()) { if(sw.eventfd >= 0) { @@ -3997,6 +4008,7 @@ bool signal_handler_exit = rocprofiler::tool::get_env("ROCPROF_INTERNAL_TEST_SIGNAL_HANDLER_VIA_EXIT", false); int signal_handler_timeout = rocprofiler::tool::get_env("ROCPROF_SIGNAL_HANDLER_TIMEOUT", 10); + } // namespace #define ROCPROFV3_INTERNAL_API __attribute__((visibility("internal"))); @@ -4152,24 +4164,8 @@ diagnose_status(pid_t _pid, int _status) } void -signal_finalization_worker() +wait_for_children(pid_t this_pid, pid_t this_ppid, uint64_t this_tid, std::string_view context) { - 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); - - uint64_t val = 0; - if(read(sw.eventfd, &val, sizeof(val)) < 0) return; - - auto this_pid = getpid(); - auto this_ppid = getppid(); - auto this_tid = common::get_tid(); - auto this_func = std::string_view{__FUNCTION__}; - auto get_children = [&this_pid]() { auto fname = fmt::format("/proc/{}/task/{}/children", this_pid, this_pid); auto ifs = std::ifstream{fname}; @@ -4185,11 +4181,11 @@ signal_finalization_worker() auto _children = get_children(); ROCP_WARNING << fmt::format( - "[PPID={}][PID={}][TID={}][{}] rocprofv3 will wait for {} children to exit", + "[PPID={}][PID={}][TID={}][{}] rocprofv3 waiting for {} children to exit", this_ppid, this_pid, this_tid, - this_func, + context, _children.size()); for(auto itr : _children) @@ -4197,6 +4193,26 @@ signal_finalization_worker() auto status = wait_pid(itr, WUNTRACED | WNOHANG); if(status) diagnose_status(itr, status.value()); } +} + +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); + + uint64_t val = 0; + if(read(sw.eventfd, &val, sizeof(val)) < 0) return; + + auto this_pid = getpid(); + auto this_ppid = getppid(); + auto this_tid = common::get_tid(); + auto this_func = std::string_view{__FUNCTION__}; ROCP_WARNING << fmt::format( "[PPID={}][PID={}][TID={}][{}] rocprofv3 finalizing after signal {}...", @@ -4210,6 +4226,11 @@ signal_finalization_worker() if(tool::get_config().enable_process_sync) wait_peer_finished(this_pid, this_ppid); + // If we're handling signals, the app likely doesn't coordinate child shutdown. + // Wait for children so they can flush their profiling data via atexit before + // the parent re-raises and dies. + 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, @@ -4229,6 +4250,10 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext) auto& sw = get_signal_worker(); + // We're in the bad path — app doesn't coordinate shutdown properly. + // All processes (parent and children) need to flush and die. + // For well-behaved apps, use --disable-signal-handlers (atexit handles everything). + // Re-entry guard (lock-free atomic, async-signal-safe) bool expected = false; if(!sw.handling.compare_exchange_strong(expected, true, std::memory_order_acquire)) return; @@ -4283,9 +4308,8 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext) get_sigaction_function()(signo, &sa, nullptr); } - // Unblock the signal (it's auto-blocked during handler execution) and re-raise. - // This ensures immediate delivery to the just-restored handler. - sigset_t unblock; + // Unblock and re-raise. App's handler (or SIG_DFL) runs in a fresh delivery. + sigset_t unblock {}; sigemptyset(&unblock); sigaddset(&unblock, signo); sigprocmask(SIG_UNBLOCK, &unblock, nullptr); @@ -4411,7 +4435,7 @@ rocprofv3_signal(int signum, sighandler_t handler) { if(!get_signal_function()) get_signal_function() = (signal_func_t) dlsym(RTLD_NEXT, "signal"); - if(get_signal_worker().handling.load(std::memory_order_relaxed) == true) + if(get_signal_worker().handling.load(std::memory_order_relaxed)) return get_signal_function()(signum, handler); if(!is_handled_signal(signum) || !tool::get_config().enable_signal_handlers) @@ -4431,7 +4455,7 @@ rocprofv3_sigaction(int signum, if(!get_sigaction_function()) get_sigaction_function() = (sigaction_func_t) dlsym(RTLD_NEXT, "sigaction"); - if(get_signal_worker().handling.load(std::memory_order_relaxed) == true) + if(get_signal_worker().handling.load(std::memory_order_relaxed)) return get_sigaction_function()(signum, act, oldact); if(!is_handled_signal(signum) || !act || !tool::get_config().enable_signal_handlers) @@ -4493,15 +4517,36 @@ rocprofv3_main(int argc, char** argv, char** envp) if(!get_sigaction_function()) get_sigaction_function() = (sigaction_func_t) dlsym(RTLD_NEXT, "sigaction"); - // Spawn signal finalization worker thread before installing signal handlers + // Determine root process for signal handling. + // Only the root process installs signal handlers and spawns the finalization worker. + // Children (fork or fork+exec) inherit ROCPROF_ROOT_PID and skip signal handling — + // they rely on atexit for finalization after the parent coordinates their shutdown. + // For apps that correctly handle signals themselves, --disable-signal-handlers + // bypasses all of this and relies purely on atexit in all processes. { + auto* root_pid_env = getenv("ROCPROF_ROOT_PID"); + bool is_child = (root_pid_env != nullptr); + + if(is_child) + { + // Reset the signal worker state — the inherited state references + // the parent's eventfd/thread which are inaccessible from here. + get_signal_worker(true); + ROCPROF_ROOT_PID = static_cast(std::stol(root_pid_env)); + } + else + { + ROCPROF_ROOT_PID = getpid(); + setenv("ROCPROF_ROOT_PID", std::to_string(ROCPROF_ROOT_PID).c_str(), 0); + } + + // Spawn finalization worker in ALL processes (parent and children). + // On the bad path (signal handlers enabled), every process needs to + // flush its own profiling data async-signal-safely via a worker thread. auto& sw = get_signal_worker(); sw.eventfd = eventfd(0, EFD_CLOEXEC); sw.timeout_sec = signal_handler_timeout; - // TODO: evaluate if tools need pre/post thread create callbacks - // if(auto* cb = get_thread_create_callback()) cb->pre("signal_finalization_worker"); if(sw.eventfd >= 0) sw.thread = std::thread{signal_finalization_worker}; - // if(auto* cb = get_thread_create_callback()) cb->post("signal_finalization_worker"); } initialize_signal_handler(get_sigaction_function()); From 9bf3a274468307bffddf604cff90b2e65f858361 Mon Sep 17 00:00:00 2001 From: "Mythreya, Kuricheti" Date: Fri, 5 Jun 2026 20:50:44 +0000 Subject: [PATCH 03/18] handle at-fork --- .../source/lib/rocprofiler-sdk-tool/tool.cpp | 57 ++++++++++++------- 1 file changed, 38 insertions(+), 19 deletions(-) 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 537bd11be9c..2f041c4adf5 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -2230,9 +2230,7 @@ finalize_rocprofv3(std::string_view context) } // Join the worker thread if called from a non-worker context (atexit / rocprofv3_main). - // Only in the root process — forked children have a stale thread object. - if(ROCPROF_ROOT_PID == getpid() && sw.thread.joinable() && - sw.thread.get_id() != std::this_thread::get_id()) + if(sw.thread.joinable() && sw.thread.get_id() != std::this_thread::get_id()) { if(sw.eventfd >= 0) { @@ -4226,9 +4224,7 @@ signal_finalization_worker() if(tool::get_config().enable_process_sync) wait_peer_finished(this_pid, this_ppid); - // If we're handling signals, the app likely doesn't coordinate child shutdown. - // Wait for children so they can flush their profiling data via atexit before - // the parent re-raises and dies. + // Wait for children to finalize before we re-raise. wait_for_children(this_pid, this_ppid, this_tid, this_func); ROCP_INFO << fmt::format( @@ -4250,11 +4246,11 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext) auto& sw = get_signal_worker(); - // We're in the bad path — app doesn't coordinate shutdown properly. + // We're handling signals = assume app doesn't coordinate shutdown properly. // All processes (parent and children) need to flush and die. // For well-behaved apps, use --disable-signal-handlers (atexit handles everything). - // Re-entry guard (lock-free atomic, async-signal-safe) + // Re-entry guard bool expected = false; if(!sw.handling.compare_exchange_strong(expected, true, std::memory_order_acquire)) return; @@ -4264,7 +4260,7 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext) // Store signal number for the worker thread to log (serialized by eventfd write/read) sw.signo = signo; - // Wake the worker thread (write is async-signal-safe) + // Wake the worker thread bool worker_notified = false; if(sw.eventfd >= 0) { @@ -4272,7 +4268,8 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext) worker_notified = (write(sw.eventfd, &val, sizeof(val)) == sizeof(val)); } - // Wait for worker with bounded timeout (futex syscall is async-signal-safe) + // Wait for worker with bounded timeout via futex syscall instead of a std::mutex + // to guarantee async signal-safe communication if(worker_notified) { struct timespec timeout = {}; @@ -4308,11 +4305,20 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext) get_sigaction_function()(signo, &sa, nullptr); } - // Unblock and re-raise. App's handler (or SIG_DFL) runs in a fresh delivery. - sigset_t unblock {}; + // Unblock and re-raise. The restored handler (or SIG_DFL) runs in a fresh delivery. + sigset_t unblock{}; sigemptyset(&unblock); sigaddset(&unblock, signo); - sigprocmask(SIG_UNBLOCK, &unblock, nullptr); + pthread_sigmask(SIG_UNBLOCK, &unblock, nullptr); + raise(signo); + + // If the chained handler returned (didn't kill the process), force exit. + { + struct sigaction sa = {}; + sa.sa_handler = SIG_DFL; + sigemptyset(&sa.sa_mask); + get_sigaction_function()(signo, &sa, nullptr); + } raise(signo); } @@ -4517,12 +4523,9 @@ rocprofv3_main(int argc, char** argv, char** envp) if(!get_sigaction_function()) get_sigaction_function() = (sigaction_func_t) dlsym(RTLD_NEXT, "sigaction"); - // Determine root process for signal handling. - // Only the root process installs signal handlers and spawns the finalization worker. - // Children (fork or fork+exec) inherit ROCPROF_ROOT_PID and skip signal handling — - // they rely on atexit for finalization after the parent coordinates their shutdown. - // For apps that correctly handle signals themselves, --disable-signal-handlers - // bypasses all of this and relies purely on atexit in all processes. + // Track root process via ROCPROF_ROOT_PID env var. + // All processes (root + children) get their own worker thread and signal handler. + // For well-behaved apps, use --disable-signal-handlers (atexit handles everything). { auto* root_pid_env = getenv("ROCPROF_ROOT_PID"); bool is_child = (root_pid_env != nullptr); @@ -4547,6 +4550,22 @@ rocprofv3_main(int argc, char** argv, char** envp) sw.eventfd = eventfd(0, EFD_CLOEXEC); sw.timeout_sec = signal_handler_timeout; 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. + get_signal_worker(true); + auto& child_sw = get_signal_worker(); + child_sw.eventfd = eventfd(0, EFD_CLOEXEC); + child_sw.timeout_sec = signal_handler_timeout; + if(child_sw.eventfd >= 0) child_sw.thread = std::thread{signal_finalization_worker}; + }); + } } initialize_signal_handler(get_sigaction_function()); From b65ba33105b63ea86e0917b4c1e6d9e3cdd3895e Mon Sep 17 00:00:00 2001 From: "Mythreya, Kuricheti" Date: Sat, 6 Jun 2026 00:17:31 +0000 Subject: [PATCH 04/18] formatting --- .../source/lib/rocprofiler-sdk-tool/tool.cpp | 28 +- .../signal-handler/signal-handler-test.cpp | 470 ++++++++++++++++++ 2 files changed, 493 insertions(+), 5 deletions(-) create mode 100644 projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp 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 2f041c4adf5..d3e9eb03492 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -4005,7 +4005,7 @@ get_sigaction_function() bool signal_handler_exit = rocprofiler::tool::get_env("ROCPROF_INTERNAL_TEST_SIGNAL_HANDLER_VIA_EXIT", false); -int signal_handler_timeout = rocprofiler::tool::get_env("ROCPROF_SIGNAL_HANDLER_TIMEOUT", 10); +int signal_handler_timeout = rocprofiler::tool::get_env("ROCPROF_SIGNAL_HANDLER_TIMEOUT", 0); } // namespace @@ -4269,15 +4269,33 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext) } // Wait for worker with bounded timeout via futex syscall instead of a std::mutex - // to guarantee async signal-safe communication + // Wait for worker to complete finalization. + // Uses exponential backoff (5, 10, 30, 60s) with periodic logging. if(worker_notified) { - struct timespec timeout = {}; - timeout.tv_sec = sw.timeout_sec; + static constexpr auto backoff_sec = std::array{5, 10, 30, 60}; + + int attempt = 0; + while(sw.finalize_done.load(std::memory_order_acquire) == 0) { + auto timeout = timespec{}; + timeout.tv_sec = backoff_sec[attempt < 4 ? attempt : 3]; + syscall(SYS_futex, &sw.finalize_done, FUTEX_WAIT, 0, &timeout, nullptr, 0); - break; + + if(sw.finalize_done.load(std::memory_order_acquire) != 0) break; + + // Still waiting — log progress + auto elapsed = backoff_sec[0]; + for(int j = 1; j <= attempt && j < 4; j++) + elapsed += backoff_sec[j]; + + // technically unsafe, but otherwise the process looks hung to users + ROCP_WARNING << fmt::format( + "signalhangler: waiting for worker thread to flush profiling data: ~{}s elapsed...", + elapsed); + attempt++; } } diff --git a/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp b/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp new file mode 100644 index 00000000000..7d0987e76d4 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp @@ -0,0 +1,470 @@ +// MIT License +// +// Copyright (c) 2023-2025 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 +#include + +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +namespace +{ +std::atomic g_shutdown{false}; + +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(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) +{ + char range_name[128]; + snprintf(range_name, sizeof(range_name), "%s_pid_%d", label, getpid()); + roctxRangePush(range_name); + + float* d_buf = nullptr; + HIP_CHECK(hipMalloc(&d_buf, 1024 * sizeof(float))); + + int iter = 0; + while(!g_shutdown.load(std::memory_order_relaxed)) + { + char iter_name[128]; + snprintf(iter_name, sizeof(iter_name), "%s_iter_%d", label, iter); + + roctxRangePush(iter_name); + 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))); + + char range_name[128]; + snprintf(range_name, sizeof(range_name), "child_%d_pid_%d", id, getpid()); + roctxRangePush(range_name); + + int iter = 0; + while(true) + { + char iter_name[64]; + snprintf(iter_name, sizeof(iter_name), "child_%d_iter_%d", id, iter); + roctxRangePush(iter_name); + + 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++) + { + pipe(pipes[i]); + 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) + { + char label[32]; + snprintf(label, sizeof(label), "child_%d", i); + run_kernels(label); + 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(strcmp(argv[i], "--single-process") == 0 || strcmp(argv[i], "--fork") == 0 || + strcmp(argv[i], "--fork-exec") == 0 || strcmp(argv[i], "--spawn") == 0) + { + mode = argv[i]; + } + else if(strcmp(argv[i], "--app-signal-handler") == 0) + { + app_handles_signals = true; + } + else if(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(app_handles_signals), + getpid()); + + if(app_handles_signals) + { + if(strcmp(mode, "--single-process") == 0) return mode_good_single_process(); + if(strcmp(mode, "--fork") == 0) return mode_good_fork(); + if(strcmp(mode, "--fork-exec") == 0) return mode_good_fork_exec(argv[0]); + if(strcmp(mode, "--spawn") == 0) return mode_good_spawn(argv[0]); + } + else + { + if(strcmp(mode, "--single-process") == 0) return mode_bad_single_process(); + if(strcmp(mode, "--fork") == 0) return mode_bad_fork(); + if(strcmp(mode, "--fork-exec") == 0) return mode_bad_fork_exec(argv[0]); + if(strcmp(mode, "--spawn") == 0) return mode_bad_spawn(argv[0]); + } + + fprintf(stderr, "Unknown mode: %s\n", mode); + return 1; +} From ca21403b744aecfb6ee4b21e9341717ddff75b5a Mon Sep 17 00:00:00 2001 From: "Mythreya, Kuricheti" Date: Sat, 6 Jun 2026 00:26:50 +0000 Subject: [PATCH 05/18] Add initial tests for signal handling Co-authored-by: Cursor --- .../rocprofiler-sdk/tests/bin/CMakeLists.txt | 1 + .../tests/bin/signal-handler/CMakeLists.txt | 45 ++++++++ .../tests/rocprofv3/CMakeLists.txt | 1 + .../rocprofv3/signal-handler/CMakeLists.txt | 91 ++++++++++++++++ .../rocprofv3/signal-handler/conftest.py | 33 ++++++ .../tests/rocprofv3/signal-handler/pytest.ini | 2 + .../rocprofv3/signal-handler/validate.py | 102 ++++++++++++++++++ 7 files changed, 275 insertions(+) create mode 100644 projects/rocprofiler-sdk/tests/bin/signal-handler/CMakeLists.txt create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/CMakeLists.txt create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/conftest.py create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/pytest.ini create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/validate.py 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..d05af88f96f --- /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-test.cpp PROPERTIES LANGUAGE HIP) +add_executable(signal-handler-test) +target_sources(signal-handler-test PRIVATE signal-handler-test.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/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..febe1fecf18 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/CMakeLists.txt @@ -0,0 +1,91 @@ +# +# 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}) + +# Process types to test +set(PROCESS_TYPES single-process fork fork-exec spawn) + +# =========================================================================== +# Generate all 8 test combos: {good, bad} x {single-process, fork, fork-exec, spawn} +# =========================================================================== + +foreach(PROC_TYPE ${PROCESS_TYPES}) + # --- GOOD CASE: app handles signals, profiler uses --disable-signal-handlers --- + set(WITH_HANDLING_NAME "app-with-signal-handling-${PROC_TYPE}") + + rocprofiler_add_integration_execute_test( + ${WITH_HANDLING_NAME} + COMMAND + ${TIMEOUT_CMD} $ --hip-trace + --marker-trace --output-format json --disable-signal-handlers -d + ${BASE_OUT_DIR}/good-${PROC_TYPE} -o out_%pid% -- + $ --${PROC_TYPE} --app-signal-handler + DEPENDS signal-handler-test + TIMEOUT ${CTEST_TIMEOUT} + LABELS "integration-tests;signal-handler" + PRELOAD "${PRELOAD_ENV}" + PASS_REGULAR_EXPRESSION "clean exit" + FIXTURES_SETUP ${WITH_HANDLING_NAME}) + + rocprofiler_add_integration_validate_test( + ${WITH_HANDLING_NAME} + TEST_PATHS validate.py + COPY conftest.py + CONFIG pytest.ini + TIMEOUT 30 + LABELS "integration-tests;signal-handler" + FIXTURES_REQUIRED ${WITH_HANDLING_NAME} + ARGS --output-dir ${BASE_OUT_DIR}/good-${PROC_TYPE} --mode good --process-type + ${PROC_TYPE}) + + # --- BAD CASE: app does NOT handle signals, profiler handles them --- + set(NO_HANDLING_NAME "app-without-signal-handling-${PROC_TYPE}") + + rocprofiler_add_integration_execute_test( + ${NO_HANDLING_NAME} + COMMAND + ${TIMEOUT_CMD} $ --hip-trace + --marker-trace --output-format json -d ${BASE_OUT_DIR}/bad-${PROC_TYPE} -o + out_%pid% -- $ --${PROC_TYPE} + --no-app-signal-handler + DEPENDS signal-handler-test + TIMEOUT ${CTEST_TIMEOUT} + LABELS "integration-tests;signal-handler" + PRELOAD "${PRELOAD_ENV}" + PASS_REGULAR_EXPRESSION "tool finalization" + FIXTURES_SETUP ${NO_HANDLING_NAME}) + + rocprofiler_add_integration_validate_test( + ${NO_HANDLING_NAME} + TEST_PATHS validate.py + COPY conftest.py + CONFIG pytest.ini + TIMEOUT ${CTEST_TIMEOUT} + LABELS "integration-tests;signal-handler" + FIXTURES_REQUIRED ${NO_HANDLING_NAME} + ARGS --output-dir ${BASE_OUT_DIR}/bad-${PROC_TYPE} --mode bad --process-type + ${PROC_TYPE}) +endforeach() 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..47b6d63fb9f --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/conftest.py @@ -0,0 +1,33 @@ +import pytest + + +def pytest_addoption(parser): + parser.addoption( + "--output-dir", action="store", help="Output directory from rocprofv3" + ) + parser.addoption("--mode", action="store", help="Test mode: good or bad") + 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.skip("--output-dir not provided") + return val + + +@pytest.fixture +def mode(request): + val = request.config.getoption("--mode") + if val is None: + pytest.skip("--mode not provided") + return val + + +@pytest.fixture +def process_type(request): + val = request.config.getoption("--process-type") + if val is None: + pytest.skip("--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..8fce174e45d --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/pytest.ini @@ -0,0 +1,2 @@ +[pytest] +addopts = -rA --tb=short 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..9a0ff547504 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/validate.py @@ -0,0 +1,102 @@ +""" +Validation for signal handler integration tests. + +Both good and bad cases must produce VALID, COMPLETE JSON output. +The profiler intercepts signals to flush cleanly — truncated JSON is a bug. + +Good case: exit_marker present in JSON (proves clean atexit finalization). +Bad case: valid JSON with >10 marker entries (proves profiler flushed before death). +""" + +import json +import os +import glob + + +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, mode, process_type): + """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, mode, process_type): + """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_good_case_exit_marker(output_dir, mode, process_type): + """Good case: mode-specific exit_marker must be present.""" + if mode != "good": + return + + 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_bad_case_markers_flushed(output_dir, mode, process_type): + """Bad case: JSON must contain >10 marker entries (proves profiler flushed).""" + if mode != "bad": + return + + 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}" + ) From f048bf83a2fe018052405b7fe3b46e3f46220b50 Mon Sep 17 00:00:00 2001 From: "Mythreya, Kuricheti" Date: Fri, 12 Jun 2026 21:58:42 +0000 Subject: [PATCH 06/18] selective logging for dangling IDs --- .../context/correlation_id.cpp | 21 ++++++++++++++++++- 1 file changed, 20 insertions(+), 1 deletion(-) 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..13e4dcb457d 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 @@ -253,7 +253,26 @@ 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) + << "retired dangling correlation IDs: " << ndangling; + } + else + { + ROCP_INFO << "retired dangling correlation IDs: " << ndangling + << " (no retirement consumer active)"; + } + } }); } } // namespace context From 42073b74de7ff0a36dc7026b5f5ff94eb25187fa Mon Sep 17 00:00:00 2001 From: "Mythreya, Kuricheti" Date: Tue, 16 Jun 2026 19:52:29 +0000 Subject: [PATCH 07/18] correct copyright notices --- .../signal-handler/signal-handler-test.cpp | 2 +- .../rocprofv3/signal-handler/conftest.py | 24 +++++++++++++++++++ .../rocprofv3/signal-handler/validate.py | 24 +++++++++++++++++++ 3 files changed, 49 insertions(+), 1 deletion(-) diff --git a/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp b/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp index 7d0987e76d4..0836d9f13e1 100644 --- a/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp +++ b/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp @@ -1,6 +1,6 @@ // MIT License // -// Copyright (c) 2023-2025 Advanced Micro Devices, Inc. All rights reserved. +// 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 diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/conftest.py b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/conftest.py index 47b6d63fb9f..f2b569e611b 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/conftest.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/conftest.py @@ -1,3 +1,27 @@ +#!/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 diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/validate.py index 9a0ff547504..266f834f5bc 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/validate.py @@ -1,3 +1,27 @@ +#!/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. From 32ab5d4d22414d6969b4b9dda5ab622873c58cb9 Mon Sep 17 00:00:00 2001 From: "Mythreya, Kuricheti" Date: Tue, 16 Jun 2026 20:05:33 +0000 Subject: [PATCH 08/18] fix typo Co-authored-by: Mark Meserve --- .../rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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 d3e9eb03492..4614a903cd3 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -4293,7 +4293,7 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext) // technically unsafe, but otherwise the process looks hung to users ROCP_WARNING << fmt::format( - "signalhangler: waiting for worker thread to flush profiling data: ~{}s elapsed...", + "signalhandler: waiting for worker thread to flush profiling data: ~{}s elapsed...", elapsed); attempt++; } From f62219e508888aaa7e774b8d77bea6bc63e34b04 Mon Sep 17 00:00:00 2001 From: "Mythreya, Kuricheti" Date: Tue, 16 Jun 2026 20:39:06 +0000 Subject: [PATCH 09/18] address review comments - use atomic - add error checking for pipe - re-add previous removed INFO call for ignored finalization - re-init on fork after memset --- .../source/lib/rocprofiler-sdk-tool/tool.cpp | 82 +++++++++++++------ .../signal-handler/signal-handler-test.cpp | 15 +++- 2 files changed, 73 insertions(+), 24 deletions(-) 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 4614a903cd3..8f05de5e4ae 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 #include #include +#include #include #include #include @@ -213,26 +214,49 @@ is_handled_signal(int signum) struct signal_worker_state { - int eventfd = -1; - std::atomic finalize_done = {0}; - std::atomic handling = {false}; - std::atomic finalized = {false}; - int signo = 0; - std::thread thread = {}; - int timeout_sec = 10; + int eventfd = -1; + std::atomic finalize_done = {}; + // set to 1 when `rocprofv3_error_signal_handler` owns the signal path (re-entry guard + + // interceptor bypass). + std::atomic handling = {0}; + std::atomic_flag finalized = ATOMIC_FLAG_INIT; + int signo = {0}; + std::thread thread = {}; + int timeout_sec = {10}; }; +static_assert(std::atomic::is_always_lock_free, + "rocprofv3 signal path requires lock-free finalize_done atomic support"); + auto& get_signal_worker(bool reset = false) { static auto*& _v = common::static_object::construct(); + auto& sw = *CHECK_NOTNULL(_v); + + const auto reset_signal_worker_state = [](signal_worker_state& obj) { + if(obj.eventfd >= 0) + { + if(::close(obj.eventfd) != 0) + { + ROCP_WARNING << "signal worker: close(eventfd) after fork failed: " + << strerror(errno); + } + } + + std::memset(static_cast(&obj), 0, sizeof(obj)); + ::new(static_cast(&obj)) signal_worker_state{}; + }; + if(reset) { - // After fork, parent's state (eventfd, thread, atomics) is stale. - // Zero everything and reconstruct in place. - std::memset(static_cast(_v), 0, sizeof(signal_worker_state)); + // Post-fork child: parent's eventfd duplicate and std::thread are invalid. Close the fd + // first, then reset storage and value-construct so defaults live only in + // signal_worker_state (we do not run ~std::thread() on a possibly-joinable parent thread). + reset_signal_worker_state(sw); } - return *CHECK_NOTNULL(_v); + + return sw; } struct buffer_ids @@ -2218,7 +2242,11 @@ void finalize_rocprofv3(std::string_view context) { auto& sw = get_signal_worker(); - if(sw.finalized.exchange(true)) return; + 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) @@ -4168,6 +4196,11 @@ wait_for_children(pid_t this_pid, pid_t this_ppid, uint64_t this_tid, std::strin auto fname = fmt::format("/proc/{}/task/{}/children", this_pid, this_pid); auto ifs = std::ifstream{fname}; auto children = std::vector{}; + if(!ifs.is_open()) + { + ROCP_WARNING << "signal worker: failed to open " << fname << ": " << strerror(errno); + return children; + } while(ifs) { pid_t child_pid = 0; @@ -4234,7 +4267,7 @@ signal_finalization_worker() this_tid, this_func); - sw.finalize_done.store(1, std::memory_order_release); + sw.finalize_done.store(1u, std::memory_order_release); syscall(SYS_futex, &sw.finalize_done, FUTEX_WAKE, 1, nullptr, nullptr, 0); } @@ -4250,9 +4283,9 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext) // All processes (parent and children) need to flush and die. // For well-behaved apps, use --disable-signal-handlers (atexit handles everything). - // Re-entry guard - bool expected = false; - if(!sw.handling.compare_exchange_strong(expected, true, std::memory_order_acquire)) return; + // Re-entry guard: lock-free compare-exchange (see static_assert on `handling`). + uint32_t expected = 0; + if(!sw.handling.compare_exchange_strong(expected, 1u, std::memory_order_acquire)) return; // For testing: allow quick_exit path if(signal_handler_exit) ::quick_exit(signo); @@ -4291,7 +4324,12 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext) for(int j = 1; j <= attempt && j < 4; j++) elapsed += backoff_sec[j]; - // technically unsafe, but otherwise the process looks hung to users + // Logging here is not async-signal-safe in the strict POSIX sense (Abseil may take + // locks). In practice this path is still safe enough for diagnostics: the handler is + // one-shot (re-entry guard above), SA_NODEFER is not set on our installed action so the + // same thread cannot be re-entered for another signal while we're here, and other + // threads may log concurrently (Abseil is thread-safe). We keep this message so users + // see progress during long flushes. ROCP_WARNING << fmt::format( "signalhandler: waiting for worker thread to flush profiling data: ~{}s elapsed...", elapsed); @@ -4379,9 +4417,7 @@ rocprofiler_configure(uint32_t version, add_destructor(execution_profile); // in case main wrapper is not used - ::atexit([]() { - if(!get_signal_worker().finalized.load()) finalize_rocprofv3("atexit"); - }); + ::atexit([]() { finalize_rocprofv3("atexit"); }); tool::get_tmp_file_name_callback() = [](domain_type type) -> std::string { return compose_tmp_file_name(tool::get_config(), type); @@ -4459,7 +4495,7 @@ rocprofv3_signal(int signum, sighandler_t handler) { if(!get_signal_function()) get_signal_function() = (signal_func_t) dlsym(RTLD_NEXT, "signal"); - if(get_signal_worker().handling.load(std::memory_order_relaxed)) + 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) @@ -4479,7 +4515,7 @@ rocprofv3_sigaction(int signum, if(!get_sigaction_function()) get_sigaction_function() = (sigaction_func_t) dlsym(RTLD_NEXT, "sigaction"); - if(get_signal_worker().handling.load(std::memory_order_relaxed)) + 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) @@ -4536,7 +4572,7 @@ rocprofv3_main(int argc, char** argv, char** envp) initialize_rocprofv3(); - // Resolve dlsym'd function pointers at init time (not in call_once later) + // 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"); diff --git a/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp b/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp index 0836d9f13e1..87a98a3f900 100644 --- a/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp +++ b/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp @@ -48,10 +48,12 @@ #include #include #include +#include #include #include #include +#include namespace { @@ -211,7 +213,18 @@ mode_good_fork() for(int i = 0; i < NUM_CHILDREN; i++) { - pipe(pipes[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) { From c35e013c5d09d28beb8b5a073ecec50193d9dfca Mon Sep 17 00:00:00 2001 From: "Mythreya, Kuricheti" Date: Tue, 4 Aug 2026 20:50:00 +0000 Subject: [PATCH 10/18] [rocprofiler-sdk] make error signal handler non-blocking The handler used to wake the finalization worker and then block the signaled thread until the flush finished, re-raising on that same thread. If the signal interrupted that thread mid-malloc (holding the arena lock), the worker's flush would block on the same lock -> deadlock; it also stalled the app's own coordinated-shutdown handlers. Now the handler hands off and gets out of the way: - Async signals (SIGINT/SIGQUIT/SIGTERM): record signo, wake the worker via eventfd, and return immediately. Returning lets the interrupted thread resume and release any lock it holds, so the worker's (non-async-signal- safe) flush cannot deadlock against it. - The worker owns termination: after finalize + reaping children it restores the wrapped disposition (the app's chained handler if one was saved, else SIG_DFL) and re-raises process-directed via kill(getpid(), signo) so it lands on an application thread (the worker itself blocks these signals). With no chained handler it unblocks first so SIG_DFL is guaranteed to exit. - SIGABRT stays synchronous: abort() resets to SIG_DFL and re-raises the instant the handler returns, so we must wait on the finalize_done futex for the flush before returning. This is the one path that can still block on a held lock -- acceptable since the process is already aborting. - Re-entry escalates instead of swallowing: if a chained handler (e.g. LLVM/comgr) restores our handler and re-raises into us, force SIG_DFL, unblock, and re-raise to guarantee termination rather than hang. --- .../source/lib/rocprofiler-sdk-tool/tool.cpp | 191 +++++++++--------- 1 file changed, 100 insertions(+), 91 deletions(-) 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 8f05de5e4ae..5bab7c061b7 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -214,15 +214,13 @@ is_handled_signal(int signum) struct signal_worker_state { - int eventfd = -1; - std::atomic finalize_done = {}; - // set to 1 when `rocprofv3_error_signal_handler` owns the signal path (re-entry guard + - // interceptor bypass). - std::atomic handling = {0}; - std::atomic_flag finalized = ATOMIC_FLAG_INIT; - int signo = {0}; - std::thread thread = {}; - int timeout_sec = {10}; + int eventfd = -1; // handler -> worker wakeup fd + std::atomic finalize_done = {}; // worker sets when flush done (futex word) + std::atomic 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 + int timeout_sec = {10}; }; static_assert(std::atomic::is_always_lock_free, @@ -250,9 +248,8 @@ get_signal_worker(bool reset = false) if(reset) { - // Post-fork child: parent's eventfd duplicate and std::thread are invalid. Close the fd - // first, then reset storage and value-construct so defaults live only in - // signal_worker_state (we do not run ~std::thread() on a possibly-joinable parent thread). + // Post-fork child: the parent's eventfd/thread are invalid here. Close the fd and reset + // in place (never run ~std::thread() on the parent's still-joinable thread). reset_signal_worker_state(sw); } @@ -4226,6 +4223,34 @@ wait_for_children(pid_t this_pid, pid_t this_ppid, uint64_t this_tid, std::strin } } +// 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() { @@ -4257,7 +4282,7 @@ signal_finalization_worker() if(tool::get_config().enable_process_sync) wait_peer_finished(this_pid, this_ppid); - // Wait for children to finalize before we re-raise. + // Wait for children to finalize before we terminate. wait_for_children(this_pid, this_ppid, this_tid, this_func); ROCP_INFO << fmt::format( @@ -4267,8 +4292,30 @@ signal_finalization_worker() this_tid, this_func); + // 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); + + // Terminate the process now that the flush is complete, but only when we were woken by an + // async signal. signo == 0 is normal-exit finalization (main/atexit woke us to be joined); + // SIGABRT terminates itself once its (synchronously waiting) handler returns into abort(). + if(sw.signo != 0 && sw.signo != SIGABRT) + { + const bool have_chained = static_cast(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) + { + // 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); + } } void @@ -4279,21 +4326,34 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext) auto& sw = get_signal_worker(); - // We're handling signals = assume app doesn't coordinate shutdown properly. - // All processes (parent and children) need to flush and die. - // For well-behaved apps, use --disable-signal-handlers (atexit handles everything). - - // Re-entry guard: lock-free compare-exchange (see static_assert on `handling`). + // Our handler being installed means the app doesn't coordinate shutdown. + // Well-behaved apps use --disable-signal-handlers (atexit handles everything). + // Re-entry guard: On re-entry (e.g. a chained handler like LLVM/comgr re-raises into us), + // don't swallow the signal as that would hang. Escalate: force SIG_DFL and re-raise so the + // process dies. uint32_t expected = 0; - if(!sw.handling.compare_exchange_strong(expected, 1u, std::memory_order_acquire)) return; + if(!sw.handling.compare_exchange_strong(expected, 1u, std::memory_order_acquire)) + { + 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; + } // For testing: allow quick_exit path if(signal_handler_exit) ::quick_exit(signo); - // Store signal number for the worker thread to log (serialized by eventfd write/read) + // Hand the signal number to the worker (ordered by the eventfd write/read below). sw.signo = signo; - // Wake the worker thread + // Wake the finalization worker: the flush is NOT async-signal-safe, so it runs there. bool worker_notified = false; if(sw.eventfd >= 0) { @@ -4301,81 +4361,30 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext) worker_notified = (write(sw.eventfd, &val, sizeof(val)) == sizeof(val)); } - // Wait for worker with bounded timeout via futex syscall instead of a std::mutex - // Wait for worker to complete finalization. - // Uses exponential backoff (5, 10, 30, 60s) with periodic logging. - if(worker_notified) - { - static constexpr auto backoff_sec = std::array{5, 10, 30, 60}; - - int attempt = 0; - - while(sw.finalize_done.load(std::memory_order_acquire) == 0) - { - auto timeout = timespec{}; - timeout.tv_sec = backoff_sec[attempt < 4 ? attempt : 3]; - - syscall(SYS_futex, &sw.finalize_done, FUTEX_WAIT, 0, &timeout, nullptr, 0); - - if(sw.finalize_done.load(std::memory_order_acquire) != 0) break; - - // Still waiting — log progress - auto elapsed = backoff_sec[0]; - for(int j = 1; j <= attempt && j < 4; j++) - elapsed += backoff_sec[j]; - - // Logging here is not async-signal-safe in the strict POSIX sense (Abseil may take - // locks). In practice this path is still safe enough for diagnostics: the handler is - // one-shot (re-entry guard above), SA_NODEFER is not set on our installed action so the - // same thread cannot be re-entered for another signal while we're here, and other - // threads may log concurrently (Abseil is thread-safe). We keep this message so users - // see progress during long flushes. - ROCP_WARNING << fmt::format( - "signalhandler: waiting for worker thread to flush profiling data: ~{}s elapsed...", - elapsed); - attempt++; - } - } - - // Restore original handler. sigaction() is async-signal-safe per POSIX. - // sw.handling is set, so our interceptor passes through to the real sigaction. - 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 + // 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(). This is the one path that can still + // block on a lock the interrupted thread holds — acceptable since we're already aborting. + if(signo == SIGABRT) { - struct sigaction sa = {}; - sa.sa_handler = SIG_DFL; - sigemptyset(&sa.sa_mask); - get_sigaction_function()(signo, &sa, nullptr); + if(worker_notified) + while(sw.finalize_done.load(std::memory_order_acquire) == 0) + syscall(SYS_futex, &sw.finalize_done, FUTEX_WAIT, 0, nullptr, nullptr, 0); + return; } - // Unblock and re-raise. The restored handler (or SIG_DFL) runs in a fresh delivery. - sigset_t unblock{}; - sigemptyset(&unblock); - sigaddset(&unblock, signo); - pthread_sigmask(SIG_UNBLOCK, &unblock, nullptr); - raise(signo); - - // If the chained handler returned (didn't kill the process), force exit. + // 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) { - struct sigaction sa = {}; - sa.sa_handler = SIG_DFL; - sigemptyset(&sa.sa_mask); - get_sigaction_function()(signo, &sa, nullptr); + // 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); } - raise(signo); } int From 507208d00894b2fb8cf4a7f6e2d20bbcd095f523 Mon Sep 17 00:00:00 2001 From: "Mythreya, Kuricheti" Date: Thu, 6 Aug 2026 00:11:42 +0000 Subject: [PATCH 11/18] re-raise before wait_for_children --- .../source/lib/rocprofiler-sdk-tool/tool.cpp | 30 +++++++++++-------- 1 file changed, 17 insertions(+), 13 deletions(-) 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 5bab7c061b7..bfe07bdd69f 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -4282,23 +4282,15 @@ signal_finalization_worker() if(tool::get_config().enable_process_sync) wait_peer_finished(this_pid, this_ppid); - // Wait for children to finalize before we terminate. - 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); - // 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); - // Terminate the process now that the flush is complete, but only when we were woken by an - // async signal. signo == 0 is normal-exit finalization (main/atexit woke us to be joined); - // SIGABRT terminates itself once its (synchronously waiting) handler returns into abort(). + // 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(get_chained_signals().at(sw.signo)); @@ -4316,6 +4308,18 @@ signal_finalization_worker() } kill(getpid(), sw.signo); } + + // Best-effort reap now that the app is tearing its children down. May not complete if we + // terminate first -- each child finalizes independently via its own worker; this just avoids + // leaving zombies when the app keeps running (e.g. a chained handler that returns). + 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); } void From bffa6e65edbb2467af4de9cda5a69e903407cb4c Mon Sep 17 00:00:00 2001 From: "Mythreya, Kuricheti" Date: Mon, 10 Aug 2026 18:43:11 +0000 Subject: [PATCH 12/18] [rocprofiler-sdk] harden signal finalization Build on the non-blocking handler + worker so signal re-delivery is safe and the profiler stays transparent to the app's own signal disposition: - Remove `SA_RESETHAND` so a second delivery of the same signal re-enters our guard instead of resetting to `SIG_DFL` and terminating mid-flush (was silent data loss). - Gate re-entry escalation on finalize_done: swallow repeats while the flush is in progress, force `SIG_DFL` + re-raise only after it completes -- a routine double `Ctrl+C` no longer truncates the profile, and a stuck flush stays killable. - Honor `SIG_IGN` symmetrically for `signal()` and `sigaction()` (`sigaction(SIG_IGN)` was dropped, wrongly terminating apps that deliberately ignore `Ctrl+C`). - Drop the unused signal-handler timeout knob; correctness relies on lock release and the gated escalation, not a clock. - signal-handler tests: each execute wipes its own output dir first (`sh` shim, `sh` located via `find_program`) so a truncated file from a prior local run can't fail validation; CI already starts clean. --- .../source/lib/rocprofiler-sdk-tool/tool.cpp | 81 +++++++++++-------- .../signal-handler/signal-handler-test.cpp | 1 + .../rocprofv3/signal-handler/CMakeLists.txt | 53 +++++++++++- .../rocprofv3/signal-handler/validate.py | 29 +++++++ 4 files changed, 126 insertions(+), 38 deletions(-) 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 bfe07bdd69f..9ed04c71090 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -220,7 +220,6 @@ struct signal_worker_state 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 - int timeout_sec = {10}; }; static_assert(std::atomic::is_always_lock_free, @@ -2137,7 +2136,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) { @@ -4030,8 +4035,6 @@ get_sigaction_function() bool signal_handler_exit = rocprofiler::tool::get_env("ROCPROF_INTERNAL_TEST_SIGNAL_HANDLER_VIA_EXIT", false); -int signal_handler_timeout = rocprofiler::tool::get_env("ROCPROF_SIGNAL_HANDLER_TIMEOUT", 0); - } // namespace #define ROCPROFV3_INTERNAL_API __attribute__((visibility("internal"))); @@ -4309,9 +4312,9 @@ signal_finalization_worker() kill(getpid(), sw.signo); } - // Best-effort reap now that the app is tearing its children down. May not complete if we - // terminate first -- each child finalizes independently via its own worker; this just avoids - // leaving zombies when the app keeps running (e.g. a chained handler that returns). + // 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( @@ -4332,22 +4335,29 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext) // Our handler being installed means the app doesn't coordinate shutdown. // Well-behaved apps use --disable-signal-handlers (atexit handles everything). - // Re-entry guard: On re-entry (e.g. a chained handler like LLVM/comgr re-raises into us), - // don't swallow the signal as that would hang. Escalate: force SIG_DFL and re-raise so the - // process dies. + // + // 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)) + if(sw.handling.compare_exchange_strong(expected, 1u, std::memory_order_acquire) == false) { - 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); + if(sw.finalize_done.load(std::memory_order_acquire) != 0) + { + 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; } @@ -4534,11 +4544,10 @@ rocprofv3_sigaction(int signum, if(!is_handled_signal(signum) || !act || !tool::get_config().enable_signal_handlers) return CHECK_NOTNULL(get_sigaction_function())(signum, act, oldact); - // Save the app's handler so we can restore it on signal delivery. - // Only save the first non-trivial handler registered for each signal. - // Later installs (e.g., LLVM comgr) often don't chain properly — they clobber - // the disposition on re-raise. Preserving the app's handler ensures the - // application sees its signal as-if the profiler wasn't there. + // 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) @@ -4548,13 +4557,19 @@ rocprofv3_sigaction(int signum, } else { - if(act->sa_handler != SIG_DFL && act->sa_handler != SIG_IGN) + // 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); @@ -4613,9 +4628,8 @@ rocprofv3_main(int argc, char** argv, char** envp) // Spawn finalization worker in ALL processes (parent and children). // On the bad path (signal handlers enabled), every process needs to // flush its own profiling data async-signal-safely via a worker thread. - auto& sw = get_signal_worker(); - sw.eventfd = eventfd(0, EFD_CLOEXEC); - sw.timeout_sec = signal_handler_timeout; + 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. @@ -4627,9 +4641,8 @@ rocprofv3_main(int argc, char** argv, char** envp) pthread_atfork(nullptr, nullptr, []() { // Child handler: reset stale state and spawn a fresh worker. get_signal_worker(true); - auto& child_sw = get_signal_worker(); - child_sw.eventfd = eventfd(0, EFD_CLOEXEC); - child_sw.timeout_sec = signal_handler_timeout; + auto& child_sw = get_signal_worker(); + child_sw.eventfd = eventfd(0, EFD_CLOEXEC); if(child_sw.eventfd >= 0) child_sw.thread = std::thread{signal_finalization_worker}; }); } diff --git a/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp b/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp index 87a98a3f900..62d368fc6c4 100644 --- a/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp +++ b/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp @@ -429,6 +429,7 @@ mode_bad_spawn(const char* self_path) fprintf(stderr, "Parent PID=%d: exit\n", getpid()); return 0; } + } // namespace // ============================================================================ // Main diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/CMakeLists.txt index febe1fecf18..3e4dbb8a631 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/CMakeLists.txt @@ -10,6 +10,8 @@ project( find_package(rocprofiler-sdk REQUIRED) +find_program(SH_EXECUTABLE NAMES sh REQUIRED) + string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}") @@ -39,6 +41,8 @@ foreach(PROC_TYPE ${PROCESS_TYPES}) rocprofiler_add_integration_execute_test( ${WITH_HANDLING_NAME} COMMAND + ${SH_EXECUTABLE} -c + "rm -rf '${BASE_OUT_DIR}/good-${PROC_TYPE}' && exec \"$@\"" ${SH_EXECUTABLE} ${TIMEOUT_CMD} $ --hip-trace --marker-trace --output-format json --disable-signal-handlers -d ${BASE_OUT_DIR}/good-${PROC_TYPE} -o out_%pid% -- @@ -67,10 +71,11 @@ foreach(PROC_TYPE ${PROCESS_TYPES}) rocprofiler_add_integration_execute_test( ${NO_HANDLING_NAME} COMMAND - ${TIMEOUT_CMD} $ --hip-trace - --marker-trace --output-format json -d ${BASE_OUT_DIR}/bad-${PROC_TYPE} -o - out_%pid% -- $ --${PROC_TYPE} - --no-app-signal-handler + ${SH_EXECUTABLE} -c "rm -rf '${BASE_OUT_DIR}/bad-${PROC_TYPE}' && exec \"$@\"" + ${SH_EXECUTABLE} ${TIMEOUT_CMD} $ + --hip-trace --marker-trace --output-format json -d + ${BASE_OUT_DIR}/bad-${PROC_TYPE} -o out_%pid% -- + $ --${PROC_TYPE} --no-app-signal-handler DEPENDS signal-handler-test TIMEOUT ${CTEST_TIMEOUT} LABELS "integration-tests;signal-handler" @@ -89,3 +94,43 @@ foreach(PROC_TYPE ${PROCESS_TYPES}) ARGS --output-dir ${BASE_OUT_DIR}/bad-${PROC_TYPE} --mode bad --process-type ${PROC_TYPE}) endforeach() + +# =========================================================================== +# Coordinated-shutdown case: profiler signal handlers are ACTIVE 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 +# (parent waits for children that are waiting for the parent) and the test hangs -> 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 (`--fork +# --app-signal-handler`). +# =========================================================================== +set(COORD_NAME "app-coordinated-shutdown-fork") + +rocprofiler_add_integration_execute_test( + ${COORD_NAME} + COMMAND + ${SH_EXECUTABLE} -c "rm -rf '${BASE_OUT_DIR}/coordinated-fork' && exec \"$@\"" + ${SH_EXECUTABLE} ${TIMEOUT_CMD} $ + --hip-trace --marker-trace --output-format json -d + ${BASE_OUT_DIR}/coordinated-fork -o out_%pid% -- + $ --fork --app-signal-handler + DEPENDS signal-handler-test + TIMEOUT ${CTEST_TIMEOUT} + LABELS "integration-tests;signal-handler" + PRELOAD "${PRELOAD_ENV}" + PASS_REGULAR_EXPRESSION "clean exit" + FIXTURES_SETUP ${COORD_NAME}) + +rocprofiler_add_integration_validate_test( + ${COORD_NAME} + TEST_PATHS validate.py + COPY conftest.py + CONFIG pytest.ini + TIMEOUT ${CTEST_TIMEOUT} + LABELS "integration-tests;signal-handler" + FIXTURES_REQUIRED ${COORD_NAME} + ARGS --output-dir ${BASE_OUT_DIR}/coordinated-fork --mode coordinated --process-type + fork) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/validate.py index 266f834f5bc..004f5edfef5 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/validate.py @@ -124,3 +124,32 @@ def test_bad_case_markers_flushed(output_dir, mode, process_type): f"Profiler may not have flushed marker data before signal death. " f"Files: {files}" ) + + +def test_coordinated_shutdown_no_deadlock(output_dir, mode, process_type): + """Coordinated-shutdown case: profiler signal handlers are ACTIVE 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, + otherwise it deadlocks: the parent waits for children that are waiting for the parent. + Reaching this validator at all means the execute step did not hang (a regression trips the + ctest TIMEOUT); here we additionally require the profiler flushed marker data. exit_marker + is intentionally NOT required -- finalization runs on the signal, before the app's + coordinated shutdown emits it.""" + if mode != "coordinated": + return + + 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 during coordinated shutdown. " + f"Files: {files}" + ) From bb1f44a3da3ef33e296dfd017f2354587146b562 Mon Sep 17 00:00:00 2001 From: "Mythreya, Kuricheti" Date: Wed, 12 Aug 2026 05:09:19 +0000 Subject: [PATCH 13/18] remove `ROCPROF_ROOT_PID` --- .../source/lib/rocprofiler-sdk-tool/tool.cpp | 33 +++---------------- 1 file changed, 5 insertions(+), 28 deletions(-) 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 9ed04c71090..d3968ec9637 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -141,10 +141,6 @@ rocprofv3_error_signal_handler(int signo, siginfo_t*, void*); namespace { -// Initialized once in rocprofv3_main. Used by the signal handler and finalization -// to detect child processes (fork/fork+exec). -pid_t ROCPROF_ROOT_PID = 0; - // Thread for safe cleanup output generation auto output_generation_thread = common::Synchronized>{}; @@ -4605,29 +4601,11 @@ rocprofv3_main(int argc, char** argv, char** envp) if(!get_sigaction_function()) get_sigaction_function() = (sigaction_func_t) dlsym(RTLD_NEXT, "sigaction"); - // Track root process via ROCPROF_ROOT_PID env var. - // All processes (root + children) get their own worker thread and signal handler. - // For well-behaved apps, use --disable-signal-handlers (atexit handles everything). + // 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* root_pid_env = getenv("ROCPROF_ROOT_PID"); - bool is_child = (root_pid_env != nullptr); - - if(is_child) - { - // Reset the signal worker state — the inherited state references - // the parent's eventfd/thread which are inaccessible from here. - get_signal_worker(true); - ROCPROF_ROOT_PID = static_cast(std::stol(root_pid_env)); - } - else - { - ROCPROF_ROOT_PID = getpid(); - setenv("ROCPROF_ROOT_PID", std::to_string(ROCPROF_ROOT_PID).c_str(), 0); - } - - // Spawn finalization worker in ALL processes (parent and children). - // On the bad path (signal handlers enabled), every process needs to - // flush its own profiling data async-signal-safely via a worker thread. auto& sw = get_signal_worker(); sw.eventfd = eventfd(0, EFD_CLOEXEC); if(sw.eventfd >= 0) sw.thread = std::thread{signal_finalization_worker}; @@ -4640,8 +4618,7 @@ rocprofv3_main(int argc, char** argv, char** envp) atfork_registered = true; pthread_atfork(nullptr, nullptr, []() { // Child handler: reset stale state and spawn a fresh worker. - get_signal_worker(true); - auto& child_sw = get_signal_worker(); + auto& child_sw = get_signal_worker(true); child_sw.eventfd = eventfd(0, EFD_CLOEXEC); if(child_sw.eventfd >= 0) child_sw.thread = std::thread{signal_finalization_worker}; }); From 76c423c82067fafb91d5c8b2844cc5e80fe2a8ff Mon Sep 17 00:00:00 2001 From: "Mythreya, Kuricheti" Date: Thu, 13 Aug 2026 00:15:29 +0000 Subject: [PATCH 14/18] correlation id finalization fix --- .../lib/rocprofiler-sdk/context/correlation_id.cpp | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) 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 13e4dcb457d..830362989d1 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) { From 68b6abfd3522ac2ebaae7fd522850f88c2001853 Mon Sep 17 00:00:00 2001 From: "Mythreya, Kuricheti" Date: Sun, 16 Aug 2026 23:45:11 +0000 Subject: [PATCH 15/18] [rocprofiler-sdk] guard tracing interceptors against null correlation id at finalization The non-blocking signal handler finalizes on a worker thread, so application threads can keep calling instrumented APIs while finalization runs. `correlation_service::construct()` returns nullptr once finalization has begun, but most interceptors dereferenced it unguarded and crashed (`SIGSEGV` in roctx range push and HIP under signal-triggered shutdown), surfacing intermittently in CI. Fall back to the untraced path (run the underlying API and return) when `construct()` yields nullptr, via a new `RETURN_UNTRACED_ON_NULL_CORRELATION_ID` macro. Applied in hip, marker, range_marker, ompt, rccl, rocshmem, hipfile, rocjpeg and rocdecode; the HSA paths already handled this. --- .../source/lib/rocprofiler-sdk/hip/hip.cpp | 13 ++++++------ .../lib/rocprofiler-sdk/hipfile/hipfile.cpp | 13 ++++++------ .../lib/rocprofiler-sdk/marker/marker.cpp | 13 ++++++------ .../rocprofiler-sdk/marker/range_marker.cpp | 16 +++++++------- .../source/lib/rocprofiler-sdk/ompt/ompt.cpp | 21 +++++++++++++------ .../source/lib/rocprofiler-sdk/rccl/rccl.cpp | 11 +++++----- .../rocprofiler-sdk/rocdecode/rocdecode.cpp | 13 ++++++------ .../lib/rocprofiler-sdk/rocjpeg/rocjpeg.cpp | 11 +++++----- .../lib/rocprofiler-sdk/rocshmem/rocshmem.cpp | 13 ++++++------ .../lib/rocprofiler-sdk/tracing/tracing.hpp | 13 ++++++++++++ 10 files changed, 84 insertions(+), 53 deletions(-) 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..71643b607e4 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,13 @@ hip_api_impl::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, info_type::get_table_func()); + 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..97108c5f251 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,13 @@ hipfile_api_impl::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, info_type::get_table_func()); + 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..46d565476ae 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,13 @@ roctx_api_impl::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, info_type::get_table_func()); + 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..2859503f023 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,13 @@ roctx_api_impl::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, info_type::get_table_func()); + 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 +311,8 @@ roctx_api_impl::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, info_type::get_push_table_func()); 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::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::end(ompt_data_t* data, Args... args) ompt_save_state* state = nullptr; if(data != nullptr) state = static_cast(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..a09a5a88c6c 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,12 @@ rccl_api_impl::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, info_type::get_table_func()); + 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..0b6ea544dd5 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,13 @@ rocdecode_api_impl::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, info_type::get_table_func()); + 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..de3a053983e 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,12 @@ rocjpeg_api_impl::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, info_type::get_table_func()); + 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..f247ee3da2e 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,13 @@ rocshmem_api_impl::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, info_type::get_table_func()); + 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..85a6aab11fb 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tracing/tracing.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tracing/tracing.hpp @@ -35,6 +35,19 @@ #include #include +// construct() returns nullptr once finalization has begun. Rather than dereference it, run the +// underlying API untraced and return. Use inside an interceptor functor (needs `exec`, `RetT`, +// `args`). +#define RETURN_UNTRACED_ON_NULL_CORRELATION_ID(CORRELATION_ID, TABLE_FUNC) \ + if(!(CORRELATION_ID)) \ + { \ + [[maybe_unused]] auto _untraced_ret = exec((TABLE_FUNC), std::forward(args)...); \ + if constexpr(!std::is_void::value) \ + return _untraced_ret; \ + else \ + return; \ + } + namespace rocprofiler { namespace tracing From 3e7869292498f67407ff738d24186201ddd7ddfd Mon Sep 17 00:00:00 2001 From: "Mythreya, Kuricheti" Date: Fri, 21 Aug 2026 20:26:51 +0000 Subject: [PATCH 16/18] code review --- .../source/lib/rocprofiler-sdk-tool/tool.cpp | 64 +++++++++++++------ 1 file changed, 45 insertions(+), 19 deletions(-) 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 d3968ec9637..efecbf6bb73 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -227,25 +227,17 @@ get_signal_worker(bool reset = false) static auto*& _v = common::static_object::construct(); auto& sw = *CHECK_NOTNULL(_v); - const auto reset_signal_worker_state = [](signal_worker_state& obj) { - if(obj.eventfd >= 0) - { - if(::close(obj.eventfd) != 0) - { - ROCP_WARNING << "signal worker: close(eventfd) after fork failed: " - << strerror(errno); - } - } - - std::memset(static_cast(&obj), 0, sizeof(obj)); - ::new(static_cast(&obj)) signal_worker_state{}; - }; - if(reset) { // Post-fork child: the parent's eventfd/thread are invalid here. Close the fd and reset // in place (never run ~std::thread() on the parent's still-joinable thread). - reset_signal_worker_state(sw); + if(sw.eventfd >= 0 && ::close(sw.eventfd) != 0) + { + ROCP_WARNING << "signal worker: close(eventfd) after fork failed: " << strerror(errno); + } + + std::memset(static_cast(&sw), 0, sizeof(sw)); + ::new(static_cast(&sw)) signal_worker_state{}; } return sw; @@ -4031,6 +4023,10 @@ 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"))); @@ -4372,13 +4368,40 @@ rocprofv3_error_signal_handler(int signo, siginfo_t* info, void* ucontext) } // 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(). This is the one path that can still - // block on a lock the interrupted thread holds — acceptable since we're already aborting. + // 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) + { + clock_gettime(CLOCK_MONOTONIC, &deadline); + deadline.tv_sec += signal_abort_flush_timeout_sec; + } + while(sw.finalize_done.load(std::memory_order_acquire) == 0) - syscall(SYS_futex, &sw.finalize_done, FUTEX_WAIT, 0, nullptr, nullptr, 0); + { + if(syscall(SYS_futex, + &sw.finalize_done, + FUTEX_WAIT_BITSET, + 0, + bounded ? &deadline : nullptr, + nullptr, + FUTEX_BITSET_MATCH_ANY) == -1 && + errno == ETIMEDOUT) + { + break; + } + } + } return; } @@ -4520,7 +4543,10 @@ rocprofv3_signal(int signum, sighandler_t 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); }); From 54308b36df300c32aaef3f9b16cdf8317965f200 Mon Sep 17 00:00:00 2001 From: "Mythreya, Kuricheti" Date: Mon, 24 Aug 2026 22:06:25 +0000 Subject: [PATCH 17/18] fix potential read EINTR --- .../source/lib/rocprofiler-sdk-tool/tool.cpp | 17 +++++++++++++++-- 1 file changed, 15 insertions(+), 2 deletions(-) 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 efecbf6bb73..f7806eb7816 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -4257,8 +4257,21 @@ signal_finalization_worker() sigaddset(&mask, sig); pthread_sigmask(SIG_BLOCK, &mask, nullptr); - uint64_t val = 0; - if(read(sw.eventfd, &val, sizeof(val)) < 0) return; + // 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(); auto this_ppid = getppid(); From 36bfd86889b5a5dff434a8bab49dbc86959f6007 Mon Sep 17 00:00:00 2001 From: "Mythreya, Kuricheti" Date: Fri, 28 Aug 2026 00:04:16 +0000 Subject: [PATCH 18/18] code review --- projects/rocprofiler-sdk/.cmake-format.yaml | 13 ++ .../source/lib/rocprofiler-sdk-tool/tool.cpp | 77 +++---- .../context/correlation_id.cpp | 7 +- .../source/lib/rocprofiler-sdk/hip/hip.cpp | 3 +- .../lib/rocprofiler-sdk/hipfile/hipfile.cpp | 3 +- .../lib/rocprofiler-sdk/marker/marker.cpp | 3 +- .../rocprofiler-sdk/marker/range_marker.cpp | 6 +- .../source/lib/rocprofiler-sdk/rccl/rccl.cpp | 3 +- .../rocprofiler-sdk/rocdecode/rocdecode.cpp | 3 +- .../lib/rocprofiler-sdk/rocjpeg/rocjpeg.cpp | 3 +- .../lib/rocprofiler-sdk/rocshmem/rocshmem.cpp | 3 +- .../lib/rocprofiler-sdk/tracing/tracing.hpp | 12 +- .../tests/bin/signal-handler/CMakeLists.txt | 4 +- ...al-handler-test.cpp => signal-handler.cpp} | 62 +++-- .../tests/common/CMakeLists.txt | 41 ++++ .../rocprofv3/signal-handler/CMakeLists.txt | 212 ++++++++++-------- .../rocprofv3/signal-handler/conftest.py | 22 +- .../tests/rocprofv3/signal-handler/pytest.ini | 5 +- .../rocprofv3/signal-handler/validate.py | 72 +++--- 19 files changed, 322 insertions(+), 232 deletions(-) rename projects/rocprofiler-sdk/tests/bin/signal-handler/{signal-handler-test.cpp => signal-handler.cpp} (89%) 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 f7806eb7816..4f03cec7464 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -216,30 +216,44 @@ struct signal_worker_state 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::is_always_lock_free, "rocprofv3 signal path requires lock-free finalize_done atomic support"); auto& -get_signal_worker(bool reset = false) +get_signal_worker() { static auto*& _v = common::static_object::construct(); auto& sw = *CHECK_NOTNULL(_v); - - if(reset) - { - // Post-fork child: the parent's eventfd/thread are invalid here. Close the fd and reset - // in place (never run ~std::thread() on the parent's still-joinable thread). - if(sw.eventfd >= 0 && ::close(sw.eventfd) != 0) - { - ROCP_WARNING << "signal worker: close(eventfd) after fork failed: " << strerror(errno); - } - - std::memset(static_cast(&sw), 0, sizeof(sw)); - ::new(static_cast(&sw)) signal_worker_state{}; - } - return sw; } @@ -2247,28 +2261,8 @@ finalize_rocprofv3(std::string_view context) client_identifier = nullptr; } - // Join the worker thread if called from a non-worker context (atexit / rocprofv3_main). - if(sw.thread.joinable() && sw.thread.get_id() != std::this_thread::get_id()) - { - if(sw.eventfd >= 0) - { - // Write to unblock the worker's read() — closing the fd is UB while read() blocks - uint64_t val = 1; - if(write(sw.eventfd, &val, sizeof(val)) < 0) - { - ROCP_WARNING << "failed to write to signal worker eventfd: " << strerror(errno); - } - } - sw.thread.join(); - if(sw.eventfd >= 0) - { - if(close(sw.eventfd) < 0) - { - ROCP_WARNING << "failed to close signal worker eventfd: " << strerror(errno); - } - sw.eventfd = -1; - } - } + // Join the worker thread (from atexit / rocprofv3_main; a call from the worker itself no-ops). + sw.join(); } bool @@ -4657,7 +4651,14 @@ rocprofv3_main(int argc, char** argv, char** envp) atfork_registered = true; pthread_atfork(nullptr, nullptr, []() { // Child handler: reset stale state and spawn a fresh worker. - auto& child_sw = get_signal_worker(true); + 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(&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}; }); 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 830362989d1..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 @@ -270,12 +270,13 @@ correlation_id_finalize() if(!retirement_ctxs.empty()) { ROCP_CI_LOG_IF(INFO, ndangling > 0) - << "retired dangling correlation IDs: " << ndangling; + << fmt::format("retired dangling correlation IDs: {}", ndangling); } else { - ROCP_INFO << "retired dangling correlation IDs: " << ndangling - << " (no retirement consumer active)"; + ROCP_INFO << fmt::format( + "retired dangling correlation IDs: {} (no retirement consumer active)", + ndangling); } } }); 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 71643b607e4..ff79b185a83 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.cpp @@ -235,7 +235,8 @@ hip_api_impl::functor(Args... args) 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, info_type::get_table_func()); + RETURN_UNTRACED_ON_NULL_CORRELATION_ID( + corr_id, RetT, exec(info_type::get_table_func(), std::forward(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/hipfile/hipfile.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hipfile/hipfile.cpp index 97108c5f251..46dc76d5c1a 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hipfile/hipfile.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hipfile/hipfile.cpp @@ -246,7 +246,8 @@ hipfile_api_impl::functor(Args... args) 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, info_type::get_table_func()); + RETURN_UNTRACED_ON_NULL_CORRELATION_ID( + corr_id, RetT, exec(info_type::get_table_func(), std::forward(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/marker/marker.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/marker.cpp index 46d565476ae..19f8251973d 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/marker.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/marker.cpp @@ -171,7 +171,8 @@ roctx_api_impl::functor(Args... args) 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, info_type::get_table_func()); + RETURN_UNTRACED_ON_NULL_CORRELATION_ID( + corr_id, RetT, exec(info_type::get_table_func(), std::forward(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/marker/range_marker.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/range_marker.cpp index 2859503f023..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 @@ -197,7 +197,8 @@ roctx_api_impl::functor(Args... args) 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, info_type::get_table_func()); + RETURN_UNTRACED_ON_NULL_CORRELATION_ID( + corr_id, RetT, exec(info_type::get_table_func(), std::forward(args)...)); auto internal_corr_id = corr_id->internal; auto ancestor_corr_id = corr_id->ancestor; @@ -312,7 +313,8 @@ roctx_api_impl::push_functor(Args... args) auto*& corr_id = range_data.corr_id; corr_id = tracing::correlation_service::construct(ref_count); - RETURN_UNTRACED_ON_NULL_CORRELATION_ID(corr_id, info_type::get_push_table_func()); + RETURN_UNTRACED_ON_NULL_CORRELATION_ID( + corr_id, RetT, exec(info_type::get_push_table_func(), std::forward(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/rccl/rccl.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rccl/rccl.cpp index a09a5a88c6c..56e385129eb 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rccl/rccl.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rccl/rccl.cpp @@ -189,7 +189,8 @@ rccl_api_impl::functor(Args... args) 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, info_type::get_table_func()); + RETURN_UNTRACED_ON_NULL_CORRELATION_ID( + corr_id, RetT, exec(info_type::get_table_func(), std::forward(args)...)); auto internal_corr_id = corr_id->internal; auto ancestor_corr_id = corr_id->internal; 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 0b6ea544dd5..fe9e0f79a2d 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocdecode/rocdecode.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocdecode/rocdecode.cpp @@ -207,7 +207,8 @@ rocdecode_api_impl::functor(Args... args) 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, info_type::get_table_func()); + RETURN_UNTRACED_ON_NULL_CORRELATION_ID( + corr_id, RetT, exec(info_type::get_table_func(), std::forward(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/rocjpeg/rocjpeg.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocjpeg/rocjpeg.cpp index de3a053983e..aa34f91b458 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocjpeg/rocjpeg.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocjpeg/rocjpeg.cpp @@ -188,7 +188,8 @@ rocjpeg_api_impl::functor(Args... args) 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, info_type::get_table_func()); + RETURN_UNTRACED_ON_NULL_CORRELATION_ID( + corr_id, RetT, exec(info_type::get_table_func(), std::forward(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/rocshmem/rocshmem.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocshmem/rocshmem.cpp index f247ee3da2e..deea24d672d 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocshmem/rocshmem.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/rocshmem/rocshmem.cpp @@ -181,7 +181,8 @@ rocshmem_api_impl::functor(Args... args) 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, info_type::get_table_func()); + RETURN_UNTRACED_ON_NULL_CORRELATION_ID( + corr_id, RetT, exec(info_type::get_table_func(), std::forward(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/tracing/tracing.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tracing/tracing.hpp index 85a6aab11fb..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,16 +33,16 @@ #include #include +#include #include // construct() returns nullptr once finalization has begun. Rather than dereference it, run the -// underlying API untraced and return. Use inside an interceptor functor (needs `exec`, `RetT`, -// `args`). -#define RETURN_UNTRACED_ON_NULL_CORRELATION_ID(CORRELATION_ID, TABLE_FUNC) \ - if(!(CORRELATION_ID)) \ +// 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 = exec((TABLE_FUNC), std::forward(args)...); \ - if constexpr(!std::is_void::value) \ + [[maybe_unused]] auto _untraced_ret = (UNTRACED_CALL); \ + if constexpr(!std::is_void::value) \ return _untraced_ret; \ else \ return; \ diff --git a/projects/rocprofiler-sdk/tests/bin/signal-handler/CMakeLists.txt b/projects/rocprofiler-sdk/tests/bin/signal-handler/CMakeLists.txt index d05af88f96f..092011603e9 100644 --- a/projects/rocprofiler-sdk/tests/bin/signal-handler/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/bin/signal-handler/CMakeLists.txt @@ -32,9 +32,9 @@ set(CMAKE_HIP_STANDARD 17) set(CMAKE_HIP_EXTENSIONS OFF) set(CMAKE_HIP_STANDARD_REQUIRED ON) -set_source_files_properties(signal-handler-test.cpp PROPERTIES LANGUAGE HIP) +set_source_files_properties(signal-handler.cpp PROPERTIES LANGUAGE HIP) add_executable(signal-handler-test) -target_sources(signal-handler-test PRIVATE signal-handler-test.cpp) +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) diff --git a/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp b/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler.cpp similarity index 89% rename from projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp rename to projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler.cpp index 62d368fc6c4..e877ae17335 100644 --- a/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler-test.cpp +++ b/projects/rocprofiler-sdk/tests/bin/signal-handler/signal-handler.cpp @@ -43,22 +43,31 @@ #include #include -#include -#include -#include #include #include #include #include +#include +#include +#include #include #include -#include namespace { std::atomic g_shutdown{false}; +template +std::string +str_join(Args&&... _args) +{ + auto _ss = std::stringstream{}; + ((_ss << std::forward(_args)), ...); + return _ss.str(); +} + +template void emit_roctx_marker(const char* fmt, ...) { @@ -112,9 +121,7 @@ sigint_handler(int) void run_kernels(const char* label) { - char range_name[128]; - snprintf(range_name, sizeof(range_name), "%s_pid_%d", label, getpid()); - roctxRangePush(range_name); + roctxRangePush(str_join(label, "_pid_", getpid()).c_str()); float* d_buf = nullptr; HIP_CHECK(hipMalloc(&d_buf, 1024 * sizeof(float))); @@ -122,10 +129,7 @@ run_kernels(const char* label) int iter = 0; while(!g_shutdown.load(std::memory_order_relaxed)) { - char iter_name[128]; - snprintf(iter_name, sizeof(iter_name), "%s_iter_%d", label, iter); - - roctxRangePush(iter_name); + roctxRangePush(str_join(label, "_iter_", iter).c_str()); test_kernel<<<4, 256>>>(d_buf, 1024); roctxRangePop(); @@ -151,16 +155,12 @@ fork_child_worker(int id, int pipe_rd) float* d_buf = nullptr; HIP_CHECK(hipMalloc(&d_buf, 1024 * sizeof(float))); - char range_name[128]; - snprintf(range_name, sizeof(range_name), "child_%d_pid_%d", id, getpid()); - roctxRangePush(range_name); + roctxRangePush(str_join("child_", id, "_pid_", getpid()).c_str()); int iter = 0; while(true) { - char iter_name[64]; - snprintf(iter_name, sizeof(iter_name), "child_%d_iter_%d", id, iter); - roctxRangePush(iter_name); + roctxRangePush(str_join("child_", id, "_iter_", iter).c_str()); test_kernel<<<4, 256>>>(d_buf, 1024); @@ -353,9 +353,7 @@ mode_bad_fork() pid_t pid = fork(); if(pid == 0) { - char label[32]; - snprintf(label, sizeof(label), "child_%d", i); - run_kernels(label); + run_kernels(str_join("child_", i).c_str()); emit_roctx_marker("exit_marker child fork ppid:%d pid:%d", getppid(), getpid()); exit(0); } @@ -443,16 +441,16 @@ main(int argc, char** argv) for(int i = 1; i < argc; i++) { - if(strcmp(argv[i], "--single-process") == 0 || strcmp(argv[i], "--fork") == 0 || - strcmp(argv[i], "--fork-exec") == 0 || strcmp(argv[i], "--spawn") == 0) + 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(strcmp(argv[i], "--app-signal-handler") == 0) + else if(std::strcmp(argv[i], "--app-signal-handler") == 0) { app_handles_signals = true; } - else if(strcmp(argv[i], "--no-app-signal-handler") == 0) + else if(std::strcmp(argv[i], "--no-app-signal-handler") == 0) { app_handles_signals = false; } @@ -466,17 +464,17 @@ main(int argc, char** argv) if(app_handles_signals) { - if(strcmp(mode, "--single-process") == 0) return mode_good_single_process(); - if(strcmp(mode, "--fork") == 0) return mode_good_fork(); - if(strcmp(mode, "--fork-exec") == 0) return mode_good_fork_exec(argv[0]); - if(strcmp(mode, "--spawn") == 0) return mode_good_spawn(argv[0]); + 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(strcmp(mode, "--single-process") == 0) return mode_bad_single_process(); - if(strcmp(mode, "--fork") == 0) return mode_bad_fork(); - if(strcmp(mode, "--fork-exec") == 0) return mode_bad_fork_exec(argv[0]); - if(strcmp(mode, "--spawn") == 0) return mode_bad_spawn(argv[0]); + 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); 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/signal-handler/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/CMakeLists.txt index 3e4dbb8a631..0af5beaead5 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/CMakeLists.txt @@ -10,8 +10,6 @@ project( find_package(rocprofiler-sdk REQUIRED) -find_program(SH_EXECUTABLE NAMES sh REQUIRED) - string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}") @@ -27,110 +25,142 @@ set(TIMEOUT_CMD timeout --preserve-status --signal=INT ${SIGNAL_DELAY}) # Output directory base set(BASE_OUT_DIR ${CMAKE_CURRENT_BINARY_DIR}) -# Process types to test -set(PROCESS_TYPES single-process fork fork-exec spawn) - -# =========================================================================== -# Generate all 8 test combos: {good, bad} x {single-process, fork, fork-exec, spawn} -# =========================================================================== - -foreach(PROC_TYPE ${PROCESS_TYPES}) - # --- GOOD CASE: app handles signals, profiler uses --disable-signal-handlers --- - set(WITH_HANDLING_NAME "app-with-signal-handling-${PROC_TYPE}") - - rocprofiler_add_integration_execute_test( - ${WITH_HANDLING_NAME} - COMMAND - ${SH_EXECUTABLE} -c - "rm -rf '${BASE_OUT_DIR}/good-${PROC_TYPE}' && exec \"$@\"" ${SH_EXECUTABLE} - ${TIMEOUT_CMD} $ --hip-trace - --marker-trace --output-format json --disable-signal-handlers -d - ${BASE_OUT_DIR}/good-${PROC_TYPE} -o out_%pid% -- - $ --${PROC_TYPE} --app-signal-handler - DEPENDS signal-handler-test - TIMEOUT ${CTEST_TIMEOUT} - LABELS "integration-tests;signal-handler" - PRELOAD "${PRELOAD_ENV}" - PASS_REGULAR_EXPRESSION "clean exit" - FIXTURES_SETUP ${WITH_HANDLING_NAME}) - - rocprofiler_add_integration_validate_test( - ${WITH_HANDLING_NAME} - TEST_PATHS validate.py - COPY conftest.py - CONFIG pytest.ini - TIMEOUT 30 +# cmake-format: off +# add_signal_handler_test( [keyword args]) +# +# : scenario name; ctest is registered as "app--" +# PROCESS_TYPE : single-process | fork | fork-exec | spawn +# VALIDATE_EXPECT : 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 : execute PASS_REGULAR_EXPRESSION +# APP_EXTRA_FLAGS : appended to the signal-handler-test command +# (e.g. --app-signal-handler / --no-app-signal-handler) +# ROCPROFV3_EXTRA_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 -clean) +# before execute runs (FIXTURES_REQUIRED -clean), and +# validate runs after execute (FIXTURES_REQUIRED ). +# 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_REQUIRED ${WITH_HANDLING_NAME} - ARGS --output-dir ${BASE_OUT_DIR}/good-${PROC_TYPE} --mode good --process-type - ${PROC_TYPE}) - - # --- BAD CASE: app does NOT handle signals, profiler handles them --- - set(NO_HANDLING_NAME "app-without-signal-handling-${PROC_TYPE}") + FIXTURES_SETUP ${_name}-clean) rocprofiler_add_integration_execute_test( - ${NO_HANDLING_NAME} + ${_name} COMMAND - ${SH_EXECUTABLE} -c "rm -rf '${BASE_OUT_DIR}/bad-${PROC_TYPE}' && exec \"$@\"" - ${SH_EXECUTABLE} ${TIMEOUT_CMD} $ - --hip-trace --marker-trace --output-format json -d - ${BASE_OUT_DIR}/bad-${PROC_TYPE} -o out_%pid% -- - $ --${PROC_TYPE} --no-app-signal-handler + ${TIMEOUT_CMD} $ --hip-trace + --marker-trace --output-format json ${arg_ROCPROFV3_EXTRA_FLAGS} -d ${_outdir} + -o out_%pid% -- $ --${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 "tool finalization" - FIXTURES_SETUP ${NO_HANDLING_NAME}) + PASS_REGULAR_EXPRESSION "${arg_PASS_REGEX}" + FIXTURES_REQUIRED ${_name}-clean + FIXTURES_SETUP ${_name}) rocprofiler_add_integration_validate_test( - ${NO_HANDLING_NAME} + ${_name} TEST_PATHS validate.py COPY conftest.py CONFIG pytest.ini TIMEOUT ${CTEST_TIMEOUT} LABELS "integration-tests;signal-handler" - FIXTURES_REQUIRED ${NO_HANDLING_NAME} - ARGS --output-dir ${BASE_OUT_DIR}/bad-${PROC_TYPE} --mode bad --process-type - ${PROC_TYPE}) -endforeach() - -# =========================================================================== -# Coordinated-shutdown case: profiler signal handlers are ACTIVE while the app runs its -# own coordinated multi-process shutdown (the sglang/vLLM tensor-parallel pattern). Worker + 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 -# (parent waits for children that are waiting for the parent) and the test hangs -> ctest -# TIMEOUT. +# 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 (`--fork -# --app-signal-handler`). -# =========================================================================== -set(COORD_NAME "app-coordinated-shutdown-fork") - -rocprofiler_add_integration_execute_test( - ${COORD_NAME} - COMMAND - ${SH_EXECUTABLE} -c "rm -rf '${BASE_OUT_DIR}/coordinated-fork' && exec \"$@\"" - ${SH_EXECUTABLE} ${TIMEOUT_CMD} $ - --hip-trace --marker-trace --output-format json -d - ${BASE_OUT_DIR}/coordinated-fork -o out_%pid% -- - $ --fork --app-signal-handler - DEPENDS signal-handler-test - TIMEOUT ${CTEST_TIMEOUT} - LABELS "integration-tests;signal-handler" - PRELOAD "${PRELOAD_ENV}" - PASS_REGULAR_EXPRESSION "clean exit" - FIXTURES_SETUP ${COORD_NAME}) - -rocprofiler_add_integration_validate_test( - ${COORD_NAME} - TEST_PATHS validate.py - COPY conftest.py - CONFIG pytest.ini - TIMEOUT ${CTEST_TIMEOUT} - LABELS "integration-tests;signal-handler" - FIXTURES_REQUIRED ${COORD_NAME} - ARGS --output-dir ${BASE_OUT_DIR}/coordinated-fork --mode coordinated --process-type - fork) +# 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 index f2b569e611b..e085ef5cb50 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/conftest.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/conftest.py @@ -24,12 +24,20 @@ 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("--mode", action="store", help="Test mode: good or bad") + 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") @@ -37,15 +45,17 @@ def pytest_addoption(parser): def output_dir(request): val = request.config.getoption("--output-dir") if val is None: - pytest.skip("--output-dir not provided") + pytest.fail("--output-dir not provided") return val @pytest.fixture -def mode(request): - val = request.config.getoption("--mode") +def expect(request): + val = request.config.getoption("--expect") if val is None: - pytest.skip("--mode not provided") + 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 @@ -53,5 +63,5 @@ def mode(request): def process_type(request): val = request.config.getoption("--process-type") if val is None: - pytest.skip("--process-type not provided") + 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 index 8fce174e45d..8bf72b0989b 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/pytest.ini +++ b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/pytest.ini @@ -1,2 +1,5 @@ + [pytest] -addopts = -rA --tb=short +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 index 004f5edfef5..fe355480ea4 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/signal-handler/validate.py @@ -25,17 +25,25 @@ """ Validation for signal handler integration tests. -Both good and bad cases must produce VALID, COMPLETE JSON output. -The profiler intercepts signals to flush cleanly — truncated JSON is a bug. - -Good case: exit_marker present in JSON (proves clean atexit finalization). -Bad case: valid JSON with >10 marker entries (proves profiler flushed before death). +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.""" @@ -63,13 +71,13 @@ def count_markers_in_json(data): return count -def test_output_files_exist(output_dir, mode, process_type): +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, mode, process_type): +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}" @@ -85,10 +93,11 @@ def test_json_is_valid(output_dir, mode, process_type): ) -def test_good_case_exit_marker(output_dir, mode, process_type): - """Good case: mode-specific exit_marker must be present.""" - if mode != "good": - return +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 @@ -106,10 +115,14 @@ def test_good_case_exit_marker(output_dir, mode, process_type): ) -def test_bad_case_markers_flushed(output_dir, mode, process_type): - """Bad case: JSON must contain >10 marker entries (proves profiler flushed).""" - if mode != "bad": - return +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 @@ -124,32 +137,3 @@ def test_bad_case_markers_flushed(output_dir, mode, process_type): f"Profiler may not have flushed marker data before signal death. " f"Files: {files}" ) - - -def test_coordinated_shutdown_no_deadlock(output_dir, mode, process_type): - """Coordinated-shutdown case: profiler signal handlers are ACTIVE 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, - otherwise it deadlocks: the parent waits for children that are waiting for the parent. - Reaching this validator at all means the execute step did not hang (a regression trips the - ctest TIMEOUT); here we additionally require the profiler flushed marker data. exit_marker - is intentionally NOT required -- finalization runs on the signal, before the app's - coordinated shutdown emits it.""" - if mode != "coordinated": - return - - 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 during coordinated shutdown. " - f"Files: {files}" - )