From 5c77a27dfe0186ec8ce6e7fa6cb37a18ffa3734a Mon Sep 17 00:00:00 2001 From: Erik Schultheis Date: Sat, 9 May 2026 00:16:13 +0200 Subject: [PATCH 1/4] supervisor exits if no seccomp_notify support --- csrc/supervisor.cpp | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) diff --git a/csrc/supervisor.cpp b/csrc/supervisor.cpp index 385546e..79e09b9 100644 --- a/csrc/supervisor.cpp +++ b/csrc/supervisor.cpp @@ -21,6 +21,8 @@ #define dbgprint(...) #endif +extern bool supports_seccomp_notify(); + struct Config { uintptr_t sensitive_lo; uintptr_t sensitive_hi; @@ -161,7 +163,14 @@ int supervisor_main(int sock_fd) { if (prctl(PR_SET_DUMPABLE, 0) < 0) throw std::system_error(errno, std::system_category(), "prctl(PR_SET_DUMPABLE)"); - prctl(PR_SET_PDEATHSIG, SIGTERM); + if (prctl(PR_SET_PDEATHSIG, SIGTERM) < 0) + fprintf(stderr, "supervisor: PR_SET_PDEATHSIG failed: %s\n", strerror(errno)); + + if (!supports_seccomp_notify()) { + close(sock_fd); + return 0; // expected, silent + } + Config cfg; int unotify_fd = recv_setup(sock_fd, cfg); From 27587b66f19f4796aae0f88b6f83d43b7cdf2535 Mon Sep 17 00:00:00 2001 From: Erik Schultheis Date: Sat, 9 May 2026 19:53:20 +0200 Subject: [PATCH 2/4] undo mprotect on exception so shutdown works --- csrc/manager.cpp | 26 ++++++++++++++++---------- 1 file changed, 16 insertions(+), 10 deletions(-) diff --git a/csrc/manager.cpp b/csrc/manager.cpp index 10ddb63..c23dba9 100644 --- a/csrc/manager.cpp +++ b/csrc/manager.cpp @@ -397,19 +397,25 @@ nb::callable BenchmarkManager::initial_kernel_setup(double& time_estimate, const PROTECT_RANGE(lo, hi-lo, PROT_NONE); setup_seccomp(sock, install_notify, lo, hi); - nb::callable kernel = kernel_from_qualname(qualname); - CUDA_CHECK(cudaDeviceSynchronize()); - kernel(*call_args); // trigger JIT compile + try { + nb::callable kernel = kernel_from_qualname(qualname); + CUDA_CHECK(cudaDeviceSynchronize()); + kernel(*call_args); // trigger JIT compile - time_estimate = run_warmup_loop(kernel, call_args, stream, - cc_memory, l2_clear_size, discard_cache, - warmup_seconds); + time_estimate = run_warmup_loop(kernel, call_args, stream, + cc_memory, l2_clear_size, discard_cache, + warmup_seconds); - PROTECT_RANGE(lo, hi - lo, PROT_READ | PROT_WRITE); - mSupervisorSock = -1; - nvtx_pop(); + PROTECT_RANGE(lo, hi - lo, PROT_READ | PROT_WRITE); + mSupervisorSock = -1; + nvtx_pop(); - return kernel; + return kernel; + } catch (...) { + PROTECT_RANGE(lo, hi - lo, PROT_READ | PROT_WRITE); + nvtx_pop(); + throw; + } } void BenchmarkManager::randomize_before_test(int num_calls, std::mt19937& rng, cudaStream_t stream) { From 30ae4328df3071125a0e6865ee51d039158dd47d Mon Sep 17 00:00:00 2001 From: Erik Schultheis Date: Sat, 9 May 2026 21:56:15 +0200 Subject: [PATCH 3/4] fix install_notify --- csrc/manager.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/csrc/manager.cpp b/csrc/manager.cpp index c23dba9..764c088 100644 --- a/csrc/manager.cpp +++ b/csrc/manager.cpp @@ -387,7 +387,7 @@ nb::callable BenchmarkManager::initial_kernel_setup(double& time_estimate, const // snapshot all member state needed in the thread before protecting the arena const int sock = mSupervisorSock; - const bool install_notify = mSeal || supports_seccomp_notify(); + const bool install_notify = supports_seccomp_notify(); const double warmup_seconds = mWarmupSeconds; void* const cc_memory = mDeviceDummyMemory; const std::size_t l2_clear_size = mL2CacheSize; From 1b20aca0ff6cf6ca11d46d0ae49aa233db033a01 Mon Sep 17 00:00:00 2001 From: Erik Schultheis Date: Sat, 9 May 2026 22:55:10 +0200 Subject: [PATCH 4/4] bump nanobind version --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 6fdf037..dec418e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -19,7 +19,7 @@ FetchContent_Declare( nanobind QUIET GIT_REPOSITORY https://github.com/wjakob/nanobind.git - GIT_TAG v2.9.2 + GIT_TAG v2.12.0 ) find_package(CUDAToolkit REQUIRED)