fix(hrr): make the capture writer safe across fork and in the crash path

projects/hrr/DESIGN.md

@@ -104,8 +104,11 @@ processes. Every HIP-owning process writes an independent sub-archive at
 two live writers into one `events.bin` without needing an advisory lock.
 
 `writer::open()` (`hip_capture_writer.cpp`) always selects the current process's
-PID directory. A `fork()` child re-opens from the base dir in `atfork_child`, so
-the child naturally switches to its own `pid-<childpid>/` sub-archive. The root
+PID directory. A `fork()` child re-opens from the base dir on its first record,
+blob or code object after the fork, so it switches to its own `pid-<childpid>/`
+sub-archive. `atfork_child` only drops the parent's events fd and paths: the
+child of a multithreaded process may make only async-signal-safe calls until it
+execs, and a child that execs or exits without recording leaves no archive. The root
 `manifest.json` is a common aggregate index with the schema fields
 `version`, `capture_mode`, `owner_pid`, and `processes[]`; each process rewrites
 it best-effort on clean shutdown by scanning existing `pid-*/manifest.json`

projects/hrr/tests/hrr_test_common.hh

@@ -309,6 +309,12 @@ inline void hrr_capture_direct(const std::string& direct_case,
 }
 #endif  // HRR_TEST_EXE
 
+// How many children Unit_HRR_ForkWhileRecording_Direct forks. Every other
+// child records a call before it exits, starting with the first, and the
+// roundtrip expects one archive for each of those besides the parent's.
+inline constexpr int kHrrForkWhileRecordingForks = 200;
+inline constexpr int kHrrForkWhileRecordingArchives = 1 + (kHrrForkWhileRecordingForks + 1) / 2;
+
 inline std::pair<int, std::string> hrr_playback_env(
     const fs::path& cap_path,
     const std::vector<std::pair<std::string, std::string>>& env,

projects/hrr/tests/integration/CMakeLists.txt

@@ -42,6 +42,10 @@
 # fake a small file system under HIP_HRR_CAPTURE_OUTPUT when the test sets
 # HRR_TEST_FAKE_FS_TOTAL; every other call goes on to libc. Contains its own
 # hidden [.][hrr-direct] workloads. POSIX only.
+#
+# hrr_clock_hook.cc defines clock_gettime for the binary, so a workload can run
+# a hook on one thread's clock reads, as the capture writer timestamps a record.
+# POSIX only.
 
 # The workload translation units carry __global__ kernels, device built-ins
 # (blockIdx/blockDim/threadIdx), hipLaunchKernelGGL, and HIP_SYMBOL, so they must
@@ -74,6 +78,7 @@ add_executable(hrr-integration-tests
   hrr_roundtrip_test.cc
   hrr_api_matrix_test.cc
   hrr_secure_exec_test.cc
+  hrr_clock_hook.cc
 )
 set_target_properties(hrr-integration-tests PROPERTIES
   CXX_STANDARD 17

projects/hrr/tests/integration/hrr_clock_hook.cc

@@ -0,0 +1,20 @@
+/*
+ * Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+ *
+ * SPDX-License-Identifier: MIT
+ */
+
+#include "hrr_clock_hook.hh"
+
+#ifndef _WIN32
+#include <sys/syscall.h>
+#include <time.h>
+#include <unistd.h>
+
+thread_local void (*t_hrr_clock_hook)() = nullptr;
+
+extern "C" int clock_gettime(clockid_t clk, struct timespec* ts) noexcept {
+  if (auto hook = t_hrr_clock_hook) hook();
+  return static_cast<int>(syscall(SYS_clock_gettime, clk, ts));
+}
+#endif  // !_WIN32

projects/hrr/tests/integration/hrr_clock_hook.hh

@@ -0,0 +1,16 @@
+/*
+ * Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+ *
+ * SPDX-License-Identifier: MIT
+ */
+
+#pragma once
+
+// The test binary defines clock_gettime() itself, in hrr_clock_hook.cc. The
+// dynamic linker binds the HIP runtime's calls to it before libc's. The capture
+// writer calls it to timestamp each record, before it takes the buffer lock.
+//
+// A thread that sets t_hrr_clock_hook gets it called on each clock_gettime()
+// that thread makes, before the clock is read. Other threads are not affected.
+// POSIX only.
+extern thread_local void (*t_hrr_clock_hook)();

projects/hrr/tests/integration/hrr_disk_space_test.cc

@@ -194,7 +194,11 @@ extern "C" int hrr_disk_space_statvfs64(const char* path, struct statvfs* buf) {
   return fake_statvfs("statvfs64", path, buf);
 }
 
+// Defined in hrr_workload_test.cc: the fork cases there hold a thread in fsync.
+extern "C" void hrr_workload_fsync_hook(int fd);
+
 extern "C" int hrr_disk_space_fsync(int fd) {
+  hrr_workload_fsync_hook(fd);
   using Fn = int (*)(int);
   static Fn real = reinterpret_cast<Fn>(dlsym(RTLD_NEXT, "fsync"));
   if (g_fsync_hold.exchange(false)) {
@@ -265,6 +269,17 @@ void copy_to_device(void* dev, uint32_t seed, size_t len) {
   HRR_HIP_CHECK(hipMemcpy(dev, host.data(), len, hipMemcpyHostToDevice));
 }
 
+// A forked child opens its archive on its first record, not at fork. This
+// records one without touching the device, which a forked child must not use:
+// a launch configuration pushed and popped through the compiler dispatch table.
+void record_in_child() {
+  dim3 grid, block;
+  size_t shared = 0;
+  hipStream_t stream = nullptr;
+  (void)__hipPushCallConfiguration(dim3(1), dim3(1), 0, nullptr);
+  (void)__hipPopCallConfiguration(&grid, &block, &shared, &stream);
+}
+
 // Wait for a forked child for at most 60 s; -1 when it had to be killed.
 int wait_child(pid_t child) {
   const auto deadline = std::chrono::steady_clock::now() + std::chrono::seconds(60);
@@ -389,6 +404,7 @@ TEST_CASE("Unit_HRR_DiskSpace_ForkAfterStop_Direct", "[.][hrr-direct]") {
   const pid_t child = fork();
   if (child == 0) {
     g_quick_exit_pid.store(getpid());
+    record_in_child();
     std::exit(0);
   }
   REQUIRE(child > 0);
@@ -417,6 +433,8 @@ TEST_CASE("Unit_HRR_DiskSpace_RefusedForkChild_Direct", "[.][hrr-direct]") {
   const pid_t child = fork();
   if (child == 0) {
     g_quick_exit_pid.store(getpid());
+    // Its open, and so the refusal, comes with its first record.
+    record_in_child();
     std::exit(0);
   }
   REQUIRE(child > 0);