From ffaaa965a6a763a3e2750d40921473006cd1ccd0 Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Fri, 16 Oct 2026 00:32:00 -0700 Subject: [PATCH 1/8] Report activity records that arrived without timing On an L4 a whole Trace session's kernel and memcpy records arrived with start == end == 0. They were dropped, as they must be, but the session then claimed no kernel activity record had arrived. Count the records that reach the timestamp check and the ones it rejects (zero-zero separately from other invalid timing), and report kernel_timestamps_unavailable, _invalid or _partial (memcpy alike) as the capability reason with the counts. Statuses are unchanged, so the server's classification is unaffected. --- .../nvidia/capture_capability_resolver.cpp | 72 +++++++++++++++++++ .../nvidia/capture_capability_resolver.hpp | 31 ++++++++ .../gpufl/backends/nvidia/cupti_backend.cpp | 5 +- .../gpufl/backends/nvidia/cupti_backend.hpp | 34 +++++++++ .../nvidia/cupti_capture_capabilities.cpp | 18 +++++ .../backends/nvidia/kernel_launch_handler.cpp | 2 + .../backends/nvidia/mem_transfer_handler.cpp | 4 ++ include/gpufl/core/monitor.cpp | 24 +++++-- include/gpufl/core/monitor_backend.hpp | 14 ++++ .../test_capture_capability_resolver.cpp | 56 +++++++++++++++ 10 files changed, 254 insertions(+), 6 deletions(-) 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..ceb2db8 100644 --- a/include/gpufl/backends/nvidia/cupti_backend.cpp +++ b/include/gpufl/backends/nvidia/cupti_backend.cpp @@ -983,8 +983,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..1ccaf07 100644 --- a/include/gpufl/backends/nvidia/cupti_backend.hpp +++ b/include/gpufl/backends/nvidia/cupti_backend.hpp @@ -51,6 +51,26 @@ 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)}; + } +}; + /** * @brief CUPTI-based monitoring backend for NVIDIA GPUs. * @@ -123,6 +143,16 @@ 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(); + } + 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 +414,9 @@ 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 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 +458,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..73a00ba 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 = @@ -159,6 +175,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/monitor.cpp b/include/gpufl/core/monitor.cpp index eb6ec35..6be2f56 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(); } diff --git a/include/gpufl/core/monitor_backend.hpp b/include/gpufl/core/monitor_backend.hpp index 38a6d29..da0ed65 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,11 @@ class IMonitorBackend { */ virtual bool kernelActivityExpectedButMissing() const { return false; } + /** 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/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; From 9afef22775047a5a2b242d362625f9bd491cd6bd Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Fri, 16 Oct 2026 01:02:24 -0700 Subject: [PATCH 2/8] Gate the first CUDA call on checked capture readiness (Linux) The injected library started CUPTI from a detached worker while the application thread went on into context creation and module loading; on an L4 whole Trace sessions lost every kernel timestamp that way, and nothing checked whether capture was ready. CaptureGate separates the initializer lifecycle from a latched readiness result and from coverage notes. The first supported runtime call triggers driver init through the real cudaGetDeviceCount (driver calls through a checked cuInit) and waits, within one CLOCK_MONOTONIC deadline, until the trigger and the worker have both returned and the Trace subscription and every required activity kind are on. A timeout, a late success, a failed trigger, a missing callback or an unowned first touch is latched and never becomes ready. Manual mode keeps one synchronous initializer in the handshake; --warmup is reported as an unqualified capture. The library writes a per-pid bootstrap status file without the logger, and the session gets capture_readiness and capture_coverage rows. The launcher sets the protocol, warmup and status variables, polls in every mode, and kills a child whose initialization outlives the published deadline. At exit an initializer still live after the deadline ends the process with status 86 instead of running teardown under it. Design: agent_comm l4-qualification-20261015 v3 + v3.1 (Codex CLEAR). --- CMakeLists.txt | 6 + daemon/launcher/CMakeLists.txt | 1 + daemon/launcher/bootstrap_status.cpp | 129 ++++ daemon/launcher/bootstrap_status.hpp | 77 +++ daemon/launcher/trace_command.cpp | 54 +- daemon/launcher/trace_command_common.cpp | 91 ++- daemon/launcher/trace_command_common.hpp | 15 + .../gpufl/backends/nvidia/cupti_backend.cpp | 17 +- .../gpufl/backends/nvidia/cupti_backend.hpp | 21 + .../nvidia/cupti_capture_capabilities.cpp | 1 + include/gpufl/core/env_vars.hpp | 17 + include/gpufl/core/monitor.cpp | 21 + include/gpufl/core/monitor.hpp | 17 + include/gpufl/core/monitor_backend.hpp | 11 + include/gpufl/inject/capture_gate.cpp | 645 ++++++++++++++++++ include/gpufl/inject/capture_gate.hpp | 222 ++++++ include/gpufl/inject/capture_gate_linux.cpp | 247 +++++++ include/gpufl/inject/capture_gate_linux.hpp | 44 ++ include/gpufl/inject/cuda_interpose_linux.cpp | 146 +++- include/gpufl/inject/inject_entry.cpp | 233 ++----- tests/CMakeLists.txt | 10 + tests/inject/test_capture_gate.cpp | 492 +++++++++++++ tests/launcher/test_bootstrap_status.cpp | 99 +++ 23 files changed, 2405 insertions(+), 211 deletions(-) create mode 100644 daemon/launcher/bootstrap_status.cpp create mode 100644 daemon/launcher/bootstrap_status.hpp create mode 100644 include/gpufl/inject/capture_gate.cpp create mode 100644 include/gpufl/inject/capture_gate.hpp create mode 100644 include/gpufl/inject/capture_gate_linux.cpp create mode 100644 include/gpufl/inject/capture_gate_linux.hpp create mode 100644 tests/inject/test_capture_gate.cpp create mode 100644 tests/launcher/test_bootstrap_status.cpp 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..94ad753 --- /dev/null +++ b/daemon/launcher/bootstrap_status.cpp @@ -0,0 +1,129 @@ +#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"); + + BootstrapStatus status; + status.valid = true; + status.seq = doc.value("seq", 0); + status.lifecycle = doc.value("lifecycle", ""); + status.owner = doc.value("owner", ""); + status.readiness = doc.value("readiness", ""); + status.reason = doc.value("reason", ""); + status.deadline_monotonic_ns = doc.value("deadline_monotonic_ns", 0); + status.protocol = doc.value("protocol", ""); + if (doc.contains("coverage_notes") && doc["coverage_notes"].is_array()) { + for (const json::JsonValue& note : doc["coverage_notes"].get_array()) { + if (note.is_string()) status.coverage_notes.insert(note.get_string()); + } + } + if (status.lifecycle.empty() || status.readiness.empty()) { + return Invalid("malformed"); + } + 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; +} + +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..09a8994 --- /dev/null +++ b/daemon/launcher/bootstrap_status.hpp @@ -0,0 +1,77 @@ +#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_; +}; + +/** 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/trace_command.cpp b/daemon/launcher/trace_command.cpp index 11243f2..802b3d9 100644 --- a/daemon/launcher/trace_command.cpp +++ b/daemon/launcher/trace_command.cpp @@ -146,13 +146,30 @@ 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 = [&]() -> bool { + if (!watch) 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(); // the final record after exit + out.bootstrap = watchdog.seen_valid() ? watchdog.last() : last_read; + }; + + 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 +183,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..c9791e9 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) { @@ -222,8 +224,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 +781,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 +942,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(), DescribeBootstrapStatus(run.bootstrap).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(), + DescribeBootstrapStatus(run.bootstrap).c_str()); + } +#endif + if (!args.quiet) { if (run.window_stopped) { std::fprintf(stderr, @@ -973,6 +1049,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..def8631 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,11 @@ 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; std::string error; }; @@ -29,6 +39,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/cupti_backend.cpp b/include/gpufl/backends/nvidia/cupti_backend.cpp index ceb2db8..4f393fd 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,21 @@ 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; the + // injection gate reads this (captureFacilitiesReady). + if (res != CUPTI_SUCCESS && !IsBenignCuptiResult(res)) { + std::lock_guard lock(capture_facility_mu_); + if (capture_facility_failure_.empty()) { + capture_facility_failure_ = + "activity_enable_failed:" + std::to_string(static_cast(k)) + + ":" + std::to_string(static_cast(res)); + } + } + } + required_activity_enabled_.store(true, std::memory_order_release); } // Initialize and start the engine (requires CUDA context) diff --git a/include/gpufl/backends/nvidia/cupti_backend.hpp b/include/gpufl/backends/nvidia/cupti_backend.hpp index 1ccaf07..1a11033 100644 --- a/include/gpufl/backends/nvidia/cupti_backend.hpp +++ b/include/gpufl/backends/nvidia/cupti_backend.hpp @@ -146,6 +146,23 @@ class CuptiBackend : public IMonitorBackend { 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"); + } + if (!required_activity_enabled_.load(std::memory_order_acquire)) { + return fail("activity_not_enabled"); + } + std::lock_guard lock(capture_facility_mu_); + if (!capture_facility_failure_.empty()) { + return fail(capture_facility_failure_); + } + return true; + } void NoteMemTransferActivityEligible() { mem_transfer_activity_timestamps_.NoteEligible(); } @@ -415,6 +432,10 @@ class CuptiBackend : public IMonitorBackend { std::atomic kernel_activity_emitted_{0}; std::atomic kernel_activity_throttled_{0}; ActivityTimestampTally kernel_activity_timestamps_; + std::atomic cupti_subscribed_{false}; + std::atomic required_activity_enabled_{false}; + mutable std::mutex capture_facility_mu_; + std::string capture_facility_failure_; std::atomic mem_transfer_activity_emitted_{0}; ActivityTimestampTally mem_transfer_activity_timestamps_; std::atomic sync_activity_emitted_{0}; diff --git a/include/gpufl/backends/nvidia/cupti_capture_capabilities.cpp b/include/gpufl/backends/nvidia/cupti_capture_capabilities.cpp index 73a00ba..cf378e1 100644 --- a/include/gpufl/backends/nvidia/cupti_capture_capabilities.cpp +++ b/include/gpufl/backends/nvidia/cupti_capture_capabilities.cpp @@ -161,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; diff --git a/include/gpufl/core/env_vars.hpp b/include/gpufl/core/env_vars.hpp index 4cc468f..3b41faa 100644 --- a/include/gpufl/core/env_vars.hpp +++ b/include/gpufl/core/env_vars.hpp @@ -120,6 +120,23 @@ 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 6be2f56..3d98273 100644 --- a/include/gpufl/core/monitor.cpp +++ b/include/gpufl/core/monitor.cpp @@ -739,6 +739,27 @@ 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; + if (!g_state.initialized.load()) { + 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 da0ed65..c2ba825 100644 --- a/include/gpufl/core/monitor_backend.hpp +++ b/include/gpufl/core/monitor_backend.hpp @@ -127,6 +127,17 @@ 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 {}; diff --git a/include/gpufl/inject/capture_gate.cpp b/include/gpufl/inject/capture_gate.cpp new file mode 100644 index 0000000..31e295b --- /dev/null +++ b/include/gpufl/inject/capture_gate.cpp @@ -0,0 +1,645 @@ +#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; + const int64_t max = std::numeric_limits::max(); + state_.deadline_ns = now > max - config_.wait_ns ? max : now + config_.wait_ns; +} + +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); +} + +GateSnapshot CaptureGate::TakeSnapshotLocked() { + ++state_.seq; + return state_; +} + +void CaptureGate::Publish(const GateSnapshot& snapshot) { + std::lock_guard lock(publish_mu_); + // Snapshots are taken under mu_ with increasing seq but published after + // it is released; an older one arriving late must not replace a newer one. + if (snapshot.seq <= last_published_seq_) return; + last_published_seq_ = snapshot.seq; + if (hooks_.publish) hooks_.publish(snapshot); +} + +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"); + const GateSnapshot snapshot = TakeSnapshotLocked(); + lock.unlock(); + cv_.notify_all(); + Publish(snapshot); + lock.lock(); + } + 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) { + const GateSnapshot snapshot = TakeSnapshotLocked(); + lock.unlock(); + Publish(snapshot); + } + return; + } + + if (state_.lifecycle == GateLifecycle::NotRequested) { + ElectLocked(GateOwner::Gate); + state_.trigger.expected = true; + GateSnapshot snapshot = TakeSnapshotLocked(); + lock.unlock(); + Publish(snapshot); + + 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(); + snapshot = TakeSnapshotLocked(); + lock.unlock(); + cv_.notify_all(); + Publish(snapshot); + lock.lock(); + } + 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; + } + } + const GateSnapshot snapshot = TakeSnapshotLocked(); + lock.unlock(); + Publish(snapshot); + if (spawn) { + if (!hooks_.spawn || !hooks_.spawn([this] { RunWorker(); })) { + std::unique_lock relock(mu_); + state_.worker.returned = true; + LatchFailedLocked("worker_spawn_unavailable"); + ResolveLocked(); + const GateSnapshot failed = TakeSnapshotLocked(); + relock.unlock(); + cv_.notify_all(); + Publish(failed); + } + } +} + +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(); + const GateSnapshot snapshot = TakeSnapshotLocked(); + lock.unlock(); + cv_.notify_all(); + Publish(snapshot); + return; + } + if (state_.lifecycle == GateLifecycle::DriverInitializing) { + state_.lifecycle = GateLifecycle::DriverReady; + } + const GateSnapshot snapshot = TakeSnapshotLocked(); + lock.unlock(); + Publish(snapshot); + } + if (hooks_.sleep_ms) hooks_.sleep_ms(config_.delay_ms); + { + std::unique_lock lock(mu_); + if (state_.lifecycle == GateLifecycle::DriverReady) { + state_.lifecycle = GateLifecycle::CaptureInitializing; + } + const GateSnapshot snapshot = TakeSnapshotLocked(); + lock.unlock(); + Publish(snapshot); + } + 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(); + const GateSnapshot snapshot = TakeSnapshotLocked(); + lock.unlock(); + cv_.notify_all(); + Publish(snapshot); +} + +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; + GateSnapshot snapshot = TakeSnapshotLocked(); + lock.unlock(); + Publish(snapshot); + + 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; + snapshot = TakeSnapshotLocked(); + lock.unlock(); + Publish(snapshot); + 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(); + snapshot = TakeSnapshotLocked(); + lock.unlock(); + cv_.notify_all(); + Publish(snapshot); + lock.lock(); + } 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; + } + if (WaitLocked(lock, /*until_finished=*/true)) return true; + state_.coverage_notes.insert("exit_while_initializing"); + const GateSnapshot snapshot = TakeSnapshotLocked(); + lock.unlock(); + Publish(snapshot); + return false; +} + +void CaptureGate::AddCoverageNote(const std::string& note) { + std::unique_lock lock(mu_); + if (!state_.coverage_notes.insert(note).second) return; + const GateSnapshot snapshot = TakeSnapshotLocked(); + lock.unlock(); + Publish(snapshot); +} + +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..55f7da9 --- /dev/null +++ b/include/gpufl/inject/capture_gate.hpp @@ -0,0 +1,222 @@ +#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, and on an L4 +// (2026-10-16) whole sessions lost every kernel timestamp that way. The gate +// makes the application's first supported CUDA call wait until capture is +// ready, within one bounded deadline, and records every other outcome. +// +// 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; + std::function sleep_ms; + // gpufl::init plus the Trace facility check. Never throws out. + std::function initialize; + // Publishes a snapshot (the status file). Called without the gate mutex; + // 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(); + GateSnapshot TakeSnapshotLocked(); + bool WaitLocked(std::unique_lock& lock, bool until_finished); + void Publish(const GateSnapshot& snapshot); + void RunWorker(); + void StartWorkerLocked(); + + const GateConfig config_; + const GateHooks hooks_; + + mutable std::mutex mu_; + std::condition_variable cv_; + GateSnapshot state_; + bool worker_started_ = false; + std::atomic settled_{false}; + std::mutex publish_mu_; + uint64_t last_published_seq_ = 0; +}; + +/** + * 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..6c4c26e --- /dev/null +++ b/include/gpufl/inject/capture_gate_linux.cpp @@ -0,0 +1,247 @@ +// NVIDIA calls InitializeInjection from inside the driver's own cuInit, so +// CUPTI cannot start there; a worker starts it after the driver is up. On an +// L4 (2026-10-16) the application thread raced that worker into context and +// module setup, and whole sessions lost every kernel timestamp. 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.sleep_ms = [](const int ms) { + if (ms > 0) std::this_thread::sleep_for(std::chrono::milliseconds(ms)); + }; + 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..0038599 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,49 @@ 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"}; } - g_deferred_init_cv.notify_all(); -} - -// 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")); + 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}; } -#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); + return {true, ""}; } -void startDeferredInjectInit() { - registerDeferredWaitAtexit(); +void contributeGateCapabilities(gpufl::CaptureCapabilitiesEvent& evt) { + gpufl::inject::AppendGateCapabilities(gpufl::inject::LinuxCaptureGate().Snapshot(), + evt); +} - 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 +641,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 +783,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 +845,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 +855,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 +875,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 +885,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..9df42fa 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 diff --git a/tests/inject/test_capture_gate.cpp b/tests/inject/test_capture_gate.cpp new file mode 100644 index 0000000..0d54bdd --- /dev/null +++ b/tests/inject/test_capture_gate.cpp @@ -0,0 +1,492 @@ +#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.sleep_ms = [](int) {}; + 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..8f5e311 --- /dev/null +++ b/tests/launcher/test_bootstrap_status.cpp @@ -0,0 +1,99 @@ +#include + +#include +#include +#include +#include + +#include "bootstrap_status.hpp" + +namespace fs = std::filesystem; +using gpufl::launcher::BootstrapStatus; +using gpufl::launcher::InitWatchdog; +using gpufl::launcher::ReadBootstrapStatus; + +namespace { + +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; + } + + static std::string Record(const std::string& nonce, long pid, int seq, + const std::string& lifecycle, int64_t deadline) { + return "{\"schema\":1,\"nonce\":\"" + nonce + "\",\"pid\":" + std::to_string(pid) + + ",\"seq\":" + std::to_string(seq) + ",\"lifecycle\":\"" + lifecycle + + "\",\"owner\":\"gate\",\"readiness\":\"pending\",\"reason\":\"\"," + + "\"coverage_notes\":[\"unsupported_first_touch\"]," + + "\"deadline_monotonic_ns\":" + std::to_string(deadline) + "}"; + } + + 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 = "pending"; + s.deadline_monotonic_ns = deadline; + return s; +} + +} // namespace + +TEST_F(BootstrapStatusTest, ReadsOnlyThisChildsOwnRecord) { + Write(42, Record("n1", 42, 3, "capture_initializing", 1000)); + const BootstrapStatus s = ReadBootstrapStatus(dir_, 42, "n1"); + 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(ReadBootstrapStatus(dir_, 42, "other").error, "foreign_nonce"); + Write(43, Record("n1", 42, 3, "finished", 1000)); + EXPECT_EQ(ReadBootstrapStatus(dir_, 43, "n1").error, "foreign_pid"); + EXPECT_EQ(ReadBootstrapStatus(dir_, 44, "n1").error, "missing"); + Write(45, "{not json"); + EXPECT_EQ(ReadBootstrapStatus(dir_, 45, "n1").error, "malformed"); + Write(46, std::string(5000, ' ')); + EXPECT_EQ(ReadBootstrapStatus(dir_, 46, "n1").error, "oversized"); +} + +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()); +} From f660cf9ddab3fcfe2374191d05d08f212be1e62b Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Fri, 16 Oct 2026 01:09:00 -0700 Subject: [PATCH 3/8] Test the capture gate end to end with fake CUDA libraries A stand-in libgpufl_inject.so built from the real gate, Linux glue and typed CUDA wrappers, with only gpufl::init faked, runs a target linked against a fake libcuda.so.1 and libcudart.so.13 that call InitializeInjection from inside cuInit like the driver does. The scenarios cover the runtime and cudaSetDevice routes, an unwrapped and a Driver API first touch, the Rust bootstrap order in manual mode, trigger failure after the worker succeeded, a missing callback, a timeout, exit with a wedged initializer (status 86, no teardown) and with a late one (teardown after it), a process without CUDA and an accidental preload, plus the real `gpufl trace` watchdog in natural-exit and windowed runs and for a trigger blocked after its worker succeeded. They found that the worker published Ready after waking the application, so a process could exit with a stale pending record and the worker still used gate state after Finished. Snapshots are now published under the gate mutex. --- include/gpufl/inject/capture_gate.cpp | 73 +++----- include/gpufl/inject/capture_gate.hpp | 10 +- tests/CMakeLists.txt | 63 +++++++ tests/inject/gate_e2e/fake_cuda.cpp | 98 +++++++++++ tests/inject/gate_e2e/fake_cudart.cpp | 51 ++++++ tests/inject/gate_e2e/fake_log.h | 24 +++ tests/inject/gate_e2e/harness.cpp | 68 ++++++++ tests/inject/gate_e2e/run_gate_e2e.py | 233 ++++++++++++++++++++++++++ tests/inject/gate_e2e/target.cpp | 86 ++++++++++ 9 files changed, 650 insertions(+), 56 deletions(-) create mode 100644 tests/inject/gate_e2e/fake_cuda.cpp create mode 100644 tests/inject/gate_e2e/fake_cudart.cpp create mode 100644 tests/inject/gate_e2e/fake_log.h create mode 100644 tests/inject/gate_e2e/harness.cpp create mode 100644 tests/inject/gate_e2e/run_gate_e2e.py create mode 100644 tests/inject/gate_e2e/target.cpp diff --git a/include/gpufl/inject/capture_gate.cpp b/include/gpufl/inject/capture_gate.cpp index 31e295b..948e0af 100644 --- a/include/gpufl/inject/capture_gate.cpp +++ b/include/gpufl/inject/capture_gate.cpp @@ -257,18 +257,12 @@ void CaptureGate::ResolveLocked() { settled_.store(true, std::memory_order_release); } -GateSnapshot CaptureGate::TakeSnapshotLocked() { +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; - return state_; -} - -void CaptureGate::Publish(const GateSnapshot& snapshot) { - std::lock_guard lock(publish_mu_); - // Snapshots are taken under mu_ with increasing seq but published after - // it is released; an older one arriving late must not replace a newer one. - if (snapshot.seq <= last_published_seq_) return; - last_published_seq_ = snapshot.seq; - if (hooks_.publish) hooks_.publish(snapshot); + if (hooks_.publish) hooks_.publish(state_); } bool CaptureGate::WaitLocked(std::unique_lock& lock, @@ -282,11 +276,8 @@ bool CaptureGate::WaitLocked(std::unique_lock& lock, if (now >= state_.deadline_ns) { if (state_.readiness == GateReadiness::Pending) { LatchFailedLocked("capture_wait_timeout"); - const GateSnapshot snapshot = TakeSnapshotLocked(); - lock.unlock(); + PublishLocked(); cv_.notify_all(); - Publish(snapshot); - lock.lock(); } return until_finished ? state_.lifecycle == GateLifecycle::Finished @@ -314,9 +305,8 @@ void CaptureGate::Enter(const std::function& trigger) { // The host must finish the handshake before capture-relevant work. if (state_.readiness == GateReadiness::Pending && state_.coverage_notes.insert("manual_protocol_violation").second) { - const GateSnapshot snapshot = TakeSnapshotLocked(); + PublishLocked(); lock.unlock(); - Publish(snapshot); } return; } @@ -324,9 +314,8 @@ void CaptureGate::Enter(const std::function& trigger) { if (state_.lifecycle == GateLifecycle::NotRequested) { ElectLocked(GateOwner::Gate); state_.trigger.expected = true; - GateSnapshot snapshot = TakeSnapshotLocked(); + PublishLocked(); lock.unlock(); - Publish(snapshot); GateStep step; { @@ -344,11 +333,8 @@ void CaptureGate::Enter(const std::function& trigger) { LatchFailedLocked("no_injection_callback"); } ResolveLocked(); - snapshot = TakeSnapshotLocked(); - lock.unlock(); + PublishLocked(); cv_.notify_all(); - Publish(snapshot); - lock.lock(); } WaitLocked(lock, /*until_finished=*/false); } @@ -374,19 +360,17 @@ void CaptureGate::OnInjectionCallback() { spawn = true; } } - const GateSnapshot snapshot = TakeSnapshotLocked(); + PublishLocked(); lock.unlock(); - Publish(snapshot); if (spawn) { if (!hooks_.spawn || !hooks_.spawn([this] { RunWorker(); })) { std::unique_lock relock(mu_); state_.worker.returned = true; LatchFailedLocked("worker_spawn_unavailable"); ResolveLocked(); - const GateSnapshot failed = TakeSnapshotLocked(); - relock.unlock(); + PublishLocked(); cv_.notify_all(); - Publish(failed); + relock.unlock(); } } } @@ -404,18 +388,16 @@ void CaptureGate::RunWorker() { LatchFailedLocked(barrier.reason.empty() ? "cuinit_failed" : barrier.reason); ResolveLocked(); - const GateSnapshot snapshot = TakeSnapshotLocked(); - lock.unlock(); + PublishLocked(); cv_.notify_all(); - Publish(snapshot); + lock.unlock(); return; } if (state_.lifecycle == GateLifecycle::DriverInitializing) { state_.lifecycle = GateLifecycle::DriverReady; } - const GateSnapshot snapshot = TakeSnapshotLocked(); + PublishLocked(); lock.unlock(); - Publish(snapshot); } if (hooks_.sleep_ms) hooks_.sleep_ms(config_.delay_ms); { @@ -423,9 +405,8 @@ void CaptureGate::RunWorker() { if (state_.lifecycle == GateLifecycle::DriverReady) { state_.lifecycle = GateLifecycle::CaptureInitializing; } - const GateSnapshot snapshot = TakeSnapshotLocked(); + PublishLocked(); lock.unlock(); - Publish(snapshot); } const GateStep init = Call(hooks_.initialize, "initialize_missing", "init_exception"); @@ -436,10 +417,9 @@ void CaptureGate::RunWorker() { LatchFailedLocked(init.reason.empty() ? "init_failed" : init.reason); } ResolveLocked(); - const GateSnapshot snapshot = TakeSnapshotLocked(); - lock.unlock(); + PublishLocked(); cv_.notify_all(); - Publish(snapshot); + lock.unlock(); } int CaptureGate::ManualHandshake(const std::function& driver_trigger) { @@ -457,9 +437,8 @@ int CaptureGate::ManualHandshake(const std::function& driver_trigger if (state_.lifecycle == GateLifecycle::NotRequested) { ElectLocked(GateOwner::Manual); state_.handshake.expected = true; - GateSnapshot snapshot = TakeSnapshotLocked(); + PublishLocked(); lock.unlock(); - Publish(snapshot); GateStep result; { @@ -469,9 +448,8 @@ int CaptureGate::ManualHandshake(const std::function& driver_trigger if (result.ok) { lock.lock(); state_.lifecycle = GateLifecycle::CaptureInitializing; - snapshot = TakeSnapshotLocked(); + PublishLocked(); lock.unlock(); - Publish(snapshot); result = Call(hooks_.initialize, "initialize_missing", "init_exception"); } @@ -484,11 +462,8 @@ int CaptureGate::ManualHandshake(const std::function& driver_trigger LatchFailedLocked(result.reason.empty() ? "init_failed" : result.reason); } ResolveLocked(); - snapshot = TakeSnapshotLocked(); - lock.unlock(); + PublishLocked(); cv_.notify_all(); - Publish(snapshot); - lock.lock(); } else { WaitLocked(lock, /*until_finished=*/false); } @@ -506,18 +481,16 @@ bool CaptureGate::AwaitAtExit() { } if (WaitLocked(lock, /*until_finished=*/true)) return true; state_.coverage_notes.insert("exit_while_initializing"); - const GateSnapshot snapshot = TakeSnapshotLocked(); + PublishLocked(); lock.unlock(); - Publish(snapshot); return false; } void CaptureGate::AddCoverageNote(const std::string& note) { std::unique_lock lock(mu_); if (!state_.coverage_notes.insert(note).second) return; - const GateSnapshot snapshot = TakeSnapshotLocked(); + PublishLocked(); lock.unlock(); - Publish(snapshot); } WrapperScope::WrapperScope() { ++t_wrapper_depth; } diff --git a/include/gpufl/inject/capture_gate.hpp b/include/gpufl/inject/capture_gate.hpp index 55f7da9..fdd564e 100644 --- a/include/gpufl/inject/capture_gate.hpp +++ b/include/gpufl/inject/capture_gate.hpp @@ -111,8 +111,9 @@ struct GateHooks { std::function sleep_ms; // gpufl::init plus the Trace facility check. Never throws out. std::function initialize; - // Publishes a snapshot (the status file). Called without the gate mutex; - // snapshots carry a strictly increasing seq. + // 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; }; @@ -168,9 +169,8 @@ class CaptureGate { void ElectLocked(GateOwner owner); void LatchFailedLocked(const std::string& reason); void ResolveLocked(); - GateSnapshot TakeSnapshotLocked(); + void PublishLocked(); bool WaitLocked(std::unique_lock& lock, bool until_finished); - void Publish(const GateSnapshot& snapshot); void RunWorker(); void StartWorkerLocked(); @@ -182,8 +182,6 @@ class CaptureGate { GateSnapshot state_; bool worker_started_ = false; std::atomic settled_{false}; - std::mutex publish_mu_; - uint64_t last_published_seq_ = 0; }; /** diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 9df42fa..1db43ae 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -300,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/inject/gate_e2e/fake_cuda.cpp b/tests/inject/gate_e2e/fake_cuda.cpp new file mode 100644 index 0000000..1ad04ac --- /dev/null +++ b/tests/inject/gate_e2e/fake_cuda.cpp @@ -0,0 +1,98 @@ +// 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 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..1750520 --- /dev/null +++ b/tests/inject/gate_e2e/fake_cudart.cpp @@ -0,0 +1,51 @@ +// 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)) +// 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..ad57293 --- /dev/null +++ b/tests/inject/gate_e2e/harness.cpp @@ -0,0 +1,68 @@ +// 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. + +#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..70116ff --- /dev/null +++ b/tests/inject/gate_e2e/run_gate_e2e.py @@ -0,0 +1,233 @@ +"""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 + + +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 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_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_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..d429290 --- /dev/null +++ b/tests/inject/gate_e2e/target.cpp @@ -0,0 +1,86 @@ +// 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); +} + +// 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 (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 at once. + cudaDriverGetVersion(&version); + atexit(app_late_atexit); // registered after the gate's guard + } else if (strcmp(scenario, "none") == 0) { + fake_log("app:no_cuda"); + } + fake_log("app:return"); + return 0; +} From fc11d83a4bf3cca8d31b4936ab658d96f67affa4 Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Fri, 16 Oct 2026 01:09:51 -0700 Subject: [PATCH 4/8] Document the Linux capture gate, its variables and limits --- CHANGELOG.md | 66 ++++++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 66 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 9deb34d..c16d53f 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -168,6 +168,36 @@ 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. 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: ...`. + - **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 +206,42 @@ 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. + `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 From 5902f069cd43e2f930f86ce1e4fd8162ec032c2a Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Fri, 16 Oct 2026 01:12:08 -0700 Subject: [PATCH 5/8] Stop reading the bootstrap status once initialization finished; tidy names --- daemon/launcher/cli_parse.hpp | 2 +- daemon/launcher/trace_command.cpp | 8 +++++--- include/gpufl/core/env_vars.hpp | 1 - 3 files changed, 6 insertions(+), 5 deletions(-) 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 802b3d9..8c59589 100644 --- a/daemon/launcher/trace_command.cpp +++ b/daemon/launcher/trace_command.cpp @@ -152,14 +152,16 @@ class PosixTracePlatform final : public TracePlatform { const bool watch = !opts.bootstrap_dir.empty(); InitWatchdog watchdog(opts.watchdog_grace_ms * 1'000'000LL); BootstrapStatus last_read; - auto watchdogFires = [&]() -> bool { - if (!watch) return false; + 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(); // the final record after exit + (void)watchdogFires(/*final_read=*/true); // the record after exit out.bootstrap = watchdog.seen_valid() ? watchdog.last() : last_read; }; diff --git a/include/gpufl/core/env_vars.hpp b/include/gpufl/core/env_vars.hpp index 3b41faa..195c9fa 100644 --- a/include/gpufl/core/env_vars.hpp +++ b/include/gpufl/core/env_vars.hpp @@ -116,7 +116,6 @@ 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"; From 58da3bf897d22970d8f5359e5cc115f3de29eaf8 Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Fri, 16 Oct 2026 01:30:07 -0700 Subject: [PATCH 6/8] Address the capture gate code review Warmup: the initialization deadline now includes the requested delay once, so a healthy warmup longer than the init allowance is not killed, and an exit that starts during the delay stops the worker before CUPTI (exit_before_capture) instead of waiting it out. The delay hook is gone. Facilities: every required activity kind's enable result is recorded, at start and at the re-enable after engine start. A non-success latches, including the results CUPTI logs as benign, and a required kind lost after readiness adds trace_facility_lost at shutdown. The facility query no longer depends on the monitor's initialized flag, which shutdown clears before the final capability report. Launcher: the read after exit is combined with the last accepted record; a missing, malformed, foreign or older final record is reported as such next to the accepted one instead of being replaced by it. The status parser requires every field the library writes and refuses states the library cannot be in. Tests: every first-touch runtime wrapper and the copy and sync wrappers end to end, the driver wrappers after an unwrapped cuInit, exit during a warmup, a healthy long warmup and a wedge after it under the real launcher, the required-kind tally, and the final-read combination. The comments no longer state the attach overlap as the established cause. --- CHANGELOG.md | 16 +++- daemon/launcher/bootstrap_status.cpp | 81 +++++++++++++--- daemon/launcher/bootstrap_status.hpp | 14 +++ daemon/launcher/trace_command.cpp | 5 +- daemon/launcher/trace_command_common.cpp | 17 +++- daemon/launcher/trace_command_common.hpp | 3 + .../gpufl/backends/nvidia/cupti_backend.cpp | 23 +++-- .../gpufl/backends/nvidia/cupti_backend.hpp | 42 ++++++--- include/gpufl/core/monitor.cpp | 4 +- include/gpufl/inject/capture_gate.cpp | 26 +++++- include/gpufl/inject/capture_gate.hpp | 17 +++- include/gpufl/inject/capture_gate_linux.cpp | 10 +- include/gpufl/inject/inject_entry.cpp | 11 ++- tests/backends/nvidia/test_nvidia_backend.cpp | 27 ++++++ tests/inject/gate_e2e/fake_cuda.cpp | 8 ++ tests/inject/gate_e2e/fake_cudart.cpp | 12 +++ tests/inject/gate_e2e/harness.cpp | 5 + tests/inject/gate_e2e/run_gate_e2e.py | 61 ++++++++++++ tests/inject/gate_e2e/target.cpp | 83 ++++++++++++++++- tests/inject/test_capture_gate.cpp | 1 - tests/launcher/test_bootstrap_status.cpp | 93 +++++++++++++++---- 21 files changed, 476 insertions(+), 83 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index c16d53f..b70f7a9 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -171,7 +171,10 @@ types are additive. - **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. A new library launched with + 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 @@ -180,7 +183,11 @@ types are additive. - **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: ...`. + 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 @@ -215,7 +222,10 @@ types are additive. - **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. + 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`, diff --git a/daemon/launcher/bootstrap_status.cpp b/daemon/launcher/bootstrap_status.cpp index 94ad753..62f485c 100644 --- a/daemon/launcher/bootstrap_status.cpp +++ b/daemon/launcher/bootstrap_status.cpp @@ -46,23 +46,59 @@ BootstrapStatus ReadBootstrapStatus(const std::filesystem::path& dir, } if (doc.value("pid", 0) != pid) return Invalid("foreign_pid"); - BootstrapStatus status; - status.valid = true; - status.seq = doc.value("seq", 0); - status.lifecycle = doc.value("lifecycle", ""); - status.owner = doc.value("owner", ""); - status.readiness = doc.value("readiness", ""); - status.reason = doc.value("reason", ""); - status.deadline_monotonic_ns = doc.value("deadline_monotonic_ns", 0); - status.protocol = doc.value("protocol", ""); - if (doc.contains("coverage_notes") && doc["coverage_notes"].is_array()) { - for (const json::JsonValue& note : doc["coverage_notes"].get_array()) { - if (note.is_string()) status.coverage_notes.insert(note.get_string()); + // Every field the library writes must be there with its type, and the + // values must name a state the library can be in. + 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"}) { + if (!has(key, Type::Int)) 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 (status.lifecycle.empty() || status.readiness.empty()) { - 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"); } + status.valid = true; return status; } @@ -98,6 +134,23 @@ bool InitWatchdog::Observe(const BootstrapStatus& status, const int64_t now_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 + diff --git a/daemon/launcher/bootstrap_status.hpp b/daemon/launcher/bootstrap_status.hpp index 09a8994..0cc9884 100644 --- a/daemon/launcher/bootstrap_status.hpp +++ b/daemon/launcher/bootstrap_status.hpp @@ -68,6 +68,20 @@ class InitWatchdog { 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); diff --git a/daemon/launcher/trace_command.cpp b/daemon/launcher/trace_command.cpp index 8c59589..787296a 100644 --- a/daemon/launcher/trace_command.cpp +++ b/daemon/launcher/trace_command.cpp @@ -162,7 +162,10 @@ class PosixTracePlatform final : public TracePlatform { auto finishBootstrap = [&]() { if (!watch) return; (void)watchdogFires(/*final_read=*/true); // the record after exit - out.bootstrap = watchdog.seen_valid() ? watchdog.last() : last_read; + // 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) { diff --git a/daemon/launcher/trace_command_common.cpp b/daemon/launcher/trace_command_common.cpp index c9791e9..d9c4c44 100644 --- a/daemon/launcher/trace_command_common.cpp +++ b/daemon/launcher/trace_command_common.cpp @@ -208,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) { @@ -1000,7 +1013,7 @@ int runTraceCommon(const TraceArgs& args, const TracePlatform& platform) { std::fprintf(stderr, "[gpufl] %s: capture initialization never finished; the " "watchdog killed the target (%s)\n", - what.c_str(), DescribeBootstrapStatus(run.bootstrap).c_str()); + 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")) { @@ -1011,7 +1024,7 @@ int runTraceCommon(const TraceArgs& args, const TracePlatform& platform) { } if (!args.quiet) { std::fprintf(stderr, "[gpufl] %s capture readiness: %s\n", what.c_str(), - DescribeBootstrapStatus(run.bootstrap).c_str()); + describeRunBootstrap(run).c_str()); } #endif diff --git a/daemon/launcher/trace_command_common.hpp b/daemon/launcher/trace_command_common.hpp index def8631..4f75f00 100644 --- a/daemon/launcher/trace_command_common.hpp +++ b/daemon/launcher/trace_command_common.hpp @@ -29,6 +29,9 @@ struct TraceProcessResult { 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; }; diff --git a/include/gpufl/backends/nvidia/cupti_backend.cpp b/include/gpufl/backends/nvidia/cupti_backend.cpp index 4f393fd..ec5416b 100644 --- a/include/gpufl/backends/nvidia/cupti_backend.cpp +++ b/include/gpufl/backends/nvidia/cupti_backend.cpp @@ -508,18 +508,11 @@ void CuptiBackend::start() { for (auto k : kinds) { const CUptiResult res = cuptiActivityEnable(k); CUPTI_CHECK(res); - // Capture readiness requires every required kind to be on; the - // injection gate reads this (captureFacilitiesReady). - if (res != CUPTI_SUCCESS && !IsBenignCuptiResult(res)) { - std::lock_guard lock(capture_facility_mu_); - if (capture_facility_failure_.empty()) { - capture_facility_failure_ = - "activity_enable_failed:" + std::to_string(static_cast(k)) + - ":" + std::to_string(static_cast(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)); } - required_activity_enabled_.store(true, std::memory_order_release); } // Initialize and start the engine (requires CUDA context) @@ -550,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 diff --git a/include/gpufl/backends/nvidia/cupti_backend.hpp b/include/gpufl/backends/nvidia/cupti_backend.hpp index 1a11033..3daa57a 100644 --- a/include/gpufl/backends/nvidia/cupti_backend.hpp +++ b/include/gpufl/backends/nvidia/cupti_backend.hpp @@ -71,6 +71,35 @@ struct ActivityTimestampTally { } }; +// 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. * @@ -154,14 +183,7 @@ class CuptiBackend : public IMonitorBackend { if (!cupti_subscribed_.load(std::memory_order_acquire)) { return fail("cupti_subscribe_failed"); } - if (!required_activity_enabled_.load(std::memory_order_acquire)) { - return fail("activity_not_enabled"); - } - std::lock_guard lock(capture_facility_mu_); - if (!capture_facility_failure_.empty()) { - return fail(capture_facility_failure_); - } - return true; + return required_activity_.Ready(reason); } void NoteMemTransferActivityEligible() { mem_transfer_activity_timestamps_.NoteEligible(); @@ -433,9 +455,7 @@ class CuptiBackend : public IMonitorBackend { std::atomic kernel_activity_throttled_{0}; ActivityTimestampTally kernel_activity_timestamps_; std::atomic cupti_subscribed_{false}; - std::atomic required_activity_enabled_{false}; - mutable std::mutex capture_facility_mu_; - std::string capture_facility_failure_; + RequiredActivityTally required_activity_; std::atomic mem_transfer_activity_emitted_{0}; ActivityTimestampTally mem_transfer_activity_timestamps_; std::atomic sync_activity_emitted_{0}; diff --git a/include/gpufl/core/monitor.cpp b/include/gpufl/core/monitor.cpp index 3d98273..2c754b9 100644 --- a/include/gpufl/core/monitor.cpp +++ b/include/gpufl/core/monitor.cpp @@ -742,7 +742,9 @@ 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; - if (!g_state.initialized.load()) { + // 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; } diff --git a/include/gpufl/inject/capture_gate.cpp b/include/gpufl/inject/capture_gate.cpp index 948e0af..3b78ec6 100644 --- a/include/gpufl/inject/capture_gate.cpp +++ b/include/gpufl/inject/capture_gate.cpp @@ -216,8 +216,13 @@ 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(); - state_.deadline_ns = now > max - config_.wait_ns ? max : now + config_.wait_ns; + 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) { @@ -399,14 +404,23 @@ void CaptureGate::RunWorker() { PublishLocked(); lock.unlock(); } - if (hooks_.sleep_ms) hooks_.sleep_ms(config_.delay_ms); { + // 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_); - if (state_.lifecycle == GateLifecycle::DriverReady) { - state_.lifecycle = GateLifecycle::CaptureInitializing; + 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(); - lock.unlock(); } const GateStep init = Call(hooks_.initialize, "initialize_missing", "init_exception"); @@ -479,6 +493,8 @@ bool CaptureGate::AwaitAtExit() { 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(); diff --git a/include/gpufl/inject/capture_gate.hpp b/include/gpufl/inject/capture_gate.hpp index fdd564e..cbea6a9 100644 --- a/include/gpufl/inject/capture_gate.hpp +++ b/include/gpufl/inject/capture_gate.hpp @@ -4,10 +4,17 @@ // // 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, and on an L4 -// (2026-10-16) whole sessions lost every kernel timestamp that way. The gate -// makes the application's first supported CUDA call wait until capture is -// ready, within one bounded deadline, and records every other outcome. +// 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): @@ -108,7 +115,6 @@ struct GateHooks { std::function body)> spawn; // The worker's checked post-driver barrier (cuInit). std::function driver_barrier; - std::function sleep_ms; // 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 @@ -181,6 +187,7 @@ class CaptureGate { std::condition_variable cv_; GateSnapshot state_; bool worker_started_ = false; + bool exit_requested_ = false; std::atomic settled_{false}; }; diff --git a/include/gpufl/inject/capture_gate_linux.cpp b/include/gpufl/inject/capture_gate_linux.cpp index 6c4c26e..677d76d 100644 --- a/include/gpufl/inject/capture_gate_linux.cpp +++ b/include/gpufl/inject/capture_gate_linux.cpp @@ -1,7 +1,8 @@ // NVIDIA calls InitializeInjection from inside the driver's own cuInit, so -// CUPTI cannot start there; a worker starts it after the driver is up. On an -// L4 (2026-10-16) the application thread raced that worker into context and -// module setup, and whole sessions lost every kernel timestamp. The gate makes +// 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. @@ -202,9 +203,6 @@ CaptureGate& LinuxCaptureGate() { } }; hooks.driver_barrier = CheckedCuInit; - hooks.sleep_ms = [](const int ms) { - if (ms > 0) std::this_thread::sleep_for(std::chrono::milliseconds(ms)); - }; hooks.initialize = BoundInitialize; hooks.publish = PublishStatus; auto* created = new CaptureGate(GateConfigFromEnv(nullptr), std::move(hooks)); diff --git a/include/gpufl/inject/inject_entry.cpp b/include/gpufl/inject/inject_entry.cpp index 0038599..8f71a51 100644 --- a/include/gpufl/inject/inject_entry.cpp +++ b/include/gpufl/inject/inject_entry.cpp @@ -596,8 +596,15 @@ gpufl::inject::GateStep initializeCapture() { } void contributeGateCapabilities(gpufl::CaptureCapabilitiesEvent& evt) { - gpufl::inject::AppendGateCapabilities(gpufl::inject::LinuxCaptureGate().Snapshot(), - 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"); + } + gpufl::inject::AppendGateCapabilities(snapshot, evt); } void registerGateCapabilityContributor() { 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 index 1ad04ac..dc4916d 100644 --- a/tests/inject/gate_e2e/fake_cuda.cpp +++ b/tests/inject/gate_e2e/fake_cuda.cpp @@ -68,6 +68,14 @@ int cuModuleLoadData(void** module, const void* image) { 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"); diff --git a/tests/inject/gate_e2e/fake_cudart.cpp b/tests/inject/gate_e2e/fake_cudart.cpp index 1750520..11ba473 100644 --- a/tests/inject/gate_e2e/fake_cudart.cpp +++ b/tests/inject/gate_e2e/fake_cudart.cpp @@ -45,6 +45,18 @@ 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)) diff --git a/tests/inject/gate_e2e/harness.cpp b/tests/inject/gate_e2e/harness.cpp index ad57293..17f6cf2 100644 --- a/tests/inject/gate_e2e/harness.cpp +++ b/tests/inject/gate_e2e/harness.cpp @@ -6,6 +6,11 @@ // 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 diff --git a/tests/inject/gate_e2e/run_gate_e2e.py b/tests/inject/gate_e2e/run_gate_e2e.py index 70116ff..5f154c7 100644 --- a/tests/inject/gate_e2e/run_gate_e2e.py +++ b/tests/inject/gate_e2e/run_gate_e2e.py @@ -20,6 +20,18 @@ 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")) @@ -50,6 +62,10 @@ def setUp(self): 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_") @@ -160,6 +176,37 @@ def test_exit_waits_for_a_late_initializer_before_teardown(self): 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) @@ -211,6 +258,20 @@ def test_launcher_watchdog_also_bounds_a_windowed_run(self): 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) diff --git a/tests/inject/gate_e2e/target.cpp b/tests/inject/gate_e2e/target.cpp index d429290..448158a 100644 --- a/tests/inject/gate_e2e/target.cpp +++ b/tests/inject/gate_e2e/target.cpp @@ -16,6 +16,68 @@ 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, @@ -41,7 +103,21 @@ int main(int argc, char** argv) { unsigned long long dp = 0; int version = 0; - if (strcmp(scenario, "runtime") == 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(); @@ -75,9 +151,12 @@ int main(int argc, char** argv) { 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 at once. + // 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"); } diff --git a/tests/inject/test_capture_gate.cpp b/tests/inject/test_capture_gate.cpp index 0d54bdd..bf10bd0 100644 --- a/tests/inject/test_capture_gate.cpp +++ b/tests/inject/test_capture_gate.cpp @@ -96,7 +96,6 @@ struct Fakes { Event("barrier"); return barrier(); }; - hooks.sleep_ms = [](int) {}; hooks.initialize = [this] { ++initialized; Event("init"); diff --git a/tests/launcher/test_bootstrap_status.cpp b/tests/launcher/test_bootstrap_status.cpp index 8f5e311..e62ee8e 100644 --- a/tests/launcher/test_bootstrap_status.cpp +++ b/tests/launcher/test_bootstrap_status.cpp @@ -9,11 +9,31 @@ 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 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") { + const std::string participant = + "{\"expected\":true,\"returned\":false,\"ok\":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\":" + participant + ",\"worker\":" + participant + + ",\"handshake\":" + participant + + ",\"callbacks_seen\":1,\"handshake_reentries\":0," + + "\"wait_config_invalid\":false,\"protocol\":\"2\"}\n"; +} + class BootstrapStatusTest : public ::testing::Test { protected: void SetUp() override { @@ -27,14 +47,8 @@ class BootstrapStatusTest : public ::testing::Test { void Write(long pid, const std::string& text) { std::ofstream(gpufl::launcher::BootstrapStatusPath(dir_, pid)) << text; } - - static std::string Record(const std::string& nonce, long pid, int seq, - const std::string& lifecycle, int64_t deadline) { - return "{\"schema\":1,\"nonce\":\"" + nonce + "\",\"pid\":" + std::to_string(pid) + - ",\"seq\":" + std::to_string(seq) + ",\"lifecycle\":\"" + lifecycle + - "\",\"owner\":\"gate\",\"readiness\":\"pending\",\"reason\":\"\"," + - "\"coverage_notes\":[\"unsupported_first_touch\"]," + - "\"deadline_monotonic_ns\":" + std::to_string(deadline) + "}"; + BootstrapStatus Read(long pid, const std::string& nonce = "n1") { + return ReadBootstrapStatus(dir_, pid, nonce); } fs::path dir_; @@ -45,30 +59,48 @@ BootstrapStatus Status(int seq, const std::string& lifecycle, int64_t deadline) s.valid = true; s.seq = seq; s.lifecycle = lifecycle; - s.readiness = "pending"; + s.readiness = lifecycle == "finished" ? "ready" : "pending"; s.deadline_monotonic_ns = deadline; return s; } } // namespace -TEST_F(BootstrapStatusTest, ReadsOnlyThisChildsOwnRecord) { - Write(42, Record("n1", 42, 3, "capture_initializing", 1000)); - const BootstrapStatus s = ReadBootstrapStatus(dir_, 42, "n1"); +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(ReadBootstrapStatus(dir_, 42, "other").error, "foreign_nonce"); - Write(43, Record("n1", 42, 3, "finished", 1000)); - EXPECT_EQ(ReadBootstrapStatus(dir_, 43, "n1").error, "foreign_pid"); - EXPECT_EQ(ReadBootstrapStatus(dir_, 44, "n1").error, "missing"); + 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(ReadBootstrapStatus(dir_, 45, "n1").error, "malformed"); + EXPECT_EQ(Read(45).error, "malformed"); Write(46, std::string(5000, ' ')); - EXPECT_EQ(ReadBootstrapStatus(dir_, 46, "n1").error, "oversized"); + 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) { @@ -97,3 +129,28 @@ TEST(InitWatchdog, DisarmsWhenEveryParticipantReturnedAndIgnoresStaleRecords) { 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"); +} From 5b0989f2205a6d26a035d8cd58e62548d0101b6d Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Fri, 16 Oct 2026 01:39:44 -0700 Subject: [PATCH 7/8] Validate bootstrap participants and round-trip the gate's own snapshots --- daemon/launcher/bootstrap_status.cpp | 44 ++++++++- tests/launcher/test_bootstrap_status.cpp | 111 ++++++++++++++++++++++- 2 files changed, 148 insertions(+), 7 deletions(-) diff --git a/daemon/launcher/bootstrap_status.cpp b/daemon/launcher/bootstrap_status.cpp index 62f485c..8fff56a 100644 --- a/daemon/launcher/bootstrap_status.cpp +++ b/daemon/launcher/bootstrap_status.cpp @@ -46,15 +46,19 @@ BootstrapStatus ReadBootstrapStatus(const std::filesystem::path& dir, } if (doc.value("pid", 0) != pid) return Invalid("foreign_pid"); - // Every field the library writes must be there with its type, and the - // values must name a state the library can be in. + // 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"}) { + 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"); } @@ -98,6 +102,40 @@ BootstrapStatus ReadBootstrapStatus(const std::filesystem::path& dir, (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; } diff --git a/tests/launcher/test_bootstrap_status.cpp b/tests/launcher/test_bootstrap_status.cpp index e62ee8e..8c6040f 100644 --- a/tests/launcher/test_bootstrap_status.cpp +++ b/tests/launcher/test_bootstrap_status.cpp @@ -7,6 +7,15 @@ #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; @@ -16,20 +25,29 @@ 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") { - const std::string participant = - "{\"expected\":true,\"returned\":false,\"ok\":false}"; + // 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\":" + participant + ",\"worker\":" + participant + - ",\"handshake\":" + participant + + ",\"trigger\":" + trigger + ",\"worker\":" + worker + + ",\"handshake\":" + handshake + ",\"callbacks_seen\":1,\"handshake_reentries\":0," + "\"wait_config_invalid\":false,\"protocol\":\"2\"}\n"; } @@ -154,3 +172,88 @@ TEST(CombineFinalBootstrap, AFinalReadThatDoesNotConfirmIsReportedNotHidden) { 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 success, failure, timeout and manual paths, parses as valid. +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 From 4ae50229a5692471a36fc534567d6db257e3f8b0 Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Fri, 16 Oct 2026 01:46:00 -0700 Subject: [PATCH 8/8] Say which gate paths the bootstrap round trip covers --- tests/launcher/test_bootstrap_status.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/tests/launcher/test_bootstrap_status.cpp b/tests/launcher/test_bootstrap_status.cpp index 8c6040f..91e6fa8 100644 --- a/tests/launcher/test_bootstrap_status.cpp +++ b/tests/launcher/test_bootstrap_status.cpp @@ -175,7 +175,8 @@ TEST(CombineFinalBootstrap, AFinalReadThatDoesNotConfirmIsReportedNotHidden) { #if defined(__linux__) // Round trip against the library's own writer: every snapshot a real gate -// publishes, on success, failure, timeout and manual paths, parses as valid. +// 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 {