diff --git a/CHANGELOG.md b/CHANGELOG.md index 9deb34d..b70f7a9 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -168,6 +168,43 @@ types are additive. `--agent-drain-ms`, `--agent-cursor` and `--all-sessions` control it, and the launcher sends upload-complete only after a clean agent drain. `--upload` is scoped to a per-run folder rather than the shared `--output` directory. +- **Linux injection lifecycle (`gpufl trace`).** + - **Warmup.** `--warmup` now sets `GPUFL_INJECT_WARMUP_MS` instead of + `GPUFL_INJECT_INIT_DELAY_MS`. Nothing waits for a warmup, so a warmup run + is reported as an unqualified Trace capture. The initialization deadline + includes the warmup once, so a long healthy warmup is not killed. A + process that exits during the warmup stops the worker before CUPTI starts + (`exit_before_capture`), and normal teardown runs. A new library launched with + a long `GPUFL_INJECT_INIT_DELAY_MS` and without the launcher's + `GPUFL_INJECT_PROTOCOL` (an older launcher, or a direct injection) records + `legacy_launcher_delay`. The launcher and the library ship together in one + wheel. `GPUFL_INJECT_SYNC_WAIT_MS` is no longer read: one deadline covers + every wait. + - **Status file.** The launcher gives each child a status directory + (`GPUFL_INJECT_STATUS_DIR`, `GPUFL_INJECT_STATUS_NONCE`). The library + writes `status-.json` there without the logger, and the launcher + prints the result as `capture readiness: ...`. If the read after exit + does not confirm the last accepted record (missing, malformed, foreign or + older), the summary says so and shows that record as history. The status + is diagnostics, not qualification authority. Writing it can hold the gate + on a slow filesystem past the nominal deadline. + - **Watchdog.** In every mode, including a natural exit, the launcher kills + (SIGKILL) a child whose capture initialization outlives the published + deadline plus `GPUFL_INJECT_WATCHDOG_GRACE_MS` (default 5000). It does not + bound the application's own run time. + - **Watchdog limits.** + - It covers only a paired library after its first status record. It does + not cover an older library, a manual host's own `cuInit` before its + handshake, a failed status write, or a wedge before the first record. + - SIGKILL plus a blocking reap is not a bound on a process stuck in the + kernel. + - **Exit status 86.** If an initializer is still inside CUDA/CUPTI at exit + after the deadline, the process now ends with status 86. Without this, the + remaining exit handlers and destructors would run under that initializer. + Application exit handlers registered after the first CUDA call still run + before that check: that overlap is not controlled. + - **Driver API hosts.** A host that calls `cuInit` itself, without the + manual handshake, is reported as `unsupported_first_touch`. ### Deprecated @@ -176,6 +213,45 @@ types are additive. ### Fixed +- **Linux `gpufl trace` capture readiness.** The injected library started + CUPTI on a worker thread while the application went on into context and + module setup. On an L4, whole Trace sessions received every kernel and + memcpy record with zero timestamps. Nothing checked whether capture was + ready. Whether this race is the CUPTI-internal cause is not established, + and real-GPU requalification validates the fix. + - **The first call waits.** The application's first supported CUDA runtime + call now triggers driver initialization itself (through the real + `cudaGetDeviceCount`) and waits, within one deadline, until the driver, + the Trace subscription and every required activity kind are ready. A kind + counts only when its enable, and its re-enable after the engine starts, + returned success: a fallback result CUPTI logs as benign is not an enable. + A required kind lost after readiness adds `trace_facility_lost`. + `GPUFL_INJECT_LAUNCH_WAIT_MS` (1-600000, default 15000) sets that + deadline. The first-touch calls covered are: + - `cudaMalloc`, `cudaMallocManaged`, `cudaMallocHost`, `cudaHostAlloc`, + `cudaFree`; + - `cudaSetDevice`, `cudaGetDeviceCount`, `cudaGetDeviceProperties`, + `cudaMemGetInfo`; + - `cudaEventCreate*`, `cudaStreamCreate*`; + - the driver's `cuMemAlloc_v2` and `cuModuleLoadData*`; + - the existing launch, sync and copy wrappers. + - **New session rows.** Each session gets a `capture_readiness` row + (`ready`, `failed` or `pending`, with the reason) and a `capture_coverage` + row. Coverage is `no_known_loss` or `incomplete` with reasons, such as + `unsupported_first_touch` for driver init through an unwrapped route, + `manual_protocol_violation`, `warmup_unqualified` and + `kernel_timestamps_unavailable`. `no_known_loss` is not a completeness + proof: a kernel that produced no record cannot be counted. + - **Timeouts.** A timeout lets the application continue. The run then stays + `failed` even when initialization finishes later. + - **Manual handshake.** Hosts that use the manual handshake + (`GPUFL_INJECT_MANUAL_INIT=1`: Numba, and the backend's Rust and Python + hosts) get one synchronous initializer. `GpuFlightInitializeInjectionAfterCuda` + returns 1 only when capture is ready. +- **Activity records without timing are reported as such.** Kernel and memcpy + records that arrive with `start == end == 0` are counted, and the capability + reason says `kernel_timestamps_unavailable` (or `_invalid`, or `_partial`), + with the counts, instead of "no kernel activity record arrived". - **PM/PC sample timestamps under `gpufl trace`.** The deferred engine start (the path Windows injection always takes, since gpufl initializes before the target creates a CUDA context) built its `EngineContext` without the diff --git a/CMakeLists.txt b/CMakeLists.txt index 888e438..96a6b0f 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -804,6 +804,12 @@ if(BUILD_GPUFL_INJECT AND ((UNIX AND NOT APPLE) OR WIN32)) add_library(gpufl_inject SHARED include/gpufl/inject/inject_entry.cpp ) + if(UNIX AND NOT APPLE) + target_sources(gpufl_inject PRIVATE + include/gpufl/inject/capture_gate.cpp + include/gpufl/inject/capture_gate_linux.cpp + ) + endif() if(UNIX AND NOT APPLE AND GPUFL_HAS_CUDA) # Compile CUDA boundary wrappers against the toolkit's official ABI. # inject_entry.cpp keeps its header-free fallback for no-CUDA builds. diff --git a/daemon/launcher/CMakeLists.txt b/daemon/launcher/CMakeLists.txt index 6779560..74907db 100644 --- a/daemon/launcher/CMakeLists.txt +++ b/daemon/launcher/CMakeLists.txt @@ -21,6 +21,7 @@ endif() add_executable(gpufl_launcher main.cpp agent_launcher.cpp + bootstrap_status.cpp cli_help.cpp cli_parse.cpp cli_parse_internal.cpp diff --git a/daemon/launcher/bootstrap_status.cpp b/daemon/launcher/bootstrap_status.cpp new file mode 100644 index 0000000..8fff56a --- /dev/null +++ b/daemon/launcher/bootstrap_status.cpp @@ -0,0 +1,220 @@ +#include "bootstrap_status.hpp" + +#include +#include +#include +#include + +#include "gpufl/core/json/json.hpp" + +#if !defined(_WIN32) +#include +#endif + +namespace gpufl::launcher { +namespace { + +constexpr std::size_t kMaxStatusBytes = 4096; + +BootstrapStatus Invalid(std::string why) { + BootstrapStatus status; + status.error = std::move(why); + return status; +} + +} // namespace + +std::filesystem::path BootstrapStatusPath(const std::filesystem::path& dir, + const long pid) { + return dir / ("status-" + std::to_string(pid) + ".json"); +} + +BootstrapStatus ReadBootstrapStatus(const std::filesystem::path& dir, + const long pid, const std::string& nonce) { + std::ifstream in(BootstrapStatusPath(dir, pid), std::ios::binary); + if (!in) return Invalid("missing"); + std::string text(kMaxStatusBytes + 1, '\0'); + in.read(text.data(), static_cast(text.size())); + text.resize(static_cast(in.gcount())); + if (text.size() > kMaxStatusBytes) return Invalid("oversized"); + + const json::JsonValue doc = json::parseJson(text); + if (!doc.is_object()) return Invalid("malformed"); + if (doc.value("schema", 0) != 1) return Invalid("schema"); + if (doc.value("nonce", "") != nonce || nonce.empty()) { + return Invalid("foreign_nonce"); + } + if (doc.value("pid", 0) != pid) return Invalid("foreign_pid"); + + // Every field of schema 1 must be there with its type, and the record + // must describe a state the gate can publish (the checks below). This is + // an operational diagnostic: a valid record is a consistent report, not + // evidence that capture worked. + const auto has = [&doc](const char* key, const json::JsonValue::Type type) { + return doc.contains(key) && doc[key].type() == type; + }; + using Type = json::JsonValue::Type; + for (const char* key : {"seq", "deadline_monotonic_ns", "callbacks_seen", + "handshake_reentries"}) { + if (!has(key, Type::Int)) return Invalid("malformed"); + } + if (!has("wait_config_invalid", Type::Bool)) return Invalid("malformed"); + for (const char* key : {"lifecycle", "owner", "readiness", "reason", "protocol"}) { + if (!has(key, Type::String)) return Invalid("malformed"); + } + for (const char* key : {"trigger", "worker", "handshake"}) { + if (!has(key, Type::Object)) return Invalid("malformed"); + for (const char* field : {"expected", "returned", "ok"}) { + if (!doc[key].contains(field) || !doc[key][field].is_bool()) { + return Invalid("malformed"); + } + } + } + if (!has("coverage_notes", Type::Array)) return Invalid("malformed"); + + BootstrapStatus status; + status.seq = doc["seq"].get_int(); + status.lifecycle = doc["lifecycle"].get_string(); + status.owner = doc["owner"].get_string(); + status.readiness = doc["readiness"].get_string(); + status.reason = doc["reason"].get_string(); + status.deadline_monotonic_ns = doc["deadline_monotonic_ns"].get_int(); + status.protocol = doc["protocol"].get_string(); + for (const json::JsonValue& note : doc["coverage_notes"].get_array()) { + if (!note.is_string()) return Invalid("malformed"); + status.coverage_notes.insert(note.get_string()); + } + + const std::set lifecycles = {"not_requested", "driver_initializing", + "driver_ready", "capture_initializing", + "finished"}; + const std::set owners = {"none", "gate", "callback", "manual"}; + const std::set readiness = {"pending", "ready", "failed"}; + if (status.seq < 1 || !lifecycles.count(status.lifecycle) || + !owners.count(status.owner) || !readiness.count(status.readiness)) { + return Invalid("inconsistent"); + } + const bool requested = status.lifecycle != "not_requested"; + if (requested != (status.owner != "none") || + requested != (status.deadline_monotonic_ns > 0) || + (status.lifecycle == "finished" && status.readiness == "pending") || + (status.readiness == "ready" && status.lifecycle != "finished") || + (status.readiness == "failed") == status.reason.empty()) { + return Invalid("inconsistent"); + } + + // Participants: returned needs expected, ok needs returned; the owner + // decides which may be expected; Finished needs every expected one + // returned, and Ready needs every expected one ok. + struct Participant { bool expected, returned, ok; }; + const auto participant = [&doc](const char* key) { + const json::JsonValue& p = doc[key]; + return Participant{p["expected"].get_bool(), p["returned"].get_bool(), + p["ok"].get_bool()}; + }; + const Participant trigger = participant("trigger"); + const Participant worker = participant("worker"); + const Participant handshake = participant("handshake"); + bool any_expected = false; + for (const Participant& p : {trigger, worker, handshake}) { + if ((p.returned && !p.expected) || (p.ok && !p.returned)) { + return Invalid("inconsistent"); + } + any_expected = any_expected || p.expected; + if (status.lifecycle == "finished" && p.expected && !p.returned) { + return Invalid("inconsistent"); + } + if (status.readiness == "ready" && p.expected && !p.ok) { + return Invalid("inconsistent"); + } + } + const bool owner_fits = + status.owner == "none" ? !any_expected + : status.owner == "gate" ? trigger.expected && !handshake.expected + : status.owner == "callback" ? worker.expected && !trigger.expected && + !handshake.expected + : handshake.expected && !trigger.expected && + !worker.expected; + if (!owner_fits) return Invalid("inconsistent"); + status.valid = true; + return status; +} + +int64_t MonotonicNowNs() { +#if defined(_WIN32) + return std::chrono::duration_cast( + std::chrono::steady_clock::now().time_since_epoch()) + .count(); +#else + timespec ts{}; + clock_gettime(CLOCK_MONOTONIC, &ts); + return static_cast(ts.tv_sec) * 1'000'000'000LL + ts.tv_nsec; +#endif +} + +bool InitWatchdog::Observe(const BootstrapStatus& status, const int64_t now_ns) { + if (status.valid) { + if (status.seq >= last_seq_) { + last_seq_ = status.seq; + last_ = status; + seen_valid_ = true; + if (status.lifecycle == "finished") disarmed_ = true; + if (!armed_ && status.lifecycle != "not_requested" && + status.deadline_monotonic_ns > 0) { + armed_ = true; + deadline_ns_ = status.deadline_monotonic_ns; + } + } + } + if (!armed_ || disarmed_) return false; + const int64_t max = std::numeric_limits::max(); + const int64_t limit = deadline_ns_ > max - grace_ns_ ? max : deadline_ns_ + grace_ns_; + return now_ns > limit; +} + +FinalBootstrapObservation CombineFinalBootstrap(const InitWatchdog& watchdog, + const BootstrapStatus& final_read) { + FinalBootstrapObservation out; + if (!watchdog.seen_valid()) { + out.accepted = final_read; // invalid; its error says why + out.final_error = final_read.error; + return out; + } + out.accepted = watchdog.last(); + if (!final_read.valid) { + out.final_error = final_read.error; + } else if (final_read.seq < watchdog.last().seq) { + out.final_error = "stale_final_seq"; + } + return out; +} + +std::string DescribeBootstrapStatus(const BootstrapStatus& status) { + if (!status.valid) return "unavailable (" + status.error + ")"; + std::string out = status.readiness + " owner=" + status.owner + + " lifecycle=" + status.lifecycle; + if (!status.reason.empty()) out += " reason=" + status.reason; + if (!status.coverage_notes.empty()) { + out += " coverage="; + bool first = true; + for (const std::string& note : status.coverage_notes) { + if (!first) out += ','; + first = false; + out += note; + } + } + return out; +} + +std::string MakeBootstrapNonce() { + std::random_device device; + static const char* kHex = "0123456789abcdef"; + std::string out; + for (int i = 0; i < 4; ++i) { + const uint32_t word = device(); + for (int shift = 28; shift >= 0; shift -= 4) out += kHex[(word >> shift) & 0xf]; + } + return out; +} + +} // namespace gpufl::launcher diff --git a/daemon/launcher/bootstrap_status.hpp b/daemon/launcher/bootstrap_status.hpp new file mode 100644 index 0000000..0cc9884 --- /dev/null +++ b/daemon/launcher/bootstrap_status.hpp @@ -0,0 +1,91 @@ +#pragma once + +// The launcher's side of the injected library's bootstrap status file +// (gpufl/inject/capture_gate.hpp): reading it, and the initialization +// watchdog that kills a child whose capture initialization never finished. + +#include +#include +#include +#include + +namespace gpufl::launcher { + +struct BootstrapStatus { + // False: missing, unreadable, oversized, malformed, foreign or stale. Never + // read as success. + bool valid = false; + std::string error; + int64_t seq = 0; + std::string lifecycle; + std::string owner; + std::string readiness; + std::string reason; + std::set coverage_notes; + int64_t deadline_monotonic_ns = 0; + std::string protocol; +}; + +/** status-.json inside the launcher-owned per-child directory. */ +std::filesystem::path BootstrapStatusPath(const std::filesystem::path& dir, + long pid); + +/** + * Reads at most 4 KiB and requires schema 1, `nonce` and `pid`. A file that + * fails any of those is returned invalid with the reason in `error`. + */ +BootstrapStatus ReadBootstrapStatus(const std::filesystem::path& dir, long pid, + const std::string& nonce); + +/** CLOCK_MONOTONIC in ns: the clock the library publishes its deadline in. */ +int64_t MonotonicNowNs(); + +/** + * Arms on the first accepted record that carries a deadline, keeps that first + * deadline (a later value never renews the budget), and disarms for good once + * the record says every initialization participant returned. A record with a + * lower seq than one already seen is stale and ignored. + */ +class InitWatchdog { + public: + explicit InitWatchdog(int64_t grace_ns) : grace_ns_(grace_ns) {} + + /** Feeds one read; true means the child must be killed now. */ + bool Observe(const BootstrapStatus& status, int64_t now_ns); + + bool armed() const { return armed_; } + bool disarmed() const { return disarmed_; } + const BootstrapStatus& last() const { return last_; } + bool seen_valid() const { return seen_valid_; } + + private: + int64_t grace_ns_; + bool armed_ = false; + bool disarmed_ = false; + bool seen_valid_ = false; + int64_t deadline_ns_ = 0; + int64_t last_seq_ = -1; + BootstrapStatus last_; +}; + +/** + * The last accepted record plus whether the read after exit confirmed it. + * `final_error` is empty only when that read was valid and current; otherwise + * it names what went wrong (missing, malformed, foreign_*, stale_final_seq), + * and the accepted record is historical, not a fresh observation. + */ +struct FinalBootstrapObservation { + BootstrapStatus accepted; + std::string final_error; +}; + +FinalBootstrapObservation CombineFinalBootstrap(const InitWatchdog& watchdog, + const BootstrapStatus& final_read); + +/** One line for the run summary, e.g. "ready owner=gate" or "unavailable (...)". */ +std::string DescribeBootstrapStatus(const BootstrapStatus& status); + +/** A fresh 128-bit hex nonce. */ +std::string MakeBootstrapNonce(); + +} // namespace gpufl::launcher diff --git a/daemon/launcher/cli_parse.hpp b/daemon/launcher/cli_parse.hpp index b856c75..aeb2498 100644 --- a/daemon/launcher/cli_parse.hpp +++ b/daemon/launcher/cli_parse.hpp @@ -33,7 +33,7 @@ struct TraceArgs { int agent_drain_ms = 60000; // --agent-drain-ms: cap on waiting for the agent to finish uploading // Bounded window profiling (`gpufl trace` only): bound a capture of a // long-running target that never exits on its own. warmup defers the - // capture start (via GPUFL_INJECT_INIT_DELAY_MS); window then runs for a + // capture start (via GPUFL_INJECT_WARMUP_MS); window then runs for a // fixed wall-clock before the launcher stops the target. All in ms. int64_t warmup_ms = 0; // --warmup; 0 = start capturing immediately int64_t window_ms = 0; // --window; 0 = run to the target's natural exit diff --git a/daemon/launcher/trace_command.cpp b/daemon/launcher/trace_command.cpp index 11243f2..787296a 100644 --- a/daemon/launcher/trace_command.cpp +++ b/daemon/launcher/trace_command.cpp @@ -146,13 +146,35 @@ class PosixTracePlatform final : public TracePlatform { fillFromStatus(status); }; - if (opts.run_ms <= 0) { + // Capture-initialization watchdog: the injected library publishes its + // deadline in a per-child status file. It bounds that initialization + // only, never the application's own run time. + const bool watch = !opts.bootstrap_dir.empty(); + InitWatchdog watchdog(opts.watchdog_grace_ms * 1'000'000LL); + BootstrapStatus last_read; + auto watchdogFires = [&](const bool final_read = false) -> bool { + // Once initialization finished nothing can rearm it; a long run + // stops reading the file until the final read after exit. + if (!watch || (watchdog.disarmed() && !final_read)) return false; + last_read = ReadBootstrapStatus(opts.bootstrap_dir, pid, opts.bootstrap_nonce); + return watchdog.Observe(last_read, MonotonicNowNs()); + }; + auto finishBootstrap = [&]() { + if (!watch) return; + (void)watchdogFires(/*final_read=*/true); // the record after exit + // The last accepted record, plus whether this final read confirmed it. + const FinalBootstrapObservation final_obs = CombineFinalBootstrap(watchdog, last_read); + out.bootstrap = final_obs.accepted; + out.bootstrap_final_error = final_obs.final_error; + }; + + if (opts.run_ms <= 0 && !watch) { reapBlocking(); return out; } - // Bounded window: poll for a natural exit until the deadline, then - // stop the child (SIGTERM, escalating to SIGKILL after the grace). + // Poll for a natural exit (until the deadline, for a bounded window), + // then stop the child (SIGTERM, escalating to SIGKILL after the grace). constexpr auto poll = std::chrono::milliseconds(50); // reaped: -1 launcher error, 0 still running, 1 exited (out filled). auto pollReap = [&]() -> int { @@ -166,29 +188,48 @@ class PosixTracePlatform final : public TracePlatform { out.error = std::string("waitpid failed: ") + std::strerror(errno); return -1; }; - auto waitUntil = [&](std::chrono::steady_clock::time_point until) -> int { + // 1 exited, -1 error, 0 still running at `until`. + auto waitUntil = [&](const std::chrono::steady_clock::time_point until) -> int { while (std::chrono::steady_clock::now() < until) { const int r = pollReap(); - if (r != 0) return r; // -1 error or 1 exited + if (r != 0) return r; + if (watchdogFires()) { + // SIGKILL, never SIGTERM: the application may handle + // SIGTERM, and the wedged initializer must not run on. + out.init_watchdog_killed = true; + ::kill(pid, SIGKILL); + reapBlocking(); + return 1; + } std::this_thread::sleep_for(poll); } - return 0; // timed out, still running + return 0; }; + const auto unbounded = std::chrono::steady_clock::time_point::max(); const auto now = std::chrono::steady_clock::now(); - const int r = waitUntil(now + std::chrono::milliseconds(opts.run_ms)); - if (r != 0) return out; // natural exit (or error) before the deadline + const int r = waitUntil(opts.run_ms <= 0 + ? unbounded + : now + std::chrono::milliseconds(opts.run_ms)); + if (r != 0) { // natural exit, watchdog kill, or error + finishBootstrap(); + return out; + } // Deadline reached → stop the target. out.window_stopped = true; ::kill(pid, SIGTERM); const auto grace = std::chrono::steady_clock::now() + std::chrono::milliseconds(opts.stop_grace_ms); - if (waitUntil(grace) != 0) return out; // exited within the grace + if (waitUntil(grace) != 0) { // exited within the grace + finishBootstrap(); + return out; + } // Still alive → hard kill and reap. ::kill(pid, SIGKILL); reapBlocking(); + finishBootstrap(); return out; } }; diff --git a/daemon/launcher/trace_command_common.cpp b/daemon/launcher/trace_command_common.cpp index c8a9097..d9c4c44 100644 --- a/daemon/launcher/trace_command_common.cpp +++ b/daemon/launcher/trace_command_common.cpp @@ -153,6 +153,8 @@ struct SyntheticShutdownContext { bool signaled = false; int signal = 0; bool window_stopped = false; // intentional window-end stop, not a crash + // Capture initialization never finished; the launcher killed the target. + bool init_watchdog_killed = false; }; void inspectLifecycleLine(const std::string& line, SessionLifecycleInfo& info) { @@ -206,6 +208,19 @@ bool decompressGzipToFile(const fs::path& gz_path, const fs::path& out_path) { return n == 0 && err == Z_OK && out.good(); } +#if !defined(_WIN32) +// The bootstrap summary, saying so when the read after exit did not confirm +// the last accepted record: that record is then history, not a current state. +std::string describeRunBootstrap(const TraceProcessResult& run) { + std::string text = DescribeBootstrapStatus(run.bootstrap); + if (run.bootstrap.valid && !run.bootstrap_final_error.empty()) { + text += " [last accepted record; final status read: " + + run.bootstrap_final_error + "]"; + } + return text; +} +#endif + std::string syntheticShutdownLine(const std::string& session_id, const SessionLifecycleInfo& info, const SyntheticShutdownContext& context) { @@ -222,8 +237,9 @@ std::string syntheticShutdownLine(const std::string& session_id, << ",\"signaled\":" << (context.signaled ? "true" : "false") << ",\"signal\":" << context.signal << ",\"window_stopped\":" << (context.window_stopped ? "true" : "false") - << ",\"crashed\":" << (crashed ? "true" : "false") - << "}"; + << ",\"crashed\":" << (crashed ? "true" : "false"); + if (context.init_watchdog_killed) os << ",\"init_watchdog_killed\":true"; + os << "}"; return os.str(); } @@ -778,11 +794,22 @@ int runTraceCommon(const TraceArgs& args, const TracePlatform& platform) { return 2; } + // The injected library and this launcher ship together; the protocol + // variable tells the library which contract it was launched under. + if (!setEnvOrPrint(platform, env::kInjectProtocol, env::kInjectProtocolVersion)) { + return 2; + } + // Warmup defers the in-child capture start so the window skips cold-start - // (autotune / JIT / allocator growth). Reuses the inject init-delay knob. - if (args.warmup_ms > 0 && - !setEnvOrPrint(platform, env::kInjectInitDelayMs, - std::to_string(args.warmup_ms))) { + // (autotune / JIT / allocator growth). Nothing waits for it, so a warmup + // run is reported as an unqualified Trace capture. Without --warmup the + // variable is removed, so a parent shell's value cannot leak in. + if (args.warmup_ms > 0) { + if (!setEnvOrPrint(platform, env::kInjectWarmupMs, + std::to_string(args.warmup_ms))) { + return 2; + } + } else if (!unsetEnvOrPrint(platform, env::kInjectWarmupMs)) { return 2; } @@ -928,17 +955,79 @@ int runTraceCommon(const TraceArgs& args, const TracePlatform& platform) { } } + RunOptions pass_opts = run_opts; +#if !defined(_WIN32) + // A fresh launcher-owned status directory per child: the injected + // library writes status-.json there, and the watchdog reads only + // this child's file. Hidden from the session scan by its dot name. + pass_opts.bootstrap_nonce = MakeBootstrapNonce(); + pass_opts.bootstrap_dir = output_dir / ".gpufl-bootstrap" / + ("pass-" + std::to_string(i) + "-" + + pass_opts.bootstrap_nonce.substr(0, 8)); + { + std::error_code dir_ec; + fs::create_directories(pass_opts.bootstrap_dir, dir_ec); + fs::permissions(pass_opts.bootstrap_dir, fs::perms::owner_all, + fs::perm_options::replace, dir_ec); + if (dir_ec) { + std::fprintf(stderr, "gpufl: cannot create %s: %s\n", + pass_opts.bootstrap_dir.string().c_str(), + dir_ec.message().c_str()); + return 2; + } + } + if (!setEnvOrPrint(platform, env::kInjectStatusDir, + pass_opts.bootstrap_dir.string()) || + !setEnvOrPrint(platform, env::kInjectStatusNonce, + pass_opts.bootstrap_nonce)) { + return 2; + } + if (const char* grace = std::getenv(env::kInjectWatchdogGraceMs)) { + char* end = nullptr; + const long long parsed = std::strtoll(grace, &end, 10); + if (end && *end == '\0' && parsed >= 0 && parsed <= 3'600'000) { + pass_opts.watchdog_grace_ms = parsed; + } + } +#endif + const int64_t pass_start_ns = nowNs(); const auto t_start = std::chrono::steady_clock::now(); - const TraceProcessResult run = platform.runProcess(args.command, run_opts); + const TraceProcessResult run = platform.runProcess(args.command, pass_opts); const auto t_elapsed = std::chrono::duration_cast( std::chrono::steady_clock::now() - t_start).count(); +#if !defined(_WIN32) + { + std::error_code dir_ec; + fs::remove_all(pass_opts.bootstrap_dir, dir_ec); + } +#endif if (run.launcher_error) { std::fprintf(stderr, "gpufl: %s\n", run.error.c_str()); return run.rc; } +#if !defined(_WIN32) + if (run.init_watchdog_killed) { + std::fprintf(stderr, + "[gpufl] %s: capture initialization never finished; the " + "watchdog killed the target (%s)\n", + what.c_str(), describeRunBootstrap(run).c_str()); + } else if (!run.signaled && run.rc == kInitWedgedExitCode && + run.bootstrap.valid && + run.bootstrap.coverage_notes.count("exit_while_initializing")) { + std::fprintf(stderr, + "[gpufl] %s: exited while capture initialization was " + "still running (rc=%d)\n", + what.c_str(), run.rc); + } + if (!args.quiet) { + std::fprintf(stderr, "[gpufl] %s capture readiness: %s\n", what.c_str(), + describeRunBootstrap(run).c_str()); + } +#endif + if (!args.quiet) { if (run.window_stopped) { std::fprintf(stderr, @@ -973,6 +1062,7 @@ int runTraceCommon(const TraceArgs& args, const TracePlatform& platform) { shutdown_context.signaled = run.signaled; shutdown_context.signal = run.signal; shutdown_context.window_stopped = run.window_stopped; + shutdown_context.init_watchdog_killed = run.init_watchdog_killed; const int synthesized = ensureTraceCompletionMarkers( output_dir, pass_start_ns, shutdown_context); if (synthesized > 0) { diff --git a/daemon/launcher/trace_command_common.hpp b/daemon/launcher/trace_command_common.hpp index 2df96ac..4f75f00 100644 --- a/daemon/launcher/trace_command_common.hpp +++ b/daemon/launcher/trace_command_common.hpp @@ -1,5 +1,6 @@ #pragma once +#include "bootstrap_status.hpp" #include "cli_parse.hpp" #include @@ -11,6 +12,10 @@ namespace gpufl::launcher { namespace fs = std::filesystem; +// The injected library ends the process with this status when an initializer +// is still inside CUDA/CUPTI at exit after its deadline (capture_gate). +constexpr int kInitWedgedExitCode = 86; + struct TraceProcessResult { int rc = 1; bool signaled = false; @@ -19,6 +24,14 @@ struct TraceProcessResult { // The launcher stopped the target at the window deadline (intentional — // not a crash and not a failure). rc/signal reflect the kill, not health. bool window_stopped = false; + // The initialization watchdog killed the target: capture initialization + // never finished within the library's published deadline plus the grace. + bool init_watchdog_killed = false; + // The child's last bootstrap status (invalid when none was read). + BootstrapStatus bootstrap; + // Empty when the read after exit confirmed `bootstrap`; otherwise why it + // did not (the record above is then historical). + std::string bootstrap_final_error; std::string error; }; @@ -29,6 +42,11 @@ struct TraceProcessResult { struct RunOptions { int64_t run_ms = 0; // stop the target after this wall-clock (0 = unbounded) int stop_grace_ms = 5000; // after the stop signal, hard-kill if still alive + // Per-child bootstrap status directory and nonce (empty: no watchdog). + fs::path bootstrap_dir; + std::string bootstrap_nonce; + // Allowed past the library's initialization deadline before the kill. + int64_t watchdog_grace_ms = 5000; }; class TracePlatform { diff --git a/include/gpufl/backends/nvidia/capture_capability_resolver.cpp b/include/gpufl/backends/nvidia/capture_capability_resolver.cpp index 0542d72..465a60f 100644 --- a/include/gpufl/backends/nvidia/capture_capability_resolver.cpp +++ b/include/gpufl/backends/nvidia/capture_capability_resolver.cpp @@ -15,8 +15,70 @@ void AddCapability(CaptureCapabilitiesEvent& evt, std::string feature, std::move(reason), std::move(message)}); } +// Refines an activity capability whose records were rejected for their +// timing. The status stays what the row count decided (the server classifies +// statuses only); the reason and message say why rows are missing. Synthetic +// and intentionally disabled capture keep their own reasons. +void NoteRejectedTimestamps(CaptureCapability& cap, const std::string& kind, + const char* what, const uint64_t eligible, + const uint64_t emitted, const uint64_t zero_ts, + const uint64_t invalid_ts) { + const std::string counts = + " (" + DescribeActivityTimestamps(eligible, zero_ts, invalid_ts) + + " emitted=" + std::to_string(emitted) + ")"; + switch (ClassifyActivityTimestamps(eligible, emitted, zero_ts, invalid_ts)) { + case ActivityTimestampOutcome::Complete: + return; + case ActivityTimestampOutcome::Unavailable: + if (cap.status != "enabled_no_data") return; + cap.reason_code = kind + "_timestamps_unavailable"; + cap.message = std::string("Every ") + what + + " activity record arrived with start == end == 0, so " + "CUPTI timing was unavailable this session and no row " + "was emitted" + counts + "."; + return; + case ActivityTimestampOutcome::Invalid: + if (cap.status != "enabled_no_data") return; + cap.reason_code = kind + "_timestamps_invalid"; + cap.message = std::string("Every ") + what + + " activity record was dropped for invalid timestamps, " + "so no row was emitted" + counts + "."; + return; + case ActivityTimestampOutcome::Partial: + if (cap.status != "collected" || !cap.reason_code.empty()) return; + cap.reason_code = kind + "_timestamps_partial"; + cap.message = std::string("Some ") + what + + " activity records were dropped for invalid " + "timestamps; the timeline is incomplete" + counts + "."; + return; + } +} + } // namespace +ActivityTimestampOutcome ClassifyActivityTimestamps(const uint64_t eligible, + const uint64_t emitted, + const uint64_t zero_ts, + const uint64_t invalid_ts) { + const uint64_t rejected = zero_ts + invalid_ts; + if (rejected == 0) return ActivityTimestampOutcome::Complete; + // Observed on an L4 (2026-10-16): a whole Trace session's kernel and + // memcpy records arrived with start == end == 0. That is not "no record + // arrived", and it must not read like it. + if (emitted > 0) return ActivityTimestampOutcome::Partial; + return invalid_ts == 0 && zero_ts == eligible + ? ActivityTimestampOutcome::Unavailable + : ActivityTimestampOutcome::Invalid; +} + +std::string DescribeActivityTimestamps(const uint64_t eligible, + const uint64_t zero_ts, + const uint64_t invalid_ts) { + return "eligible=" + std::to_string(eligible) + + " zero_timestamps=" + std::to_string(zero_ts) + + " other_invalid=" + std::to_string(invalid_ts); +} + std::vector BuildCaptureCapabilityWarnings( const CaptureCapabilityInput& input) { std::vector warnings; @@ -130,6 +192,11 @@ CaptureCapabilitiesEvent BuildCaptureCapabilitiesEvent( : input.sass_metrics_only ? "Kernel activity was intentionally disabled because CUPTI SASS Metrics requires metrics-only mode on this GPU/driver to produce non-zero counters." : "Kernel tracing was enabled but emitted no kernel rows this session."))); + NoteRejectedTimestamps(evt.capabilities.back(), "kernel", "kernel", + input.counters.kernel_records_eligible, + input.counters.kernel_rows, + input.counters.kernel_records_zero_ts, + input.counters.kernel_records_invalid_ts); AddCapability( evt, "kernel_names", input.kernel_activity, kernelHasData @@ -216,6 +283,11 @@ CaptureCapabilitiesEvent BuildCaptureCapabilitiesEvent( ? "Memcpy/memset activity records were collected." : "Memcpy/memset activity was enabled but emitted no rows this session.") : "Memcpy/memset timeline activity was not requested by the selected engine domains."); + NoteRejectedTimestamps(evt.capabilities.back(), "memcpy", "memcpy/memset", + input.counters.mem_transfer_records_eligible, + input.counters.mem_transfer_rows, + input.counters.mem_transfer_records_zero_ts, + input.counters.mem_transfer_records_invalid_ts); const bool syncRequested = input.kernel_activity && input.options.enable_synchronization && input.allow_sass_sync_activity; diff --git a/include/gpufl/backends/nvidia/capture_capability_resolver.hpp b/include/gpufl/backends/nvidia/capture_capability_resolver.hpp index 717ae1f..fc05331 100644 --- a/include/gpufl/backends/nvidia/capture_capability_resolver.hpp +++ b/include/gpufl/backends/nvidia/capture_capability_resolver.hpp @@ -23,8 +23,39 @@ struct CaptureCapabilityCounters { uint64_t source_rows = 0; uint64_t function_rows = 0; uint64_t launch_count = 0; + /** + * Activity records that reached the timestamp check (intentionally + * filtered records excluded), and the ones it rejected: start == end == 0 + * (CUPTI's "timestamps unavailable") and any other invalid timing. + */ + uint64_t kernel_records_eligible = 0; + uint64_t kernel_records_zero_ts = 0; + uint64_t kernel_records_invalid_ts = 0; + uint64_t mem_transfer_records_eligible = 0; + uint64_t mem_transfer_records_zero_ts = 0; + uint64_t mem_transfer_records_invalid_ts = 0; }; +enum class ActivityTimestampOutcome { + /** Nothing was rejected for its timing (or nothing arrived). */ + Complete, + /** Records arrived, none was emitted, and every one had start == end == 0. */ + Unavailable, + /** None was emitted, and the rejected ones were not all zero-zero. */ + Invalid, + /** Some records were rejected for their timing; the rest still count. */ + Partial, +}; + +ActivityTimestampOutcome ClassifyActivityTimestamps(uint64_t eligible, + uint64_t emitted, + uint64_t zero_ts, + uint64_t invalid_ts); + +/** "eligible=N zero_timestamps=Z other_invalid=I" for capability messages and logs. */ +std::string DescribeActivityTimestamps(uint64_t eligible, uint64_t zero_ts, + uint64_t invalid_ts); + struct CaptureCapabilityOptions { bool enable_memory_tracking = true; bool enable_external_correlation = true; diff --git a/include/gpufl/backends/nvidia/cupti_backend.cpp b/include/gpufl/backends/nvidia/cupti_backend.cpp index fb7cedb..ec5416b 100644 --- a/include/gpufl/backends/nvidia/cupti_backend.cpp +++ b/include/gpufl/backends/nvidia/cupti_backend.cpp @@ -145,6 +145,7 @@ void CuptiBackend::initialize(const MonitorOptions& opts) { " - Insufficient permissions\n" " - CUDA driver issues"); GFL_LOG_DEBUG("CUPTI subscription successful"); + cupti_subscribed_.store(true, std::memory_order_release); std::set domains; std::set> callbacks; @@ -504,7 +505,14 @@ void CuptiBackend::start() { std::set kinds; for (const auto& h : handlers_) for (auto k : h->requiredActivityKinds()) kinds.insert(k); - for (auto k : kinds) CUPTI_CHECK(cuptiActivityEnable(k)); + for (auto k : kinds) { + const CUptiResult res = cuptiActivityEnable(k); + CUPTI_CHECK(res); + // Capture readiness requires every required kind to be on; a + // "benign" result is a fallback for logging, not an enable. + required_activity_.Record(static_cast(k), res == CUPTI_SUCCESS, + static_cast(res)); + } } // Initialize and start the engine (requires CUDA context) @@ -535,7 +543,13 @@ void CuptiBackend::ReenableActivityAfterEngineStart_() { std::set kinds; for (const auto& h : handlers_) for (auto k : h->requiredActivityKinds()) kinds.insert(k); - for (auto k : kinds) cuptiActivityEnable(k); + // The engine may have reset the subscription, so the re-enable result + // counts too: an earlier success does not stand for this one. + for (auto k : kinds) { + const CUptiResult res = cuptiActivityEnable(k); + required_activity_.Record(static_cast(k), res == CUPTI_SUCCESS, + static_cast(res)); + } } // Re-resolve NVTX marker policy after engine start. This is primarily for @@ -983,8 +997,11 @@ void CuptiBackend::stop() { kernel_activity_emitted_.load(std::memory_order_relaxed); const uint64_t throttled = kernel_activity_throttled_.load(std::memory_order_relaxed); + const ActivityTimestampCounts ts = kernel_activity_timestamps_.Snapshot(); GFL_LOG_DEBUG("[KernelLaunchHandler] activity summary seen=", seen, - " emitted=", emitted, " throttled=", throttled); + " emitted=", emitted, " throttled=", throttled, + " eligible=", ts.eligible, " zero_timestamps=", ts.zero_ts, + " other_invalid=", ts.invalid_ts); } void CuptiBackend::RegisterHandler( diff --git a/include/gpufl/backends/nvidia/cupti_backend.hpp b/include/gpufl/backends/nvidia/cupti_backend.hpp index 29b03ab..3daa57a 100644 --- a/include/gpufl/backends/nvidia/cupti_backend.hpp +++ b/include/gpufl/backends/nvidia/cupti_backend.hpp @@ -51,6 +51,55 @@ inline bool ShouldSuppressOrphanKernelSynthesis(bool collects_kernel_events, return collects_kernel_events || !will_emit_synthetic; } +// Activity records that reached a handler's timestamp check, and the ones it +// rejected. Process-cumulative like the other activity counters; callers take +// segment deltas from Snapshot(). +struct ActivityTimestampTally { + std::atomic eligible{0}; + std::atomic zero_ts{0}; + std::atomic invalid_ts{0}; + + void NoteEligible() { eligible.fetch_add(1, std::memory_order_relaxed); } + void NoteRejected(const uint64_t start, const uint64_t end) { + (start == 0 && end == 0 ? zero_ts : invalid_ts) + .fetch_add(1, std::memory_order_relaxed); + } + ActivityTimestampCounts Snapshot() const { + return {eligible.load(std::memory_order_relaxed), + zero_ts.load(std::memory_order_relaxed), + invalid_ts.load(std::memory_order_relaxed)}; + } +}; + +// Results of enabling the required activity kinds, at start and at every +// re-enable. Ready only when at least one enable was recorded and none +// failed: a failure latches, so a later success never hides it, and a +// non-success the logger calls benign (a fallback) is still not an enable. +class RequiredActivityTally { + public: + void Record(const int kind, const bool success, const int result) { + std::lock_guard lock(mu_); + recorded_ = true; + if (!success && failure_.empty()) { + failure_ = "activity_enable_failed:" + std::to_string(kind) + ":" + + std::to_string(result); + } + } + bool Ready(std::string* reason) const { + std::lock_guard lock(mu_); + if (!failure_.empty() || !recorded_) { + if (reason) *reason = recorded_ ? failure_ : "activity_not_enabled"; + return false; + } + return true; + } + + private: + mutable std::mutex mu_; + bool recorded_ = false; + std::string failure_; +}; + /** * @brief CUPTI-based monitoring backend for NVIDIA GPUs. * @@ -123,6 +172,26 @@ class CuptiBackend : public IMonitorBackend { kernel_launch_callback_count_.load(std::memory_order_acquire), kernel_activity_emitted_.load(std::memory_order_relaxed)); } + ActivityTimestampCounts kernelActivityTimestamps() const override { + return kernel_activity_timestamps_.Snapshot(); + } + bool captureFacilitiesReady(std::string* reason) const override { + const auto fail = [reason](std::string why) { + if (reason) *reason = std::move(why); + return false; + }; + if (!cupti_subscribed_.load(std::memory_order_acquire)) { + return fail("cupti_subscribe_failed"); + } + return required_activity_.Ready(reason); + } + void NoteMemTransferActivityEligible() { + mem_transfer_activity_timestamps_.NoteEligible(); + } + void NoteMemTransferActivityRejected(const uint64_t start, + const uint64_t end) { + mem_transfer_activity_timestamps_.NoteRejected(start, end); + } bool AllowSassMarkerActivity() const { return resolved_plan_.allow_sass_marker_activity; } @@ -384,7 +453,11 @@ class CuptiBackend : public IMonitorBackend { std::atomic kernel_activity_seen_{0}; std::atomic kernel_activity_emitted_{0}; std::atomic kernel_activity_throttled_{0}; + ActivityTimestampTally kernel_activity_timestamps_; + std::atomic cupti_subscribed_{false}; + RequiredActivityTally required_activity_; std::atomic mem_transfer_activity_emitted_{0}; + ActivityTimestampTally mem_transfer_activity_timestamps_; std::atomic sync_activity_emitted_{0}; std::atomic nvtx_marker_emitted_{0}; std::atomic graph_activity_emitted_{0}; @@ -426,6 +499,8 @@ class CuptiBackend : public IMonitorBackend { mutable uint64_t capability_launch_count_baseline_ = 0; mutable uint64_t capability_scope_truncated_baseline_ = 0; mutable uint64_t capability_pm_rows_baseline_ = 0; + mutable ActivityTimestampCounts capability_kernel_timestamps_baseline_; + mutable ActivityTimestampCounts capability_mem_transfer_timestamps_baseline_; uint32_t device_id_ = 0; std::string chip_name_; diff --git a/include/gpufl/backends/nvidia/cupti_capture_capabilities.cpp b/include/gpufl/backends/nvidia/cupti_capture_capabilities.cpp index 7436c1a..cf378e1 100644 --- a/include/gpufl/backends/nvidia/cupti_capture_capabilities.cpp +++ b/include/gpufl/backends/nvidia/cupti_capture_capabilities.cpp @@ -47,6 +47,10 @@ void CuptiBackend::EmitCaptureCapabilities_() const { function_record_seen_.load(std::memory_order_relaxed); const uint64_t launch_count = kernel_launch_callback_count_.load(std::memory_order_acquire); + const ActivityTimestampCounts kernel_ts = + kernel_activity_timestamps_.Snapshot(); + const ActivityTimestampCounts transfer_ts = + mem_transfer_activity_timestamps_.Snapshot(); const EngineRequestSet requests = BuildEngineRequestSet(opts_.profiling_engine, combo_); @@ -104,6 +108,18 @@ void CuptiBackend::EmitCaptureCapabilities_() const { delta(function_rows, capability_function_rows_baseline_); input.counters.launch_count = delta(launch_count, capability_launch_count_baseline_); + const auto& kb = capability_kernel_timestamps_baseline_; + input.counters.kernel_records_eligible = delta(kernel_ts.eligible, kb.eligible); + input.counters.kernel_records_zero_ts = delta(kernel_ts.zero_ts, kb.zero_ts); + input.counters.kernel_records_invalid_ts = + delta(kernel_ts.invalid_ts, kb.invalid_ts); + const auto& tb = capability_mem_transfer_timestamps_baseline_; + input.counters.mem_transfer_records_eligible = + delta(transfer_ts.eligible, tb.eligible); + input.counters.mem_transfer_records_zero_ts = + delta(transfer_ts.zero_ts, tb.zero_ts); + input.counters.mem_transfer_records_invalid_ts = + delta(transfer_ts.invalid_ts, tb.invalid_ts); input.pc_insufficient_privileges = engine_ && engine_->hasInsufficientPrivileges(); input.pc_stall_reasons_unavailable = @@ -145,6 +161,7 @@ void CuptiBackend::EmitCaptureCapabilities_() const { evt.capabilities.push_back(std::move(cap)); } + Monitor::ContributeCaptureCapabilities(evt); segment->logger->write(model::CaptureCapabilitiesModel(evt)); capability_kernel_rows_baseline_ = kernel_rows; capability_memory_rows_baseline_ = memory_rows; @@ -159,6 +176,8 @@ void CuptiBackend::EmitCaptureCapabilities_() const { capability_launch_count_baseline_ = launch_count; capability_scope_truncated_baseline_ = truncated_total; capability_pm_rows_baseline_ = pm_rows_total; + capability_kernel_timestamps_baseline_ = kernel_ts; + capability_mem_transfer_timestamps_baseline_ = transfer_ts; capture_capabilities_session_id_ = segment->session_id; } diff --git a/include/gpufl/backends/nvidia/kernel_launch_handler.cpp b/include/gpufl/backends/nvidia/kernel_launch_handler.cpp index 845a6d1..3d53fc8 100644 --- a/include/gpufl/backends/nvidia/kernel_launch_handler.cpp +++ b/include/gpufl/backends/nvidia/kernel_launch_handler.cpp @@ -484,6 +484,7 @@ bool KernelLaunchHandler::handleActivityRecord(const CUpti_Activity* record, "internal wait kernel corr=", k->correlationId); return false; } + backend_->kernel_activity_timestamps_.NoteEligible(); // Diagnostic: log every real activity record as it arrives. GFL_LOG_DEBUG("[KernelLaunchHandler] ACTIVITY_RECORD corr=", @@ -515,6 +516,7 @@ bool KernelLaunchHandler::handleActivityRecord(const CUpti_Activity* record, GFL_LOG_DEBUG("[KernelLaunchHandler] dropping kernel activity with " "invalid timestamps corr=", k->correlationId, " start=", k->start, " end=", k->end); + backend_->kernel_activity_timestamps_.NoteRejected(k->start, k->end); ActivityRecord drop{}; drop.type = TraceType::KERNEL_META_DISCARD; drop.corr_id = k->correlationId; diff --git a/include/gpufl/backends/nvidia/mem_transfer_handler.cpp b/include/gpufl/backends/nvidia/mem_transfer_handler.cpp index 831229e..d230985 100644 --- a/include/gpufl/backends/nvidia/mem_transfer_handler.cpp +++ b/include/gpufl/backends/nvidia/mem_transfer_handler.cpp @@ -253,6 +253,7 @@ bool MemTransferHandler::handleActivityRecord(const CUpti_Activity* record, if (record->kind == CUPTI_ACTIVITY_KIND_MEMCPY || record->kind == CUPTI_ACTIVITY_KIND_MEMCPY2) { const auto* m = reinterpret_cast(record); + backend_->NoteMemTransferActivityEligible(); const int64_t wallStartNs = baseCpuNs + static_cast(m->start - baseCuptiTs); if (m->start == 0 || m->end < m->start || @@ -261,6 +262,7 @@ bool MemTransferHandler::handleActivityRecord(const CUpti_Activity* record, GFL_LOG_DEBUG("[MemTransferHandler] dropping memcpy activity with " "invalid timestamps corr=", m->correlationId, " start=", m->start, " end=", m->end); + backend_->NoteMemTransferActivityRejected(m->start, m->end); ActivityRecord drop{}; drop.type = TraceType::KERNEL_META_DISCARD; drop.corr_id = m->correlationId; @@ -289,6 +291,7 @@ bool MemTransferHandler::handleActivityRecord(const CUpti_Activity* record, if (record->kind == CUPTI_ACTIVITY_KIND_MEMSET) { const auto* m = reinterpret_cast(record); + backend_->NoteMemTransferActivityEligible(); const int64_t wallStartNs = baseCpuNs + static_cast(m->start - baseCuptiTs); if (m->start == 0 || m->end < m->start || @@ -297,6 +300,7 @@ bool MemTransferHandler::handleActivityRecord(const CUpti_Activity* record, GFL_LOG_DEBUG("[MemTransferHandler] dropping memset activity with " "invalid timestamps corr=", m->correlationId, " start=", m->start, " end=", m->end); + backend_->NoteMemTransferActivityRejected(m->start, m->end); ActivityRecord drop{}; drop.type = TraceType::KERNEL_META_DISCARD; drop.corr_id = m->correlationId; diff --git a/include/gpufl/core/env_vars.hpp b/include/gpufl/core/env_vars.hpp index 4cc468f..195c9fa 100644 --- a/include/gpufl/core/env_vars.hpp +++ b/include/gpufl/core/env_vars.hpp @@ -116,10 +116,26 @@ constexpr const char* kFlushLogsAlways = "GPUFL_FLUSH_LOGS_ALWAYS"; // ── Inject-lib timing knobs (tuning / debugging the injection lifecycle) ──── constexpr const char* kInjectShutdownDelayMs = "GPUFL_INJECT_SHUTDOWN_DELAY_MS"; -constexpr const char* kInjectSyncWaitMs = "GPUFL_INJECT_SYNC_WAIT_MS"; constexpr const char* kInjectLaunchWaitMs = "GPUFL_INJECT_LAUNCH_WAIT_MS"; constexpr const char* kInjectProcessScope = "GPUFL_INJECT_PROCESS_SCOPE"; constexpr const char* kInjectInitDelayMs = "GPUFL_INJECT_INIT_DELAY_MS"; +// --warmup: a deliberate capture delay. Unlike the init delay above, the first +// CUDA calls do not wait for it, so a warmup run is never a qualified Trace +// capture (coverage note warmup_unqualified). +constexpr const char* kInjectWarmupMs = "GPUFL_INJECT_WARMUP_MS"; +// Launcher/library pairing. The launcher sets it; the library records what it +// saw. Absent means an older launcher or a direct injection. +constexpr const char* kInjectProtocol = "GPUFL_INJECT_PROTOCOL"; +constexpr const char* kInjectProtocolVersion = "2"; +// Launcher-owned per-child bootstrap status directory and its nonce. Each +// injected process writes status-.json there without needing the logger; +// the launcher reads only its own child's file, so a grandchild that inherited +// the environment cannot overwrite it. +constexpr const char* kInjectStatusDir = "GPUFL_INJECT_STATUS_DIR"; +constexpr const char* kInjectStatusNonce = "GPUFL_INJECT_STATUS_NONCE"; +// Extra time the launcher allows past the library's published initialization +// deadline before it kills a child whose initialization never finished. +constexpr const char* kInjectWatchdogGraceMs = "GPUFL_INJECT_WATCHDOG_GRACE_MS"; // ── Profiling engine selection ────────────────────────────────────────────── // Override the resolved ProfilingEngine; gpufl::init() is the single string→enum diff --git a/include/gpufl/core/monitor.cpp b/include/gpufl/core/monitor.cpp index eb6ec35..2c754b9 100644 --- a/include/gpufl/core/monitor.cpp +++ b/include/gpufl/core/monitor.cpp @@ -608,11 +608,25 @@ void Monitor::Shutdown() { // enabled_but_no_records, so the session stays self-describing. if (IMonitorBackend* b = g_state.adapter->backend(); b && b->kernelActivityExpectedButMissing()) { - GFL_LOG_ERROR( - "[Monitor] Kernel activity was enabled and launches happened, " - "but no kernel activity record arrived this session. Kernel " - "rows are omitted (a synthetic fallback would carry host-gap " - "durations, not kernel time)."); + // Records that arrived without timing are a different failure + // from records that never arrived; say which one this was. + const ActivityTimestampCounts ts = b->kernelActivityTimestamps(); + if (ts.eligible == 0) { + GFL_LOG_ERROR( + "[Monitor] Kernel activity was enabled and launches " + "happened, but no kernel activity record arrived this " + "session. Kernel rows are omitted (a synthetic fallback " + "would carry host-gap durations, not kernel time)."); + } else { + GFL_LOG_ERROR( + "[Monitor] Kernel activity was enabled and launches " + "happened, but no kernel activity record had valid timing " + "this session (eligible=", ts.eligible, + " zero_timestamps=", ts.zero_ts, + " other_invalid=", ts.invalid_ts, + "). Kernel rows are omitted (a synthetic fallback would " + "carry host-gap durations, not kernel time)."); + } } g_state.adapter->shutdown(); } @@ -725,6 +739,29 @@ void Monitor::RequestSyntheticDrainAndWait() { } void Monitor::Start() { if (g_state.adapter) g_state.adapter->start(); } +bool Monitor::CaptureFacilitiesReady(std::string* reason) { + IMonitorBackend* backend = + g_state.adapter ? g_state.adapter->backend() : nullptr; + // The adapter, not `initialized`: Shutdown clears that flag before the + // final capability report, which asks this again. + if (!g_state.adapter) { + if (reason) *reason = "monitor_not_started"; + return false; + } + // A monitoring-only session has no capture backend and needs none. + return !backend || backend->captureFacilitiesReady(reason); +} +namespace { +std::atomic g_capability_contributor{nullptr}; +} // namespace +void Monitor::SetCaptureCapabilityContributor(CaptureCapabilityContributor fn) { + g_capability_contributor.store(fn, std::memory_order_release); +} +void Monitor::ContributeCaptureCapabilities(CaptureCapabilitiesEvent& evt) { + if (const auto fn = g_capability_contributor.load(std::memory_order_acquire)) { + fn(evt); + } +} void Monitor::Stop() { if (g_state.adapter) g_state.adapter->stop(); } void Monitor::PushRange(const char* name) { diff --git a/include/gpufl/core/monitor.hpp b/include/gpufl/core/monitor.hpp index 86cb080..5c0e3d7 100644 --- a/include/gpufl/core/monitor.hpp +++ b/include/gpufl/core/monitor.hpp @@ -308,6 +308,23 @@ class Monitor { */ static void Start(); + /** + * @brief Whether the started backend set up every capture facility it + * requires (see IMonitorBackend::captureFacilitiesReady). False + * with "monitor_not_started" when no backend is running. + */ + static bool CaptureFacilitiesReady(std::string* reason); + + /** + * @brief A process-wide hook that appends rows to every capture + * capabilities report before it is written. The Linux injection + * library uses it for capture readiness and coverage. A plain + * function pointer, so nothing is destroyed under a late call. + */ + using CaptureCapabilityContributor = void (*)(CaptureCapabilitiesEvent& evt); + static void SetCaptureCapabilityContributor(CaptureCapabilityContributor fn); + static void ContributeCaptureCapabilities(CaptureCapabilitiesEvent& evt); + /** * @brief Stops global collection. */ diff --git a/include/gpufl/core/monitor_backend.hpp b/include/gpufl/core/monitor_backend.hpp index 38a6d29..c2ba825 100644 --- a/include/gpufl/core/monitor_backend.hpp +++ b/include/gpufl/core/monitor_backend.hpp @@ -1,5 +1,6 @@ #pragma once +#include #include #include #include @@ -9,6 +10,14 @@ namespace gpufl { +/** Activity records that reached the timestamp check, and the ones rejected. */ +struct ActivityTimestampCounts { + uint64_t eligible = 0; + /** start == end == 0: CUPTI could not collect timing for the record. */ + uint64_t zero_ts = 0; + uint64_t invalid_ts = 0; +}; + /** * @brief Interface for backend-specific monitoring implementations. * @@ -118,6 +127,22 @@ class IMonitorBackend { */ virtual bool kernelActivityExpectedButMissing() const { return false; } + /** + * @brief True when every capture facility this session requires was set + * up successfully (for CUPTI: the subscription and every required + * activity kind). The Linux injection gate reports capture ready + * only on true. On false, @p reason names the first failure. + */ + virtual bool captureFacilitiesReady(std::string* reason) const { + (void)reason; + return true; + } + + /** Kernel activity records seen at the timestamp check, process-wide. */ + virtual ActivityTimestampCounts kernelActivityTimestamps() const { + return {}; + } + virtual void OnScopeStart(const char* name) {} virtual void OnScopeStop(const char* name) {} diff --git a/include/gpufl/inject/capture_gate.cpp b/include/gpufl/inject/capture_gate.cpp new file mode 100644 index 0000000..3b78ec6 --- /dev/null +++ b/include/gpufl/inject/capture_gate.cpp @@ -0,0 +1,634 @@ +#include "gpufl/inject/capture_gate.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include + +#include "gpufl/core/env_vars.hpp" +#include "gpufl/core/events/lifecycle_events.hpp" + +#if !defined(_WIN32) +#include +#include +#endif + +namespace gpufl::inject { +namespace { + +constexpr int64_t kDefaultWaitMs = 15'000; +constexpr int64_t kMaxWaitMs = 600'000; +constexpr int kDefaultDelayMs = 1; +constexpr std::size_t kMaxReasonBytes = 200; + +// Which gate, if any, this thread belongs to. A gate never waits on its own +// initializer, inside the driver's injection callback, or inside the driver +// init it triggered: each of those would wait for itself. +thread_local const CaptureGate* t_initializer = nullptr; +thread_local const CaptureGate* t_in_callback = nullptr; +thread_local const CaptureGate* t_in_trigger = nullptr; +thread_local int t_wrapper_depth = 0; + +class ThreadMark { + public: + ThreadMark(const CaptureGate*& slot, const CaptureGate* gate) + : slot_(slot), previous_(slot) { + slot_ = gate; + } + ~ThreadMark() { slot_ = previous_; } + ThreadMark(const ThreadMark&) = delete; + ThreadMark& operator=(const ThreadMark&) = delete; + + private: + const CaptureGate*& slot_; + const CaptureGate* previous_; +}; + +bool ParseInt(const char* text, int64_t& out) { + if (!text || !*text) return false; + char* end = nullptr; + errno = 0; + const long long value = std::strtoll(text, &end, 10); + if (errno != 0 || !end || *end != '\0') return false; + out = value; + return true; +} + +GateStep Call(const std::function& fn, const char* missing, + const char* exception) { + if (!fn) return {false, missing}; + try { + return fn(); + } catch (...) { + return {false, exception}; + } +} + +std::string Bounded(std::string text) { + if (text.size() > kMaxReasonBytes) text.resize(kMaxReasonBytes); + return text; +} + +void AppendJsonString(std::string& out, const std::string& value) { + out += '"'; + for (const unsigned char c : value) { + switch (c) { + case '"': out += "\\\""; break; + case '\\': out += "\\\\"; break; + case '\n': out += "\\n"; break; + case '\r': out += "\\r"; break; + case '\t': out += "\\t"; break; + default: + if (c < 0x20) { + static const char* kHex = "0123456789abcdef"; + out += "\\u00"; + out += kHex[c >> 4]; + out += kHex[c & 0xf]; + } else { + out += static_cast(c); + } + } + } + out += '"'; +} + +void AppendParticipant(std::string& out, const char* name, + const GateParticipant& p) { + out += ",\""; + out += name; + out += "\":{\"expected\":"; + out += p.expected ? "true" : "false"; + out += ",\"returned\":"; + out += p.returned ? "true" : "false"; + out += ",\"ok\":"; + out += p.ok ? "true" : "false"; + out += '}'; +} + +} // namespace + +const char* GateLifecycleName(const GateLifecycle value) { + switch (value) { + case GateLifecycle::NotRequested: return "not_requested"; + case GateLifecycle::DriverInitializing: return "driver_initializing"; + case GateLifecycle::DriverReady: return "driver_ready"; + case GateLifecycle::CaptureInitializing: return "capture_initializing"; + case GateLifecycle::Finished: return "finished"; + } + return "unknown"; +} + +const char* GateOwnerName(const GateOwner value) { + switch (value) { + case GateOwner::None: return "none"; + case GateOwner::Gate: return "gate"; + case GateOwner::Callback: return "callback"; + case GateOwner::Manual: return "manual"; + } + return "unknown"; +} + +const char* GateReadinessName(const GateReadiness value) { + switch (value) { + case GateReadiness::Pending: return "pending"; + case GateReadiness::Ready: return "ready"; + case GateReadiness::Failed: return "failed"; + } + return "unknown"; +} + +GateConfig GateConfigFromEnv( + const std::function& getenv_fn) { + const auto get = [&](const char* name) -> const char* { + const char* value = getenv_fn ? getenv_fn(name) : std::getenv(name); + return value && *value ? value : nullptr; + }; + GateConfig config; + + int64_t delay = kDefaultDelayMs; + bool delay_set = false; + if (int64_t parsed = 0; + ParseInt(get(env::kInjectInitDelayMs), parsed) && parsed >= 0 && + parsed <= std::numeric_limits::max()) { + delay = parsed; + delay_set = true; + } + config.delay_ms = static_cast(delay); + + if (const char* manual = get(env::kInjectManualInit); + manual && std::strcmp(manual, "1") == 0) { + config.mode = GateMode::Manual; + } else if (int64_t warmup = 0; + ParseInt(get(env::kInjectWarmupMs), warmup) && warmup > 0 && + warmup <= std::numeric_limits::max()) { + config.mode = GateMode::Warmup; + config.delay_ms = static_cast(warmup); + } + + if (const char* wait = get(env::kInjectLaunchWaitMs)) { + int64_t parsed = 0; + if (ParseInt(wait, parsed) && parsed >= 1 && parsed <= kMaxWaitMs) { + config.wait_ns = parsed * 1'000'000LL; + } else { + config.wait_ns = kDefaultWaitMs * 1'000'000LL; + config.wait_config_invalid = true; + } + } + + if (const char* protocol = get(env::kInjectProtocol)) { + config.protocol = protocol; + } + // An older launcher wrote --warmup into the init delay. Without the + // protocol variable a long delay may have been one, and a warmup run is + // not a qualified capture. + config.legacy_launcher_delay = config.protocol.empty() && + config.mode == GateMode::Automatic && + delay_set && delay > kDefaultDelayMs; + return config; +} + +CaptureGate::CaptureGate(GateConfig config, GateHooks hooks) + : config_(std::move(config)), hooks_(std::move(hooks)) { + state_.wait_config_invalid = config_.wait_config_invalid; + state_.protocol = config_.protocol; + if (config_.mode == GateMode::Warmup) { + state_.coverage_notes.insert("warmup_unqualified"); + } + if (config_.legacy_launcher_delay) { + state_.coverage_notes.insert("legacy_launcher_delay"); + } +} + +bool CaptureGate::Settled() const { + return settled_.load(std::memory_order_acquire); +} + +GateSnapshot CaptureGate::Snapshot() const { + std::lock_guard lock(mu_); + return state_; +} + +void CaptureGate::ElectLocked(const GateOwner owner) { + state_.owner = owner; + state_.lifecycle = GateLifecycle::DriverInitializing; + const int64_t now = hooks_.now_ns ? hooks_.now_ns() : 0; + // One immutable budget: the worker's requested delay (the readiness delay, + // or a --warmup) once, plus the initialization allowance. Nothing renews it. + const int64_t max = std::numeric_limits::max(); + const int64_t delay_ns = + owner == GateOwner::Manual ? 0 : static_cast(config_.delay_ms) * 1'000'000LL; + const int64_t budget = delay_ns > max - config_.wait_ns ? max : delay_ns + config_.wait_ns; + state_.deadline_ns = now > max - budget ? max : now + budget; +} + +void CaptureGate::LatchFailedLocked(const std::string& reason) { + if (state_.readiness != GateReadiness::Pending) return; + state_.readiness = GateReadiness::Failed; + state_.reason = Bounded(reason); + settled_.store(true, std::memory_order_release); +} + +void CaptureGate::ResolveLocked() { + const GateParticipant* participants[] = {&state_.trigger, &state_.worker, + &state_.handshake}; + bool any = false; + bool all_returned = true; + bool all_ok = true; + for (const GateParticipant* p : participants) { + if (!p->expected) continue; + any = true; + all_returned = all_returned && p->returned; + all_ok = all_ok && p->ok; + } + if (!any || !all_returned) return; + // Every participant has stopped using initialization-owned state. + state_.lifecycle = GateLifecycle::Finished; + if (state_.readiness != GateReadiness::Pending) return; + if (!all_ok) { + LatchFailedLocked("init_failed"); + return; + } + const int64_t now = hooks_.now_ns ? hooks_.now_ns() : 0; + if (now > state_.deadline_ns) { + // No waiter has to be present for the budget to apply. + LatchFailedLocked("capture_init_deadline_exceeded"); + return; + } + state_.readiness = GateReadiness::Ready; + settled_.store(true, std::memory_order_release); +} + +void CaptureGate::PublishLocked() { + // Under mu_: a state is on disk before any thread can act on it, a newer + // record can never be replaced by an older one, and nothing is published + // after the process saw Finished and went on into teardown. + ++state_.seq; + if (hooks_.publish) hooks_.publish(state_); +} + +bool CaptureGate::WaitLocked(std::unique_lock& lock, + const bool until_finished) { + for (;;) { + const bool done = until_finished + ? state_.lifecycle == GateLifecycle::Finished + : state_.readiness != GateReadiness::Pending; + if (done) return true; + const int64_t now = hooks_.now_ns ? hooks_.now_ns() : 0; + if (now >= state_.deadline_ns) { + if (state_.readiness == GateReadiness::Pending) { + LatchFailedLocked("capture_wait_timeout"); + PublishLocked(); + cv_.notify_all(); + } + return until_finished + ? state_.lifecycle == GateLifecycle::Finished + : false; + } + cv_.wait_for(lock, std::chrono::nanoseconds(state_.deadline_ns - now)); + } +} + +void CaptureGate::StartWorkerLocked() { + worker_started_ = true; + state_.worker.expected = true; +} + +void CaptureGate::Enter(const std::function& trigger) { + if (Settled()) return; + if (t_initializer == this || t_in_callback == this || t_in_trigger == this) { + return; + } + if (t_wrapper_depth > 1) return; // an inner wrapper of an outer one + if (config_.mode == GateMode::Warmup) return; + + std::unique_lock lock(mu_); + if (config_.mode == GateMode::Manual) { + // The host must finish the handshake before capture-relevant work. + if (state_.readiness == GateReadiness::Pending && + state_.coverage_notes.insert("manual_protocol_violation").second) { + PublishLocked(); + lock.unlock(); + } + return; + } + + if (state_.lifecycle == GateLifecycle::NotRequested) { + ElectLocked(GateOwner::Gate); + state_.trigger.expected = true; + PublishLocked(); + lock.unlock(); + + GateStep step; + { + ThreadMark mark(t_in_trigger, this); + step = Call(trigger, "trigger_missing", "trigger_exception"); + } + + lock.lock(); + state_.trigger.returned = true; + state_.trigger.ok = step.ok; + if (!step.ok) { + LatchFailedLocked(step.reason.empty() ? "trigger_failed" : step.reason); + } else if (!worker_started_) { + // Driver init succeeded but never reached InitializeInjection. + LatchFailedLocked("no_injection_callback"); + } + ResolveLocked(); + PublishLocked(); + cv_.notify_all(); + } + WaitLocked(lock, /*until_finished=*/false); +} + +void CaptureGate::OnInjectionCallback() { + ThreadMark mark(t_in_callback, this); + std::unique_lock lock(mu_); + ++state_.callbacks_seen; + bool spawn = false; + if (config_.mode != GateMode::Manual) { + if (state_.lifecycle == GateLifecycle::NotRequested) { + // Driver init came from a route no gate owns. + ElectLocked(GateOwner::Callback); + if (config_.mode == GateMode::Automatic) { + state_.coverage_notes.insert("unsupported_first_touch"); + } + StartWorkerLocked(); + spawn = true; + } else if (state_.owner == GateOwner::Gate && !worker_started_ && + !state_.trigger.returned) { + // The gate owner's trigger is on the stack: arrange its worker. + StartWorkerLocked(); + spawn = true; + } + } + PublishLocked(); + lock.unlock(); + if (spawn) { + if (!hooks_.spawn || !hooks_.spawn([this] { RunWorker(); })) { + std::unique_lock relock(mu_); + state_.worker.returned = true; + LatchFailedLocked("worker_spawn_unavailable"); + ResolveLocked(); + PublishLocked(); + cv_.notify_all(); + relock.unlock(); + } + } +} + +void CaptureGate::RunWorker() { + ThreadMark mark(t_initializer, this); + const GateStep barrier = + Call(hooks_.driver_barrier, "driver_barrier_missing", + "driver_barrier_exception"); + { + std::unique_lock lock(mu_); + if (!barrier.ok) { + state_.worker.returned = true; + state_.worker.ok = false; + LatchFailedLocked(barrier.reason.empty() ? "cuinit_failed" + : barrier.reason); + ResolveLocked(); + PublishLocked(); + cv_.notify_all(); + lock.unlock(); + return; + } + if (state_.lifecycle == GateLifecycle::DriverInitializing) { + state_.lifecycle = GateLifecycle::DriverReady; + } + PublishLocked(); + lock.unlock(); + } + { + // The delay is not CUDA work: an exit that starts during it stops the + // worker here, before CUPTI, so ordinary teardown stays safe. + std::unique_lock lock(mu_); + cv_.wait_for(lock, std::chrono::milliseconds(config_.delay_ms), + [this] { return exit_requested_; }); + if (exit_requested_) { + state_.worker.returned = true; + state_.worker.ok = false; + LatchFailedLocked("exit_before_capture"); + ResolveLocked(); + PublishLocked(); + cv_.notify_all(); + return; + } + state_.lifecycle = GateLifecycle::CaptureInitializing; + PublishLocked(); + } + const GateStep init = + Call(hooks_.initialize, "initialize_missing", "init_exception"); + std::unique_lock lock(mu_); + state_.worker.returned = true; + state_.worker.ok = init.ok; + if (!init.ok) { + LatchFailedLocked(init.reason.empty() ? "init_failed" : init.reason); + } + ResolveLocked(); + PublishLocked(); + cv_.notify_all(); + lock.unlock(); +} + +int CaptureGate::ManualHandshake(const std::function& driver_trigger) { + if (config_.mode != GateMode::Manual) { + // Without manual mode the handshake is one more outer entry. + Enter(driver_trigger); + std::lock_guard lock(mu_); + return state_.readiness == GateReadiness::Ready ? 1 : 0; + } + std::unique_lock lock(mu_); + if (t_initializer == this) { + ++state_.handshake_reentries; + return 0; + } + if (state_.lifecycle == GateLifecycle::NotRequested) { + ElectLocked(GateOwner::Manual); + state_.handshake.expected = true; + PublishLocked(); + lock.unlock(); + + GateStep result; + { + ThreadMark mark(t_initializer, this); + result = Call(hooks_.driver_barrier, "driver_barrier_missing", + "driver_barrier_exception"); + if (result.ok) { + lock.lock(); + state_.lifecycle = GateLifecycle::CaptureInitializing; + PublishLocked(); + lock.unlock(); + result = Call(hooks_.initialize, "initialize_missing", + "init_exception"); + } + } + + lock.lock(); + state_.handshake.returned = true; + state_.handshake.ok = result.ok; + if (!result.ok) { + LatchFailedLocked(result.reason.empty() ? "init_failed" : result.reason); + } + ResolveLocked(); + PublishLocked(); + cv_.notify_all(); + } else { + WaitLocked(lock, /*until_finished=*/false); + } + return state_.readiness == GateReadiness::Ready && + state_.coverage_notes.count("manual_protocol_violation") == 0 + ? 1 + : 0; +} + +bool CaptureGate::AwaitAtExit() { + std::unique_lock lock(mu_); + if (state_.lifecycle == GateLifecycle::NotRequested || + state_.lifecycle == GateLifecycle::Finished) { + return true; + } + exit_requested_ = true; // a worker still in its delay stands down + cv_.notify_all(); + if (WaitLocked(lock, /*until_finished=*/true)) return true; + state_.coverage_notes.insert("exit_while_initializing"); + PublishLocked(); + lock.unlock(); + return false; +} + +void CaptureGate::AddCoverageNote(const std::string& note) { + std::unique_lock lock(mu_); + if (!state_.coverage_notes.insert(note).second) return; + PublishLocked(); + lock.unlock(); +} + +WrapperScope::WrapperScope() { ++t_wrapper_depth; } +WrapperScope::~WrapperScope() { --t_wrapper_depth; } + +void AppendGateCapabilities(const GateSnapshot& s, CaptureCapabilitiesEvent& evt) { + std::vector reasons; + if (s.readiness == GateReadiness::Pending) { + reasons.emplace_back("readiness_pending"); + } else if (s.readiness == GateReadiness::Failed) { + reasons.emplace_back("readiness_failed"); + } + for (const std::string& note : s.coverage_notes) reasons.push_back(note); + // Rows the timestamp counters marked (see capture_capability_resolver). + for (const CaptureCapability& cap : evt.capabilities) { + const std::string& code = cap.reason_code; + const auto ends_with = [&code](const char* suffix) { + const std::size_t n = std::strlen(suffix); + return code.size() >= n && code.compare(code.size() - n, n, suffix) == 0; + }; + if (ends_with("_timestamps_unavailable") || ends_with("_timestamps_invalid") || + ends_with("_timestamps_partial")) { + reasons.push_back(code); + } + } + + CaptureCapability readiness; + readiness.feature = "capture_readiness"; + readiness.requested = true; + readiness.status = GateReadinessName(s.readiness); + readiness.mode = GateOwnerName(s.owner); + readiness.reason_code = s.reason; + readiness.message = std::string("lifecycle=") + GateLifecycleName(s.lifecycle) + + " callbacks=" + std::to_string(s.callbacks_seen) + + " protocol=" + (s.protocol.empty() ? "absent" : s.protocol); + evt.capabilities.push_back(std::move(readiness)); + + CaptureCapability coverage; + coverage.feature = "capture_coverage"; + coverage.requested = true; + coverage.mode = "injection_gate"; + if (reasons.empty()) { + coverage.status = "no_known_loss"; + coverage.message = + "No known capture loss. This is not a completeness proof: a kernel " + "that produced no record cannot be counted."; + } else { + coverage.status = "incomplete"; + for (const std::string& reason : reasons) { + if (!coverage.reason_code.empty()) coverage.reason_code += ','; + coverage.reason_code += reason; + } + coverage.message = "The timeline of this session may be incomplete."; + } + evt.capabilities.push_back(std::move(coverage)); +} + +std::string GateStatusJson(const GateSnapshot& s, const std::string& nonce, + const long pid) { + std::string out; + out.reserve(1024); + out += "{\"schema\":1,\"nonce\":"; + AppendJsonString(out, nonce); + out += ",\"pid\":" + std::to_string(pid); + out += ",\"seq\":" + std::to_string(s.seq); + out += ",\"lifecycle\":\""; + out += GateLifecycleName(s.lifecycle); + out += "\",\"owner\":\""; + out += GateOwnerName(s.owner); + out += "\",\"readiness\":\""; + out += GateReadinessName(s.readiness); + out += "\",\"reason\":"; + AppendJsonString(out, s.reason); + out += ",\"coverage_notes\":["; + bool first = true; + for (const std::string& note : s.coverage_notes) { + if (!first) out += ','; + first = false; + AppendJsonString(out, note); + } + out += "],\"deadline_monotonic_ns\":" + std::to_string(s.deadline_ns); + AppendParticipant(out, "trigger", s.trigger); + AppendParticipant(out, "worker", s.worker); + AppendParticipant(out, "handshake", s.handshake); + out += ",\"callbacks_seen\":" + std::to_string(s.callbacks_seen); + out += ",\"handshake_reentries\":" + std::to_string(s.handshake_reentries); + out += ",\"wait_config_invalid\":"; + out += s.wait_config_invalid ? "true" : "false"; + out += ",\"protocol\":"; + AppendJsonString(out, s.protocol.substr(0, 16)); + out += "}\n"; + return out; +} + +bool WriteGateStatusFile(const std::string& path, const std::string& json) { +#if defined(_WIN32) + (void)path; + (void)json; + return false; +#else + if (path.empty() || json.size() > 4096) return false; + const std::string temp = path + ".tmp"; + const int fd = ::open(temp.c_str(), O_WRONLY | O_CREAT | O_TRUNC | O_CLOEXEC, + 0600); + if (fd < 0) return false; + std::size_t written = 0; + bool ok = true; + while (written < json.size()) { + const ssize_t n = ::write(fd, json.data() + written, json.size() - written); + if (n < 0) { + if (errno == EINTR) continue; + ok = false; + break; + } + written += static_cast(n); + } + ok = ok && ::fsync(fd) == 0; + ok = ::close(fd) == 0 && ok; + ok = ok && ::rename(temp.c_str(), path.c_str()) == 0; + if (!ok) ::unlink(temp.c_str()); + return ok; +#endif +} + +} // namespace gpufl::inject diff --git a/include/gpufl/inject/capture_gate.hpp b/include/gpufl/inject/capture_gate.hpp new file mode 100644 index 0000000..cbea6a9 --- /dev/null +++ b/include/gpufl/inject/capture_gate.hpp @@ -0,0 +1,227 @@ +#pragma once + +// Capture readiness for Linux CUDA injection. +// +// The injected library starts CUPTI from a worker thread after the driver's +// own initialization. Without a gate, the application thread continues into +// context creation and module loading while that worker attaches. On an L4 +// (2026-10-16) whole Trace sessions received every kernel record with zero +// timestamps; that is consistent with this overlap, but the CUPTI-internal +// cause is not established. The gate makes the application's first supported +// CUDA call wait until capture is ready, within one bounded deadline, and +// records every other outcome. +// +// Limits: status publication runs under the gate mutex, so a slow filesystem +// can hold waiters past the nominal deadline (the launcher watchdog bounds +// that only after its first accepted record). Coverage is a weaker claim +// than completeness, and nothing here is evidence about real driver locks. +// +// Three separate facts, see the design (agent_comm l4-qualification-20261015, +// client-capture-gate-design-v3.md + v3.1): +// * lifecycle - what the initializer is doing (diagnostics only); +// * readiness - latched once: Ready or Failed(reason); +// * coverage notes - reasons the run cannot be reported as no_known_loss. +// +// The component is pure: driver calls, the worker thread, the clock and the +// status file are hooks, so the state machine runs on CPU in unit tests. + +#include +#include +#include +#include +#include +#include +#include + +namespace gpufl { +struct CaptureCapabilitiesEvent; +} // namespace gpufl + +namespace gpufl::inject { + +enum class GateLifecycle : int { + NotRequested = 0, + DriverInitializing, + DriverReady, + CaptureInitializing, + Finished, +}; + +enum class GateOwner : int { None = 0, Gate, Callback, Manual }; + +enum class GateReadiness : int { Pending = 0, Ready, Failed }; + +enum class GateMode : int { + // The first supported CUDA call triggers driver init and waits. + Automatic = 0, + // The host calls GpuFlightInitializeInjectionAfterCuda after cuInit. + Manual, + // --warmup: capture starts late on purpose; nothing waits. + Warmup, +}; + +struct GateConfig { + GateMode mode = GateMode::Automatic; + int64_t wait_ns = 15'000'000'000LL; + bool wait_config_invalid = false; + // Worker sleep between the driver barrier and CUPTI setup: the readiness + // delay, or the warmup in Warmup mode. + int delay_ms = 1; + // The launcher's protocol variable was absent while a legacy delay above + // the default was set: it may have been a warmup. + bool legacy_launcher_delay = false; + // Protocol the launcher announced ("" when absent). + std::string protocol; +}; + +/** Reads the gate configuration from the environment (getenv by default). */ +GateConfig GateConfigFromEnv( + const std::function& getenv_fn); + +struct GateStep { + bool ok = false; + // Machine-readable cause when !ok, e.g. "cuinit_failed:100". + std::string reason; +}; + +struct GateParticipant { + bool expected = false; + bool returned = false; + bool ok = false; +}; + +struct GateSnapshot { + uint64_t seq = 0; + GateLifecycle lifecycle = GateLifecycle::NotRequested; + GateOwner owner = GateOwner::None; + GateReadiness readiness = GateReadiness::Pending; + std::string reason; + std::set coverage_notes; + // CLOCK_MONOTONIC ns; 0 until an owner is elected. + int64_t deadline_ns = 0; + GateParticipant trigger; + GateParticipant worker; + GateParticipant handshake; + uint64_t callbacks_seen = 0; + uint64_t handshake_reentries = 0; + bool wait_config_invalid = false; + std::string protocol; +}; + +struct GateHooks { + // CLOCK_MONOTONIC in ns (clock_gettime), the value the launcher compares. + std::function now_ns; + // Starts `body` on a new detached thread; false when no thread could start. + std::function body)> spawn; + // The worker's checked post-driver barrier (cuInit). + std::function driver_barrier; + // gpufl::init plus the Trace facility check. Never throws out. + std::function initialize; + // Publishes the state (the status file), with the gate mutex held so a + // record is written before anyone acts on it. Must not call back into the + // gate or into CUDA; snapshots carry a strictly increasing seq. + std::function publish; +}; + +const char* GateLifecycleName(GateLifecycle value); +const char* GateOwnerName(GateOwner value); +const char* GateReadinessName(GateReadiness value); + +class CaptureGate { + public: + CaptureGate(GateConfig config, GateHooks hooks); + CaptureGate(const CaptureGate&) = delete; + CaptureGate& operator=(const CaptureGate&) = delete; + + /** + * A supported CUDA entry point, before its real call. `trigger` performs + * that route's driver initialization (runtime: the real + * cudaGetDeviceCount; driver: the checked cuInit) and is called at most + * once per process, by the elected gate owner. + * + * Returns at once when readiness is settled, on gpufl's own threads, inside + * the injection callback, inside the trigger, and in Manual or Warmup mode. + * Otherwise waits until readiness settles or the deadline passes. + */ + void Enter(const std::function& trigger); + + /** NVIDIA's InitializeInjection (Linux). Never waits. */ + void OnInjectionCallback(); + + /** + * GpuFlightInitializeInjectionAfterCuda. In Manual mode it is the only + * initializer; returns 1 only for Ready with no protocol violation. + * `driver_trigger` is used only in Automatic mode, as an outer entry. + */ + int ManualHandshake(const std::function& driver_trigger); + + /** + * The atexit guard. true: ordinary teardown may continue (nothing ran, or + * every participant returned). false: an initializer is still live after + * the deadline; the caller must end the process without teardown. + */ + bool AwaitAtExit(); + + /** Adds a coverage note decided outside the gate (e.g. at exit). */ + void AddCoverageNote(const std::string& note); + + GateSnapshot Snapshot() const; + + /** One acquire load: readiness settled or the deadline already passed. */ + bool Settled() const; + + private: + // All under mu_. + void ElectLocked(GateOwner owner); + void LatchFailedLocked(const std::string& reason); + void ResolveLocked(); + void PublishLocked(); + bool WaitLocked(std::unique_lock& lock, bool until_finished); + void RunWorker(); + void StartWorkerLocked(); + + const GateConfig config_; + const GateHooks hooks_; + + mutable std::mutex mu_; + std::condition_variable cv_; + GateSnapshot state_; + bool worker_started_ = false; + bool exit_requested_ = false; + std::atomic settled_{false}; +}; + +/** + * Counts nested gpufl CUDA wrappers on this thread. Only the outermost one + * gates: an inner one is the real call of an outer wrapper reaching another + * wrapped symbol. This cannot see CUDA work an unwrapped entry started. + */ +class WrapperScope { + public: + WrapperScope(); + ~WrapperScope(); + WrapperScope(const WrapperScope&) = delete; + WrapperScope& operator=(const WrapperScope&) = delete; +}; + +/** + * Appends capture_readiness and capture_coverage rows. Coverage is + * "no_known_loss" only when readiness is Ready, no coverage note is set and no + * kernel or memcpy row reports rejected timestamps. That is weaker than + * completeness: a kernel that produced no record cannot be counted. + */ +void AppendGateCapabilities(const GateSnapshot& snapshot, + CaptureCapabilitiesEvent& evt); + +/** The snapshot as the bootstrap status JSON (schema 1, at most 4 KiB). */ +std::string GateStatusJson(const GateSnapshot& snapshot, const std::string& nonce, + long pid); + +/** + * Writes the status JSON to `path` with temp file + fsync + rename. Only + * open/write/fsync/rename/close, so it needs neither the logger nor the + * driver. Returns false when any step failed. + */ +bool WriteGateStatusFile(const std::string& path, const std::string& json); + +} // namespace gpufl::inject diff --git a/include/gpufl/inject/capture_gate_linux.cpp b/include/gpufl/inject/capture_gate_linux.cpp new file mode 100644 index 0000000..677d76d --- /dev/null +++ b/include/gpufl/inject/capture_gate_linux.cpp @@ -0,0 +1,245 @@ +// NVIDIA calls InitializeInjection from inside the driver's own cuInit, so +// CUPTI cannot start there; a worker starts it after the driver is up, while +// the application thread continues into context and module setup. The L4 +// sessions that lost every kernel timestamp (2026-10-16) are consistent with +// that overlap; the CUPTI-internal cause is not established. The gate makes +// the first supported CUDA call trigger driver init itself and wait, within +// one deadline, until capture is ready, and it records every route and failure +// it cannot cover. + +#if !defined(_WIN32) + +#include "gpufl/inject/capture_gate_linux.hpp" + +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "gpufl/core/env_vars.hpp" + +namespace gpufl::inject { +namespace { + +std::atomic g_initialize{nullptr}; +std::atomic g_after_initialization_exit{nullptr}; +std::atomic g_on_gate_created{nullptr}; + +const char* EnvOrNull(const char* name) { + const char* value = std::getenv(name); + return value && *value ? value : nullptr; +} + +int64_t MonotonicNowNs() { + timespec ts{}; + clock_gettime(CLOCK_MONOTONIC, &ts); + return static_cast(ts.tv_sec) * 1'000'000'000LL + ts.tv_nsec; +} + +std::string DlerrorText() { + const char* error = dlerror(); + return error ? error : "unknown"; +} + +bool ResolvesIntoThisLibrary(void* symbol) { + Dl_info target{}; + Dl_info self{}; + if (dladdr(symbol, &target) == 0 || + dladdr(reinterpret_cast(&MonotonicNowNs), &self) == 0) { + return false; + } + return target.dli_fbase == self.dli_fbase; +} + +using CuInitFn = int (*)(unsigned int); + +// The driver's own cuInit: the already-loaded libcuda.so.1 first, then the +// global scope (libcuda may sit ahead of this library). Never loads a driver. +GateStep ResolveCuInit(CuInitFn& out) { + dlerror(); + void* symbol = nullptr; + if (void* handle = dlopen("libcuda.so.1", RTLD_NOLOAD | RTLD_LAZY)) { + dlerror(); + symbol = dlsym(handle, "cuInit"); + dlclose(handle); // drops only the reference RTLD_NOLOAD took + } + if (!symbol) { + dlerror(); + symbol = dlsym(RTLD_DEFAULT, "cuInit"); + if (!symbol) return {false, "cuinit_symbol_missing:" + DlerrorText()}; + } + if (ResolvesIntoThisLibrary(symbol)) return {false, "cuinit_resolves_to_self"}; + out = reinterpret_cast(symbol); + return {true, ""}; +} + +// Driver route trigger, and the worker's post-driver barrier: cuInit is +// idempotent and blocks on the driver's init lock. +GateStep CheckedCuInit() { + CuInitFn fn = nullptr; + GateStep found = ResolveCuInit(fn); + if (!found.ok) return found; + const int rc = fn(0); + if (rc != 0) return {false, "cuinit_failed:" + std::to_string(rc)}; + return {true, ""}; +} + +// Runtime route trigger: the real cudaGetDeviceCount from the same RTLD_NEXT +// chain the wrapper forwards to, so the driver is the one cudart uses. It runs +// cuInit without creating a context. +GateStep RuntimeTrigger() { + using CudaGetDeviceCountFn = int (*)(int*); + static auto* fn = + reinterpret_cast(dlsym(RTLD_NEXT, "cudaGetDeviceCount")); + if (!fn) return {false, "runtime_symbol_missing"}; + int count = 0; + const int rc = fn(&count); + if (rc != 0) return {false, "runtime_init_failed:" + std::to_string(rc)}; + return {true, ""}; +} + +GateStep BoundInitialize() { + GateStep (*initialize)() = g_initialize.load(std::memory_order_acquire); + if (!initialize) return {false, "initializer_unbound"}; + return initialize(); +} + +struct StatusTarget { + std::string dir; + std::string nonce; + bool debug = false; + std::atomic printed_readiness{-1}; +}; + +StatusTarget& Status() { + static auto* target = [] { + auto* t = new StatusTarget(); + if (const char* dir = EnvOrNull(env::kInjectStatusDir)) t->dir = dir; + if (const char* nonce = EnvOrNull(env::kInjectStatusNonce)) t->nonce = nonce; + if (const char* debug = EnvOrNull(env::kDebugOutput)) { + t->debug = std::strcmp(debug, "1") == 0; + } + return t; + }(); + return *target; +} + +// Logger-independent: the status file and one stderr line under GPUFL_DEBUG. +void PublishStatus(const GateSnapshot& snapshot) { + StatusTarget& target = Status(); + const long pid = static_cast(::getpid()); + if (!target.dir.empty()) { + // Per pid: a child that inherited the environment writes its own file. + const std::string path = + target.dir + "/status-" + std::to_string(pid) + ".json"; + (void)WriteGateStatusFile(path, GateStatusJson(snapshot, target.nonce, pid)); + } + const int readiness = static_cast(snapshot.readiness); + if (target.debug && snapshot.readiness != GateReadiness::Pending && + target.printed_readiness.exchange(readiness) != readiness) { + std::fprintf(stderr, "[gpufl] capture readiness: %s owner=%s reason=%s\n", + GateReadinessName(snapshot.readiness), + GateOwnerName(snapshot.owner), + snapshot.reason.empty() ? "-" : snapshot.reason.c_str()); + } +} + +// atexit: either every initialization participant has returned and ordinary +// teardown may run, or one is still live after the deadline and the process +// ends here, without the handlers and destructors that would run under it. +// Handlers the application registered after this one already ran before it. +void ExitGuard() { + if (!LinuxCaptureGate().AwaitAtExit()) std::_Exit(kInitWedgedExitCode); + if (void (*after)() = g_after_initialization_exit.load(std::memory_order_acquire)) { + after(); + } +} + +void Enter(const bool driver_route) { + if (!LinuxInjectionRequested()) return; // an accidental preload stays inert + CaptureGate& gate = LinuxCaptureGate(); + if (gate.Settled()) return; + gate.Enter(driver_route ? CheckedCuInit : RuntimeTrigger); +} + +} // namespace + +void BindLinuxCaptureGate(const LinuxGateBindings& bindings) { + g_initialize.store(bindings.initialize, std::memory_order_release); + g_after_initialization_exit.store(bindings.after_initialization_exit, + std::memory_order_release); + g_on_gate_created.store(bindings.on_gate_created, std::memory_order_release); +} + +bool LinuxInjectionRequested() { + static const bool requested = [] { + const char* sentinel = EnvOrNull(env::kInject); + return sentinel && std::strcmp(sentinel, "1") == 0; + }(); + return requested; +} + +CaptureGate& LinuxCaptureGate() { + // Process lifetime: a late worker, or a call at exit, never touches a + // destroyed gate. + static auto* gate = [] { + GateHooks hooks; + hooks.now_ns = MonotonicNowNs; + hooks.spawn = [](std::function body) { + try { + std::thread(std::move(body)).detach(); + return true; + } catch (...) { + return false; + } + }; + hooks.driver_barrier = CheckedCuInit; + hooks.initialize = BoundInitialize; + hooks.publish = PublishStatus; + auto* created = new CaptureGate(GateConfigFromEnv(nullptr), std::move(hooks)); + if (void (*notify)() = g_on_gate_created.load(std::memory_order_acquire)) { + notify(); + } + // Registered before init can register shutdown, so shutdown runs first. + std::atexit(ExitGuard); + return created; + }(); + return *gate; +} + +void LinuxInjectionCallback() { + if (LinuxInjectionRequested()) LinuxCaptureGate().OnInjectionCallback(); +} + +int LinuxManualHandshake() { + if (!LinuxInjectionRequested()) return 0; + return LinuxCaptureGate().ManualHandshake(CheckedCuInit); +} + +} // namespace gpufl::inject + +extern "C" { + +// Cross-TU capture-gate hooks used by the typed CUDA interpose translation +// unit, one per route. Hidden inside the library; only CUDA ABI symbols are +// public. +__attribute__((visibility("hidden"))) void GpuFlightCaptureGateRuntime() { + gpufl::inject::Enter(/*driver_route=*/false); +} + +__attribute__((visibility("hidden"))) void GpuFlightCaptureGateDriver() { + gpufl::inject::Enter(/*driver_route=*/true); +} + +} // extern "C" + +#endif // !defined(_WIN32) diff --git a/include/gpufl/inject/capture_gate_linux.hpp b/include/gpufl/inject/capture_gate_linux.hpp new file mode 100644 index 0000000..e12c453 --- /dev/null +++ b/include/gpufl/inject/capture_gate_linux.hpp @@ -0,0 +1,44 @@ +#pragma once + +// The Linux injection glue around CaptureGate: driver-init triggers, the +// status file, the exit guard and the process-wide gate. Separate from +// inject_entry.cpp so the CPU end-to-end tests can load it, with the real +// typed wrappers, into a process with fake CUDA libraries. + +#include "gpufl/inject/capture_gate.hpp" + +namespace gpufl::inject { + +// The documented exit status when the process ends at exit with an +// initializer still inside CUDA/CUPTI after the deadline. +constexpr int kInitWedgedExitCode = 86; + +struct LinuxGateBindings { + // gpufl::init plus the Trace facility check, on the initializer thread. + GateStep (*initialize)() = nullptr; + // At exit, after every initialization participant returned. + void (*after_initialization_exit)() = nullptr; + // Called once when the gate is created (e.g. to register a capability + // contributor). + void (*on_gate_created)() = nullptr; +}; + +/** + * Binds the initializer. Called from the injection library's static + * initialization, before any CUDA call can reach the gate. + */ +void BindLinuxCaptureGate(const LinuxGateBindings& bindings); + +/** GPUFL_INJECT=1: anything else is an accidental preload and stays inert. */ +bool LinuxInjectionRequested(); + +/** The process-lifetime gate (created on first use). */ +CaptureGate& LinuxCaptureGate(); + +/** InitializeInjection's Linux body. */ +void LinuxInjectionCallback(); + +/** GpuFlightInitializeInjectionAfterCuda's Linux body. */ +int LinuxManualHandshake(); + +} // namespace gpufl::inject diff --git a/include/gpufl/inject/cuda_interpose_linux.cpp b/include/gpufl/inject/cuda_interpose_linux.cpp index ed81e29..9eeb967 100644 --- a/include/gpufl/inject/cuda_interpose_linux.cpp +++ b/include/gpufl/inject/cuda_interpose_linux.cpp @@ -13,10 +13,13 @@ #include #include +#include "gpufl/inject/capture_gate.hpp" + extern "C" { -void GpuFlightWaitAtCudaLaunchBoundary(); -void GpuFlightWaitAtCudaSyncBoundary(); -void GpuFlightWaitAtCudaMemoryBoundary(); +// The capture gate, by route: runtime wrappers trigger driver init through the +// real cudaGetDeviceCount, driver wrappers through the checked cuInit. +void GpuFlightCaptureGateRuntime(); +void GpuFlightCaptureGateDriver(); } namespace { @@ -34,6 +37,7 @@ Fn ResolveNext(const char* symbol) noexcept { #define GPUFL_CUDA_INTERPOSE(WAIT, RETURN_TYPE, NAME, PARAMS, ARGS, FAILURE) \ extern "C" __attribute__((visibility("default"))) \ RETURN_TYPE NAME PARAMS { \ + gpufl::inject::WrapperScope gpufl_wrapper_scope; \ WAIT(); \ using Function = RETURN_TYPE (*) PARAMS; \ static Function resolved_function = ResolveNext(#NAME); \ @@ -44,27 +48,27 @@ Fn ResolveNext(const char* symbol) noexcept { // with ABI-compatible opaque types. Keeping them here means the same official // headers now validate every interposed CUDA function. GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaLaunchBoundary, cudaError_t, __cudaLaunchKernel, + GpuFlightCaptureGateRuntime, cudaError_t, __cudaLaunchKernel, (const void* function_address, dim3 grid_dim, dim3 block_dim, void** args, std::size_t shared_mem, cudaStream_t stream), (function_address, grid_dim, block_dim, args, shared_mem, stream), cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaLaunchBoundary, cudaError_t, cudaLaunchKernel, + GpuFlightCaptureGateRuntime, cudaError_t, cudaLaunchKernel, (const void* function_address, dim3 grid_dim, dim3 block_dim, void** args, std::size_t shared_mem, cudaStream_t stream), (function_address, grid_dim, block_dim, args, shared_mem, stream), cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaLaunchBoundary, cudaError_t, cudaLaunchKernelExC, + GpuFlightCaptureGateRuntime, cudaError_t, cudaLaunchKernelExC, (const cudaLaunchConfig_t* config, const void* function_address, void** args), (config, function_address, args), cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaLaunchBoundary, CUresult, cuLaunchKernel, + GpuFlightCaptureGateDriver, CUresult, cuLaunchKernel, (CUfunction function, unsigned int grid_x, unsigned int grid_y, unsigned int grid_z, unsigned int block_x, unsigned int block_y, unsigned int block_z, unsigned int shared_mem, CUstream stream, @@ -74,43 +78,43 @@ GPUFL_CUDA_INTERPOSE( CUDA_ERROR_UNKNOWN) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaSyncBoundary, cudaError_t, cudaDeviceSynchronize, + GpuFlightCaptureGateRuntime, cudaError_t, cudaDeviceSynchronize, (), (), cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaSyncBoundary, cudaError_t, cudaStreamSynchronize, + GpuFlightCaptureGateRuntime, cudaError_t, cudaStreamSynchronize, (cudaStream_t stream), (stream), cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaSyncBoundary, CUresult, cuCtxSynchronize, + GpuFlightCaptureGateDriver, CUresult, cuCtxSynchronize, (), (), CUDA_ERROR_UNKNOWN) // CUDA Runtime API: synchronous, asynchronous, pitched, 3D, and peer copies. GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, cudaError_t, cudaMemcpy, + GpuFlightCaptureGateRuntime, cudaError_t, cudaMemcpy, (void* dst, const void* src, std::size_t bytes, cudaMemcpyKind kind), (dst, src, bytes, kind), cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, cudaError_t, cudaMemcpyAsync, + GpuFlightCaptureGateRuntime, cudaError_t, cudaMemcpyAsync, (void* dst, const void* src, std::size_t bytes, cudaMemcpyKind kind, cudaStream_t stream), (dst, src, bytes, kind, stream), cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, cudaError_t, cudaMemcpyAsync_ptsz, + GpuFlightCaptureGateRuntime, cudaError_t, cudaMemcpyAsync_ptsz, (void* dst, const void* src, std::size_t bytes, cudaMemcpyKind kind, cudaStream_t stream), (dst, src, bytes, kind, stream), cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, cudaError_t, cudaMemcpy2D, + GpuFlightCaptureGateRuntime, cudaError_t, cudaMemcpy2D, (void* dst, std::size_t dst_pitch, const void* src, std::size_t src_pitch, std::size_t width, std::size_t height, cudaMemcpyKind kind), (dst, dst_pitch, src, src_pitch, width, height, kind), cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, cudaError_t, cudaMemcpy2DAsync, + GpuFlightCaptureGateRuntime, cudaError_t, cudaMemcpy2DAsync, (void* dst, std::size_t dst_pitch, const void* src, std::size_t src_pitch, std::size_t width, std::size_t height, cudaMemcpyKind kind, cudaStream_t stream), @@ -118,7 +122,7 @@ GPUFL_CUDA_INTERPOSE( cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, cudaError_t, cudaMemcpy2DAsync_ptsz, + GpuFlightCaptureGateRuntime, cudaError_t, cudaMemcpy2DAsync_ptsz, (void* dst, std::size_t dst_pitch, const void* src, std::size_t src_pitch, std::size_t width, std::size_t height, cudaMemcpyKind kind, cudaStream_t stream), @@ -126,48 +130,48 @@ GPUFL_CUDA_INTERPOSE( cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, cudaError_t, cudaMemcpy3D, + GpuFlightCaptureGateRuntime, cudaError_t, cudaMemcpy3D, (const cudaMemcpy3DParms* params), (params), cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, cudaError_t, cudaMemcpy3DAsync, + GpuFlightCaptureGateRuntime, cudaError_t, cudaMemcpy3DAsync, (const cudaMemcpy3DParms* params, cudaStream_t stream), (params, stream), cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, cudaError_t, cudaMemcpy3DAsync_ptsz, + GpuFlightCaptureGateRuntime, cudaError_t, cudaMemcpy3DAsync_ptsz, (const cudaMemcpy3DParms* params, cudaStream_t stream), (params, stream), cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, cudaError_t, cudaMemcpyPeer, + GpuFlightCaptureGateRuntime, cudaError_t, cudaMemcpyPeer, (void* dst, int dst_device, const void* src, int src_device, std::size_t bytes), (dst, dst_device, src, src_device, bytes), cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, cudaError_t, cudaMemcpyPeerAsync, + GpuFlightCaptureGateRuntime, cudaError_t, cudaMemcpyPeerAsync, (void* dst, int dst_device, const void* src, int src_device, std::size_t bytes, cudaStream_t stream), (dst, dst_device, src, src_device, bytes, stream), cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, cudaError_t, cudaMemcpyPeerAsync_ptsz, + GpuFlightCaptureGateRuntime, cudaError_t, cudaMemcpyPeerAsync_ptsz, (void* dst, int dst_device, const void* src, int src_device, std::size_t bytes, cudaStream_t stream), (dst, dst_device, src, src_device, bytes, stream), cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, cudaError_t, cudaMemcpy3DPeer, + GpuFlightCaptureGateRuntime, cudaError_t, cudaMemcpy3DPeer, (const cudaMemcpy3DPeerParms* params), (params), cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, cudaError_t, cudaMemcpy3DPeerAsync, + GpuFlightCaptureGateRuntime, cudaError_t, cudaMemcpy3DPeerAsync, (const cudaMemcpy3DPeerParms* params, cudaStream_t stream), (params, stream), cudaErrorUnknown) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, cudaError_t, cudaMemcpy3DPeerAsync_ptsz, + GpuFlightCaptureGateRuntime, cudaError_t, cudaMemcpy3DPeerAsync_ptsz, (const cudaMemcpy3DPeerParms* params, cudaStream_t stream), (params, stream), cudaErrorUnknown) @@ -176,39 +180,39 @@ GPUFL_CUDA_INTERPOSE( // streams. Export both spellings so the readiness barrier cannot be bypassed // by that compile-time mode. GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, CUresult, cuMemcpy, + GpuFlightCaptureGateDriver, CUresult, cuMemcpy, (CUdeviceptr dst, CUdeviceptr src, std::size_t bytes), (dst, src, bytes), CUDA_ERROR_UNKNOWN) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, CUresult, cuMemcpy_ptds, + GpuFlightCaptureGateDriver, CUresult, cuMemcpy_ptds, (CUdeviceptr dst, CUdeviceptr src, std::size_t bytes), (dst, src, bytes), CUDA_ERROR_UNKNOWN) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, CUresult, cuMemcpyAsync, + GpuFlightCaptureGateDriver, CUresult, cuMemcpyAsync, (CUdeviceptr dst, CUdeviceptr src, std::size_t bytes, CUstream stream), (dst, src, bytes, stream), CUDA_ERROR_UNKNOWN) GPUFL_CUDA_INTERPOSE( - GpuFlightWaitAtCudaMemoryBoundary, CUresult, cuMemcpyAsync_ptsz, + GpuFlightCaptureGateDriver, CUresult, cuMemcpyAsync_ptsz, (CUdeviceptr dst, CUdeviceptr src, std::size_t bytes, CUstream stream), (dst, src, bytes, stream), CUDA_ERROR_UNKNOWN) #define GPUFL_DRIVER_COPY_SYNC_PAIR(NAME, PARAMS, ARGS) \ GPUFL_CUDA_INTERPOSE( \ - GpuFlightWaitAtCudaMemoryBoundary, CUresult, NAME, PARAMS, ARGS, \ + GpuFlightCaptureGateDriver, CUresult, NAME, PARAMS, ARGS, \ CUDA_ERROR_UNKNOWN) \ GPUFL_CUDA_INTERPOSE( \ - GpuFlightWaitAtCudaMemoryBoundary, CUresult, NAME##_ptds, PARAMS, \ + GpuFlightCaptureGateDriver, CUresult, NAME##_ptds, PARAMS, \ ARGS, CUDA_ERROR_UNKNOWN) #define GPUFL_DRIVER_COPY_ASYNC_PAIR(NAME, PARAMS, ARGS) \ GPUFL_CUDA_INTERPOSE( \ - GpuFlightWaitAtCudaMemoryBoundary, CUresult, NAME, PARAMS, ARGS, \ + GpuFlightCaptureGateDriver, CUresult, NAME, PARAMS, ARGS, \ CUDA_ERROR_UNKNOWN) \ GPUFL_CUDA_INTERPOSE( \ - GpuFlightWaitAtCudaMemoryBoundary, CUresult, NAME##_ptsz, PARAMS, \ + GpuFlightCaptureGateDriver, CUresult, NAME##_ptsz, PARAMS, \ ARGS, CUDA_ERROR_UNKNOWN) GPUFL_DRIVER_COPY_SYNC_PAIR( @@ -283,6 +287,82 @@ GPUFL_DRIVER_COPY_ASYNC_PAIR( (const CUDA_MEMCPY3D_PEER* params, CUstream stream), (params, stream)) +// First-touch calls: real first CUDA calls of the backend's lab sources and of +// common hosts. Gating them makes driver init, CUPTI subscription and activity +// enable complete before the application's own context, allocation or module +// work, instead of racing it. +GPUFL_CUDA_INTERPOSE( + GpuFlightCaptureGateRuntime, cudaError_t, cudaMalloc, + (void** pointer, std::size_t bytes), (pointer, bytes), cudaErrorUnknown) + +GPUFL_CUDA_INTERPOSE( + GpuFlightCaptureGateRuntime, cudaError_t, cudaMallocManaged, + (void** pointer, std::size_t bytes, unsigned int flags), + (pointer, bytes, flags), cudaErrorUnknown) + +GPUFL_CUDA_INTERPOSE( + GpuFlightCaptureGateRuntime, cudaError_t, cudaMallocHost, + (void** pointer, std::size_t bytes), (pointer, bytes), cudaErrorUnknown) + +GPUFL_CUDA_INTERPOSE( + GpuFlightCaptureGateRuntime, cudaError_t, cudaHostAlloc, + (void** pointer, std::size_t bytes, unsigned int flags), + (pointer, bytes, flags), cudaErrorUnknown) + +GPUFL_CUDA_INTERPOSE( + GpuFlightCaptureGateRuntime, cudaError_t, cudaFree, + (void* pointer), (pointer), cudaErrorUnknown) + +GPUFL_CUDA_INTERPOSE( + GpuFlightCaptureGateRuntime, cudaError_t, cudaSetDevice, + (int device), (device), cudaErrorUnknown) + +GPUFL_CUDA_INTERPOSE( + GpuFlightCaptureGateRuntime, cudaError_t, cudaGetDeviceCount, + (int* count), (count), cudaErrorUnknown) + +GPUFL_CUDA_INTERPOSE( + GpuFlightCaptureGateRuntime, cudaError_t, cudaGetDeviceProperties_v2, + (cudaDeviceProp* properties, int device), (properties, device), + cudaErrorUnknown) + +GPUFL_CUDA_INTERPOSE( + GpuFlightCaptureGateRuntime, cudaError_t, cudaMemGetInfo, + (std::size_t* free_bytes, std::size_t* total_bytes), + (free_bytes, total_bytes), cudaErrorUnknown) + +GPUFL_CUDA_INTERPOSE( + GpuFlightCaptureGateRuntime, cudaError_t, cudaEventCreate, + (cudaEvent_t* event), (event), cudaErrorUnknown) + +GPUFL_CUDA_INTERPOSE( + GpuFlightCaptureGateRuntime, cudaError_t, cudaEventCreateWithFlags, + (cudaEvent_t* event, unsigned int flags), (event, flags), cudaErrorUnknown) + +GPUFL_CUDA_INTERPOSE( + GpuFlightCaptureGateRuntime, cudaError_t, cudaStreamCreate, + (cudaStream_t* stream), (stream), cudaErrorUnknown) + +GPUFL_CUDA_INTERPOSE( + GpuFlightCaptureGateRuntime, cudaError_t, cudaStreamCreateWithFlags, + (cudaStream_t* stream, unsigned int flags), (stream, flags), + cudaErrorUnknown) + +GPUFL_CUDA_INTERPOSE( + GpuFlightCaptureGateDriver, CUresult, cuMemAlloc_v2, + (CUdeviceptr* pointer, std::size_t bytes), (pointer, bytes), + CUDA_ERROR_UNKNOWN) + +GPUFL_CUDA_INTERPOSE( + GpuFlightCaptureGateDriver, CUresult, cuModuleLoadData, + (CUmodule* module, const void* image), (module, image), CUDA_ERROR_UNKNOWN) + +GPUFL_CUDA_INTERPOSE( + GpuFlightCaptureGateDriver, CUresult, cuModuleLoadDataEx, + (CUmodule* module, const void* image, unsigned int option_count, + CUjit_option* options, void** option_values), + (module, image, option_count, options, option_values), CUDA_ERROR_UNKNOWN) + #undef GPUFL_DRIVER_COPY_ASYNC_PAIR #undef GPUFL_DRIVER_COPY_SYNC_PAIR #undef GPUFL_CUDA_INTERPOSE diff --git a/include/gpufl/inject/inject_entry.cpp b/include/gpufl/inject/inject_entry.cpp index 6676b3b..8f71a51 100644 --- a/include/gpufl/inject/inject_entry.cpp +++ b/include/gpufl/inject/inject_entry.cpp @@ -18,6 +18,7 @@ // CMake gates this TU on UNIX AND NOT APPLE. #include "gpufl/inject/inject_entry.hpp" +#include "gpufl/inject/capture_gate_linux.hpp" #include "gpufl/core/env_vars.hpp" @@ -29,6 +30,8 @@ #include #include #include +#include +#include #include #include #include @@ -88,19 +91,9 @@ namespace { std::once_flag g_init_once; -std::once_flag g_deferred_wait_atexit_once; std::once_flag g_shutdown_atexit_once; std::atomic g_init_ok{false}; -std::atomic g_deferred_init_started{false}; -std::atomic g_deferred_init_finished{false}; std::atomic g_shutdown_started{false}; -std::mutex g_deferred_init_mutex; -std::condition_variable g_deferred_init_cv; -// A CUDA call made by gpufl's own deferred-init worker must never wait for -// that worker to finish. Most CUPTI setup does not call the runtime copy -// APIs, but the interpose layer must be safe if a driver/toolkit revision -// does: waiting here would be a self-deadlock. -thread_local bool g_is_deferred_init_worker = false; // Captured during init for the atexit upload path (--upload). g_log_path // mirrors InitOptions.log_path; both are read only in shutdownAndSignal, @@ -175,8 +168,6 @@ void endProcessScope() { state.active = false; } -void waitForDeferredInit(); - // Non-empty env value, or std::nullopt if unset / empty. const char* envOrNull(const char* name) { const char* v = std::getenv(name); @@ -254,50 +245,6 @@ void shutdownAndSignal() { writeCompletionByteIfRequested(); } -void waitForDeferredInit() { - if (g_deferred_init_started.load(std::memory_order_acquire) && - !g_deferred_init_finished.load(std::memory_order_acquire)) { - std::unique_lock lock(g_deferred_init_mutex); - g_deferred_init_cv.wait(lock, [] { - return g_deferred_init_finished.load(std::memory_order_acquire); - }); - } -} - -#ifndef _WIN32 -// Boundary waits exist only to back the launch/sync interpose wrappers, -// which are LD_PRELOAD-only (Linux). On Windows there is no interpose, so -// these are compiled out to avoid unused-function diagnostics. -void waitForDeferredInitForMs(const int wait_ms) { - if (wait_ms <= 0) return; - if (g_is_deferred_init_worker) return; - if (g_deferred_init_started.load(std::memory_order_acquire) && - !g_deferred_init_finished.load(std::memory_order_acquire)) { - std::unique_lock lock(g_deferred_init_mutex); - g_deferred_init_cv.wait_for( - lock, std::chrono::milliseconds(wait_ms), [] { - return g_deferred_init_finished.load(std::memory_order_acquire); - }); - } -} - -void waitAtCudaSyncBoundary() { - waitForDeferredInitForMs(envIntOrDefault(gpufl::env::kInjectSyncWaitMs, 15000)); -} - -void waitAtCudaLaunchBoundary() { - waitForDeferredInitForMs(envIntOrDefault(gpufl::env::kInjectLaunchWaitMs, 15000)); -} - -void waitAtCudaMemoryBoundary() { - // A memory transfer is an activity-capture boundary for the same reason a - // launch is: if it crosses the target before CUPTI subscription completes, - // it is gone permanently. Reuse the launch timeout rather than introduce a - // second knob with identical semantics. - waitForDeferredInitForMs(envIntOrDefault(gpufl::env::kInjectLaunchWaitMs, 15000)); -} -#endif // !_WIN32 - struct NvtxDomainStorage { std::string name; @@ -526,15 +473,6 @@ void NVTX_API CounterSampleNoValue(nvtxDomainHandle_t, const uint64_t id, } // namespace nvtx_counters_impl #endif // GPUFL_HAS_NVTX_COUNTERS -void registerDeferredWaitAtexit() { - std::call_once(g_deferred_wait_atexit_once, [] { - // atexit handlers run in reverse registration order. Register this - // wait guard early, then register shutdown only after init succeeds, - // so shutdown runs first and this guard only covers failed/slow init. - std::atexit(waitForDeferredInit); - }); -} - void registerShutdownAtexit() { std::call_once(g_shutdown_atexit_once, [] { std::atexit(shutdownAndSignal); @@ -639,71 +577,56 @@ void doInjectInit() { } -void markDeferredInitFinished() { - { - std::lock_guard lock(g_deferred_init_mutex); - g_deferred_init_finished.store(true, std::memory_order_release); +#ifndef _WIN32 +// Capture readiness on Linux: capture_gate_linux.cpp owns the gate, its +// driver-init triggers, status file and exit guard. This library binds the +// real initializer to it (gpufl::init plus the Trace facility check). +gpufl::inject::GateStep initializeCapture() { + try { + std::call_once(g_init_once, doInjectInit); + } catch (...) { + return {false, "init_exception"}; + } + if (!g_init_ok.load(std::memory_order_acquire)) return {false, "init_failed"}; + std::string why; + if (!gpufl::Monitor::CaptureFacilitiesReady(&why)) { + return {false, "trace_facility_unavailable:" + why}; } - g_deferred_init_cv.notify_all(); + return {true, ""}; } -// Blocks until the CUDA driver has finished its own initialization. -// -// cuInit is idempotent and thread-safe, so calling it here just waits on the -// driver's init lock. Resolved at runtime rather than linked: this file -// interposes CUDA symbols and cannot include cuda.h without colliding with -// its own declarations. -void waitForCudaDriverInit() { - using CuInitFn = int (*)(unsigned int); - CuInitFn fn = nullptr; -#ifdef _WIN32 - if (HMODULE cuda = GetModuleHandleA("nvcuda.dll")) { - fn = reinterpret_cast(GetProcAddress(cuda, "cuInit")); +void contributeGateCapabilities(gpufl::CaptureCapabilitiesEvent& evt) { + gpufl::inject::GateSnapshot snapshot = gpufl::inject::LinuxCaptureGate().Snapshot(); + // A required kind that failed to re-enable after readiness (for example at + // a deferred engine start) still makes this session incomplete. + std::string why; + if (snapshot.readiness == gpufl::inject::GateReadiness::Ready && + !gpufl::Monitor::CaptureFacilitiesReady(&why)) { + snapshot.coverage_notes.insert("trace_facility_lost"); } -#else - // RTLD_DEFAULT, not RTLD_NEXT: libcuda loaded this library and may sit - // ahead of it in the search order. - fn = reinterpret_cast(dlsym(RTLD_DEFAULT, "cuInit")); -#endif - if (fn) fn(0); + gpufl::inject::AppendGateCapabilities(snapshot, evt); } -void startDeferredInjectInit() { - registerDeferredWaitAtexit(); - - bool expected = false; - if (!g_deferred_init_started.compare_exchange_strong( - expected, true, std::memory_order_acq_rel)) { - return; - } +void registerGateCapabilityContributor() { + gpufl::Monitor::SetCaptureCapabilityContributor(contributeGateCapabilities); +} - std::thread([] { - g_is_deferred_init_worker = true; - // NVIDIA calls InitializeInjection from inside the CUDA driver's own - // initialization. cuptiSubscribe probes driver state, so reaching it - // while that initialization is still running faults inside libcuda. - // - // A sleep cannot fix this. Measured on an RTX 3090 / driver 610.43, - // `gpufl trace --passes Trace` crashed nondeterministically at every - // length tried - 3/10 at 5ms, 5/10 at 40ms, 0/10 at 0, 10, 20, 30 and - // 50ms - so the delay only moves the odds. cuInit is idempotent and - // thread-safe, and blocks on the driver's own init lock, which is the - // barrier the sleep was standing in for. Engines other than Trace hid - // this by doing slow PerfWorks setup before subscribing. - waitForCudaDriverInit(); - // The delay knob stays: it is what --warmup uses to skip cold start, - // and it is now measured from a driver that is actually up. - sleepMs(envIntOrDefault(gpufl::env::kInjectInitDelayMs, 1)); - try { - std::call_once(g_init_once, doInjectInit); - } catch (...) { - // Injection must never throw through libcuda's callback path. - } - markDeferredInitFinished(); - g_is_deferred_init_worker = false; - }).detach(); +void shutdownAfterInitialization() { + if (g_init_ok.load(std::memory_order_acquire)) shutdownAndSignal(); } +// Bound at this library's static initialization, before any CUDA call can +// reach the gate. +const bool g_capture_gate_bound = [] { + gpufl::inject::LinuxGateBindings bindings; + bindings.initialize = initializeCapture; + bindings.after_initialization_exit = shutdownAfterInitialization; + bindings.on_gate_created = registerGateCapabilityContributor; + gpufl::inject::BindLinuxCaptureGate(bindings); + return true; +}(); +#endif // !_WIN32 + } // namespace #if defined(GPUFL_AMD_PRELOAD_STARTUP) @@ -725,22 +648,9 @@ void initializeAmdInjection() noexcept { extern "C" { #ifndef _WIN32 -// Cross-TU boundary hooks used by the typed CUDA interpose translation unit. -// They remain hidden inside libgpufl_inject; only CUDA ABI symbols are public. -__attribute__((visibility("hidden"))) -void GpuFlightWaitAtCudaLaunchBoundary() { - waitAtCudaLaunchBoundary(); -} - -__attribute__((visibility("hidden"))) -void GpuFlightWaitAtCudaSyncBoundary() { - waitAtCudaSyncBoundary(); -} - -__attribute__((visibility("hidden"))) -void GpuFlightWaitAtCudaMemoryBoundary() { - waitAtCudaMemoryBoundary(); -} +// Capture-gate hooks, one per route (capture_gate_linux.cpp). Hidden. +void GpuFlightCaptureGateRuntime(); +void GpuFlightCaptureGateDriver(); #endif GPUFL_INJECT_EXPORT int InitializeInjectionNvtx2( @@ -880,35 +790,36 @@ int InitializeInjection(void* /*funcTable*/) { // injection pattern - so init fully completes before any kernel runs. std::call_once(g_init_once, doInjectInit); #else - // Numba loads the CUDA Driver API through cuda-python. Starting CUPTI on a - // detached thread while that first cuInit call is still on the stack races - // the driver's initialization and can segfault inside numba-cuda's - // safe_cuda_api_call. Its generated harness opts into a deterministic - // handshake: let cuInit return, then call the exported function below - // synchronously before any allocation or launch. Every other workload keeps - // the existing automatic deferred behavior. - const char* manual_init = envOrNull(gpufl::env::kInjectManualInit); - if (!manual_init || std::strcmp(manual_init, "1") != 0) { - startDeferredInjectInit(); - } + // CUPTI never starts inside this callback: the driver's cuInit is still on + // the stack. The gate arranges the worker for the call that triggered + // driver init, records an unowned (unsupported) first touch, or, in manual + // mode (Numba, the Rust and Python hosts), only counts the callback and + // leaves initialization to GpuFlightInitializeInjectionAfterCuda. + gpufl::inject::LinuxInjectionCallback(); #endif return 0; } // Deterministic post-cuInit handshake for Driver-API hosts such as Numba. -// PRECONDITION: the caller has completed CUDA Driver initialization. The -// function is idempotent with every other injection entry through g_init_once. -// Returning 1 means capture is ready; 0 means initialization was disabled or -// failed. Exported with a stable C name so Python can resolve it from the -// already-LD_PRELOADed process image via ctypes.CDLL(None). +// PRECONDITION: the caller has completed CUDA Driver initialization, and makes +// no allocation, module or launch call before this returns. On Linux in manual +// mode this is the one synchronous initializer; concurrent and repeated calls +// get the same latched result. Returning 1 means capture is ready; 0 means it +// was disabled, failed, timed out, or the host broke the protocol. Exported +// with a stable C name so Python can resolve it from the already-LD_PRELOADed +// process image via ctypes.CDLL(None). GPUFL_INJECT_EXPORT int GpuFlightInitializeInjectionAfterCuda() { +#ifdef _WIN32 try { std::call_once(g_init_once, doInjectInit); } catch (...) { return 0; } return g_init_ok.load(std::memory_order_acquire) ? 1 : 0; +#else + return gpufl::inject::LinuxManualHandshake(); +#endif } #if !defined(_WIN32) && !defined(GPUFL_TYPED_CUDA_INTERPOSE) @@ -941,7 +852,8 @@ using CuLaunchKernelFn = int (*)(void*, unsigned int, unsigned int, unsigned int __attribute__((visibility("default"))) int __cudaLaunchKernel(const void* func, const GpuflDim3 gridDim, const GpuflDim3 blockDim, void** args, const std::size_t sharedMem, void* stream) { - waitAtCudaLaunchBoundary(); + gpufl::inject::WrapperScope scope; + GpuFlightCaptureGateRuntime(); static auto* fn = reinterpret_cast(dlsym(RTLD_NEXT, "__cudaLaunchKernel")); return fn ? fn(func, gridDim, blockDim, args, sharedMem, stream) : 0; } @@ -950,14 +862,16 @@ __attribute__((visibility("default"))) int cudaLaunchKernel(const void* func, const GpuflDim3 gridDim, GpuflDim3 blockDim, void** args, const std::size_t sharedMem, void* stream) { - waitAtCudaLaunchBoundary(); + gpufl::inject::WrapperScope scope; + GpuFlightCaptureGateRuntime(); static auto* fn = reinterpret_cast(dlsym(RTLD_NEXT, "cudaLaunchKernel")); return fn ? fn(func, gridDim, blockDim, args, sharedMem, stream) : 0; } __attribute__((visibility("default"))) int cudaLaunchKernelExC(const void* config, const void* func, void** args) { - waitAtCudaLaunchBoundary(); + gpufl::inject::WrapperScope scope; + GpuFlightCaptureGateRuntime(); static auto* fn = reinterpret_cast(dlsym(RTLD_NEXT, "cudaLaunchKernelExC")); return fn ? fn(config, func, args) : 0; } @@ -968,7 +882,8 @@ int cuLaunchKernel(void* f, const unsigned int gridDimX, const unsigned int grid const unsigned int blockDimY, const unsigned int blockDimZ, const unsigned int sharedMemBytes, void* hStream, void** kernelParams, void** extra) { - waitAtCudaLaunchBoundary(); + gpufl::inject::WrapperScope scope; + GpuFlightCaptureGateDriver(); static auto* fn = reinterpret_cast(dlsym(RTLD_NEXT, "cuLaunchKernel")); return fn ? fn(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra) @@ -977,21 +892,24 @@ int cuLaunchKernel(void* f, const unsigned int gridDimX, const unsigned int grid __attribute__((visibility("default"))) int cudaDeviceSynchronize() { - waitAtCudaSyncBoundary(); + gpufl::inject::WrapperScope scope; + GpuFlightCaptureGateRuntime(); static auto* fn = reinterpret_cast(dlsym(RTLD_NEXT, "cudaDeviceSynchronize")); return fn ? fn() : 0; } __attribute__((visibility("default"))) int cuCtxSynchronize() { - waitAtCudaSyncBoundary(); + gpufl::inject::WrapperScope scope; + GpuFlightCaptureGateDriver(); static auto* fn = reinterpret_cast(dlsym(RTLD_NEXT, "cuCtxSynchronize")); return fn ? fn() : 0; } __attribute__((visibility("default"))) int cudaStreamSynchronize(void* stream) { - waitAtCudaSyncBoundary(); + gpufl::inject::WrapperScope scope; + GpuFlightCaptureGateRuntime(); static auto* fn = reinterpret_cast(dlsym(RTLD_NEXT, "cudaStreamSynchronize")); return fn ? fn(stream) : 0; } diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index ed1cbb4..1db43ae 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -60,6 +60,8 @@ set(GPUFL_TEST_SOURCES ${CMAKE_SOURCE_DIR}/daemon/launcher/segmentation_env.cpp ${CMAKE_SOURCE_DIR}/daemon/launcher/trace_run_plan.cpp ${CMAKE_SOURCE_DIR}/daemon/launcher/agent_launcher.cpp + launcher/test_bootstrap_status.cpp + ${CMAKE_SOURCE_DIR}/daemon/launcher/bootstrap_status.cpp ) if(GPUFL_ENABLE_NVIDIA AND GPUFL_HAS_CUPTI) @@ -102,6 +104,14 @@ if(APPLE AND GPUFL_ENABLE_METAL AND GPUFL_HAS_METAL) ) endif() +# The Linux injection capture gate is a pure state machine (fake driver hooks). +if(UNIX AND NOT APPLE) + list(APPEND GPUFL_TEST_SOURCES + inject/test_capture_gate.cpp + ${CMAKE_SOURCE_DIR}/include/gpufl/inject/capture_gate.cpp + ) +endif() + add_executable(gpufl_tests ${GPUFL_TEST_SOURCES}) # Standalone target for the Linux injection readiness regression. It @@ -290,3 +300,66 @@ else() DISCOVERY_MODE POST_BUILD ) endif() + +# 8. Capture-gate end-to-end tests on CPU (Linux). Real dynamic loading: a +# fake libcuda.so.1 / libcudart.so.13, a stand-in libgpufl_inject.so built from +# the real gate, Linux glue and typed CUDA wrappers with a fake gpufl::init, +# and the real `gpufl trace` launcher for the watchdog. Built with gpufl_tests +# so the CI test step runs them. +if(UNIX AND NOT APPLE AND GPUFL_HAS_CUDA AND CUDAToolkit_INCLUDE_DIRS) + find_package(Python3 COMPONENTS Interpreter) + set(GATE_E2E_DIR ${CMAKE_CURRENT_BINARY_DIR}/gate_e2e) + set(GATE_E2E_SRC ${CMAKE_CURRENT_SOURCE_DIR}/inject/gate_e2e) + + add_library(gate_e2e_fake_cuda SHARED ${GATE_E2E_SRC}/fake_cuda.cpp) + set_target_properties(gate_e2e_fake_cuda PROPERTIES + OUTPUT_NAME cuda SOVERSION 1 + LIBRARY_OUTPUT_DIRECTORY ${GATE_E2E_DIR}/fake) + target_link_libraries(gate_e2e_fake_cuda PRIVATE ${CMAKE_DL_LIBS} pthread) + + add_library(gate_e2e_fake_cudart SHARED ${GATE_E2E_SRC}/fake_cudart.cpp) + set_target_properties(gate_e2e_fake_cudart PROPERTIES + OUTPUT_NAME cudart SOVERSION 13 + LIBRARY_OUTPUT_DIRECTORY ${GATE_E2E_DIR}/fake) + target_link_libraries(gate_e2e_fake_cudart PRIVATE ${CMAKE_DL_LIBS}) + + add_library(gate_e2e_inject SHARED + ${GATE_E2E_SRC}/harness.cpp + ${CMAKE_SOURCE_DIR}/include/gpufl/inject/capture_gate.cpp + ${CMAKE_SOURCE_DIR}/include/gpufl/inject/capture_gate_linux.cpp + ${CMAKE_SOURCE_DIR}/include/gpufl/inject/cuda_interpose_linux.cpp) + set_target_properties(gate_e2e_inject PROPERTIES + OUTPUT_NAME gpufl_inject + LIBRARY_OUTPUT_DIRECTORY ${GATE_E2E_DIR}/inject + CXX_VISIBILITY_PRESET hidden) + target_include_directories(gate_e2e_inject PRIVATE + ${CMAKE_SOURCE_DIR}/include ${GATE_E2E_SRC} ${CUDAToolkit_INCLUDE_DIRS}) + target_link_libraries(gate_e2e_inject PRIVATE ${CMAKE_DL_LIBS} pthread) + + add_executable(gate_e2e_target ${GATE_E2E_SRC}/target.cpp) + set_target_properties(gate_e2e_target PROPERTIES + RUNTIME_OUTPUT_DIRECTORY ${GATE_E2E_DIR} + ENABLE_EXPORTS ON) # the harness resolves TestOwnedTouch + target_include_directories(gate_e2e_target PRIVATE ${GATE_E2E_SRC}) + target_link_libraries(gate_e2e_target PRIVATE + gate_e2e_fake_cudart gate_e2e_fake_cuda ${CMAKE_DL_LIBS}) + # DT_RPATH, searched before LD_LIBRARY_PATH: the target must load the fakes + # even where the environment points at a real or stub libcudart. + target_link_options(gate_e2e_target PRIVATE -Wl,--disable-new-dtags) + + add_dependencies(gpufl_tests gate_e2e_target gate_e2e_inject) + set(GATE_E2E_LAUNCHER_ARGS) + if(BUILD_GPUFL_LAUNCHER) + add_dependencies(gpufl_tests gpufl_launcher) + set(GATE_E2E_LAUNCHER_ARGS --launcher $) + endif() + if(Python3_Interpreter_FOUND) + add_test(NAME CaptureGateEndToEnd + COMMAND ${Python3_EXECUTABLE} -B ${GATE_E2E_SRC}/run_gate_e2e.py + --target $ + --inject $ + --fake-dir ${GATE_E2E_DIR}/fake + ${GATE_E2E_LAUNCHER_ARGS}) + set_tests_properties(CaptureGateEndToEnd PROPERTIES TIMEOUT 300) + endif() +endif() diff --git a/tests/backends/nvidia/test_capture_capability_resolver.cpp b/tests/backends/nvidia/test_capture_capability_resolver.cpp index 360fd1e..5bacb25 100644 --- a/tests/backends/nvidia/test_capture_capability_resolver.cpp +++ b/tests/backends/nvidia/test_capture_capability_resolver.cpp @@ -122,6 +122,62 @@ TEST(CaptureCapabilityResolver, EXPECT_EQ(graph->reason_code, "invalid_timestamps_dropped"); } +// The L4 failure (2026-10-16): every kernel and memcpy record arrived with +// start == end == 0. That is not "no record arrived". +TEST(CaptureCapabilityResolver, AllZeroTimestampsReportTimestampsUnavailable) { + auto input = BaseInput(gpufl::ProfilingEngine::Trace); + input.kernel_activity = true; + input.counters.launch_count = 4096; + input.counters.kernel_records_eligible = 4096; + input.counters.kernel_records_zero_ts = 4096; + input.counters.mem_transfer_records_eligible = 2; + input.counters.mem_transfer_records_zero_ts = 2; + + const auto evt = gpufl::BuildCaptureCapabilitiesEvent(input); + + const auto* kernel = FindCapability(evt, "kernel_events"); + ASSERT_NE(kernel, nullptr); + EXPECT_EQ(kernel->status, "enabled_no_data"); + EXPECT_EQ(kernel->reason_code, "kernel_timestamps_unavailable"); + EXPECT_NE(kernel->message.find("eligible=4096 zero_timestamps=4096"), + std::string::npos); + const auto* memcpy = FindCapability(evt, "memcpy_activity"); + ASSERT_NE(memcpy, nullptr); + EXPECT_EQ(memcpy->status, "enabled_no_data"); + EXPECT_EQ(memcpy->reason_code, "memcpy_timestamps_unavailable"); +} + +TEST(CaptureCapabilityResolver, SomeRejectedTimestampsReportPartialTimeline) { + auto input = BaseInput(gpufl::ProfilingEngine::Trace); + input.kernel_activity = true; + input.counters.kernel_rows = 10; + input.counters.kernel_records_eligible = 12; + input.counters.kernel_records_zero_ts = 1; + input.counters.kernel_records_invalid_ts = 1; + + const auto evt = gpufl::BuildCaptureCapabilitiesEvent(input); + + const auto* kernel = FindCapability(evt, "kernel_events"); + ASSERT_NE(kernel, nullptr); + EXPECT_EQ(kernel->status, "collected"); + EXPECT_EQ(kernel->reason_code, "kernel_timestamps_partial"); + EXPECT_EQ(FindCapability(evt, "memcpy_activity")->reason_code, + "enabled_but_no_records"); +} + +TEST(CaptureCapabilityResolver, RecordsThatNeverArrivedKeepNoRecordsReason) { + auto input = BaseInput(gpufl::ProfilingEngine::Trace); + input.kernel_activity = true; + input.counters.launch_count = 3; + + const auto evt = gpufl::BuildCaptureCapabilitiesEvent(input); + + const auto* kernel = FindCapability(evt, "kernel_events"); + ASSERT_NE(kernel, nullptr); + EXPECT_EQ(kernel->status, "enabled_no_data"); + EXPECT_EQ(kernel->reason_code, "enabled_but_no_records"); +} + TEST(CaptureCapabilityResolver, PcSamplingSyntheticMajorityReportsFallback) { auto input = BaseInput(gpufl::ProfilingEngine::PcSampling); input.counters.kernel_rows = 2; diff --git a/tests/backends/nvidia/test_nvidia_backend.cpp b/tests/backends/nvidia/test_nvidia_backend.cpp index af4bef3..40695d8 100644 --- a/tests/backends/nvidia/test_nvidia_backend.cpp +++ b/tests/backends/nvidia/test_nvidia_backend.cpp @@ -11,6 +11,33 @@ class CuptiBackendTest : public ::testing::Test { void SetUp() override { SKIP_IF_NO_CUDA(); } }; +// Capture readiness needs every required activity kind enabled, at start and +// after the engine re-enable. A "benign" CUPTI result (NOT_COMPATIBLE, a +// fallback the logger keeps quiet about) is not an enable. +TEST(RequiredActivityTally, ABenignResultOrALaterFailedReEnableIsNotReady) { + std::string reason; + gpufl::RequiredActivityTally none; + EXPECT_FALSE(none.Ready(&reason)); + EXPECT_EQ(reason, "activity_not_enabled"); + + gpufl::RequiredActivityTally benign; + benign.Record(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL, true, CUPTI_SUCCESS); + benign.Record(CUPTI_ACTIVITY_KIND_MEMCPY, false, CUPTI_ERROR_NOT_COMPATIBLE); + ASSERT_TRUE(gpufl::IsBenignCuptiResult(CUPTI_ERROR_NOT_COMPATIBLE)); + EXPECT_FALSE(benign.Ready(&reason)); + EXPECT_EQ(reason, "activity_enable_failed:" + + std::to_string(CUPTI_ACTIVITY_KIND_MEMCPY) + ":" + + std::to_string(CUPTI_ERROR_NOT_COMPATIBLE)); + + gpufl::RequiredActivityTally reenable; + reenable.Record(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL, true, CUPTI_SUCCESS); + EXPECT_TRUE(reenable.Ready(&reason)); + reenable.Record(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL, false, + CUPTI_ERROR_NOT_INITIALIZED); + reenable.Record(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL, true, CUPTI_SUCCESS); + EXPECT_FALSE(reenable.Ready(&reason)); // the failure latches +} + // Pure decision for Monitor::Shutdown's orphan-synthesis suppression. No CUDA // needed. The case this exists for: a real-record session (kernel activity // enabled, not PC/SASS synthesize-by-design) where launches happened but zero diff --git a/tests/inject/gate_e2e/fake_cuda.cpp b/tests/inject/gate_e2e/fake_cuda.cpp new file mode 100644 index 0000000..dc4916d --- /dev/null +++ b/tests/inject/gate_e2e/fake_cuda.cpp @@ -0,0 +1,106 @@ +// A fake libcuda.so.1 for the CPU capture-gate end-to-end tests. Like the real +// driver, its first cuInit loads CUDA_INJECTION64_PATH and calls +// InitializeInjection from inside cuInit. FAKE_CUINIT selects a failure: +// ok (default) | fail | nocallback | block (after the callback, forever) +// | block_first (before anything, forever) + +#include +#include +#include +#include +#include +#include + +#include "fake_log.h" + +extern "C" { + +static pthread_mutex_t g_init_lock = PTHREAD_MUTEX_INITIALIZER; +static int g_initialized = 0; + +static const char* mode(void) { + const char* value = getenv("FAKE_CUINIT"); + return value && *value ? value : "ok"; +} + +int cuInit(unsigned int flags) { + (void)flags; + if (strcmp(mode(), "block_first") == 0) { + fake_log("cuInit:block_first"); + for (;;) pause(); + } + pthread_mutex_lock(&g_init_lock); + int first = !g_initialized; + g_initialized = 1; + if (first) { + fake_log("cuInit:first"); + const char* inject = getenv("CUDA_INJECTION64_PATH"); + if (inject && *inject && strcmp(mode(), "nocallback") != 0) { + void* handle = dlopen(inject, RTLD_NOW); + int (*initialize)(void*) = + handle ? (int (*)(void*))dlsym(handle, "InitializeInjection") : NULL; + if (initialize) { + fake_log("cuInit:callback"); + initialize(NULL); + } + } + } + pthread_mutex_unlock(&g_init_lock); + if (strcmp(mode(), "fail") == 0) return 100; + if (first && strcmp(mode(), "block") == 0) { + fake_log("cuInit:block"); + for (;;) pause(); + } + return 0; +} + +int cuMemAlloc_v2(unsigned long long* pointer, size_t bytes) { + (void)bytes; + fake_log("cuMemAlloc_v2"); + *pointer = 0x1000; + return 0; +} + +int cuModuleLoadData(void** module, const void* image) { + (void)image; + fake_log("cuModuleLoadData"); + *module = (void*)0x2000; + return 0; +} + +int cuModuleLoadDataEx(void** module, const void* image, unsigned int count, + void* options, void** values) { + (void)image; (void)count; (void)options; (void)values; + fake_log("cuModuleLoadDataEx"); + *module = (void*)0x2000; + return 0; +} + +int cuDevicePrimaryCtxRetain(void** context, int device) { + (void)device; + fake_log("cuDevicePrimaryCtxRetain"); + *context = (void*)0x3000; + return 0; +} + +int cuCtxSetCurrent(void* context) { + (void)context; + fake_log("cuCtxSetCurrent"); + return 0; +} + +int cuCtxSynchronize(void) { + fake_log("cuCtxSynchronize"); + return 0; +} + +int cuLaunchKernel(void* f, unsigned gx, unsigned gy, unsigned gz, unsigned bx, + unsigned by, unsigned bz, unsigned shared, void* stream, + void** params, void** extra) { + (void)f; (void)gx; (void)gy; (void)gz; (void)bx; (void)by; (void)bz; + (void)shared; (void)stream; (void)params; (void)extra; + fake_log("cuLaunchKernel"); + return 0; +} + +} // extern "C" diff --git a/tests/inject/gate_e2e/fake_cudart.cpp b/tests/inject/gate_e2e/fake_cudart.cpp new file mode 100644 index 0000000..11ba473 --- /dev/null +++ b/tests/inject/gate_e2e/fake_cudart.cpp @@ -0,0 +1,63 @@ +// A fake libcudart for the CPU capture-gate end-to-end tests. Like the real +// runtime, its first call loads libcuda.so.1 and runs cuInit. FAKE_CUDART +// selects a failure of cudaGetDeviceCount after that cuInit: +// ok (default) | fail (35) | block (forever) + +#include +#include +#include +#include +#include + +#include "fake_log.h" + +extern "C" { + +static int runtime_init(void) { + void* driver = dlopen("libcuda.so.1", RTLD_NOW | RTLD_GLOBAL); + int (*cu_init)(unsigned int) = + driver ? (int (*)(unsigned int))dlsym(driver, "cuInit") : NULL; + return cu_init ? cu_init(0) : 3; +} + +int cudaGetDeviceCount(int* count) { + fake_log("cudaGetDeviceCount"); + const int rc = runtime_init(); + const char* mode = getenv("FAKE_CUDART"); + if (mode && strcmp(mode, "fail") == 0) return 35; + if (mode && strcmp(mode, "block") == 0) { + fake_log("cudaGetDeviceCount:block"); + for (;;) pause(); + } + if (rc != 0) return 3; + *count = 1; + return 0; +} + +#define FAKE_RUNTIME_CALL(NAME, PARAMS) \ + int NAME PARAMS { \ + fake_log(#NAME); \ + return runtime_init() == 0 ? 0 : 3; \ + } + +FAKE_RUNTIME_CALL(cudaMalloc, (void** pointer, size_t bytes)) +FAKE_RUNTIME_CALL(cudaFree, (void* pointer)) +FAKE_RUNTIME_CALL(cudaSetDevice, (int device)) +FAKE_RUNTIME_CALL(cudaDeviceSynchronize, (void)) +FAKE_RUNTIME_CALL(cudaMemcpy, (void* dst, const void* src, size_t bytes, int kind)) +FAKE_RUNTIME_CALL(cudaMallocManaged, (void** pointer, size_t bytes, unsigned int flags)) +FAKE_RUNTIME_CALL(cudaMallocHost, (void** pointer, size_t bytes)) +FAKE_RUNTIME_CALL(cudaHostAlloc, (void** pointer, size_t bytes, unsigned int flags)) +FAKE_RUNTIME_CALL(cudaGetDeviceProperties_v2, (void* properties, int device)) +FAKE_RUNTIME_CALL(cudaMemGetInfo, (size_t* free_bytes, size_t* total_bytes)) +FAKE_RUNTIME_CALL(cudaEventCreate, (void** event)) +FAKE_RUNTIME_CALL(cudaEventCreateWithFlags, (void** event, unsigned int flags)) +FAKE_RUNTIME_CALL(cudaStreamCreate, (void** stream)) +FAKE_RUNTIME_CALL(cudaStreamCreateWithFlags, (void** stream, unsigned int flags)) +FAKE_RUNTIME_CALL(cudaStreamSynchronize, (void* stream)) +FAKE_RUNTIME_CALL(cudaMemcpyAsync, (void* dst, const void* src, size_t bytes, int kind, + void* stream)) +// A call the gate does not wrap: it initializes the driver unowned. +FAKE_RUNTIME_CALL(cudaDriverGetVersion, (int* version)) + +} // extern "C" diff --git a/tests/inject/gate_e2e/fake_log.h b/tests/inject/gate_e2e/fake_log.h new file mode 100644 index 0000000..3b708c8 --- /dev/null +++ b/tests/inject/gate_e2e/fake_log.h @@ -0,0 +1,24 @@ +// Shared event log for the capture-gate end-to-end fakes: one line per event, +// appended to FAKE_LOG with O_APPEND so lines from threads and processes do +// not interleave. +#pragma once + +#include +#include +#include +#include +#include + +static inline void fake_log(const char* event) { + const char* path = getenv("FAKE_LOG"); + if (!path || !*path) return; + int fd = open(path, O_WRONLY | O_CREAT | O_APPEND | O_CLOEXEC, 0600); + if (fd < 0) return; + char line[256]; + int n = snprintf(line, sizeof(line), "%s\n", event); + if (n > 0) { + ssize_t written = write(fd, line, (size_t)n); + (void)written; + } + close(fd); +} diff --git a/tests/inject/gate_e2e/harness.cpp b/tests/inject/gate_e2e/harness.cpp new file mode 100644 index 0000000..17f6cf2 --- /dev/null +++ b/tests/inject/gate_e2e/harness.cpp @@ -0,0 +1,73 @@ +// The capture-gate end-to-end harness: built as a stand-in libgpufl_inject.so +// from the real capture_gate.cpp, capture_gate_linux.cpp and the typed CUDA +// wrappers (cuda_interpose_linux.cpp), with only gpufl::init replaced by a +// fake. FAKE_INIT selects its outcome: +// ok (default) | fail | block (until FAKE_INIT_RELEASE_FILE exists, or forever) +// When the target exports TestOwnedTouch, the fake calls it after it returns +// from the block, so a test can prove it never runs after the target's own +// teardown. +// +// The fake does not register the real shutdownAndSignal atexit handler, so +// these tests say nothing about production shutdown registration, and they +// are not teardown-safety evidence for the timeout path (unsupported by +// design: application handlers registered later still run before the guard). + +#include +#include + +#include +#include +#include +#include + +#include "fake_log.h" +#include "gpufl/inject/capture_gate_linux.hpp" + +namespace { + +gpufl::inject::GateStep FakeInitialize() { + fake_log("init:begin"); + const char* mode = std::getenv("FAKE_INIT"); + if (mode && std::strcmp(mode, "block") == 0) { + const char* release = std::getenv("FAKE_INIT_RELEASE_FILE"); + for (;;) { + if (release && access(release, F_OK) == 0) break; + std::this_thread::sleep_for(std::chrono::milliseconds(10)); + } + fake_log("init:released"); + if (auto touch = reinterpret_cast(dlsym(RTLD_DEFAULT, "TestOwnedTouch"))) { + touch(); + } + } + if (mode && std::strcmp(mode, "fail") == 0) { + fake_log("init:fail"); + return {false, "init_failed"}; + } + fake_log("init:ok"); + return {true, ""}; +} + +void FakeShutdown() { fake_log("shutdown"); } + +const bool g_bound = [] { + gpufl::inject::LinuxGateBindings bindings; + bindings.initialize = FakeInitialize; + bindings.after_initialization_exit = FakeShutdown; + gpufl::inject::BindLinuxCaptureGate(bindings); + return true; +}(); + +} // namespace + +extern "C" { + +__attribute__((visibility("default"))) int InitializeInjection(void*) { + gpufl::inject::LinuxInjectionCallback(); + return 0; +} + +__attribute__((visibility("default"))) int GpuFlightInitializeInjectionAfterCuda() { + return gpufl::inject::LinuxManualHandshake(); +} + +} // extern "C" diff --git a/tests/inject/gate_e2e/run_gate_e2e.py b/tests/inject/gate_e2e/run_gate_e2e.py new file mode 100644 index 0000000..5f154c7 --- /dev/null +++ b/tests/inject/gate_e2e/run_gate_e2e.py @@ -0,0 +1,294 @@ +"""Capture-gate end-to-end tests on CPU. + +Runs the target under the stand-in injection library (real gate, Linux glue +and typed CUDA wrappers; fake gpufl::init) with fake libcuda/libcudart, and +optionally under the real `gpufl trace` launcher. Fake functions prove the +ordering and state behaviour, not the absence of real driver locks. +""" + +import argparse +import json +import os +import shutil +import subprocess +import sys +import tempfile +import threading +import time +import unittest +from pathlib import Path + +ARGS = None + +# Every first-touch runtime wrapper the gate supports, plus waiting copy and +# sync wrappers (launches are covered by the unit tests: their by-value dim3 +# ABI is not faked here). +RUNTIME_WRAPPERS = [ + "cudaMalloc", "cudaMallocManaged", "cudaMallocHost", "cudaHostAlloc", "cudaFree", + "cudaSetDevice", "cudaGetDeviceCount", "cudaGetDeviceProperties_v2", "cudaMemGetInfo", + "cudaEventCreate", "cudaEventCreateWithFlags", "cudaStreamCreate", + "cudaStreamCreateWithFlags", "cudaDeviceSynchronize", "cudaStreamSynchronize", + "cudaMemcpy", "cudaMemcpyAsync", +] +DRIVER_WRAPPERS = ["cuMemAlloc_v2", "cuModuleLoadData", "cuModuleLoadDataEx", "cuCtxSynchronize"] + + +def read_status(status_dir): + files = sorted(Path(status_dir).glob("status-*.json")) + if not files: + return None + return json.loads(files[0].read_text()) + + +class Run: + def __init__(self, rc, log, status, elapsed, stderr=""): + self.rc = rc + self.log = log + self.status = status + self.elapsed = elapsed + self.stderr = stderr + + def index(self, event): + return self.log.index(event) + + def before(self, first, second): + return self.index(first) < self.index(second) + + +class GateEndToEnd(unittest.TestCase): + def setUp(self): + self.work = Path(tempfile.mkdtemp(prefix="gpufl-gate-e2e-")) + + def tearDown(self): + shutil.rmtree(self.work, ignore_errors=True) + + def work_reset(self): + shutil.rmtree(self.work, ignore_errors=True) + self.work.mkdir() + + def base_env(self, **extra): + env = {k: v for k, v in os.environ.items() + if not k.startswith("GPUFL_") and not k.startswith("FAKE_") + and k not in ("LD_PRELOAD", "CUDA_INJECTION64_PATH")} + env.update({ + "FAKE_LOG": str(self.work / "events.log"), + "GPUFL_INJECT_LAUNCH_WAIT_MS": "5000", + }) + env.update({k: str(v) for k, v in extra.items()}) + return env + + def direct(self, scenario, release_after=None, **extra): + status_dir = self.work / "status" + status_dir.mkdir(exist_ok=True) + env = self.base_env(**extra) + env.update({ + "LD_PRELOAD": ARGS.inject, + "CUDA_INJECTION64_PATH": ARGS.inject, + "GPUFL_INJECT": "1", + "GPUFL_INJECT_PROTOCOL": "2", + "GPUFL_INJECT_STATUS_DIR": str(status_dir), + "GPUFL_INJECT_STATUS_NONCE": "e2e", + }) + release = self.work / "release" + if release_after is not None: + env["FAKE_INIT_RELEASE_FILE"] = str(release) + threading.Timer(release_after, release.touch).start() + start = time.monotonic() + proc = subprocess.run([ARGS.target, scenario], env=env, timeout=60, + capture_output=True, text=True) + elapsed = time.monotonic() - start + log_path = self.work / "events.log" + log = log_path.read_text().split() if log_path.exists() else [] + return Run(proc.returncode, log, read_status(status_dir), elapsed, proc.stderr) + + def assert_status(self, run, readiness, owner, reason="", notes=()): + self.assertIsNotNone(run.status, run.log) + self.assertEqual(run.status["readiness"], readiness, run.status) + self.assertEqual(run.status["owner"], owner, run.status) + self.assertEqual(run.status["reason"], reason, run.status) + self.assertEqual(sorted(run.status["coverage_notes"]), sorted(notes), run.status) + + def test_first_runtime_call_waits_for_capture(self): + run = self.direct("runtime") + self.assertEqual(run.rc, 0, run.stderr) + self.assert_status(run, "ready", "gate") + self.assertEqual(run.status["lifecycle"], "finished") + self.assertTrue(run.before("cudaGetDeviceCount", "cuInit:callback")) + self.assertTrue(run.before("init:ok", "cudaMalloc")) + self.assertTrue(run.before("cudaMalloc", "app:after_first_call")) + self.assertTrue(run.before("shutdown", "destroyed")) + + def test_set_device_first_is_a_supported_route(self): + run = self.direct("setdevice") + self.assert_status(run, "ready", "gate") + self.assertTrue(run.before("init:ok", "cudaSetDevice")) + + def test_unwrapped_first_touch_is_recorded(self): + run = self.direct("unwrapped") + self.assert_status(run, "ready", "callback", notes=["unsupported_first_touch"]) + self.assertTrue(run.before("init:ok", "cudaMalloc")) + + def test_driver_api_host_without_handshake_is_unsupported(self): + run = self.direct("driver") + self.assert_status(run, "ready", "callback", notes=["unsupported_first_touch"]) + self.assertTrue(run.before("init:ok", "cuMemAlloc_v2")) + + def test_rust_bootstrap_order_uses_one_synchronous_initializer(self): + run = self.direct("manual") + self.assertIn("app:handshake_1", run.log) + self.assert_status(run, "ready", "manual") + self.assertTrue(run.before("cuInit:callback", "cuCtxSetCurrent")) + self.assertTrue(run.before("cuCtxSetCurrent", "init:begin")) + self.assertEqual(run.log.count("init:begin"), 1) + + def test_trigger_failure_after_the_worker_ran_stays_failed(self): + run = self.direct("runtime", FAKE_CUDART="fail") + self.assert_status(run, "failed", "gate", reason="runtime_init_failed:35") + self.assertEqual(run.status["lifecycle"], "finished") + + def test_driver_init_without_callback_fails(self): + run = self.direct("runtime", FAKE_CUINIT="nocallback") + self.assert_status(run, "failed", "gate", reason="no_injection_callback") + self.assertNotIn("init:begin", run.log) + + def test_timeout_lets_the_application_go_on_and_exit_does_not_tear_down(self): + run = self.direct("runtime", FAKE_INIT="block", GPUFL_INJECT_LAUNCH_WAIT_MS=300) + self.assertIn("app:after_first_call", run.log) + self.assertNotIn("init:released", run.log) + self.assertEqual(run.rc, 86) + self.assertNotIn("destroyed", run.log) + self.assertEqual(run.status["reason"], "capture_wait_timeout") + self.assertIn("exit_while_initializing", run.status["coverage_notes"]) + + def test_exit_with_a_wedged_initializer_skips_teardown(self): + run = self.direct("exit_early", FAKE_INIT="block", GPUFL_INJECT_LAUNCH_WAIT_MS=300) + self.assertEqual(run.rc, 86) + self.assertLess(run.elapsed, 10) + self.assertNotIn("destroyed", run.log) + # Documented overlap: a handler registered after the guard runs first. + self.assertIn("app_late_atexit", run.log) + self.assertIn("exit_while_initializing", run.status["coverage_notes"]) + + def test_exit_waits_for_a_late_initializer_before_teardown(self): + run = self.direct("exit_early", release_after=0.3, FAKE_INIT="block") + self.assertEqual(run.rc, 0, run.stderr) + self.assertTrue(run.before("touch:alive", "destroyed")) + self.assertNotIn("touch:after_destroy", run.log) + self.assertEqual(run.log.count("destroyed"), 1) + + def test_every_first_touch_runtime_wrapper_waits_for_capture(self): + for name in RUNTIME_WRAPPERS: + with self.subTest(name): + self.work_reset() + run = self.direct("first:" + name) + self.assertEqual(run.rc, 0, run.stderr) + self.assert_status(run, "ready", "gate") + # The application's own call is the last one with that name + # (cudaGetDeviceCount is also the trigger). + last = len(run.log) - 1 - run.log[::-1].index(name) + self.assertLess(run.index("init:ok"), last, run.log) + + def test_every_driver_wrapper_waits_after_an_unwrapped_cuinit(self): + for name in DRIVER_WRAPPERS: + with self.subTest(name): + self.work_reset() + run = self.direct("driver:" + name) + self.assertEqual(run.rc, 0, run.stderr) + self.assert_status(run, "ready", "callback", + notes=["unsupported_first_touch"]) + self.assertTrue(run.before("init:ok", name), run.log) + + def test_exit_during_the_warmup_stands_the_worker_down(self): + run = self.direct("exit_early", GPUFL_INJECT_WARMUP_MS=5000) + self.assertEqual(run.rc, 0, run.stderr) + self.assertLess(run.elapsed, 3) + self.assertNotIn("init:begin", run.log) + self.assertIn("destroyed", run.log) + self.assertEqual(run.status["reason"], "exit_before_capture") + self.assertIn("warmup_unqualified", run.status["coverage_notes"]) + + def test_a_process_without_cuda_writes_nothing(self): + run = self.direct("none") + self.assertEqual(run.rc, 0) + self.assertIsNone(run.status) + + def test_accidental_preload_stays_inert(self): + status_dir = self.work / "status" + env = self.base_env(LD_PRELOAD=ARGS.inject, CUDA_INJECTION64_PATH=ARGS.inject, + GPUFL_INJECT_STATUS_DIR=status_dir) + status_dir.mkdir() + proc = subprocess.run([ARGS.target, "runtime"], env=env, timeout=60) + self.assertEqual(proc.returncode, 0) + self.assertIsNone(read_status(status_dir)) + + # -- The real launcher ------------------------------------------------- + + def launch(self, scenario, *trace_args, **extra): + if not ARGS.launcher: + self.skipTest("launcher not built") + bin_dir = self.work / "bin" + bin_dir.mkdir() + launcher = bin_dir / "gpufl" + shutil.copy2(ARGS.launcher, launcher) + shutil.copy2(ARGS.inject, bin_dir / "libgpufl_inject.so") + env = self.base_env(GPUFL_INJECT_WATCHDOG_GRACE_MS=200, **extra) + start = time.monotonic() + proc = subprocess.run( + [str(launcher), "trace", "-o", str(self.work / "out"), "--passes", "Trace", + *trace_args, "--", ARGS.target, scenario], + env=env, timeout=120, capture_output=True, text=True) + elapsed = time.monotonic() - start + log_path = self.work / "events.log" + log = log_path.read_text().split() if log_path.exists() else [] + return Run(proc.returncode, log, None, elapsed, proc.stderr) + + def test_launcher_reports_readiness(self): + run = self.launch("runtime") + self.assertIn("capture readiness: ready owner=gate", run.stderr) + + def test_launcher_watchdog_kills_a_wedged_driver_init_on_natural_exit(self): + run = self.launch("runtime", FAKE_CUINIT="block_first", + GPUFL_INJECT_LAUNCH_WAIT_MS=300) + self.assertLess(run.elapsed, 60) + self.assertIn("watchdog killed the target", run.stderr) + + def test_launcher_watchdog_also_bounds_a_windowed_run(self): + run = self.launch("runtime", "--window-timeout", "60s", + FAKE_CUINIT="block_first", GPUFL_INJECT_LAUNCH_WAIT_MS=300) + self.assertLess(run.elapsed, 30) + self.assertIn("watchdog killed the target", run.stderr) + + def test_launcher_keeps_a_healthy_warmup_longer_than_the_init_allowance(self): + run = self.launch("sleep:2500", "--warmup", "1s", GPUFL_INJECT_LAUNCH_WAIT_MS=300) + self.assertEqual(run.rc, 0, run.stderr) + self.assertNotIn("watchdog killed", run.stderr) + self.assertTrue(run.before("init:ok", "app:after_sleep"), run.log) + self.assertIn("coverage=warmup_unqualified", run.stderr) + + def test_launcher_kills_an_initializer_that_wedges_after_the_warmup(self): + run = self.launch("sleep:20000", "--warmup", "1s", FAKE_INIT="block", + GPUFL_INJECT_LAUNCH_WAIT_MS=300) + self.assertIn("init:begin", run.log) # the warmup itself was allowed + self.assertLess(run.elapsed, 10) + self.assertIn("watchdog killed the target", run.stderr) + + def test_launcher_watchdog_covers_a_blocked_trigger_after_worker_success(self): + run = self.launch("runtime", FAKE_CUDART="block", GPUFL_INJECT_LAUNCH_WAIT_MS=300) + self.assertIn("init:ok", run.log) + self.assertLess(run.elapsed, 60) + self.assertIn("watchdog killed the target", run.stderr) + + +def main(): + global ARGS + parser = argparse.ArgumentParser() + parser.add_argument("--target", required=True) + parser.add_argument("--inject", required=True) + parser.add_argument("--fake-dir", required=True) + parser.add_argument("--launcher") + ARGS, rest = parser.parse_known_args() + unittest.main(argv=[sys.argv[0], "-v", *rest]) + + +if __name__ == "__main__": + main() diff --git a/tests/inject/gate_e2e/target.cpp b/tests/inject/gate_e2e/target.cpp new file mode 100644 index 0000000..448158a --- /dev/null +++ b/tests/inject/gate_e2e/target.cpp @@ -0,0 +1,165 @@ +// Target program for the capture-gate end-to-end tests. Linked against the +// fake libcudart and libcuda; the scenario is argv[1]. + +#include +#include +#include +#include +#include + +#include "fake_log.h" + +extern "C" { +int cudaMalloc(void** pointer, size_t bytes); +int cudaSetDevice(int device); +int cudaDeviceSynchronize(void); +int cudaDriverGetVersion(int* version); +int cuInit(unsigned int flags); +int cuMemAlloc_v2(unsigned long long* pointer, size_t bytes); +int cudaMallocManaged(void** pointer, size_t bytes, unsigned int flags); +int cudaMallocHost(void** pointer, size_t bytes); +int cudaHostAlloc(void** pointer, size_t bytes, unsigned int flags); +int cudaFree(void* pointer); +int cudaGetDeviceCount(int* count); +int cudaGetDeviceProperties_v2(void* properties, int device); +int cudaMemGetInfo(size_t* free_bytes, size_t* total_bytes); +int cudaEventCreate(void** event); +int cudaEventCreateWithFlags(void** event, unsigned int flags); +int cudaStreamCreate(void** stream); +int cudaStreamCreateWithFlags(void** stream, unsigned int flags); +int cudaStreamSynchronize(void* stream); +int cudaMemcpy(void* dst, const void* src, size_t bytes, int kind); +int cudaMemcpyAsync(void* dst, const void* src, size_t bytes, int kind, void* stream); +int cuModuleLoadData(void** module, const void* image); +int cuModuleLoadDataEx(void** module, const void* image, unsigned int count, + void* options, void** values); +int cuCtxSynchronize(void); +} + +// "first:": the named wrapped runtime call is the application's first +// CUDA call. Returns false for an unknown name. +static bool call_first(const char* name) { + void* p = NULL; + void* handle = NULL; + size_t a = 0, b = 0; + int count = 0; + char properties[4096]; + char host[16] = {0}; + if (!strcmp(name, "cudaMalloc")) cudaMalloc(&p, 16); + else if (!strcmp(name, "cudaMallocManaged")) cudaMallocManaged(&p, 16, 1); + else if (!strcmp(name, "cudaMallocHost")) cudaMallocHost(&p, 16); + else if (!strcmp(name, "cudaHostAlloc")) cudaHostAlloc(&p, 16, 0); + else if (!strcmp(name, "cudaFree")) cudaFree(NULL); + else if (!strcmp(name, "cudaSetDevice")) cudaSetDevice(0); + else if (!strcmp(name, "cudaGetDeviceCount")) cudaGetDeviceCount(&count); + else if (!strcmp(name, "cudaGetDeviceProperties_v2")) cudaGetDeviceProperties_v2(properties, 0); + else if (!strcmp(name, "cudaMemGetInfo")) cudaMemGetInfo(&a, &b); + else if (!strcmp(name, "cudaEventCreate")) cudaEventCreate(&handle); + else if (!strcmp(name, "cudaEventCreateWithFlags")) cudaEventCreateWithFlags(&handle, 0); + else if (!strcmp(name, "cudaStreamCreate")) cudaStreamCreate(&handle); + else if (!strcmp(name, "cudaStreamCreateWithFlags")) cudaStreamCreateWithFlags(&handle, 0); + else if (!strcmp(name, "cudaDeviceSynchronize")) cudaDeviceSynchronize(); + else if (!strcmp(name, "cudaStreamSynchronize")) cudaStreamSynchronize(NULL); + else if (!strcmp(name, "cudaMemcpy")) cudaMemcpy(host, host + 8, 8, 0); + else if (!strcmp(name, "cudaMemcpyAsync")) cudaMemcpyAsync(host, host + 8, 8, 0, NULL); + else return false; + return true; +} + +// "driver:": the host initializes the driver itself (unwrapped cuInit), +// then the named wrapped driver call follows and must wait for capture. +static bool call_driver(const char* name) { + unsigned long long dp = 0; + void* module = NULL; + const char image[8] = {0}; + if (!strcmp(name, "cuMemAlloc_v2")) cuMemAlloc_v2(&dp, 16); + else if (!strcmp(name, "cuModuleLoadData")) cuModuleLoadData(&module, image); + else if (!strcmp(name, "cuModuleLoadDataEx")) cuModuleLoadDataEx(&module, image, 0, NULL, NULL); + else if (!strcmp(name, "cuCtxSynchronize")) cuCtxSynchronize(); + else return false; + return true; +} + +// A test-owned object: its "destructor" (an atexit handler registered first, +// so it runs last) marks it destroyed, and the late initializer touches it. +static int g_destroyed = 0; + +extern "C" __attribute__((visibility("default"))) void TestOwnedTouch(void) { + fake_log(g_destroyed ? "touch:after_destroy" : "touch:alive"); +} + +static void destroy_test_owned(void) { + g_destroyed = 1; + fake_log("destroyed"); +} + +static void app_late_atexit(void) { fake_log("app_late_atexit"); } + +int main(int argc, char** argv) { + // Registered before any CUDA call: runs after the gate's exit guard. + atexit(destroy_test_owned); + const char* scenario = argc > 1 ? argv[1] : "runtime"; + void* p = NULL; + unsigned long long dp = 0; + int version = 0; + + if (strncmp(scenario, "first:", 6) == 0) { + if (!call_first(scenario + 6)) return 2; + fake_log("app:after_first_call"); + } else if (strncmp(scenario, "driver:", 7) == 0) { + cuInit(0); + if (!call_driver(scenario + 7)) return 2; + fake_log("app:after_first_call"); + } else if (strncmp(scenario, "sleep:", 6) == 0) { + // A long-running host: driver init through an unwrapped call, then + // it works for the given milliseconds before its next CUDA call. + cudaDriverGetVersion(&version); + usleep(static_cast(atoi(scenario + 6)) * 1000u); + cudaMalloc(&p, 16); + fake_log("app:after_sleep"); + } else if (strcmp(scenario, "runtime") == 0) { + int rc = cudaMalloc(&p, 16); + fake_log(rc == 0 ? "app:after_first_call" : "app:first_call_failed"); + cudaDeviceSynchronize(); + } else if (strcmp(scenario, "setdevice") == 0) { + cudaSetDevice(0); + fake_log("app:after_first_call"); + } else if (strcmp(scenario, "unwrapped") == 0) { + // An unwrapped first call initializes the driver; a wrapped call + // follows. + cudaDriverGetVersion(&version); + cudaMalloc(&p, 16); + fake_log("app:after_first_call"); + } else if (strcmp(scenario, "driver") == 0) { + cuInit(0); // unwrapped by design + cuMemAlloc_v2(&dp, 16); + fake_log("app:after_first_call"); + } else if (strcmp(scenario, "manual") == 0) { + // The Rust bootstrap order: MANUAL=1, cuInit (callback fires), + // primary context, then the handshake. + setenv("GPUFL_INJECT_MANUAL_INIT", "1", 1); + void* driver = dlopen("libcuda.so.1", RTLD_NOW); + int (*cu_init)(unsigned int) = (int (*)(unsigned int))dlsym(driver, "cuInit"); + int (*retain)(void**, int) = (int (*)(void**, int))dlsym(driver, "cuDevicePrimaryCtxRetain"); + int (*set_current)(void*) = (int (*)(void*))dlsym(driver, "cuCtxSetCurrent"); + void* context = NULL; + cu_init(0); + retain(&context, 0); + set_current(context); + int (*handshake)(void) = + (int (*)(void))dlsym(RTLD_DEFAULT, "GpuFlightInitializeInjectionAfterCuda"); + int ok = handshake ? handshake() : -1; + fake_log(ok == 1 ? "app:handshake_1" : "app:handshake_0"); + } else if (strcmp(scenario, "exit_early") == 0) { + // Starts initialization and leaves main before it finishes: shortly + // after the worker's delay (so it is inside init), or during a long + // warmup (so it is still in its delay). + cudaDriverGetVersion(&version); + atexit(app_late_atexit); // registered after the gate's guard + usleep(200 * 1000); + } else if (strcmp(scenario, "none") == 0) { + fake_log("app:no_cuda"); + } + fake_log("app:return"); + return 0; +} diff --git a/tests/inject/test_capture_gate.cpp b/tests/inject/test_capture_gate.cpp new file mode 100644 index 0000000..bf10bd0 --- /dev/null +++ b/tests/inject/test_capture_gate.cpp @@ -0,0 +1,491 @@ +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "gpufl/inject/capture_gate.hpp" + +using gpufl::inject::CaptureGate; +using gpufl::inject::GateConfig; +using gpufl::inject::GateHooks; +using gpufl::inject::GateLifecycle; +using gpufl::inject::GateMode; +using gpufl::inject::GateOwner; +using gpufl::inject::GateReadiness; +using gpufl::inject::GateSnapshot; +using gpufl::inject::GateStep; + +namespace { + +using Clock = std::chrono::steady_clock; + +int64_t RealNowNs() { + return std::chrono::duration_cast( + Clock::now().time_since_epoch()) + .count(); +} + +// A one-shot latch the fakes block on. +class Latch { + public: + void Open() { + { + std::lock_guard lock(mu_); + open_ = true; + } + cv_.notify_all(); + } + void Wait() { + std::unique_lock lock(mu_); + cv_.wait(lock, [this] { return open_; }); + } + + private: + std::mutex mu_; + std::condition_variable cv_; + bool open_ = false; +}; + +// Real threads and the real clock (plus an offset a test can push forward), +// fake driver and CUPTI steps. +struct Fakes { + std::atomic clock_offset{0}; + std::atomic spawned{0}; + std::atomic initialized{0}; + std::function barrier = [] { return GateStep{true, ""}; }; + std::function init = [] { return GateStep{true, ""}; }; + std::mutex mu; + std::vector events; + std::vector published; + std::vector threads; + + ~Fakes() { JoinAll(); } + + void JoinAll() { + std::vector pending; + { + std::lock_guard lock(mu); + pending.swap(threads); + } + for (auto& t : pending) t.join(); + } + + void Event(const std::string& e) { + std::lock_guard lock(mu); + events.push_back(e); + } + + GateHooks Hooks() { + GateHooks hooks; + hooks.now_ns = [this] { return RealNowNs() + clock_offset.load(); }; + hooks.spawn = [this](std::function body) { + ++spawned; + std::lock_guard lock(mu); + threads.emplace_back(std::move(body)); + return true; + }; + hooks.driver_barrier = [this] { + Event("barrier"); + return barrier(); + }; + hooks.initialize = [this] { + ++initialized; + Event("init"); + return init(); + }; + hooks.publish = [this](const GateSnapshot& s) { + std::lock_guard lock(mu); + published.push_back(s); + }; + return hooks; + } + + bool EverPublishedReady() { + std::lock_guard lock(mu); + for (const auto& s : published) { + if (s.readiness == GateReadiness::Ready) return true; + } + return false; + } +}; + +GateConfig Config(GateMode mode = GateMode::Automatic, + int64_t wait_ms = 10'000) { + GateConfig config; + config.mode = mode; + config.wait_ns = wait_ms * 1'000'000LL; + config.protocol = "2"; + return config; +} + +template +void WaitFor(Predicate predicate) { + const auto until = Clock::now() + std::chrono::seconds(10); + while (!predicate()) { + ASSERT_LT(Clock::now(), until) << "condition never became true"; + std::this_thread::sleep_for(std::chrono::milliseconds(1)); + } +} + +} // namespace + +TEST(CaptureGate, FirstCallInitializesCaptureBeforeTheApplicationProceeds) { + Fakes fakes; + CaptureGate gate(Config(), fakes.Hooks()); + + gate.Enter([&] { + fakes.Event("driver_init"); + gate.OnInjectionCallback(); // the driver calls InitializeInjection + return GateStep{true, ""}; + }); + fakes.Event("application_call"); + + fakes.JoinAll(); + const auto s = gate.Snapshot(); + EXPECT_EQ(s.readiness, GateReadiness::Ready); + EXPECT_EQ(s.lifecycle, GateLifecycle::Finished); + EXPECT_EQ(s.owner, GateOwner::Gate); + EXPECT_TRUE(s.coverage_notes.empty()); + EXPECT_EQ(fakes.events, (std::vector{ + "driver_init", "barrier", "init", "application_call"})); +} + +TEST(CaptureGate, TriggerFailureBeforeAnyCallbackSettlesAtOnce) { + Fakes fakes; + CaptureGate gate(Config(), fakes.Hooks()); + const auto start = Clock::now(); + + gate.Enter([] { return GateStep{false, "runtime_init_failed:35"}; }); + + EXPECT_LT(Clock::now() - start, std::chrono::seconds(1)); + const auto s = gate.Snapshot(); + EXPECT_EQ(s.readiness, GateReadiness::Failed); + EXPECT_EQ(s.reason, "runtime_init_failed:35"); + EXPECT_EQ(s.lifecycle, GateLifecycle::Finished); // no worker ever started + EXPECT_FALSE(s.worker.expected); + EXPECT_TRUE(gate.AwaitAtExit()); +} + +TEST(CaptureGate, DriverInitWithoutTheInjectionCallbackFails) { + Fakes fakes; + CaptureGate gate(Config(), fakes.Hooks()); + + gate.Enter([] { return GateStep{true, ""}; }); + + const auto s = gate.Snapshot(); + EXPECT_EQ(s.readiness, GateReadiness::Failed); + EXPECT_EQ(s.reason, "no_injection_callback"); + EXPECT_EQ(s.lifecycle, GateLifecycle::Finished); + EXPECT_EQ(fakes.spawned.load(), 0); +} + +TEST(CaptureGate, WorkerSuccessDoesNotHideALaterTriggerFailure) { + Fakes fakes; + CaptureGate gate(Config(), fakes.Hooks()); + + gate.Enter([&] { + gate.OnInjectionCallback(); + WaitFor([&] { return gate.Snapshot().worker.returned; }); + EXPECT_EQ(gate.Snapshot().lifecycle, GateLifecycle::CaptureInitializing); + return GateStep{false, "runtime_init_failed:100"}; + }); + + fakes.JoinAll(); + const auto s = gate.Snapshot(); + EXPECT_EQ(s.readiness, GateReadiness::Failed); + EXPECT_EQ(s.reason, "runtime_init_failed:100"); + EXPECT_EQ(s.lifecycle, GateLifecycle::Finished); + EXPECT_FALSE(fakes.EverPublishedReady()); +} + +TEST(CaptureGate, ABlockedTriggerKeepsInitializationUnfinished) { + Fakes fakes; + CaptureGate gate(Config(GateMode::Automatic, 200), fakes.Hooks()); + Latch release; + + std::thread trigger_thread([&] { + gate.Enter([&] { + gate.OnInjectionCallback(); + release.Wait(); // never returns until the test lets it + return GateStep{true, ""}; + }); + }); + WaitFor([&] { return gate.Snapshot().worker.returned; }); + + // An independent waiter is what observes the deadline in-process. + gate.Enter([] { return GateStep{true, ""}; }); + auto s = gate.Snapshot(); + EXPECT_EQ(s.readiness, GateReadiness::Failed); + EXPECT_EQ(s.reason, "capture_wait_timeout"); + EXPECT_NE(s.lifecycle, GateLifecycle::Finished); + EXPECT_FALSE(gate.AwaitAtExit()); // the process must not tear down + EXPECT_EQ(gate.Snapshot().coverage_notes.count("exit_while_initializing"), 1u); + + release.Open(); + trigger_thread.join(); + s = gate.Snapshot(); + EXPECT_EQ(s.lifecycle, GateLifecycle::Finished); + EXPECT_EQ(s.reason, "capture_wait_timeout"); + EXPECT_FALSE(fakes.EverPublishedReady()); +} + +TEST(CaptureGate, TimeoutStaysFailedAfterALateSuccessAndNothingWaitsAgain) { + Fakes fakes; + Latch release; + fakes.init = [&] { + release.Wait(); + return GateStep{true, ""}; + }; + CaptureGate gate(Config(GateMode::Automatic, 200), fakes.Hooks()); + + gate.Enter([&] { + gate.OnInjectionCallback(); + return GateStep{true, ""}; + }); + EXPECT_EQ(gate.Snapshot().reason, "capture_wait_timeout"); + + release.Open(); + fakes.JoinAll(); + EXPECT_EQ(gate.Snapshot().readiness, GateReadiness::Failed); + EXPECT_EQ(gate.Snapshot().lifecycle, GateLifecycle::Finished); + + const auto start = Clock::now(); + gate.Enter([] { return GateStep{true, ""}; }); + EXPECT_LT(Clock::now() - start, std::chrono::milliseconds(50)); +} + +TEST(CaptureGate, SuccessAfterTheDeadlineFailsWithoutAnyWaiter) { + Fakes fakes; + CaptureGate gate(Config(), fakes.Hooks()); + fakes.init = [&] { + fakes.clock_offset = 60'000'000'000LL; // a minute later + return GateStep{true, ""}; + }; + + gate.OnInjectionCallback(); // driver init from an unwrapped route + fakes.JoinAll(); + + const auto s = gate.Snapshot(); + EXPECT_EQ(s.owner, GateOwner::Callback); + EXPECT_EQ(s.readiness, GateReadiness::Failed); + EXPECT_EQ(s.reason, "capture_init_deadline_exceeded"); + EXPECT_EQ(s.coverage_notes.count("unsupported_first_touch"), 1u); +} + +TEST(CaptureGate, ConcurrentFirstCallersShareOneInitializer) { + Fakes fakes; + CaptureGate gate(Config(), fakes.Hooks()); + std::atomic triggers{0}; + std::vector callers; + for (int i = 0; i < 8; ++i) { + callers.emplace_back([&] { + gate.Enter([&] { + ++triggers; + gate.OnInjectionCallback(); + gate.OnInjectionCallback(); // duplicate callback + return GateStep{true, ""}; + }); + EXPECT_EQ(gate.Snapshot().readiness, GateReadiness::Ready); + }); + } + for (auto& t : callers) t.join(); + fakes.JoinAll(); + EXPECT_EQ(triggers.load(), 1); + EXPECT_EQ(fakes.spawned.load(), 1); + EXPECT_EQ(fakes.initialized.load(), 1); +} + +// The Rust bootstrap order: MANUAL=1, cuInit (its callback fires), primary +// context preparation, then the handshake. +TEST(CaptureGate, ManualModeRunsOneSynchronousInitializer) { + Fakes fakes; + CaptureGate gate(Config(GateMode::Manual), fakes.Hooks()); + + gate.OnInjectionCallback(); + gate.OnInjectionCallback(); + EXPECT_EQ(gate.Snapshot().lifecycle, GateLifecycle::NotRequested); + EXPECT_EQ(gate.ManualHandshake(nullptr), 1); + + int concurrent = -1; + std::thread other([&] { concurrent = gate.ManualHandshake(nullptr); }); + other.join(); + EXPECT_EQ(concurrent, 1); + + const auto s = gate.Snapshot(); + EXPECT_EQ(fakes.spawned.load(), 0); + EXPECT_EQ(fakes.initialized.load(), 1); + EXPECT_EQ(s.owner, GateOwner::Manual); + EXPECT_EQ(s.callbacks_seen, 2u); + EXPECT_TRUE(s.coverage_notes.empty()); +} + +TEST(CaptureGate, ManualHandshakeReentryAndOwnCallsDoNotWait) { + Fakes fakes; + CaptureGate* gate_ptr = nullptr; + int reentry = -1; + fakes.init = [&] { + reentry = gate_ptr->ManualHandshake(nullptr); + gate_ptr->Enter(nullptr); // gpufl's own CUDA call during init + return GateStep{true, ""}; + }; + CaptureGate gate(Config(GateMode::Manual), fakes.Hooks()); + gate_ptr = &gate; + + EXPECT_EQ(gate.ManualHandshake(nullptr), 1); + EXPECT_EQ(reentry, 0); + EXPECT_EQ(gate.Snapshot().handshake_reentries, 1u); + EXPECT_TRUE(gate.Snapshot().coverage_notes.empty()); +} + +TEST(CaptureGate, ManualCallFromAnotherThreadBeforeReadyIsAViolation) { + Fakes fakes; + Latch in_init; + Latch release; + fakes.init = [&] { + in_init.Open(); + release.Wait(); + return GateStep{true, ""}; + }; + CaptureGate gate(Config(GateMode::Manual), fakes.Hooks()); + + int result = -1; + std::thread host([&] { result = gate.ManualHandshake(nullptr); }); + in_init.Wait(); + gate.Enter(nullptr); // application work while the handshake runs + release.Open(); + host.join(); + + EXPECT_EQ(result, 0); + EXPECT_EQ(gate.Snapshot().readiness, GateReadiness::Ready); + EXPECT_EQ(gate.Snapshot().coverage_notes.count("manual_protocol_violation"), 1u); +} + +TEST(CaptureGate, WarmupNeverWaitsAndIsUnqualified) { + Fakes fakes; + CaptureGate gate(Config(GateMode::Warmup), fakes.Hooks()); + bool triggered = false; + + gate.Enter([&] { + triggered = true; + return GateStep{true, ""}; + }); + + EXPECT_FALSE(triggered); + EXPECT_EQ(gate.Snapshot().coverage_notes.count("warmup_unqualified"), 1u); +} + +// An unwrapped entry holds a lock the initializer needs and then reaches a +// wrapper. The TLS flags cannot see that; the deadline bounds the stall. +TEST(CaptureGate, AHeldLockStallsOnlyUntilTheDeadline) { + Fakes fakes; + std::mutex driver_lock; + fakes.init = [&] { + std::lock_guard lock(driver_lock); + return GateStep{true, ""}; + }; + CaptureGate gate(Config(GateMode::Automatic, 200), fakes.Hooks()); + + const auto start = Clock::now(); + { + std::lock_guard held(driver_lock); + gate.OnInjectionCallback(); // the unwrapped entry ran cuInit + gate.Enter([] { return GateStep{true, ""}; }); + } + EXPECT_LT(Clock::now() - start, std::chrono::seconds(5)); + fakes.JoinAll(); + + const auto s = gate.Snapshot(); + EXPECT_EQ(s.reason, "capture_wait_timeout"); + EXPECT_EQ(s.lifecycle, GateLifecycle::Finished); + EXPECT_EQ(s.coverage_notes.count("unsupported_first_touch"), 1u); +} + +TEST(CaptureGate, ExitWithNothingRequestedOrFinishedTearsDownNormally) { + Fakes fakes; + CaptureGate idle(Config(), fakes.Hooks()); + EXPECT_TRUE(idle.AwaitAtExit()); + + CaptureGate done(Config(), fakes.Hooks()); + done.Enter([&] { + done.OnInjectionCallback(); + return GateStep{true, ""}; + }); + EXPECT_TRUE(done.AwaitAtExit()); +} + +TEST(CaptureGateConfig, ReadsModesDelaysAndRejectsBadWaits) { + const auto config = [](std::vector> vars) { + return gpufl::inject::GateConfigFromEnv([vars](const char* name) -> const char* { + for (const auto& [k, v] : vars) { + if (std::string(k) == name) return v; + } + return nullptr; + }); + }; + + auto c = config({{"GPUFL_INJECT_PROTOCOL", "2"}, {"GPUFL_INJECT_INIT_DELAY_MS", "100"}}); + EXPECT_EQ(c.mode, GateMode::Automatic); + EXPECT_EQ(c.delay_ms, 100); + EXPECT_FALSE(c.legacy_launcher_delay); + EXPECT_EQ(c.wait_ns, 15'000'000'000LL); + + c = config({{"GPUFL_INJECT_INIT_DELAY_MS", "5000"}}); + EXPECT_TRUE(c.legacy_launcher_delay); + + c = config({{"GPUFL_INJECT_PROTOCOL", "2"}, {"GPUFL_INJECT_WARMUP_MS", "3000"}, + {"GPUFL_INJECT_INIT_DELAY_MS", "100"}}); + EXPECT_EQ(c.mode, GateMode::Warmup); + EXPECT_EQ(c.delay_ms, 3000); + + c = config({{"GPUFL_INJECT_MANUAL_INIT", "1"}, {"GPUFL_INJECT_WARMUP_MS", "3000"}}); + EXPECT_EQ(c.mode, GateMode::Manual); + + for (const char* bad : {"0", "-5", "abc", "600001", "99999999999999999999"}) { + c = config({{"GPUFL_INJECT_LAUNCH_WAIT_MS", bad}}); + EXPECT_TRUE(c.wait_config_invalid) << bad; + EXPECT_EQ(c.wait_ns, 15'000'000'000LL) << bad; + } + c = config({{"GPUFL_INJECT_LAUNCH_WAIT_MS", "250"}}); + EXPECT_FALSE(c.wait_config_invalid); + EXPECT_EQ(c.wait_ns, 250'000'000LL); +} + +TEST(CaptureGateStatus, WritesTheSnapshotAtomically) { + GateSnapshot s; + s.seq = 7; + s.lifecycle = GateLifecycle::Finished; + s.owner = GateOwner::Gate; + s.readiness = GateReadiness::Failed; + s.reason = "cuinit_symbol_missing:\"quoted\"\n"; + s.coverage_notes = {"unsupported_first_touch"}; + s.deadline_ns = 123; + const std::string json = gpufl::inject::GateStatusJson(s, "abc", 42); + EXPECT_NE(json.find("\"schema\":1,\"nonce\":\"abc\",\"pid\":42,\"seq\":7"), + std::string::npos); + EXPECT_NE(json.find("\"readiness\":\"failed\""), std::string::npos); + EXPECT_NE(json.find("\\\"quoted\\\"\\n"), std::string::npos); + EXPECT_NE(json.find("\"coverage_notes\":[\"unsupported_first_touch\"]"), + std::string::npos); + EXPECT_NE(json.find("\"deadline_monotonic_ns\":123"), std::string::npos); + + const auto dir = std::filesystem::temp_directory_path() / + ("gpufl-gate-status-" + std::to_string(RealNowNs())); + std::filesystem::create_directories(dir); + const std::string path = (dir / "status-42.json").string(); + ASSERT_TRUE(gpufl::inject::WriteGateStatusFile(path, json)); + std::ifstream in(path); + const std::string read((std::istreambuf_iterator(in)), + std::istreambuf_iterator()); + EXPECT_EQ(read, json); + EXPECT_FALSE(std::filesystem::exists(path + ".tmp")); + EXPECT_FALSE(gpufl::inject::WriteGateStatusFile(path, std::string(5000, 'x'))); + std::filesystem::remove_all(dir); +} diff --git a/tests/launcher/test_bootstrap_status.cpp b/tests/launcher/test_bootstrap_status.cpp new file mode 100644 index 0000000..91e6fa8 --- /dev/null +++ b/tests/launcher/test_bootstrap_status.cpp @@ -0,0 +1,260 @@ +#include + +#include +#include +#include +#include + +#include "bootstrap_status.hpp" + +#if defined(__linux__) +#include +#include +#include +#include + +#include "gpufl/inject/capture_gate.hpp" +#endif + +namespace fs = std::filesystem; +using gpufl::launcher::BootstrapStatus; +using gpufl::launcher::CombineFinalBootstrap; +using gpufl::launcher::InitWatchdog; +using gpufl::launcher::ReadBootstrapStatus; + +namespace { + +// A record in the library's exact shape (capture_gate.cpp GateStatusJson). +std::string Participant(bool expected, bool returned, bool ok) { + const auto b = [](bool v) { return v ? std::string("true") : std::string("false"); }; + return "{\"expected\":" + b(expected) + ",\"returned\":" + b(returned) + + ",\"ok\":" + b(ok) + "}"; +} + +std::string Record(const std::string& nonce, long pid, int seq, + const std::string& lifecycle, const std::string& readiness, + const std::string& reason = "", int64_t deadline = 1000, + const std::string& owner = "gate") { + // Gate-owned participants consistent with the lifecycle and readiness. + const bool finished = lifecycle == "finished"; + const std::string trigger = Participant(true, finished, readiness == "ready"); + const std::string worker = Participant(true, finished, readiness == "ready"); + const std::string handshake = Participant(false, false, false); + return "{\"schema\":1,\"nonce\":\"" + nonce + "\",\"pid\":" + std::to_string(pid) + + ",\"seq\":" + std::to_string(seq) + ",\"lifecycle\":\"" + lifecycle + + "\",\"owner\":\"" + owner + "\",\"readiness\":\"" + readiness + + "\",\"reason\":\"" + reason + "\"," + + "\"coverage_notes\":[\"unsupported_first_touch\"]," + + "\"deadline_monotonic_ns\":" + std::to_string(deadline) + + ",\"trigger\":" + trigger + ",\"worker\":" + worker + + ",\"handshake\":" + handshake + + ",\"callbacks_seen\":1,\"handshake_reentries\":0," + + "\"wait_config_invalid\":false,\"protocol\":\"2\"}\n"; +} + +class BootstrapStatusTest : public ::testing::Test { + protected: + void SetUp() override { + dir_ = fs::temp_directory_path() / + ("gpufl-bootstrap-test-" + + std::to_string(std::chrono::steady_clock::now().time_since_epoch().count())); + fs::create_directories(dir_); + } + void TearDown() override { fs::remove_all(dir_); } + + void Write(long pid, const std::string& text) { + std::ofstream(gpufl::launcher::BootstrapStatusPath(dir_, pid)) << text; + } + BootstrapStatus Read(long pid, const std::string& nonce = "n1") { + return ReadBootstrapStatus(dir_, pid, nonce); + } + + fs::path dir_; +}; + +BootstrapStatus Status(int seq, const std::string& lifecycle, int64_t deadline) { + BootstrapStatus s; + s.valid = true; + s.seq = seq; + s.lifecycle = lifecycle; + s.readiness = lifecycle == "finished" ? "ready" : "pending"; + s.deadline_monotonic_ns = deadline; + return s; +} + +} // namespace + +TEST_F(BootstrapStatusTest, ReadsOnlyThisChildsOwnConsistentRecord) { + Write(42, Record("n1", 42, 3, "capture_initializing", "pending")); + const BootstrapStatus s = Read(42); + ASSERT_TRUE(s.valid) << s.error; + EXPECT_EQ(s.seq, 3); + EXPECT_EQ(s.lifecycle, "capture_initializing"); + EXPECT_EQ(s.deadline_monotonic_ns, 1000); + EXPECT_EQ(s.coverage_notes.count("unsupported_first_touch"), 1u); + + EXPECT_EQ(Read(42, "other").error, "foreign_nonce"); + Write(43, Record("n1", 42, 3, "finished", "ready")); + EXPECT_EQ(Read(43).error, "foreign_pid"); + EXPECT_EQ(Read(44).error, "missing"); + Write(45, "{not json"); + EXPECT_EQ(Read(45).error, "malformed"); + Write(46, std::string(5000, ' ')); + EXPECT_EQ(Read(46).error, "oversized"); + Write(47, "{\"schema\":1,\"nonce\":\"n1\",\"pid\":47,\"lifecycle\":\"finished\"," + "\"readiness\":\"ready\"}"); + EXPECT_EQ(Read(47).error, "malformed"); // fields missing +} + +TEST_F(BootstrapStatusTest, RefusesStatesTheLibraryCannotBeIn) { + Write(50, Record("n1", 50, 2, "finished", "pending")); + EXPECT_EQ(Read(50).error, "inconsistent"); + Write(51, Record("n1", 51, 2, "capture_initializing", "ready")); + EXPECT_EQ(Read(51).error, "inconsistent"); + Write(52, Record("n1", 52, 2, "finished", "failed")); // failed without a reason + EXPECT_EQ(Read(52).error, "inconsistent"); + Write(53, Record("n1", 53, 2, "driver_ready", "pending", "", 0)); // no deadline + EXPECT_EQ(Read(53).error, "inconsistent"); + Write(54, Record("n1", 54, 2, "sleeping", "pending")); + EXPECT_EQ(Read(54).error, "inconsistent"); + Write(55, Record("n1", 55, 2, "finished", "failed", "capture_wait_timeout")); + EXPECT_TRUE(Read(55).valid); +} + +TEST(InitWatchdog, KeepsTheFirstDeadlineAndFiresAfterTheGrace) { + InitWatchdog watchdog(/*grace_ns=*/100); + EXPECT_FALSE(watchdog.Observe(Status(1, "not_requested", 0), 10'000)); + EXPECT_FALSE(watchdog.armed()); + + EXPECT_FALSE(watchdog.Observe(Status(2, "driver_initializing", 1'000), 500)); + EXPECT_TRUE(watchdog.armed()); + // A later record cannot renew the budget. + EXPECT_FALSE(watchdog.Observe(Status(3, "capture_initializing", 900'000), 1'100)); + EXPECT_TRUE(watchdog.Observe(Status(3, "capture_initializing", 900'000), 1'101)); +} + +TEST(InitWatchdog, DisarmsWhenEveryParticipantReturnedAndIgnoresStaleRecords) { + InitWatchdog watchdog(0); + EXPECT_FALSE(watchdog.Observe(Status(5, "driver_initializing", 1'000), 0)); + EXPECT_FALSE(watchdog.Observe(Status(9, "finished", 1'000), 2'000)); + EXPECT_TRUE(watchdog.disarmed()); + // An older record after Finished changes nothing. + EXPECT_FALSE(watchdog.Observe(Status(6, "capture_initializing", 1'000), 5'000)); + EXPECT_EQ(watchdog.last().lifecycle, "finished"); + + // An unreadable file never arms or disarms. + InitWatchdog unread(0); + EXPECT_FALSE(unread.Observe(BootstrapStatus{}, 1'000'000)); + EXPECT_FALSE(unread.seen_valid()); +} + +TEST(CombineFinalBootstrap, AFinalReadThatDoesNotConfirmIsReportedNotHidden) { + InitWatchdog watchdog(0); + watchdog.Observe(Status(9, "finished", 1'000), 0); + + auto confirmed = CombineFinalBootstrap(watchdog, Status(9, "finished", 1'000)); + EXPECT_TRUE(confirmed.final_error.empty()); + + for (const char* error : {"missing", "malformed", "foreign_nonce"}) { + BootstrapStatus lost; + lost.error = error; + const auto observed = CombineFinalBootstrap(watchdog, lost); + EXPECT_EQ(observed.final_error, error); + EXPECT_EQ(observed.accepted.seq, 9); // history, flagged as such + } + EXPECT_EQ(CombineFinalBootstrap(watchdog, Status(4, "driver_ready", 1'000)).final_error, + "stale_final_seq"); + + InitWatchdog never(0); + BootstrapStatus missing; + missing.error = "missing"; + const auto none = CombineFinalBootstrap(never, missing); + EXPECT_FALSE(none.accepted.valid); + EXPECT_EQ(none.final_error, "missing"); +} + +#if defined(__linux__) +// Round trip against the library's own writer: every snapshot a real gate +// publishes on the success, trigger-failure, unowned-init-failure and manual +// paths parses as valid. A timeout path is not induced here. +namespace { + +struct Published { + std::mutex mu; + std::vector json; + std::vector threads; + ~Published() { + for (auto& t : threads) t.join(); + } +}; + +gpufl::inject::GateHooks RoundTripHooks(Published& out, gpufl::inject::GateStep init) { + gpufl::inject::GateHooks hooks; + hooks.now_ns = gpufl::launcher::MonotonicNowNs; + hooks.spawn = [&out](std::function body) { + std::lock_guard lock(out.mu); + out.threads.emplace_back(std::move(body)); + return true; + }; + hooks.driver_barrier = [] { return gpufl::inject::GateStep{true, ""}; }; + hooks.initialize = [init] { return init; }; + hooks.publish = [&out](const gpufl::inject::GateSnapshot& s) { + std::lock_guard lock(out.mu); + out.json.push_back(gpufl::inject::GateStatusJson(s, "n1", 60)); + }; + return hooks; +} + +} // namespace + +TEST_F(BootstrapStatusTest, EverySnapshotTheGatePublishesParsesAsValid) { + using gpufl::inject::CaptureGate; + using gpufl::inject::GateMode; + using gpufl::inject::GateStep; + const auto config = [](GateMode mode, int64_t wait_ms) { + gpufl::inject::GateConfig c; + c.mode = mode; + c.wait_ns = wait_ms * 1'000'000LL; + c.protocol = "2"; + return c; + }; + std::vector all; + const auto collect = [&all](Published& p) { + for (auto& t : p.threads) t.join(); + p.threads.clear(); + all.insert(all.end(), p.json.begin(), p.json.end()); + }; + { + Published p; // gate owner, success + CaptureGate gate(config(GateMode::Automatic, 5'000), RoundTripHooks(p, {true, ""})); + gate.Enter([&] { gate.OnInjectionCallback(); return GateStep{true, ""}; }); + collect(p); + } + { + Published p; // gate owner, trigger failure without a callback + CaptureGate gate(config(GateMode::Automatic, 5'000), RoundTripHooks(p, {true, ""})); + gate.Enter([] { return GateStep{false, "runtime_init_failed:35"}; }); + collect(p); + } + { + Published p; // unowned first touch, init failure + CaptureGate gate(config(GateMode::Automatic, 5'000), + RoundTripHooks(p, {false, "init_failed"})); + gate.OnInjectionCallback(); + collect(p); + } + { + Published p; // manual handshake after a counted callback + CaptureGate gate(config(GateMode::Manual, 5'000), RoundTripHooks(p, {true, ""})); + gate.OnInjectionCallback(); + EXPECT_EQ(gate.ManualHandshake(nullptr), 1); + collect(p); + } + ASSERT_GT(all.size(), 8u); + for (const std::string& json : all) { + Write(60, json); + const BootstrapStatus s = Read(60); + EXPECT_TRUE(s.valid) << s.error << ": " << json; + } +} +#endif