supervisor exits if no seccomp_notify support

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)

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;
@@ -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) {

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);