From 3670c89f593ff2c196a1c10c87cbe938b5e6275d Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Fri, 24 Jul 2026 22:19:57 -0700 Subject: [PATCH 01/10] fix(capture): carry the clock anchor into the deferred engine start, bounded deep-profiling window --- CHANGELOG.md | 47 ++ CMakeLists.txt | 2 + daemon/launcher/cli_parse.cpp | 50 +++ daemon/launcher/cli_parse.hpp | 11 + daemon/launcher/trace_command_common.cpp | 28 ++ example/cuda/CMakeLists.txt | 12 + example/cuda/deep_window_demo.cu | 227 ++++++++++ .../gpufl/backends/nvidia/cupti_backend.cpp | 21 +- .../gpufl/backends/nvidia/cupti_backend.hpp | 29 +- .../nvidia/engine/pc_sampling_engine.cpp | 36 +- .../nvidia/engine/pm_sampling_engine.cpp | 8 +- .../nvidia/engine/pm_sampling_engine.hpp | 7 + .../nvidia/engine/range_profiler_engine.cpp | 55 ++- .../nvidia/engine/range_profiler_engine.hpp | 3 + include/gpufl/core/deep_window.cpp | 404 ++++++++++++++++++ include/gpufl/core/deep_window.hpp | 140 ++++++ include/gpufl/core/env_vars.hpp | 16 + include/gpufl/core/events.hpp | 27 ++ include/gpufl/core/gpufl.cpp | 70 +-- .../gpufl/core/model/deep_window_model.cpp | 28 ++ .../gpufl/core/model/deep_window_model.hpp | 23 + include/gpufl/core/monitor.cpp | 4 + include/gpufl/core/monitor.hpp | 28 ++ include/gpufl/core/monitor_backend.hpp | 17 + include/gpufl/gpufl.hpp | 67 +++ python/bindings.cpp | 12 + python/gpufl/__init__.py | 62 +++ scripts/deep_window_check.py | 157 +++++++ scripts/deep_window_e2e.sh | 208 +++++++++ scripts/deep_window_target.cu | 72 ++++ tests/CMakeLists.txt | 1 + tests/core/test_deep_window.cpp | 325 ++++++++++++++ tests/launcher/test_cli_parse.cpp | 52 +++ 33 files changed, 2186 insertions(+), 63 deletions(-) create mode 100644 example/cuda/deep_window_demo.cu create mode 100644 include/gpufl/core/deep_window.cpp create mode 100644 include/gpufl/core/deep_window.hpp create mode 100644 include/gpufl/core/model/deep_window_model.cpp create mode 100644 include/gpufl/core/model/deep_window_model.hpp create mode 100644 scripts/deep_window_check.py create mode 100644 scripts/deep_window_e2e.sh create mode 100644 scripts/deep_window_target.cu create mode 100644 tests/core/test_deep_window.cpp diff --git a/CHANGELOG.md b/CHANGELOG.md index 232fa22..ab76184 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -9,6 +9,53 @@ versioning follows PEP 440 for the Python wheel and semver-style ### Added +- **Bounded deep-profiling windows.** `gpufl::deepWindow(ms, max_launches)` + (Python: `gpufl.deep_window(seconds, max_launches)`) arms the deep engines + for a short region and disarms them automatically, so a long-running job + can profile the moment it went wrong without carrying replay cost for its + whole lifetime: + + ```cpp + if (tokens_per_sec < 1000) gpufl::deepWindow(3000); + ``` + + Calling it again while a window is open is ignored rather than treated as + an extension, so a check inside a training loop can run every step. The + duration and launch bounds combine with OR; prefer a launch budget for the + replay engines, where wall time and work done diverge sharply. Each window + emits a `deep_window` event carrying the launches it actually covered and + which bound closed it. + + Set `InitOptions::deep_window_only` (Python `deep_window_only=True`, env + `GPUFL_DEEP_ARM=window`) to keep PC sampling, SASS metrics, PM sampling and + the Range profiler idle outside windows. `GPUFL_DEEP_WINDOW_MS` and + `GPUFL_DEEP_WINDOW_MAX_LAUNCHES` supply defaults for bounds left at 0. + + `DeepWindowSpec::cooldown_ms` (`GPUFL_DEEP_WINDOW_COOLDOWN_MS`) sets the + minimum quiet time between windows. Without it a condition that stays true + reopens a window the instant the last one expired; only the library knows + when that was, so the bound lives there rather than in your trigger. +- **`gpufl trace --deep-after` / `--deep-for` / `--deep-launches` / + `--deep-cooldown`.** Arms a deep window inside a target whose source you + can't edit, which otherwise has no way to call `deepWindow()`. Unlike + `--window`, which bounds the target's LIFETIME, these leave it running and + bound only how long the deep engines stay armed: + + ```bash + gpufl trace --deep-after 30s --deep-for 3s --passes PmSampling -- python train.py + ``` + + Any `--deep-*` flag implies `GPUFL_DEEP_ARM=window`. A window with neither + a duration nor a launch bound is rejected. + +### Fixed + +- **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 + CUPTI-to-wall-clock anchor. Engines that stamp their own samples emitted + raw CUPTI timestamps, putting every PM sample days away from the kernel + timeline it should line up with. - **Shared-memory bank-conflict profiling.** `RangeProfilerKernelReplay` now emits per-kernel shared load/store/total conflict counts, shared wavefronts, conflict overhead, and average N-way diff --git a/CMakeLists.txt b/CMakeLists.txt index d5bc18c..fb0fc1c 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -101,6 +101,8 @@ target_sources(gpufl PRIVATE include/gpufl/core/model/memory_alloc_event_model.cpp include/gpufl/core/model/graph_launch_event_model.cpp include/gpufl/core/model/system_event_model.cpp + include/gpufl/core/model/deep_window_model.cpp + include/gpufl/core/deep_window.cpp include/gpufl/core/sampler.cpp include/gpufl/core/runtime.cpp include/gpufl/core/backend_factory.cpp diff --git a/daemon/launcher/cli_parse.cpp b/daemon/launcher/cli_parse.cpp index fbcd901..37bf6c9 100644 --- a/daemon/launcher/cli_parse.cpp +++ b/daemon/launcher/cli_parse.cpp @@ -168,6 +168,18 @@ const char* traceHelp() { " Hard cap on total target runtime (safety).\n" " --after-window=\n" " What to do at window end. Only 'stop' today.\n" + " --deep-after= Arm the DEEP engines this long into the run,\n" + " then disarm. Unlike --window the target keeps\n" + " running. Needs a bound below. Default: 0 (arm\n" + " at the first kernel launch).\n" + " --deep-for= How long the deep window stays armed.\n" + " --deep-launches= Kernel-launch bound on the deep window; ends it\n" + " at whichever bound is hit first. PREFER THIS for\n" + " SASS / Range: replay re-runs every kernel, so a\n" + " second of wall time covers ~25x less work there\n" + " than under PM sampling.\n" + " --deep-cooldown=\n" + " Quiet time before another window may open.\n" " --pc-sample-period=\n" " PC sampling period: log2 of GPU cycles per sample\n" " (5..31; default 10). Lower = more frequent — for\n" @@ -369,6 +381,35 @@ TraceParseResult parseTraceArgs(const std::vector& argv) { " (expected a duration like 30s, 5m, 1h, " "or a bare number of seconds)"}; } + } else if (key == "--deep-after" || key == "--deep-for" || + key == "--deep-cooldown") { + std::string v; + auto err = take_value(v); + if (!err.empty()) return {std::nullopt, err}; + int64_t ms = 0; + if (!parseDurationMs(v, ms)) { + return {std::nullopt, + "invalid " + key + " value: " + v + + " (expected a duration like 30s, 500ms, 5m, 1h, " + "or a bare number of seconds)"}; + } + if (key == "--deep-after") out.deep_after_ms = ms; + else if (key == "--deep-for") out.deep_for_ms = ms; + else out.deep_cooldown_ms = ms; + out.deep_requested = true; + } else if (key == "--deep-launches") { + std::string v; + auto err = take_value(v); + if (!err.empty()) return {std::nullopt, err}; + char* end = nullptr; + const unsigned long long n = std::strtoull(v.c_str(), &end, 10); + if (end == v.c_str() || (end && *end != '\0') || n == 0) { + return {std::nullopt, + "invalid --deep-launches value: " + v + + " (expected a positive number of kernel launches)"}; + } + out.deep_launches = static_cast(n); + out.deep_requested = true; } else if (key == "--after-window") { auto err = take_value(out.after_window); if (!err.empty()) return {std::nullopt, err}; @@ -398,6 +439,15 @@ TraceParseResult parseTraceArgs(const std::vector& argv) { if (out.command.empty()) { return {std::nullopt, "no command specified after `--`"}; } + // A deep window with neither bound would arm and never disarm, which is + // just "profile deeply for the whole run" with extra steps. + if (out.deep_requested && out.deep_for_ms == 0 && out.deep_launches == 0) { + return {std::nullopt, + "a deep window needs a bound: pass --deep-for " + "or --deep-launches (prefer --deep-launches for the " + "replay engines, where a second of wall time covers far " + "less work)"}; + } return {out, ""}; } diff --git a/daemon/launcher/cli_parse.hpp b/daemon/launcher/cli_parse.hpp index 2e125b3..f8a35d1 100644 --- a/daemon/launcher/cli_parse.hpp +++ b/daemon/launcher/cli_parse.hpp @@ -37,6 +37,17 @@ struct TraceArgs { int64_t window_ms = 0; // --window; 0 = run to the target's natural exit int64_t window_timeout_ms = 0; // --window-timeout; hard cap on total runtime (0 = warmup+window) std::string after_window = "stop"; // --after-window; "stop" is the only value today + // Bounded DEEP window - unrelated to --window above, which bounds the + // target's LIFETIME. These keep the target running and instead bound how + // long the deep engines (PC sampling / SASS / PM / Range) stay armed + // inside it. A target whose source can't be edited has no way to call + // gpufl::deepWindow(), so time is the trigger the launcher can offer. + // Any of them turns on window-only arming (GPUFL_DEEP_ARM=window). + int64_t deep_after_ms = 0; // --deep-after; 0 = arm at the first launch + int64_t deep_for_ms = 0; // --deep-for; duration bound, 0 = none + uint64_t deep_launches = 0; // --deep-launches; launch bound, 0 = none + int64_t deep_cooldown_ms = 0; // --deep-cooldown; quiet time between windows + bool deep_requested = false; // any --deep-* flag was given // PC sampling period as a log2 exponent (2^N GPU cycles/sample, valid 5..31; // lower = more frequent → catches shorter kernels). 0 = leave the engine // default. Plumbed to the injected target via GPUFL_PC_SAMPLING_PERIOD. diff --git a/daemon/launcher/trace_command_common.cpp b/daemon/launcher/trace_command_common.cpp index 77e648f..9cb9b88 100644 --- a/daemon/launcher/trace_command_common.cpp +++ b/daemon/launcher/trace_command_common.cpp @@ -616,6 +616,34 @@ int runTraceCommon(const TraceArgs& args, const TracePlatform& platform) { return 2; } + // --deep-*: bound how long the DEEP engines stay armed inside a target + // that keeps running. Distinct from --window above, which bounds the + // target's lifetime. Asking for a deep window implies window-only + // arming, or the engines would be armed from the first kernel and the + // window would bound nothing. + if (args.deep_requested) { + if (!setEnvOrPrint(platform, env::kDeepArm, "window") || + !setEnvOrPrint(platform, env::kDeepAfterMs, + std::to_string(args.deep_after_ms))) { + return 2; + } + if (args.deep_for_ms > 0 && + !setEnvOrPrint(platform, env::kDeepWindowMs, + std::to_string(args.deep_for_ms))) { + return 2; + } + if (args.deep_launches > 0 && + !setEnvOrPrint(platform, env::kDeepWindowMaxLaunches, + std::to_string(args.deep_launches))) { + return 2; + } + if (args.deep_cooldown_ms > 0 && + !setEnvOrPrint(platform, env::kDeepWindowCooldownMs, + std::to_string(args.deep_cooldown_ms))) { + return 2; + } + } + // A bounded window stops the target after warmup+window wall-clock; // run_ms == 0 keeps the historical "run until the target exits" behavior. RunOptions run_opts; diff --git a/example/cuda/CMakeLists.txt b/example/cuda/CMakeLists.txt index 610ed32..4b6e896 100644 --- a/example/cuda/CMakeLists.txt +++ b/example/cuda/CMakeLists.txt @@ -42,6 +42,7 @@ add_executable(memory_coalescing_demo memory_coalescing_demo.cu) add_executable(shared_bank_conflicts_demo shared_bank_conflicts_demo.cu) add_executable(deep_deadlock_repro deep_deadlock_repro.cu) add_executable(multi_engine_demo multi_engine_demo.cu) +add_executable(deep_window_demo deep_window_demo.cu) target_compile_options(gfl_block_example PRIVATE $<$: @@ -183,6 +184,17 @@ target_link_libraries(multi_engine_demo PRIVATE CUDA::cudart ) +target_link_libraries(deep_window_demo PRIVATE + gpufl::gpufl + CUDA::cupti + CUDA::cudart +) +# PC / SASS want the cubin embedded for source correlation. +target_compile_options(deep_window_demo PRIVATE + $<$: + -lineinfo + >) + # Properties set_target_properties(gfl_block_example PROPERTIES diff --git a/example/cuda/deep_window_demo.cu b/example/cuda/deep_window_demo.cu new file mode 100644 index 0000000..48db5a0 --- /dev/null +++ b/example/cuda/deep_window_demo.cu @@ -0,0 +1,227 @@ +// deep_window_demo.cu +// +// Exercises the bounded deep-profiling window end to end: a long-running +// kernel loop that arms deep capture for a few seconds when a metric drops, +// then keeps running once the window closes itself. +// +// It stands in for the real case - a training loop that notices its +// throughput fell and wants to know why, without paying replay cost for the +// whole run. +// +// ── WHAT TO CHECK ──────────────────────────────────────────────────────────── +// 1. Before the trigger: the engine is armed but IDLE, so the session +// carries no PC / PM / SASS samples from those iterations. +// 2. During the window: samples appear, bounded to the window. +// 3. After the window: the process keeps running to completion. A window +// ending is not a session ending. +// 4. The trigger fires 50 times in a row and still opens ONE window - a +// re-fire while a window is open is ignored, not an extension. (Note +// that a trigger left firing across many iterations would correctly +// open a NEW window each time the previous one expired; that is the +// caller's cooldown to own, not the library's.) +// 5. The deep_window event records how many launches the window actually +// covered and which bound closed it. Under a replay engine that count +// is small; that is the engine, not a failure. +// +// ── RUN ────────────────────────────────────────────────────────────────────── +// PowerShell: +// $env:GPUFL_PROFILING_ENGINE = "PmSampling" # or PcSampling / SassMetrics +// .\deep_window_demo.exe +// +// Bound the window by launches instead of time (better for replay engines): +// $env:GPUFL_DEEP_WINDOW_MAX_LAUNCHES = "200" +// +// PerfWorks engines (PmSampling / RangeProfiler*) need GPU performance-counter +// access: run elevated, or NVIDIA Control Panel -> Developer -> "Manage GPU +// Performance Counters" -> allow all users. + +#include + +#include // std::getenv +#include +#include + +#include "gpufl/gpufl.hpp" + +static bool CheckCuda(const cudaError_t err, const char* call, const char* file, + const int line) { + if (err == cudaSuccess) return true; + std::cerr << "[CUDA ERROR] " << file << ":" << line << " " << call + << " failed: " << cudaGetErrorString(err) << " (" + << err << ")" << std::endl; + return false; +} + +#define CHECK_CUDA(call) \ + do { \ + if (!CheckCuda((call), #call, __FILE__, __LINE__)) return 2; \ + } while (0) + +// Long FMA chain - keeps the SMs busy long enough for PC sampling to land +// stall samples and for PM to read steady-state counters. +__global__ void computeHeavy(float* out, const float* in, int n, int iters) { + const int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx >= n) return; + float val = in[idx]; + for (int i = 0; i < iters; ++i) { + val = val * 1.0009f + 0.0001f; + val = fmaf(val, 0.9991f, 0.0002f); + } + out[idx] = val; +} + +namespace { + +gpufl::ProfilingEngine EngineFromEnv() { + const char* v = std::getenv("GPUFL_PROFILING_ENGINE"); + // gpufl::init() parses this env var itself; we mirror it only so the + // banner below reports the right thing. + const std::string name = v ? v : "PmSampling"; + if (name == "PcSampling") return gpufl::ProfilingEngine::PcSampling; + if (name == "SassMetrics") return gpufl::ProfilingEngine::SassMetrics; + if (name == "RangeProfilerKernelReplay") + return gpufl::ProfilingEngine::RangeProfilerKernelReplay; + if (name == "Trace") return gpufl::ProfilingEngine::Trace; + return gpufl::ProfilingEngine::PmSampling; +} + +// The loop has to outlast the window by a wide margin or the run ends first +// and the window closes at session stop, proving nothing. How many +// iterations that takes depends on the GPU and on the engine (kernel replay +// costs ~40x per iteration), so both are tunable for a faster card. +int EnvIntOr(const char* name, const int fallback) { + const char* v = std::getenv(name); + if (!v || v[0] == '\0') return fallback; + const int n = std::atoi(v); + return n > 0 ? n : fallback; +} + +} // namespace + +int main() { + gpufl::InitOptions opts; + opts.app_name = "deep_window_demo"; + opts.log_path = "deep_window"; + opts.enable_debug_output = true; + opts.profiling_engine = EngineFromEnv(); + // The point of the demo: engines stay idle until a window opens. + opts.deep_window_only = true; + + if (!gpufl::init(opts)) { + std::cerr << "Failed to initialize gpufl" << std::endl; + return 1; + } + + std::cout << "=== Bounded deep window demo ===\n" + << "engine : " + << gpufl::ProfilingEngineWireName(opts.profiling_engine) << "\n" + << "deep_arm_mode : window-only\n" + << std::endl; + + constexpr int n = 1 << 20; + constexpr size_t bytes = n * sizeof(float); + + float* d_in = nullptr; + float* d_out = nullptr; + CHECK_CUDA(cudaMalloc(&d_in, bytes)); + CHECK_CUDA(cudaMalloc(&d_out, bytes)); + CHECK_CUDA(cudaMemset(d_in, 0, bytes)); + + const int threads = 256; + const int blocks = (n + threads - 1) / threads; + + // The trigger fires as a tight burst at one iteration, not spread over + // many. Spreading it across iterations would (correctly) open a fresh + // window every time the previous one expired, and how many you get + // would just track how slow the engine is - under kernel replay an + // iteration costs ~40x what it does under PM sampling. A burst inside + // a single iteration is over in microseconds, so every call after the + // first is guaranteed to land while the window is open, which is the + // invariant worth testing: a re-fire is ignored, never an extension. + const int kIterations = EnvIntOr("GPUFL_DEMO_ITERATIONS", 3000); + const int kDropsAt = kIterations / 6; + const int kTriggerBurst = 50; + const int kWindowMs = EnvIntOr("GPUFL_DEMO_WINDOW_MS", 1000); + + int trigger_calls = 0; + int windows_opened = 0; + int iterations_with_window_open = 0; + int iterations_after_close = 0; + bool was_open = false; + bool fired = false; + + for (int iter = 0; iter < kIterations; ++iter) { + const double tokens_per_sec = (iter < kDropsAt) ? 1800.0 : 850.0; + + if (tokens_per_sec < 1000.0 && !fired) { + fired = true; + for (int i = 0; i < kTriggerBurst; ++i) { + gpufl::deepWindow(kWindowMs); + ++trigger_calls; + } + } + + const bool window_open = gpufl::deepWindowActive(); + if (window_open && !was_open) { + ++windows_opened; + std::cout << "[iter " << iter << "] window opened (" << kWindowMs + << "ms)" << std::endl; + } + if (!window_open && was_open) { + std::cout << "[iter " << iter + << "] window closed on its own; the loop keeps running" + << std::endl; + } + was_open = window_open; + + if (window_open) ++iterations_with_window_open; + else if (windows_opened > 0) ++iterations_after_close; + + computeHeavy<<>>(d_out, d_in, n, 4000); + CHECK_CUDA(cudaGetLastError()); + CHECK_CUDA(cudaDeviceSynchronize()); + } + + const bool open_at_end = gpufl::deepWindowActive(); + std::cout << "\n--- summary ---\n" + << "trigger calls : " << trigger_calls + << " (a burst inside one iteration)\n" + << "windows opened : " << windows_opened + << " (expected 1)\n" + << "iterations with window open: " + << iterations_with_window_open << " of " << kIterations << "\n" + << "iterations after close : " << iterations_after_close + << "\n" + << "window still open at end : " << (open_at_end ? "yes" : "no") + << "\n\nCheck the session's deep_window event for the launches " + "it covered and the bound that closed it." + << std::endl; + + int failures = 0; + if (windows_opened != 1) { + std::cerr << "[FAIL] expected exactly 1 window, saw " << windows_opened + << " - a re-fired trigger opened extra windows." << std::endl; + ++failures; + } + if (open_at_end) { + std::cerr << "[FAIL] the window never closed - it should have hit its " + << kWindowMs << "ms deadline." << std::endl; + ++failures; + } + if (iterations_after_close == 0) { + std::cerr << "[FAIL] no iterations ran after the window closed." + << std::endl; + ++failures; + } + if (failures == 0) { + std::cout << "\n[OK] window opened once, closed on its own bound, and " + "the workload outlived it." + << std::endl; + } + + CHECK_CUDA(cudaFree(d_in)); + CHECK_CUDA(cudaFree(d_out)); + + gpufl::shutdown(); + return failures == 0 ? 0 : 1; +} diff --git a/include/gpufl/backends/nvidia/cupti_backend.cpp b/include/gpufl/backends/nvidia/cupti_backend.cpp index 3a85d2e..7975526 100644 --- a/include/gpufl/backends/nvidia/cupti_backend.cpp +++ b/include/gpufl/backends/nvidia/cupti_backend.cpp @@ -22,6 +22,7 @@ #include "gpufl/backends/nvidia/synchronization_handler.hpp" #include "gpufl/core/common.hpp" #include "gpufl/core/debug_logger.hpp" +#include "gpufl/core/deep_window.hpp" #include "gpufl/core/logger/logger.hpp" #include "gpufl/core/monitor.hpp" // Monitor::RequestSyntheticDrainAndWait #include "gpufl/core/model/perf_metric_model.hpp" @@ -670,8 +671,14 @@ void CuptiBackend::FinishDeferredEngineStart_() { profiling_request_, device_facts_, EnvOverrides::FromProcess()); ApplyComboPlanOverrides(resolved_plan_, combo_); - EngineContext ectx{ctx_, device_id_, chip_name_, &cubin_mu_, - &cubin_by_crc_}; + // Carry the CUPTI->wall clock anchor, same as the non-deferred start. + // Omitting it left base_cpu_ns/base_cupti_ts at 0, so engines that stamp + // their own samples (PM, PC) emitted raw CUPTI timestamps - a different + // clock domain from the kernel timeline, which put every sample days + // away from the session. start() captures the anchor before flagging the + // start pending, so it is already valid here. + EngineContext ectx{ctx_, device_id_, chip_name_, &cubin_mu_, + &cubin_by_crc_, base_cpu_ns_, base_cupti_ts_}; engine_->initialize(opts_, ectx); engine_->start(); ReenableActivityAfterEngineStart_(); @@ -734,11 +741,21 @@ void CuptiBackend::FlushProfilingDataBeforeCudaTeardown(const char* reason) { void CuptiBackend::EngineLaunchTick() { if (!initialized_ || !active_.load(std::memory_order_relaxed)) return; + // Bound check first: this is the app thread at launch ENTER, the only + // reliably scheduled, context-current place to run the window's + // stop/collect. Closing before the engine tick also spares the engine a + // beat it would only spend on a window that is already over. + DeepWindow::OnLaunch(); if (engine_) engine_->onLaunchTick(); } void CuptiBackend::DrainProfilingData() { if (!initialized_ || !active_.load(std::memory_order_relaxed)) return; + // Fallback for a window whose workload stopped launching before the + // deadline. On a Windows-injected target the CUPTI teardown must not run + // from this collector thread, so there it only flags the close and the + // next launch performs it. + DeepWindow::OnPeriodicTick(/*may_close_here=*/!WindowsInjectedProcess()); if (engine_) { engine_->drainData(); } diff --git a/include/gpufl/backends/nvidia/cupti_backend.hpp b/include/gpufl/backends/nvidia/cupti_backend.hpp index df07330..9a15483 100644 --- a/include/gpufl/backends/nvidia/cupti_backend.hpp +++ b/include/gpufl/backends/nvidia/cupti_backend.hpp @@ -181,13 +181,30 @@ class CuptiBackend : public IMonitorBackend { void EmitCaptureCapabilities_() const; + // In WindowOnly mode an ordinary user scope must not arm anything - + // otherwise a per-step GFL_SCOPE in a training loop keeps the engines + // armed for the whole run and the window means nothing. Deep windows + // come in through OnDeepWindowStart/Stop, which never gate. + bool ScopeArmsEngines_() const { + return opts_.deep_arm_mode != DeepArmMode::WindowOnly; + } + void OnScopeStart(const char* name) override { GFL_LOG_DEBUG("OnScopeStart"); - if (engine_) engine_->onScopeStart(name); + if (!ScopeArmsEngines_()) return; + OnDeepWindowStart(name); } void DrainProfilingData() override; void OnScopeStop(const char* name) override { GFL_LOG_DEBUG("OnScopeStop"); + if (!ScopeArmsEngines_()) return; + OnDeepWindowStop(name); + } + + void OnDeepWindowStart(const char* name) override { + if (engine_) engine_->onScopeStart(name); + } + void OnDeepWindowStop(const char* name) override { if (engine_) engine_->onScopeStop(name); // cuptiActivityFlushAll(1) permanently kills the CUPTI subscriber // callback when the SamplingAPI is armed (enableStartStopControl=0, @@ -200,9 +217,17 @@ class CuptiBackend : public IMonitorBackend { } } void OnPerfScopeStart(const char* name) override { - if (engine_) engine_->onPerfScopeStart(name); + if (!ScopeArmsEngines_()) return; + OnDeepWindowPerfStart(name); } void OnPerfScopeStop(const char* name) override { + if (!ScopeArmsEngines_()) return; + OnDeepWindowPerfStop(name); + } + void OnDeepWindowPerfStart(const char* name) override { + if (engine_) engine_->onPerfScopeStart(name); + } + void OnDeepWindowPerfStop(const char* name) override { if (engine_) engine_->onPerfScopeStop(name); } std::optional TakeLastPerfEvent() override { diff --git a/include/gpufl/backends/nvidia/engine/pc_sampling_engine.cpp b/include/gpufl/backends/nvidia/engine/pc_sampling_engine.cpp index 5aae753..2f51e74 100644 --- a/include/gpufl/backends/nvidia/engine/pc_sampling_engine.cpp +++ b/include/gpufl/backends/nvidia/engine/pc_sampling_engine.cpp @@ -200,14 +200,24 @@ void PcSamplingEngine::start() { // still degrades to a kernel trace. Synthetic-kernel fallback stays // suppressed (cupti_backend.cpp start()), so only REAL records show. cuptiActivityEnable(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL); - // Arm now — start() runs in the CONTEXT_CREATED callback, before the - // app's first kernel, while the GPU is quiet. Enable/config/Start - // return INVALID_OPERATION when kernels run concurrently (verified - // live), so pre-first-kernel is the only reliable window. profiler- - // init already ran pre-context, so stall enumeration succeeds here. + // Enable/config/Start return INVALID_OPERATION when kernels run + // concurrently (verified live), and start() runs in the + // CONTEXT_CREATED callback before the app's first kernel, while the + // GPU is quiet - the only reliable moment for them. profiler-init + // already ran pre-context, so stall enumeration succeeds here. + // + // WindowOnly splits that: enable + configure now (they must happen + // while quiet), but leave the sampler unarmed until a deep window + // opens. Only cuptiPCSamplingStart is deferred, and that one does + // succeed with kernels running - the mid-run drain-restart below + // has always relied on exactly that. { std::lock_guard lk(sampling_lifecycle_mu_); - StartPcSampling_(); + if (opts_.deep_arm_mode == DeepArmMode::WindowOnly) { + EnableSamplingFeatures_(); + } else { + StartPcSampling_(); + } } // Enable can internally disable kernel activity — re-assert it. cuptiActivityEnable(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL); @@ -391,13 +401,23 @@ void PcSamplingEngine::shutdown() { } void PcSamplingEngine::onScopeStart(const char* /*name*/) { - // Idempotent re-arm - start() already arms the whole session; this - // only matters if a prior arm attempt failed (e.g. context raced). + // Under WindowOnly this IS the arm: start() only enabled and configured + // the sampler. Otherwise it's an idempotent re-arm that matters only if + // a prior attempt failed (e.g. context raced). std::lock_guard lk(sampling_lifecycle_mu_); StartPcSampling_(); } void PcSamplingEngine::onScopeStop(const char* /*name*/) { + if (opts_.deep_arm_mode == DeepArmMode::WindowOnly) { + // Disarm for real. A forced collect would leave the sampler running + // past the window, which is the whole thing WindowOnly exists to + // avoid. The ref count keeps a user scope nested inside the window + // from disarming it early. + std::lock_guard lk(sampling_lifecycle_mu_); + StopAndCollectPcSampling_(); + return; + } // Forced collect at scope end. For the process-wide scope this is the // last healthy moment before Windows process-exit teardown breaks // cuptiPCSamplingStop with CUPTI_ERROR_UNKNOWN. Re-arms afterwards, so diff --git a/include/gpufl/backends/nvidia/engine/pm_sampling_engine.cpp b/include/gpufl/backends/nvidia/engine/pm_sampling_engine.cpp index de48197..90cc9bf 100644 --- a/include/gpufl/backends/nvidia/engine/pm_sampling_engine.cpp +++ b/include/gpufl/backends/nvidia/engine/pm_sampling_engine.cpp @@ -39,13 +39,13 @@ bool PmSamplingEngine::initialize(const MonitorOptions& opts, " metrics=", metrics_.size(), " interval_us=", opts_.pm_sampling_interval_us, " max_samples=", opts_.pm_sampling_max_samples, - " scope_only=", opts_.pm_sampling_scope_only ? "true" : "false"); + " scope_gated=", ScopeGated_() ? "true" : "false"); return true; } void PmSamplingEngine::start() { #if GPUFL_HAS_PERFWORKS - if (!opts_.pm_sampling_scope_only) { + if (!ScopeGated_()) { StartPmSampling_(); } else { attempted_.store(true, std::memory_order_relaxed); @@ -84,13 +84,13 @@ void PmSamplingEngine::shutdown() { void PmSamplingEngine::onScopeStart(const char*) { #if GPUFL_HAS_PERFWORKS - if (opts_.pm_sampling_scope_only) StartPmSampling_(); + if (ScopeGated_()) StartPmSampling_(); #endif } void PmSamplingEngine::onScopeStop(const char*) { #if GPUFL_HAS_PERFWORKS - if (opts_.pm_sampling_scope_only) StopPmSampling_(); + if (ScopeGated_()) StopPmSampling_(); #endif } diff --git a/include/gpufl/backends/nvidia/engine/pm_sampling_engine.hpp b/include/gpufl/backends/nvidia/engine/pm_sampling_engine.hpp index 79ea45f..1b66941 100644 --- a/include/gpufl/backends/nvidia/engine/pm_sampling_engine.hpp +++ b/include/gpufl/backends/nvidia/engine/pm_sampling_engine.hpp @@ -42,6 +42,13 @@ class PmSamplingEngine final : public IProfilingEngine { } private: + // True when sampling arms on scope start rather than at session start. + // A deep window IS a scope, so WindowOnly reduces to the same gate the + // PM engine has always had. + bool ScopeGated_() const { + return opts_.pm_sampling_scope_only || + opts_.deep_arm_mode == DeepArmMode::WindowOnly; + } std::vector ResolveMetrics_() const; void EmitConfig_() const; diff --git a/include/gpufl/backends/nvidia/engine/range_profiler_engine.cpp b/include/gpufl/backends/nvidia/engine/range_profiler_engine.cpp index 3a8e066..16eab19 100644 --- a/include/gpufl/backends/nvidia/engine/range_profiler_engine.cpp +++ b/include/gpufl/backends/nvidia/engine/range_profiler_engine.cpp @@ -113,27 +113,38 @@ void RangeProfilerEngine::start() { if (!perf_session_active_) { InitPerfworksSession_(mode_ == Mode::Scope); } - if (mode_ == Mode::KernelReplay && perf_session_active_) { - if (kernel_replay_running_) { - GFL_LOG_DEBUG("[RangeProfilerKernelReplay] start skipped: already running"); - return; - } - CUpti_RangeProfiler_Start_Params p = { - CUpti_RangeProfiler_Start_Params_STRUCT_SIZE}; - p.pRangeProfilerObject = range_profiler_object_; - if (LogCuptiErrorIfFailed(this->name(), "cuptiRangeProfilerStart", - cuptiRangeProfilerStart(&p))) { - return; - } - kernel_replay_running_ = true; - kernel_replay_decoded_ = false; - GFL_LOG_DEBUG("[RangeProfilerKernelReplay] started"); + // WindowOnly keeps the Perfworks session built but the replay stopped: + // kernel replay re-runs every kernel, so leaving it running for the + // whole session is exactly the cost a window exists to avoid. It arms + // in onPerfScopeStart instead. + if (mode_ == Mode::KernelReplay && perf_session_active_ && + opts_.deep_arm_mode != DeepArmMode::WindowOnly) { + StartKernelReplay_(); } #else GFL_LOG_ERROR("[RangeProfilerEngine] Not built with GPUFL_HAS_PERFWORKS"); #endif } +#if GPUFL_HAS_PERFWORKS +void RangeProfilerEngine::StartKernelReplay_() { + if (kernel_replay_running_) { + GFL_LOG_DEBUG("[RangeProfilerKernelReplay] start skipped: already running"); + return; + } + CUpti_RangeProfiler_Start_Params p = { + CUpti_RangeProfiler_Start_Params_STRUCT_SIZE}; + p.pRangeProfilerObject = range_profiler_object_; + if (LogCuptiErrorIfFailed(this->name(), "cuptiRangeProfilerStart", + cuptiRangeProfilerStart(&p))) { + return; + } + kernel_replay_running_ = true; + kernel_replay_decoded_ = false; + GFL_LOG_DEBUG("[RangeProfilerKernelReplay] started"); +} +#endif + void RangeProfilerEngine::stop() { #if GPUFL_HAS_PERFWORKS if (mode_ != Mode::KernelReplay || !perf_session_active_) return; @@ -195,6 +206,14 @@ void RangeProfilerEngine::shutdown() { void RangeProfilerEngine::onPerfScopeStart(const char* name) { #if GPUFL_HAS_PERFWORKS + if (mode_ == Mode::KernelReplay) { + // WindowOnly deferred the replay arm out of start() to here. + if (opts_.deep_arm_mode == DeepArmMode::WindowOnly && + perf_session_active_) { + StartKernelReplay_(); + } + return; + } if (mode_ != Mode::Scope) return; attempted_.store(true, std::memory_order_relaxed); GFL_LOG_DEBUG("[RangeProfilerEngine] onPerfScopeStart name=", @@ -237,6 +256,12 @@ void RangeProfilerEngine::onPerfScopeStart(const char* name) { void RangeProfilerEngine::onPerfScopeStop(const char* name) { #if GPUFL_HAS_PERFWORKS + if (mode_ == Mode::KernelReplay) { + // Stops the replay and decodes what it collected; the Perfworks + // session stays built so the next window can arm again. + if (opts_.deep_arm_mode == DeepArmMode::WindowOnly) stop(); + return; + } if (mode_ != Mode::Scope) return; GFL_LOG_DEBUG("[RangeProfilerEngine] onPerfScopeStop name=", (name ? name : "(null)"), " active=", perf_session_active_); diff --git a/include/gpufl/backends/nvidia/engine/range_profiler_engine.hpp b/include/gpufl/backends/nvidia/engine/range_profiler_engine.hpp index 26d9073..03a0766 100644 --- a/include/gpufl/backends/nvidia/engine/range_profiler_engine.hpp +++ b/include/gpufl/backends/nvidia/engine/range_profiler_engine.hpp @@ -55,6 +55,9 @@ class RangeProfilerEngine final : public IProfilingEngine { bool InitPerfworksSession_(bool require_single_pass); void EndPerfPassAndDecode_(); void DecodeKernelReplayEvents_(); + // Arms kernel replay on an already-built Perfworks session. Called from + // start(), or from onPerfScopeStart when a deep window defers the arm. + void StartKernelReplay_(); bool perf_session_active_ = false; mutable std::mutex perf_mu_; diff --git a/include/gpufl/core/deep_window.cpp b/include/gpufl/core/deep_window.cpp new file mode 100644 index 0000000..7eabe37 --- /dev/null +++ b/include/gpufl/core/deep_window.cpp @@ -0,0 +1,404 @@ +#include "gpufl/core/deep_window.hpp" + +#include +#include +#include + +#include "gpufl.hpp" +#include "gpufl/core/common.hpp" +#include "gpufl/core/debug_logger.hpp" +#include "gpufl/core/env_vars.hpp" +#include "gpufl/core/events.hpp" +#include "gpufl/core/logger/logger.hpp" +#include "gpufl/core/model/deep_window_model.hpp" +#include "gpufl/core/model/perf_metric_model.hpp" +#include "gpufl/core/monitor.hpp" +#include "gpufl/core/monitor_backend.hpp" +#include "gpufl/core/runtime.hpp" +#include "gpufl/core/teardown_flag.hpp" // detail::isProcessExitTeardown + +namespace gpufl { +namespace { + +// Serializes open/close transitions. Held across the engine arm/disarm so a +// window can't be closed between "state says open" and "engines are armed". +std::mutex g_mu; + +// Read lock-free from the launch callback on every launch while a window is +// open, so the hot path never touches g_mu. Also the re-entrancy guard: +// Close() clears it before the engine teardown, and that teardown can +// synchronize the device and re-enter the launch callback. +std::atomic g_active{false}; + +std::atomic g_deadline_ns{0}; // 0 = no time bound +std::atomic g_launches_remaining{0}; // 0 = no launch bound +std::atomic g_launches_covered{0}; +std::atomic g_close_requested{false}; + +// An open asked for by a thread that can't arm one itself. Checked lock-free +// on the launch beat; g_pending_spec is only read once this is set, so the +// hot path pays for it exactly when a trigger is waiting. +std::atomic g_open_requested{false}; +std::atomic g_pending_open_at_ns{0}; // 0 = at the next launch +DeepWindowSpec g_pending_spec; // guarded by g_mu + +// When the last window closed, so a cooldown can be enforced. 0 = never. +std::atomic g_last_close_ns{0}; + +int64_t g_opened_ns = 0; +int64_t g_requested_duration_ms = 0; +uint64_t g_requested_max_launches = 0; +std::string g_name; + +bool ComboActive() { + const char* combo = std::getenv(env::kEngineCombo); + return combo && combo[0] != '\0'; +} + +// Non-negative integer from env, or `fallback` when unset or malformed. +int64_t EnvUnsignedOr(const char* name, const int64_t fallback) { + const char* v = std::getenv(name); + if (!v || v[0] == '\0') return fallback; + char* end = nullptr; + const long long n = std::strtoll(v, &end, 10); + if (end == v || *end != '\0' || n < 0) { + GFL_LOG_ERROR(name, "='", v, + "' is not a non-negative integer. Ignoring."); + return fallback; + } + return n; +} + +// Fills bounds left at 0 from the environment, so an operator can size a +// window the application code already asks for - and so the injected path, +// which can't set a DeepWindowSpec, can set one at all. +void ApplyEnvDefaults(DeepWindowSpec& spec) { + if (spec.max_duration_ms == 0) { + spec.max_duration_ms = EnvUnsignedOr(env::kDeepWindowMs, 0); + } + if (spec.max_launches == 0) { + spec.max_launches = + static_cast(EnvUnsignedOr(env::kDeepWindowMaxLaunches, 0)); + } + if (spec.cooldown_ms == 0) { + spec.cooldown_ms = EnvUnsignedOr(env::kDeepWindowCooldownMs, 0); + } +} + +// True while a just-closed window's cooldown is still running. +bool InCooldown(const DeepWindowSpec& spec) { + if (spec.cooldown_ms <= 0) return false; + const int64_t last = g_last_close_ns.load(std::memory_order_relaxed); + if (last == 0) return false; + return detail::GetTimestampNs() - last < spec.cooldown_ms * 1000000; +} + +} // namespace + +const char* DeepWindowCloseName(const DeepWindowClose reason) { + switch (reason) { + case DeepWindowClose::Deadline: return "deadline"; + case DeepWindowClose::LaunchBudget: return "launch_budget"; + case DeepWindowClose::Manual: return "manual"; + case DeepWindowClose::SessionStop: return "session_stop"; + } + return "unknown"; +} + +namespace detail { + +bool PerfScopeEnabled() { + // Also fire for an engine combo with a Trace base - otherwise a + // Trace+RangeProfiler combo would never trigger Range's perf scope. + return g_opts.profiling_engine != ProfilingEngine::Monitor && + (g_opts.profiling_engine != ProfilingEngine::Trace || ComboActive()); +} + +void BeginPerfScopeIfEnabled(const char* name, const bool is_deep_window) { + if (!PerfScopeEnabled()) return; + if (is_deep_window) { + Monitor::BeginDeepWindowPerfScope(name); + } else { + Monitor::BeginPerfScope(name); + } +} + +void EndPerfScopeIfEnabled(const char* name, const int pid, + const int64_t start_ns, const int64_t end_ns, + const bool is_deep_window) { + if (!PerfScopeEnabled()) return; + // Triggers EndPerfPassAndDecode first. + if (is_deep_window) { + Monitor::EndDeepWindowPerfScope(name); + } else { + Monitor::EndPerfScope(name); + } + + const Runtime* rt = runtime(); + if (!rt || !rt->logger) return; + IMonitorBackend* backend = Monitor::GetBackend(); + if (!backend) return; + auto event_opt = backend->TakeLastPerfEvent(); + if (!event_opt) return; + + PerfMetricEvent& pe = *event_opt; + pe.pid = pid; + pe.app = rt->app_name; + pe.session_id = rt->session_id; + pe.name = name ? name : ""; + pe.start_ns = start_ns; + pe.end_ns = end_ns; + rt->logger->write(model::PerfMetricModel(pe)); +} + +} // namespace detail + +bool DeepWindow::Active() { + return g_active.load(std::memory_order_acquire); +} + +bool DeepWindow::Open(const DeepWindowSpec& spec) { + if (const Runtime* rt = runtime(); !rt || !rt->logger) return false; + + std::string name; + { + std::lock_guard lk(g_mu); + if (g_active.load(std::memory_order_relaxed)) { + // Not an extension. A trigger that fires every step would + // otherwise hold the window open for the rest of the run. + return false; + } + if (InCooldown(spec)) { + // A condition that stays true would otherwise reopen a window the + // moment the last one expired, and the run never stops paying. + GFL_LOG_DEBUG("[DeepWindow] open suppressed: cooldown ", + spec.cooldown_ms, "ms not elapsed"); + return false; + } + + g_opened_ns = detail::GetTimestampNs(); + g_name = spec.name.empty() ? "deep_window" : spec.name; + g_requested_duration_ms = spec.max_duration_ms; + g_requested_max_launches = spec.max_launches; + g_deadline_ns.store(spec.max_duration_ms > 0 + ? g_opened_ns + spec.max_duration_ms * 1000000 + : 0, + std::memory_order_relaxed); + g_launches_remaining.store(spec.max_launches, std::memory_order_relaxed); + g_launches_covered.store(0, std::memory_order_relaxed); + g_close_requested.store(false, std::memory_order_relaxed); + name = g_name; + + // Publish last: once this is true the launch callback starts + // consuming budget, and everything it reads is already set. + g_active.store(true, std::memory_order_release); + + // Arms the deep engines. Runs under the lock so a concurrent close + // can't disarm engines this call hasn't armed yet; safe because the + // arm path doesn't re-enter DeepWindow. + Monitor::BeginDeepWindowScope(name.c_str()); + detail::BeginPerfScopeIfEnabled(name.c_str(), /*is_deep_window=*/true); + } + + GFL_LOG_DEBUG("[DeepWindow] opened name=", name, + " duration_ms=", spec.max_duration_ms, + " max_launches=", spec.max_launches); + return true; +} + +void DeepWindow::Close(const DeepWindowClose reason) { + if (!g_active.load(std::memory_order_acquire)) return; + + DeepWindowEvent ev; + std::string name; + { + int64_t start_ns = 0; + std::lock_guard lk(g_mu); + if (!g_active.load(std::memory_order_relaxed)) return; + // Clear before the engine teardown below. That teardown can + // synchronize the device and re-enter the launch callback, and + // OnLaunch's lock-free check turns the re-entry into a no-op + // instead of a deadlock on g_mu. + g_active.store(false, std::memory_order_release); + g_close_requested.store(false, std::memory_order_relaxed); + + name = g_name; + start_ns = g_opened_ns; + const int64_t end_ns = detail::GetTimestampNs(); + // Starts the cooldown clock. Set from the decision point, not after + // the engine teardown, so a slow disarm doesn't shorten the quiet time. + g_last_close_ns.store(end_ns, std::memory_order_relaxed); + + ev.pid = detail::GetPid(); + ev.name = g_name; + ev.close_reason = DeepWindowCloseName(reason); + ev.engine = ProfilingEngineWireName(g_opts.profiling_engine); + ev.start_ns = start_ns; + ev.end_ns = end_ns; + ev.duration_ns = end_ns - start_ns; + ev.launches_covered = g_launches_covered.load(std::memory_order_relaxed); + ev.requested_duration_ms = g_requested_duration_ms; + ev.requested_max_launches = g_requested_max_launches; + + // Disarms the deep engines and drains whatever they collected. + // Skipped on process-exit teardown, where cudart has already + // destroyed the context and the scope-stop path would fault against + // it - the engines' own exit handling flushes there instead. The + // event below is still written so the window is on the record. + if (!detail::isProcessExitTeardown()) { + Monitor::EndDeepWindowScope(name.c_str()); + detail::EndPerfScopeIfEnabled(name.c_str(), ev.pid, start_ns, end_ns, + /*is_deep_window=*/true); + } + } + + if (const Runtime* rt = runtime(); rt && rt->logger) { + ev.app = rt->app_name; + ev.session_id = rt->session_id; + rt->logger->write(model::DeepWindowModel(ev)); + } + + GFL_LOG_DEBUG("[DeepWindow] closed name=", name, + " reason=", ev.close_reason, + " duration_ns=", ev.duration_ns, + " launches_covered=", ev.launches_covered); +} + +void DeepWindow::RequestOpen(const DeepWindowSpec& spec) { + ScheduleOpenAfter(0, spec); +} + +void DeepWindow::ScheduleOpenAfter(const int64_t delay_ms, + const DeepWindowSpec& spec) { + { + std::lock_guard lk(g_mu); + g_pending_spec = spec; + g_pending_open_at_ns.store( + delay_ms > 0 ? detail::GetTimestampNs() + delay_ms * 1000000 : 0, + std::memory_order_relaxed); + } + // Published last: the launch beat reads the spec only once this is set. + g_open_requested.store(true, std::memory_order_release); + GFL_LOG_DEBUG("[DeepWindow] open requested delay_ms=", delay_ms, + " duration_ms=", spec.max_duration_ms, + " max_launches=", spec.max_launches); +} + +// Runs on the app thread at launch ENTER - the only place the arm's CUPTI +// calls are safe. Claims the request before opening so two launch threads +// can't both act on it. +void DeepWindow::TakePendingOpen_() { + if (!g_open_requested.load(std::memory_order_acquire)) return; + const int64_t due = g_pending_open_at_ns.load(std::memory_order_relaxed); + if (due > 0 && detail::GetTimestampNs() < due) return; + if (!g_open_requested.exchange(false, std::memory_order_acq_rel)) return; + + DeepWindowSpec spec; + { + std::lock_guard lk(g_mu); + spec = g_pending_spec; + } + // Outside the lock: Open takes it too. + Open(spec); +} + +void DeepWindow::OnLaunch() { + if (!g_active.load(std::memory_order_acquire)) { + TakePendingOpen_(); + return; + } + + g_launches_covered.fetch_add(1, std::memory_order_relaxed); + + if (g_launches_remaining.load(std::memory_order_relaxed) > 0) { + // fetch_sub returns the PREVIOUS value, so 1 means this launch + // consumed the last of the budget. + if (g_launches_remaining.fetch_sub(1, std::memory_order_relaxed) <= 1) { + Close(DeepWindowClose::LaunchBudget); + return; + } + } + + // A tick on a thread that couldn't run the teardown left this set. + if (g_close_requested.load(std::memory_order_acquire)) { + Close(DeepWindowClose::Deadline); + return; + } + + const int64_t deadline = g_deadline_ns.load(std::memory_order_relaxed); + if (deadline > 0 && detail::GetTimestampNs() >= deadline) { + Close(DeepWindowClose::Deadline); + } +} + +void DeepWindow::OnPeriodicTick(const bool may_close_here) { + if (!g_active.load(std::memory_order_acquire)) return; + + const int64_t deadline = g_deadline_ns.load(std::memory_order_relaxed); + if (deadline <= 0 || detail::GetTimestampNs() < deadline) return; + + if (may_close_here) { + Close(DeepWindowClose::Deadline); + return; + } + // Hand the close to the next launch: the CUPTI stop has to run on the + // application thread that owns the context. + g_close_requested.store(true, std::memory_order_release); +} + +void DeepWindow::ResetForTesting() { + std::lock_guard lk(g_mu); + g_active.store(false, std::memory_order_release); + g_deadline_ns.store(0, std::memory_order_relaxed); + g_launches_remaining.store(0, std::memory_order_relaxed); + g_launches_covered.store(0, std::memory_order_relaxed); + g_close_requested.store(false, std::memory_order_relaxed); + g_open_requested.store(false, std::memory_order_relaxed); + g_pending_open_at_ns.store(0, std::memory_order_relaxed); + g_pending_spec = DeepWindowSpec{}; + g_last_close_ns.store(0, std::memory_order_relaxed); + g_opened_ns = 0; + g_requested_duration_ms = 0; + g_requested_max_launches = 0; + g_name.clear(); +} + +// ---- Public API ---- + +void deepWindow(const DeepWindowSpec& spec) { + DeepWindowSpec resolved = spec; + ApplyEnvDefaults(resolved); + DeepWindow::Open(resolved); +} + +void deepWindow(const int64_t max_duration_ms, const uint64_t max_launches) { + DeepWindowSpec spec; + spec.max_duration_ms = max_duration_ms; + spec.max_launches = max_launches; + deepWindow(spec); +} + +// Reads the launcher's time-based trigger, the only one available when the +// target's source can't be edited. Called once from init(); a no-op unless +// GPUFL_DEEP_AFTER_MS is set. +void scheduleEnvDeepWindow() { + const int64_t after_ms = EnvUnsignedOr(env::kDeepAfterMs, -1); + if (after_ms < 0) return; + + DeepWindowSpec spec; + ApplyEnvDefaults(spec); + if (spec.max_duration_ms == 0 && spec.max_launches == 0) { + GFL_LOG_ERROR( + env::kDeepAfterMs, + " is set but no window bound is - the window would never close. " + "Set ", env::kDeepWindowMs, " or ", env::kDeepWindowMaxLaunches, "."); + return; + } + DeepWindow::ScheduleOpenAfter(after_ms, spec); +} + +void deepWindowClose() { DeepWindow::Close(DeepWindowClose::Manual); } + +bool deepWindowActive() { return DeepWindow::Active(); } + +} // namespace gpufl diff --git a/include/gpufl/core/deep_window.hpp b/include/gpufl/core/deep_window.hpp new file mode 100644 index 0000000..bb64e1c --- /dev/null +++ b/include/gpufl/core/deep_window.hpp @@ -0,0 +1,140 @@ +#pragma once + +#include +#include + +namespace gpufl { + +/** + * @brief Why a deep window closed. + * + * Reported on every deep_window event. A window that exhausted its launch + * budget covered exactly what the caller asked for; one that hit the + * deadline may have covered far less, because kernel replay stretches wall + * time without advancing the application. Without this field a short + * window looks like a bug rather than the bound doing its job. + */ +enum class DeepWindowClose { Deadline, LaunchBudget, Manual, SessionStop }; + +/** Wire name for a close reason (goes into the deep_window event). */ +const char* DeepWindowCloseName(DeepWindowClose reason); + +/** + * @brief Bounds for one deep window. + * + * Both bounds are optional and combine with OR - whichever is reached + * first closes the window. Wall time alone is a poor bound for the replay + * engines (SASS / Range), where a three-second window can cover a handful + * of launches; a launch budget says what the caller actually meant. + */ +struct DeepWindowSpec { + int64_t max_duration_ms = 0; // 0 = no time bound + uint64_t max_launches = 0; // 0 = no launch bound + // Minimum quiet time after a window closes before another may open. + // 0 = none. Without it a condition that stays true reopens a window the + // instant the last one expired, and the run pays deep cost forever. Only + // the library knows when the last window closed, so this bound belongs + // here rather than in the caller's trigger. + int64_t cooldown_ms = 0; + std::string name = "deep_window"; +}; + +/** + * @brief The profiler scope that deep engines arm on, opened by a trigger + * and closed by a bound instead of by a destructor. + * + * The deep engines are already scope-gated (PM sampling arms in + * onScopeStart, SASS arms on scope start and flushes on scope stop, the + * Range profiler's scope mode initializes Perfworks lazily), so this adds + * no new arming mechanism. What it adds is a scope that nobody has to + * close by hand. + * + * Process-wide singleton: a window is a property of the session, not of a + * thread, and the CUPTI calls behind it are context-bound. + */ +class DeepWindow { + public: + /** + * @brief Open a window if none is active. + * + * A second call while a window is open is IGNORED, not an extension, + * so a trigger that fires every training step cannot hold the window + * open indefinitely. Returns true only if this call opened it. + * + * Arms the engines on the calling thread. Call it from the application + * thread that runs the workload - that thread is context-current. + */ + static bool Open(const DeepWindowSpec& spec); + + /** + * @brief Ask for a window from a thread that must not arm one itself. + * + * Arming runs CUPTI calls that are only safe on the application thread + * at a launch boundary, so a trigger living anywhere else - the sampler + * evaluating a rule, a listener woken by an external signal, a timer - + * records the request here and the next launch performs the open. The + * mirror of how a deadline reached off the app thread defers its close. + * + * A pending request is replaced, not queued: the newest spec wins. + */ + static void RequestOpen(const DeepWindowSpec& spec); + + /** + * @brief Same, but not before `delay_ms` have passed. + * + * Backs the launcher's time-based trigger, which is the only trigger + * available when the target's source can't be edited. Costs no thread: + * the due-time is checked on the launch beat that is already running. + */ + static void ScheduleOpenAfter(int64_t delay_ms, const DeepWindowSpec& spec); + + /** @brief Close an active window. No-op when none is open. */ + static void Close(DeepWindowClose reason); + + static bool Active(); + + /** + * @brief Per-launch bound check, driven from the CUPTI launch callback. + * + * Consumes one launch of budget and closes the window once either + * bound is reached. The close has to land here: the launch callback on + * the application thread is the only reliably scheduled, + * context-current place to run a mid-session CUPTI stop/collect. + */ + static void OnLaunch(); + + /** + * @brief Periodic bound check for when launches stop before the deadline. + * + * `may_close_here` is false on threads that must not run the CUPTI + * teardown (the collector thread against a Windows-injected target). + * There this only records that a close is due and the next launch + * performs it, so a window whose workload stops launching entirely + * stays open until session stop. + */ + static void OnPeriodicTick(bool may_close_here); + + /** @brief Test seam: drop all state without touching a backend. */ + static void ResetForTesting(); + + private: + // Claims a pending request and opens it. Called from OnLaunch only. + static void TakePendingOpen_(); +}; + +namespace detail { + +// Perf scopes (Range Profiler / Perfworks) only mean something for engines +// that use them. Shared by ScopedMonitor and DeepWindow so the two paths +// can't drift apart on which engines get one. +// `is_deep_window` routes to the backend's deep-window hooks, which stay +// live under DeepArmMode::WindowOnly while ordinary user scopes go quiet. +bool PerfScopeEnabled(); +void BeginPerfScopeIfEnabled(const char* name, bool is_deep_window); +// Ends the perf scope and writes the decoded PerfMetricEvent, if the engine +// produced one. Mirrors the tail of ScopedMonitor's destructor. +void EndPerfScopeIfEnabled(const char* name, int pid, int64_t start_ns, + int64_t end_ns, bool is_deep_window); + +} // namespace detail +} // namespace gpufl diff --git a/include/gpufl/core/env_vars.hpp b/include/gpufl/core/env_vars.hpp index b5a14aa..9978b96 100644 --- a/include/gpufl/core/env_vars.hpp +++ b/include/gpufl/core/env_vars.hpp @@ -132,6 +132,22 @@ constexpr const char* kPcSamplingPeriod = "GPUFL_PC_SAMPLING_PERIOD"; constexpr const char* kDeepPcOnly = "GPUFL_DEEP_PC_ONLY"; constexpr const char* kDeepTryBoth = "GPUFL_DEEP_TRY_BOTH"; +// ── Bounded deep window ───────────────────────────────────────────────────── +// When the deep engines hold their session armed: "always" (default) or +// "window" (idle until gpufl::deepWindow opens one). Reaches the injection +// path, which has no other way to set MonitorOptions::deep_arm_mode. +constexpr const char* kDeepArm = "GPUFL_DEEP_ARM"; +// Default bounds for a window opened with that bound left at 0. Wall time +// alone under-describes the replay engines, so both exist. +constexpr const char* kDeepWindowMs = "GPUFL_DEEP_WINDOW_MS"; +constexpr const char* kDeepWindowMaxLaunches = "GPUFL_DEEP_WINDOW_MAX_LAUNCHES"; +// Quiet time after a window closes before another may open. +constexpr const char* kDeepWindowCooldownMs = "GPUFL_DEEP_WINDOW_COOLDOWN_MS"; +// Time-based trigger: open a window this long after init. The only trigger +// that reaches a target whose source can't be edited, so `gpufl trace` sets +// it from --deep-after. Unset = no scheduled window. +constexpr const char* kDeepAfterMs = "GPUFL_DEEP_AFTER_MS"; + // ── SASS metrics knobs ────────────────────────────────────────────────────── constexpr const char* kSassMetricsOnly = "GPUFL_SASS_METRICS_ONLY"; constexpr const char* kSassForceSafeActivity = "GPUFL_SASS_FORCE_SAFE_ACTIVITY"; diff --git a/include/gpufl/core/events.hpp b/include/gpufl/core/events.hpp index ff5577e..cb48af8 100644 --- a/include/gpufl/core/events.hpp +++ b/include/gpufl/core/events.hpp @@ -638,6 +638,33 @@ struct NvtxMarkerEvent { uint32_t marker_id = 0; // CUPTI marker ID (for debug / dedup) }; +/** + * One bounded deep-profiling window: the region between a + * gpufl::deepWindow() trigger and the bound that closed it. + * + * Deep engines arm on open and disarm on close, so this is the only + * record of what the window actually covered. Both the requested bounds + * and the outcome are carried, because they routinely disagree: under + * kernel replay a three-second window can cover a dozen launches, and + * `close_reason` is what tells the reader that was the deadline expiring + * rather than the profiler failing. + */ +struct DeepWindowEvent { + int pid = 0; + std::string app; + std::string session_id; + std::string name; + std::string close_reason; // DeepWindowCloseName(): "deadline" | ... + std::string engine; // ProfilingEngineWireName of the armed engine + int64_t start_ns = 0; + int64_t end_ns = 0; + int64_t duration_ns = 0; + uint64_t launches_covered = 0; + // What the caller asked for, so a short window is self-explanatory. + int64_t requested_duration_ms = 0; + uint64_t requested_max_launches = 0; +}; + /** * One CUDA graph launch event captured by CUPTI's * CUPTI_ACTIVITY_KIND_GRAPH_TRACE stream. diff --git a/include/gpufl/core/gpufl.cpp b/include/gpufl/core/gpufl.cpp index a66d5e9..f8e64f3 100644 --- a/include/gpufl/core/gpufl.cpp +++ b/include/gpufl/core/gpufl.cpp @@ -21,6 +21,7 @@ #include "gpufl/core/teardown_flag.hpp" // detail::isProcessExitTeardown #include "gpufl/core/config_file_loader.hpp" #include "gpufl/core/debug_logger.hpp" +#include "gpufl/core/deep_window.hpp" #include "gpufl/core/events.hpp" #include "gpufl/core/logger/logger.hpp" #include "gpufl/core/remote_config.hpp" @@ -33,7 +34,6 @@ // lives in remote_config.cpp, which includes httplib first and avoids // windows.h entirely. #include "gpufl/core/model/lifecycle_model.hpp" -#include "gpufl/core/model/perf_metric_model.hpp" #include "gpufl/core/model/system_event_model.hpp" #include "gpufl/core/monitor.hpp" #include "gpufl/core/monitor_backend.hpp" @@ -425,6 +425,26 @@ bool init(const InitOptions& opts) { mOpts.pm_sampling_preset = g_opts.pm_sampling_preset; mOpts.pm_sampling_metrics = g_opts.pm_sampling_metrics; mOpts.pm_sampling_scope_only = g_opts.pm_sampling_scope_only; + mOpts.deep_arm_mode = g_opts.deep_window_only ? DeepArmMode::WindowOnly + : DeepArmMode::Always; + // GPUFL_DEEP_ARM reaches the injection path, which can't set InitOptions. + if (const char* v = std::getenv(gpufl::env::kDeepArm)) { + const std::string val(v); + if (val == "window") { + mOpts.deep_arm_mode = DeepArmMode::WindowOnly; + } else if (val == "always") { + mOpts.deep_arm_mode = DeepArmMode::Always; + } else { + GFL_LOG_ERROR("GPUFL_DEEP_ARM='", val, + "' is not recognized. Valid values: always, window. " + "Keeping current deep arm mode."); + } + } + // WindowOnly subsumes the PM-specific gate: PM sampling already arms on + // scope start, which is exactly what a window is. + if (mOpts.deep_arm_mode == DeepArmMode::WindowOnly) { + mOpts.pm_sampling_scope_only = true; + } mOpts.backend_kind = ToMonitorBackendKind(g_opts.backend); // EAGER module loading is OPT-IN. By default we leave CUDA on its normal @@ -587,6 +607,12 @@ bool init(const InitOptions& opts) { g_nvtx_available.store(true, std::memory_order_release); #endif + // Arm the launcher's time-based deep window, if one was asked for. Last, + // so the delay is measured from a fully-initialized runtime, and safe + // from here: it only records the request; the arm itself happens on the + // app thread at the first launch past the delay. + scheduleEnvDeepWindow(); + GFL_LOG_DEBUG("Initialization complete!"); return true; } @@ -648,6 +674,11 @@ void shutdown() { Runtime* rt = runtime(); if (!rt) return; + // Close a still-open deep window before anything is torn down, so its + // engines disarm through the normal scope-stop path and the window + // lands in the log rather than vanishing with the session. + DeepWindow::Close(DeepWindowClose::SessionStop); + GFL_LOG_DEBUG("Shutdown: begin -> sampler.shutdown()"); // Stop the system sampler before CUPTI/backend teardown. The sampler can // be inside NVML while shutdown begins, especially in injection mode where @@ -787,16 +818,10 @@ void ScopedMonitor::init_(const ScopeMeta& meta) { // Scope callbacks are useful for both tracing and profiling backends. Monitor::BeginProfilerScope(name_.c_str()); - // Perf scope (Range Profiler / Perfworks). Also fire it when an engine combo - // (GPUFL_ENGINE_COMBO) is active even with a Trace base - otherwise a - // Trace+RangeProfiler combo would never trigger Range's perf scope. Harmless - // no-op for engines that don't use perf scopes (PC / PM). - const char* comboEnv = std::getenv(env::kEngineCombo); - const bool comboActive = comboEnv && comboEnv[0] != '\0'; - if (g_opts.profiling_engine != ProfilingEngine::Monitor && - (g_opts.profiling_engine != ProfilingEngine::Trace || comboActive)) { - Monitor::BeginPerfScope(name_.c_str()); - } + // Perf scope (Range Profiler / Perfworks). Harmless no-op for engines + // that don't use perf scopes (PC / PM). Shared with DeepWindow so both + // paths agree on which engines get one. + detail::BeginPerfScopeIfEnabled(name_.c_str(), /*is_deep_window=*/false); } ScopedMonitor::~ScopedMonitor() { @@ -837,27 +862,8 @@ ScopedMonitor::~ScopedMonitor() { // scope as an NVTX marker. That echo duplicated scope_event (the SPA // had to de-dupe it) and only the framework NVTX path remains useful. Monitor::EndProfilerScope(name_.c_str()); - const char* comboEnv = std::getenv(gpufl::env::kEngineCombo); - const bool comboActive = comboEnv && comboEnv[0] != '\0'; - if (g_opts.profiling_engine != ProfilingEngine::Monitor && - (g_opts.profiling_engine != ProfilingEngine::Trace || comboActive)) { - Monitor::EndPerfScope( - name_.c_str()); // triggers EndPerfPassAndDecode first - if (IMonitorBackend* b = Monitor::GetBackend()) { - if (auto event_opt = b->TakeLastPerfEvent()) { - PerfMetricEvent& pe = *event_opt; - pe.pid = pid_; - pe.app = rt->app_name; - pe.session_id = rt->session_id; - pe.name = name_; - pe.start_ns = start_ns_; - pe.end_ns = end_ns; - rt->logger->write(model::PerfMetricModel(pe)); - - GFL_LOG_DEBUG("Log Perf Metric Event"); - } - } - } + detail::EndPerfScopeIfEnabled(name_.c_str(), pid_, start_ns_, end_ns, + /*is_deep_window=*/false); } void generateReport(const std::string& output_path) { namespace fs = std::filesystem; diff --git a/include/gpufl/core/model/deep_window_model.cpp b/include/gpufl/core/model/deep_window_model.cpp new file mode 100644 index 0000000..a3d219e --- /dev/null +++ b/include/gpufl/core/model/deep_window_model.cpp @@ -0,0 +1,28 @@ +#include "gpufl/core/model/deep_window_model.hpp" + +#include + +#include "gpufl/core/model/model_utils.hpp" + +namespace gpufl::model { + +std::string DeepWindowModel::buildJson() const { + std::ostringstream oss; + oss << "{\"type\":\"deep_window_event\"" + << ",\"pid\":" << e_.pid + << ",\"app\":\"" << jsonEscape(e_.app) << "\"" + << ",\"session_id\":\"" << jsonEscape(e_.session_id) << "\"" + << ",\"name\":\"" << jsonEscape(e_.name) << "\"" + << ",\"close_reason\":\"" << jsonEscape(e_.close_reason) << "\"" + << ",\"engine\":\"" << jsonEscape(e_.engine) << "\"" + << ",\"start_ns\":" << e_.start_ns + << ",\"end_ns\":" << e_.end_ns + << ",\"duration_ns\":" << e_.duration_ns + << ",\"launches_covered\":" << e_.launches_covered + << ",\"requested_duration_ms\":" << e_.requested_duration_ms + << ",\"requested_max_launches\":" << e_.requested_max_launches + << "}"; + return oss.str(); +} + +} // namespace gpufl::model diff --git a/include/gpufl/core/model/deep_window_model.hpp b/include/gpufl/core/model/deep_window_model.hpp new file mode 100644 index 0000000..cafd346 --- /dev/null +++ b/include/gpufl/core/model/deep_window_model.hpp @@ -0,0 +1,23 @@ +#pragma once + +#include "gpufl/core/events.hpp" +#include "gpufl/core/model/serializable.hpp" + +namespace gpufl::model { + +/** + * JSON serializer for DeepWindowEvent. Emitted to the Scope channel: a + * deep window is a named region with start/end timestamps, so it belongs + * with scope_event_batch and nvtx_marker_event rather than in the + * per-device event stream. + */ +struct DeepWindowModel final : IJsonSerializable { + explicit DeepWindowModel(const DeepWindowEvent& e) : e_(e) {} + std::string buildJson() const override; + Channel channel() const override { return Channel::Scope; } + + private: + const DeepWindowEvent& e_; +}; + +} // namespace gpufl::model diff --git a/include/gpufl/core/monitor.cpp b/include/gpufl/core/monitor.cpp index 4390554..6a25f32 100644 --- a/include/gpufl/core/monitor.cpp +++ b/include/gpufl/core/monitor.cpp @@ -618,6 +618,10 @@ void Monitor::RecordStop(void* handle, StreamHandle) { void Monitor::BeginProfilerScope(const char* name) { if (auto* b = GetBackend()) b->OnScopeStart(name); } void Monitor::EndProfilerScope(const char* name) { if (auto* b = GetBackend()) b->OnScopeStop(name); } +void Monitor::BeginDeepWindowScope(const char* name) { if (auto* b = GetBackend()) b->OnDeepWindowStart(name); } +void Monitor::EndDeepWindowScope(const char* name) { if (auto* b = GetBackend()) b->OnDeepWindowStop(name); } +void Monitor::BeginDeepWindowPerfScope(const char* name) { if (auto* b = GetBackend()) b->OnDeepWindowPerfStart(name); } +void Monitor::EndDeepWindowPerfScope(const char* name) { if (auto* b = GetBackend()) b->OnDeepWindowPerfStop(name); } void Monitor::BeginPerfScope(const char* name) { if (auto* b = GetBackend()) b->OnPerfScopeStart(name); } void Monitor::EndPerfScope(const char* name) { if (auto* b = GetBackend()) b->OnPerfScopeStop(name); } diff --git a/include/gpufl/core/monitor.hpp b/include/gpufl/core/monitor.hpp index 999345f..e671630 100644 --- a/include/gpufl/core/monitor.hpp +++ b/include/gpufl/core/monitor.hpp @@ -133,6 +133,22 @@ inline const char* ProfilingEngineSessionKind(const ProfilingEngine engine) { return "monitor"; } +/** + * @brief When the deep engines hold their hardware resources armed. + * + * Always - armed from session start to session stop. + * WindowOnly - idle until a deep window opens (gpufl::deepWindow), and + * idle again once it closes. Lets a long-running job carry + * deep capture without paying for it outside the moments + * it asked for. + * + * Only the engines that hold a replay or sampling session care. Cubin + * capture and the CUPTI subscriber run from process start either way - + * they have to, or SASS correlation loses the modules loaded before the + * first window. + */ +enum class DeepArmMode { Always, WindowOnly }; + struct MonitorOptions { bool enable_debug_output = false; bool enable_stack_trace = false; @@ -175,6 +191,7 @@ struct MonitorOptions { std::string pm_sampling_preset = "overview"; std::vector pm_sampling_metrics; bool pm_sampling_scope_only = true; + DeepArmMode deep_arm_mode = DeepArmMode::Always; // Default Monitor: no CUPTI. The user-facing default lives on // InitOptions (gpufl.hpp); this internal default matches it so a // bare MonitorOptions (e.g. the system-monitor daemon, which only @@ -303,6 +320,17 @@ class Monitor { static void BeginProfilerScope(const char* name); static void EndProfilerScope(const char* name); + /** + * @brief Profiler scope control for a bounded deep window. + * + * Same engines, but routed separately so a WindowOnly backend can arm + * for a window without arming for every user scope. + */ + static void BeginDeepWindowScope(const char* name); + static void EndDeepWindowScope(const char* name); + static void BeginDeepWindowPerfScope(const char* name); + static void EndDeepWindowPerfScope(const char* name); + /** * @brief Hardware counter (Perfworks) scope control */ diff --git a/include/gpufl/core/monitor_backend.hpp b/include/gpufl/core/monitor_backend.hpp index 7d39033..5f95e0b 100644 --- a/include/gpufl/core/monitor_backend.hpp +++ b/include/gpufl/core/monitor_backend.hpp @@ -88,11 +88,28 @@ class IMonitorBackend { virtual void OnScopeStart(const char* name) {} virtual void OnScopeStop(const char* name) {} + /** + * @brief Scope hooks for a bounded deep window, as opposed to an + * ordinary user scope. + * + * They arm and disarm the same engines, but a backend running in + * DeepArmMode::WindowOnly has to tell the two apart: under that mode + * ordinary scopes must NOT arm anything, or a per-step GFL_SCOPE in a + * training loop would leave the engines armed for the whole run and + * the window would mean nothing. Default: identical to a user scope, + * which is right for every backend that arms unconditionally. + */ + virtual void OnDeepWindowStart(const char* name) { OnScopeStart(name); } + virtual void OnDeepWindowStop(const char* name) { OnScopeStop(name); } + /** @brief Periodically drain buffered profiling data. Thread-safe. */ virtual void DrainProfilingData() {} virtual void OnPerfScopeStart(const char* name) {} virtual void OnPerfScopeStop(const char* name) {} + // Perf-scope counterparts of OnDeepWindowStart/Stop; see those. + virtual void OnDeepWindowPerfStart(const char* name) { OnPerfScopeStart(name); } + virtual void OnDeepWindowPerfStop(const char* name) { OnPerfScopeStop(name); } virtual std::optional TakeLastPerfEvent() { return std::nullopt; } }; diff --git a/include/gpufl/gpufl.hpp b/include/gpufl/gpufl.hpp index 9a9be10..ff5f5ad 100644 --- a/include/gpufl/gpufl.hpp +++ b/include/gpufl/gpufl.hpp @@ -6,6 +6,7 @@ #include #include +#include "gpufl/core/deep_window.hpp" #include "gpufl/core/monitor.hpp" namespace gpufl { @@ -111,6 +112,13 @@ struct InitOptions { std::vector pm_sampling_metrics; bool pm_sampling_scope_only = true; + // When true the deep engines stay idle until gpufl::deepWindow() opens + // a window, instead of holding their session armed for the whole run. + // The setup that must happen at process start (CUPTI subscribe, cubin + // capture, Profiler API init) still does; only the sampling / replay + // session is deferred. Env override: GPUFL_DEEP_ARM=window. + bool deep_window_only = false; + // ── Configuration sources, in precedence order (low → high) ──────────── // // 1. InitOptions defaults (these field initializers) @@ -181,6 +189,65 @@ BackendProbeResult probeRocm(); void systemStart(std::string name = "system"); void systemStop(std::string name = "system"); +// Bounded deep-profiling window. Arms the deep engines (PC sampling, SASS +// metrics, PM sampling, Range profiler) for a short region and disarms +// them automatically, so a long-running job can profile the moment it +// went wrong without paying replay cost for its whole lifetime: +// +// if (tokens_per_sec < 1000) gpufl::deepWindow(3000); +// +// Non-blocking; the workload keeps running through the window and past +// its end. Calling it again while a window is open is IGNORED rather +// than treated as an extension, so a trigger that fires every iteration +// can't pin the window open. +// +// `max_duration_ms` and `max_launches` combine with OR - whichever is hit +// first closes the window; 0 disables that bound. +// +// Prefer a launch budget for the replay engines. Wall time and work done +// diverge sharply there: measured on one workload, a 1-second window +// covered ~1200 launches under PM sampling but ~50 under SASS metrics, +// because replay re-runs every kernel. A launch budget says what you +// actually meant; a duration says how long you are willing to be slow. +// +// `max_launches` counts kernel-launch API callbacks, not kernels - the +// runtime API path reports roughly two per launch (cudaLaunchKernel and +// the cuLaunchKernel beneath it), so budget accordingly. +// +// Bounds are only observed at a launch boundary, the one place a +// mid-session CUPTI stop is safe, so a window overruns its deadline by up +// to one kernel launch. The window's actual coverage and the bound that +// closed it are recorded in the session's deep_window event. +// +// Requires an engine that supports deep capture; a no-op under +// ProfilingEngine::Monitor and on backends without one. +void deepWindow(int64_t max_duration_ms, uint64_t max_launches = 0); + +// Full-control form. Adds `cooldown_ms`, the minimum quiet time after a +// window closes before another may open. Set it when the condition can stay +// true: without it a per-step check reopens a window the instant the last +// one expired and the run pays deep cost forever. The library is the only +// party that knows when the last window closed, which is why the bound +// lives here rather than in your trigger. +// +// gpufl::DeepWindowSpec spec; +// spec.max_duration_ms = 3000; +// spec.cooldown_ms = 60000; // at most one window a minute +// if (tokens_per_sec < 1000) gpufl::deepWindow(spec); +void deepWindow(const DeepWindowSpec& spec); + +// Close the current window early. No-op when none is open. +void deepWindowClose(); + +bool deepWindowActive(); + +// Arms a window `GPUFL_DEEP_AFTER_MS` after this call, sized by +// `GPUFL_DEEP_WINDOW_MS` / `GPUFL_DEEP_WINDOW_MAX_LAUNCHES`. A no-op when +// the first is unset. Called by init(); `gpufl trace --deep-after` is what +// normally sets those, since a target whose source can't be edited has no +// way to call deepWindow() itself. +void scheduleEnvDeepWindow(); + // F1 (External Correlation) - active push/pop for callers that want to // tag CUDA work with an op id WITHOUT relying on a framework profiler // being active. Used by `gpufl.torch.attach()` to stamp every aten diff --git a/python/bindings.cpp b/python/bindings.cpp index 7f02c5a..19628f9 100644 --- a/python/bindings.cpp +++ b/python/bindings.cpp @@ -142,6 +142,7 @@ PYBIND11_MODULE(_gpufl_client, m) { std::string pm_sampling_preset, std::vector pm_sampling_metrics, bool pm_sampling_scope_only, + bool deep_window_only, bool flush_logs_always) -> bool { gpufl::InitOptions opts; @@ -166,6 +167,7 @@ PYBIND11_MODULE(_gpufl_client, m) { opts.pm_sampling_preset = std::move(pm_sampling_preset); opts.pm_sampling_metrics = std::move(pm_sampling_metrics); opts.pm_sampling_scope_only = pm_sampling_scope_only; + opts.deep_window_only = deep_window_only; return gpufl::init(opts); }, py::arg("app_name"), @@ -188,6 +190,7 @@ PYBIND11_MODULE(_gpufl_client, m) { py::arg("pm_sampling_preset") = "overview", py::arg("pm_sampling_metrics") = std::vector{}, py::arg("pm_sampling_scope_only") = true, + py::arg("deep_window_only") = false, py::arg("flush_logs_always") = false); m.def("system_start", [](std::string name) { gpufl::systemStart(std::move(name)); }, @@ -196,6 +199,15 @@ PYBIND11_MODULE(_gpufl_client, m) { m.def("system_stop", [](std::string name) { gpufl::systemStop(std::move(name)); }, py::arg("name") = "system"); + m.def("deep_window", + [](double seconds, uint64_t max_launches) { + gpufl::deepWindow(static_cast(seconds * 1000.0), + max_launches); + }, + py::arg("seconds") = 0.0, py::arg("max_launches") = 0); + m.def("deep_window_close", &gpufl::deepWindowClose); + m.def("deep_window_active", &gpufl::deepWindowActive); + m.def("shutdown", &gpufl::shutdown); // ── Deferred bulk upload ──────────────────────────────────────────── diff --git a/python/gpufl/__init__.py b/python/gpufl/__init__.py index f55296d..41e10fb 100644 --- a/python/gpufl/__init__.py +++ b/python/gpufl/__init__.py @@ -182,6 +182,8 @@ def _dbg(msg): try: from ._gpufl_client import ( Scope as _CScope, init, shutdown, system_start, system_stop, + deep_window as _c_deep_window, deep_window_close as _c_deep_window_close, + deep_window_active as _c_deep_window_active, BackendKind, InitOptions, ProfilingEngine, upload_logs as _c_upload_logs, UploadOptions, UploadResult, ) @@ -209,6 +211,15 @@ def system_start(name="system"): def system_stop(name="system"): return None + def _c_deep_window(seconds=0.0, max_launches=0): + return None + + def _c_deep_window_close(): + return None + + def _c_deep_window_active(): + return False + class BackendKind: Auto = "Auto" Nvidia = "Nvidia" @@ -593,6 +604,56 @@ def system_stop(name="system"): return _original_system_stop(name) +def deep_window(seconds=0.0, max_launches=0): + """Arm deep profiling for a short, self-closing window. + + Lets a long-running job profile the moment it went wrong without + carrying replay cost for its whole lifetime:: + + if tokens_per_sec < 1000: + gpufl.deep_window(3.0) + + Returns immediately; the workload keeps running through the window and + past its end. Calling it again while a window is open is ignored rather + than treated as an extension, so a check inside a training loop can run + every step without pinning the window open. + + Args: + seconds: Wall-clock bound. 0 = no time bound. + max_launches: Launch-callback bound. 0 = no launch bound. Counts + kernel-launch API callbacks, not kernels - the runtime + API path reports roughly two per launch. + + The two bounds combine with OR - whichever is reached first closes the + window. Prefer a launch budget when the engine replays kernels (SASS, + Range profiler): measured on one workload, a 1-second window covered + ~1200 launches under PM sampling but ~50 under SASS metrics. Bounds are + observed at launch boundaries, so a window overruns its deadline by up + to one kernel launch. The coverage the window actually got, and which + bound ended it, are recorded in the session's deep_window event. + + Requires an engine that supports deep capture; a no-op under + ProfilingEngine.Monitor or when gpufl is disabled. + """ + if _disabled: + return None + return _c_deep_window(seconds, max_launches) + + +def deep_window_close(): + """Close the current deep window early. No-op when none is open.""" + if _disabled: + return None + return _c_deep_window_close() + + +def deep_window_active(): + """True while a deep window is open.""" + if _disabled: + return False + return _c_deep_window_active() + + # ── upload_logs: deferred bulk upload to the backend ──────────────────────── # # Thin Python wrapper around the C++ uploadLogs() so callers can use @@ -1015,6 +1076,7 @@ def clean_logs(log_path=None, log_prefix=None, *, dry_run=False): __all__ = [ "Scope", "init", "shutdown", "session", "clean_logs", "targeting", "system_start", "system_stop", + "deep_window", "deep_window_close", "deep_window_active", "BackendKind", "InitOptions", "ProfilingEngine", "upload_logs", "UploadOptions", "UploadResult", ] diff --git a/scripts/deep_window_check.py b/scripts/deep_window_check.py new file mode 100644 index 0000000..2d41f98 --- /dev/null +++ b/scripts/deep_window_check.py @@ -0,0 +1,157 @@ +#!/usr/bin/env python3 +"""Verify one deep-window session's NDJSON logs. + +Checks what the feature actually promises: + 1. exactly one window, closed by a bound rather than by session stop; + 2. deep samples exist and stop when the window closes; + 3. the light tier kept running for the whole session, not just the window. + +Two wrinkles a naive "is every sample inside [start,end]" check gets wrong, +both learned the hard way: + + * PC and SASS samples are stamped at COLLECT time, and the collect runs at + window close - just after the end_ns the window recorded. They are + in-window data carrying an out-of-window timestamp, so a short grace + after end_ns counts as inside. + * Some batches arrive on the raw CUPTI clock instead of the wall-clock + anchor, landing days from the session. Those are reported separately. + Under `gpufl trace` this used to hit every PM sample (the deferred engine + start dropped the anchor); it should now be zero there, so the count is + worth watching rather than ignoring. + +Usage: deep_window_check.py [--expect-windows N] +Exit: 0 pass, 1 fail, 2 nothing to check. +""" +import argparse +import gzip +import json +import pathlib +import sys + +GRACE_NS = 100_000_000 # 100ms: a close-time collect, not a leak +ALIEN_NS = 60_000_000_000 # >60s from the window = a different clock domain +WALL_FLOOR = 1_000_000_000_000_000_000 +DEEP_BATCHES = ("pm_sample_batch", "profile_sample_batch") + + +def find_session(root: pathlib.Path) -> pathlib.Path: + if list(root.glob("*.log.gz")) or list(root.glob("*.log")): + return root + dirs = [p for p in root.rglob("*") + if p.is_dir() and (list(p.glob("*.log.gz")) or list(p.glob("*.log")))] + if not dirs: + sys.exit(f"no session logs under {root}") + return sorted(dirs, key=lambda p: p.stat().st_mtime)[-1] + + +def read_rows(path: pathlib.Path): + opener = gzip.open if path.suffix == ".gz" else open + with opener(path, "rt", encoding="utf-8", errors="replace") as fh: + for line in fh: + line = line.strip() + if not line: + continue + try: + yield json.loads(line) + except json.JSONDecodeError: + continue + + +def main() -> int: + ap = argparse.ArgumentParser() + ap.add_argument("path") + ap.add_argument("--expect-windows", type=int, default=1) + # A window that closes at session stop means a bound never fired. + ap.add_argument("--allow-session-stop", action="store_true") + args = ap.parse_args() + + sess = find_session(pathlib.Path(args.path)) + rows = [] + for f in sorted(sess.glob("*.log*")): + # Lifecycle records fan out to several channels; de-dup by channel so + # counts mean what they look like. + channel = f.name.split(".")[0] + rows += [(channel, r) for r in read_rows(f)] + + windows = [r for ch, r in rows + if r.get("type") == "deep_window_event" and ch == "scope"] + kernels = sum(1 for ch, r in rows + if r.get("type") == "kernel_detail" and ch == "device") + + print(f"session: {sess.name}") + job = next((r for _, r in rows if r.get("type") == "job_start"), None) + if job: + print(f" engine : {job.get('profiling_engine')}") + print(f" deep windows : {len(windows)}") + for w in windows: + print(f" close_reason={w['close_reason']} " + f"duration={w['duration_ns'] / 1e6:.1f}ms " + f"launches={w['launches_covered']} " + f"requested={w['requested_duration_ms']}ms/" + f"{w['requested_max_launches']}launches") + + failures = [] + if len(windows) != args.expect_windows: + failures.append(f"expected {args.expect_windows} window(s), " + f"found {len(windows)}") + if not windows: + print("VERDICT: FAIL - no window to check") + return 1 + + w = windows[0] + if w["close_reason"] == "session_stop" and not args.allow_session_stop: + failures.append("window closed at session stop - no bound fired, so " + "the run was shorter than the window") + + start, end = w["start_ns"], w["end_ns"] + inside = late = alien = 0 + latest_late_ms = 0.0 + for _, r in rows: + if r.get("type") not in DEEP_BATCHES: + continue + base = r.get("base_time_ns", 0) + cols = r.get("columns", []) + if "dt_ns" not in cols: + continue + dt_i = cols.index("dt_ns") + for row in r.get("rows", []): + ts = base + row[dt_i] + if ts < start - ALIEN_NS or ts > end + ALIEN_NS: + alien += 1 + elif start <= ts <= end + GRACE_NS: + inside += 1 + else: + late += 1 + latest_late_ms = max(latest_late_ms, (ts - end) / 1e6) + + # The Range profiler reports per-kernel metric events, not sample batches. + perf_events = sum(1 for _, r in rows + if r.get("type") == "kernel_perf_metric_event") + + print(f" deep samples inside : {inside}") + print(f" deep samples late : {late}" + + (f" (latest +{latest_late_ms:.0f}ms after close)" if late else "")) + print(f" unanchored samples : {alien}" + + (" <- raw CUPTI clock, not the wall anchor" if alien else "")) + print(f" range perf events : {perf_events}") + print(f" kernel rows (run) : {kernels}") + + if late: + failures.append(f"{late} deep samples after the window closed - " + "the engine kept sampling") + if inside == 0 and perf_events == 0: + failures.append("no deep data collected in the window") + if kernels == 0: + failures.append("no kernel rows - the light tier did not run") + + if failures: + print("VERDICT: FAIL") + for f in failures: + print(f" - {f}") + return 1 + print("VERDICT: PASS") + return 0 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/scripts/deep_window_e2e.sh b/scripts/deep_window_e2e.sh new file mode 100644 index 0000000..b1fdb1a --- /dev/null +++ b/scripts/deep_window_e2e.sh @@ -0,0 +1,208 @@ +#!/usr/bin/env bash +# End-to-end verification for the bounded deep-profiling window. +# +# Written for a Linux box with working GPU performance-counter access (the +# RTX 3090 dev machine). The Windows 5060 cannot finish this: its PC sampling +# stops producing hardware samples and needs elevation plus a reboot, so the +# PC and SASS legs there are unverifiable. +# +# Three things get checked, in increasing order of what they prove: +# +# A. EMBED - an app that calls gpufl::deepWindow() itself, per engine. +# B. INJECT - `gpufl trace --deep-after/--deep-for` against a target with +# ZERO gpufl calls. This is the case the feature exists for, +# since a job whose source you can't edit can never call the +# API. Also the only leg that exercises the deferred engine +# start, which is where the clock-anchor bug lived. +# C. OVERHEAD - what an armed-but-idle session costs. Deep capture is only +# worth carrying through a long run if idle is close to free, +# and that number has never been measured. +# +# Usage: +# scripts/deep_window_e2e.sh [--build-dir DIR] [--out DIR] [--seconds N] +# [--skip-overhead] [--engines "A B C"] +# +# Requires: a build with GPUFL_ENABLE_NVIDIA=ON, BUILD_GPUFL_EXAMPLE=ON, +# BUILD_GPUFL_LAUNCHER=ON, BUILD_GPUFL_INJECT=ON; nvcc on PATH; python3. + +set -uo pipefail + +BUILD_DIR="build" +OUT_DIR="" +SECONDS_PER_RUN=12 +SKIP_OVERHEAD=0 +ENGINES="PmSampling PcSampling SassMetrics RangeProfilerKernelReplay" +REPS=3 + +while [[ $# -gt 0 ]]; do + case "$1" in + --build-dir) BUILD_DIR="$2"; shift 2 ;; + --out) OUT_DIR="$2"; shift 2 ;; + --seconds) SECONDS_PER_RUN="$2"; shift 2 ;; + --engines) ENGINES="$2"; shift 2 ;; + --reps) REPS="$2"; shift 2 ;; + --skip-overhead) SKIP_OVERHEAD=1; shift ;; + -h|--help) sed -n '2,30p' "$0"; exit 0 ;; + *) echo "unknown flag: $1" >&2; exit 2 ;; + esac +done + +SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +REPO_DIR="$(cd "$SCRIPT_DIR/.." && pwd)" +[[ -n "$OUT_DIR" ]] || OUT_DIR="$REPO_DIR/deep_window_e2e_$(date +%Y%m%d-%H%M%S)" +mkdir -p "$OUT_DIR" + +CHECK="$SCRIPT_DIR/deep_window_check.py" +REPORT="$OUT_DIR/report.md" + +# Build dirs differ between single- and multi-config generators. +find_bin() { + local name="$1" p + for p in "$BUILD_DIR/$2/$name" "$BUILD_DIR/$2/Release/$name" \ + "$BUILD_DIR/$name" "$BUILD_DIR/Release/$name"; do + [[ -x "$p" ]] && { echo "$p"; return 0; } + done + return 1 +} + +DEMO="$(find_bin deep_window_demo example/cuda)" || { + echo "deep_window_demo not found under $BUILD_DIR - build it first:" >&2 + echo " cmake -S . -B $BUILD_DIR -DGPUFL_ENABLE_NVIDIA=ON -DBUILD_GPUFL_EXAMPLE=ON \\" >&2 + echo " -DBUILD_GPUFL_LAUNCHER=ON -DBUILD_GPUFL_INJECT=ON" >&2 + echo " cmake --build $BUILD_DIR -j" >&2 + exit 2 +} +GPUFL="$(find_bin gpufl daemon/launcher)" || { + echo "gpufl launcher not found under $BUILD_DIR (BUILD_GPUFL_LAUNCHER=ON?)" >&2 + exit 2 +} + +TARGET="$OUT_DIR/deep_window_target" +echo "[e2e] compiling the gpufl-unaware target" +nvcc -O2 -lineinfo -o "$TARGET" "$SCRIPT_DIR/deep_window_target.cu" \ + >"$OUT_DIR/nvcc.log" 2>&1 || { cat "$OUT_DIR/nvcc.log"; exit 2; } + +{ + echo "# Deep window E2E" + echo + echo "- host: \`$(hostname)\`" + echo "- gpu: \`$(nvidia-smi --query-gpu=name,driver_version --format=csv,noheader 2>/dev/null | head -1)\`" + echo "- build: \`$BUILD_DIR\`" + echo "- run seconds: $SECONDS_PER_RUN" + echo +} > "$REPORT" + +PASS=0 +FAIL=0 +note() { echo "$@" | tee -a "$REPORT"; } +record() { # name, exit code + if [[ "$2" -eq 0 ]]; then PASS=$((PASS+1)); note "- **PASS** $1"; + else FAIL=$((FAIL+1)); note "- **FAIL** $1"; fi +} + +# ── A. embed ──────────────────────────────────────────────────────────────── +note "## A. Embedded API (app calls gpufl::deepWindow)" +note +for eng in $ENGINES; do + d="$OUT_DIR/embed_$eng"; mkdir -p "$d" + echo "[e2e] embed: $eng" + # A faster card gets through the default loop before the window expires, + # which closes it at session stop and proves nothing. Scale with + # GPUFL_DEMO_ITERATIONS if the check reports that. + ( cd "$d" && GPUFL_PROFILING_ENGINE="$eng" \ + GPUFL_DEMO_ITERATIONS="${GPUFL_DEMO_ITERATIONS:-3000}" "$DEMO" ) \ + >"$d/stdout.txt" 2>&1 + demo_rc=$? + python3 "$CHECK" "$d/deep_window" >"$d/check.txt" 2>&1 + chk_rc=$? + sed 's/^/ /' "$d/check.txt" >> "$REPORT" + [[ $demo_rc -eq 0 && $chk_rc -eq 0 ]] + record "embed $eng" $? +done +note + +# ── B. injection ──────────────────────────────────────────────────────────── +# The load-bearing leg: the target has no gpufl calls, so the only possible +# trigger is the launcher's. Also the only path that runs the deferred engine +# start, where the wall-clock anchor was being dropped. +note "## B. Injection (\`gpufl trace\`, target has zero gpufl calls)" +note +for eng in $ENGINES; do + d="$OUT_DIR/inject_$eng"; mkdir -p "$d" + echo "[e2e] inject: $eng" + "$GPUFL" trace -o "$d/trace" --passes "$eng" \ + --deep-after 3s --deep-for 2s -- "$TARGET" "$SECONDS_PER_RUN" \ + >"$d/stdout.txt" 2>&1 + trace_rc=$? + python3 "$CHECK" "$d/trace" >"$d/check.txt" 2>&1 + chk_rc=$? + sed 's/^/ /' "$d/check.txt" >> "$REPORT" + [[ $trace_rc -eq 0 && $chk_rc -eq 0 ]] + record "inject $eng" $? +done + +# Launch-bound leg: the bound that actually suits the replay engines, where a +# second of wall time buys far less work. +d="$OUT_DIR/inject_launch_bound"; mkdir -p "$d" +echo "[e2e] inject: SassMetrics bounded by launches" +"$GPUFL" trace -o "$d/trace" --passes SassMetrics \ + --deep-after 3s --deep-launches 200 -- "$TARGET" "$SECONDS_PER_RUN" \ + >"$d/stdout.txt" 2>&1 +trace_rc=$? +python3 "$CHECK" "$d/trace" >"$d/check.txt" 2>&1 +chk_rc=$? +sed 's/^/ /' "$d/check.txt" >> "$REPORT" +[[ $trace_rc -eq 0 && $chk_rc -eq 0 ]] +record "inject SassMetrics --deep-launches 200" $? +note + +# ── C. armed-but-idle overhead ────────────────────────────────────────────── +# Conditions are interleaved rather than run in blocks so thermal drift hits +# all of them equally. +if [[ "$SKIP_OVERHEAD" -eq 0 ]]; then + note "## C. Overhead (iterations/sec, median of $REPS)" + note + declare -A RESULTS + run_probe() { # label, then argv for the run + local label="$1"; shift + local out + out="$("$@" 2>/dev/null | grep -o 'ITERS_PER_SEC=[0-9.]*' | tail -1)" + RESULTS["$label"]+="${out#ITERS_PER_SEC=} " + } + for _ in $(seq 1 "$REPS"); do + run_probe "no gpufl" "$TARGET" "$SECONDS_PER_RUN" + run_probe "trace (light tier)" \ + "$GPUFL" trace -o "$OUT_DIR/ovh_light" --passes Trace -- "$TARGET" "$SECONDS_PER_RUN" + # Window-only arming with no window ever opened - the idle cost. Set + # via env because --deep-* always opens one. + export GPUFL_DEEP_ARM=window + run_probe "PM armed-but-idle" \ + "$GPUFL" trace -o "$OUT_DIR/ovh_idle" --passes PmSampling -- "$TARGET" "$SECONDS_PER_RUN" + unset GPUFL_DEEP_ARM + run_probe "PM always armed" \ + "$GPUFL" trace -o "$OUT_DIR/ovh_always" --passes PmSampling -- "$TARGET" "$SECONDS_PER_RUN" + done + + median() { tr ' ' '\n' <<<"$1" | grep -v '^$' | sort -n | awk '{a[NR]=$1} END{print (NR%2)?a[(NR+1)/2]:(a[NR/2]+a[NR/2+1])/2}'; } + base="$(median "${RESULTS["no gpufl"]:-}")" + note "| condition | iters/sec | vs no gpufl |" + note "|---|---|---|" + for label in "no gpufl" "trace (light tier)" "PM armed-but-idle" "PM always armed"; do + m="$(median "${RESULTS[$label]:-}")" + if [[ -n "$m" && -n "$base" && "$base" != "0" ]]; then + pct="$(awk -v a="$m" -v b="$base" 'BEGIN{printf "%+.1f%%", (a/b-1)*100}')" + else pct="n/a"; fi + note "| $label | ${m:-n/a} | $pct |" + done + note + note "\`PM armed-but-idle\` is the number that decides whether deep capture" + note "can ride along in a long run. \`PM always armed\` is what the window avoids." + note +fi + +note "## Summary" +note +note "passed: $PASS, failed: $FAIL" +echo +echo "[e2e] report: $REPORT" +[[ "$FAIL" -eq 0 ]] diff --git a/scripts/deep_window_target.cu b/scripts/deep_window_target.cu new file mode 100644 index 0000000..d19c803 --- /dev/null +++ b/scripts/deep_window_target.cu @@ -0,0 +1,72 @@ +// A CUDA workload that knows nothing about gpufl. +// +// Stands in for the case the deep window exists for: a long-running job whose +// source you can't edit, so it can never call gpufl::deepWindow() itself and +// the trigger has to come from `gpufl trace`. Deliberately NOT part of the +// example/ tree, which links gpufl - the point here is that it doesn't. +// +// Runs for a wall-clock duration and reports how many iterations it got +// through, so the same binary doubles as the throughput probe for measuring +// what an armed-but-idle session costs. +// +// nvcc -O2 -lineinfo -o deep_window_target deep_window_target.cu +// ./deep_window_target [seconds] + +#include + +#include +#include +#include + +__global__ void computeHeavy(float* out, const float* in, int n, int iters) { + const int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx >= n) return; + float val = in[idx]; + for (int i = 0; i < iters; ++i) { + val = val * 1.0009f + 0.0001f; + val = fmaf(val, 0.9991f, 0.0002f); + } + out[idx] = val; +} + +int main(int argc, char** argv) { + const int seconds = argc > 1 ? std::atoi(argv[1]) : 12; + const int n = 1 << 20; + + float* d_in = nullptr; + float* d_out = nullptr; + if (cudaMalloc(&d_in, n * sizeof(float)) != cudaSuccess || + cudaMalloc(&d_out, n * sizeof(float)) != cudaSuccess) { + std::fprintf(stderr, "cudaMalloc failed\n"); + return 2; + } + cudaMemset(d_in, 0, n * sizeof(float)); + + const int threads = 256; + const int blocks = (n + threads - 1) / threads; + + // One untimed pass so context creation, module load and clock ramp land + // outside the measured region. + computeHeavy<<>>(d_out, d_in, n, 4000); + cudaDeviceSynchronize(); + + const auto t0 = std::chrono::steady_clock::now(); + long iterations = 0; + for (;;) { + const auto elapsed = std::chrono::duration_cast( + std::chrono::steady_clock::now() - t0).count(); + if (elapsed >= seconds * 1000L) break; + computeHeavy<<>>(d_out, d_in, n, 4000); + cudaDeviceSynchronize(); + ++iterations; + } + const double secs = std::chrono::duration_cast( + std::chrono::steady_clock::now() - t0).count() / 1000.0; + + // The throughput line the harness parses. + std::printf("ITERATIONS=%ld SECONDS=%.3f ITERS_PER_SEC=%.2f\n", + iterations, secs, iterations / secs); + cudaFree(d_in); + cudaFree(d_out); + return 0; +} diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 3dfc9f7..684accd 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -15,6 +15,7 @@ set(GPUFL_TEST_SOURCES core/test_api_path_routing.cpp core/test_batch_models.cpp core/test_bench_invoker.cpp + core/test_deep_window.cpp core/test_disabled.cpp core/test_wire_contract.cpp core/test_monitor.cpp diff --git a/tests/core/test_deep_window.cpp b/tests/core/test_deep_window.cpp new file mode 100644 index 0000000..b334f76 --- /dev/null +++ b/tests/core/test_deep_window.cpp @@ -0,0 +1,325 @@ +// Tests for the bounded deep-profiling window. +// +// The window's job is to arm the deep engines for a short region and then +// close ITSELF, so what matters here is the state machine: which bound +// fires first, that a repeated trigger doesn't extend an open window, and +// that the close reason reported is the one that actually fired. The CUPTI +// arming behind it needs a GPU and is covered by the E2E runs. +// +// A real runtime is initialized (backend None, no sampler) because Open() +// deliberately refuses to report a window as open when gpufl isn't running. + +#include + +#include +#include +#include +#include + +#include "gpufl/core/deep_window.hpp" +#include "gpufl/core/events.hpp" +#include "gpufl/core/model/deep_window_model.hpp" +#include "gpufl/core/runtime.hpp" +#include "gpufl/gpufl.hpp" + +namespace { + +class DeepWindowTest : public ::testing::Test { + protected: + void SetUp() override { + log_dir_ = (std::filesystem::temp_directory_path() / + "gpufl_deep_window_test") + .string(); + std::filesystem::remove_all(log_dir_); + + gpufl::InitOptions opts; + opts.app_name = "deep-window-test"; + opts.log_path = log_dir_; + opts.backend = gpufl::BackendKind::None; + opts.system_sample_rate_ms = 0; + opts.continuous_system_sampling = false; + opts.enable_stack_trace = false; + opts.enable_source_collection = false; + ASSERT_TRUE(gpufl::init(opts)); + gpufl::DeepWindow::ResetForTesting(); + } + + void TearDown() override { + gpufl::DeepWindow::ResetForTesting(); + gpufl::shutdown(); + std::error_code ec; + std::filesystem::remove_all(log_dir_, ec); + } + + std::string log_dir_; +}; + +gpufl::DeepWindowSpec Spec(const int64_t ms, const uint64_t launches, + const int64_t cooldown_ms = 0) { + gpufl::DeepWindowSpec spec; + spec.max_duration_ms = ms; + spec.max_launches = launches; + spec.cooldown_ms = cooldown_ms; + return spec; +} + +} // namespace + +// ── open / close state machine ────────────────────────────────────────────── + +TEST_F(DeepWindowTest, StartsInactive) { + EXPECT_FALSE(gpufl::DeepWindow::Active()); +} + +TEST_F(DeepWindowTest, OpenActivatesAndManualCloseDeactivates) { + EXPECT_TRUE(gpufl::DeepWindow::Open(Spec(0, 0))); + EXPECT_TRUE(gpufl::DeepWindow::Active()); + + gpufl::DeepWindow::Close(gpufl::DeepWindowClose::Manual); + EXPECT_FALSE(gpufl::DeepWindow::Active()); +} + +TEST_F(DeepWindowTest, CloseWithoutAnOpenWindowIsHarmless) { + gpufl::DeepWindow::Close(gpufl::DeepWindowClose::Manual); + EXPECT_FALSE(gpufl::DeepWindow::Active()); +} + +TEST_F(DeepWindowTest, SecondOpenIsIgnoredNotAnExtension) { + // The motivating case: a trigger that re-fires every training step must + // not hold the window open past its bound. + ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(0, 3))); + EXPECT_FALSE(gpufl::DeepWindow::Open(Spec(0, 1000))); + + // Still bounded by the FIRST spec's budget of 3. + gpufl::DeepWindow::OnLaunch(); + EXPECT_FALSE(gpufl::DeepWindow::Open(Spec(0, 1000))); + gpufl::DeepWindow::OnLaunch(); + EXPECT_TRUE(gpufl::DeepWindow::Active()); + gpufl::DeepWindow::OnLaunch(); + EXPECT_FALSE(gpufl::DeepWindow::Active()); +} + +TEST_F(DeepWindowTest, ReopensAfterClosing) { + ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(0, 1))); + gpufl::DeepWindow::OnLaunch(); + ASSERT_FALSE(gpufl::DeepWindow::Active()); + + EXPECT_TRUE(gpufl::DeepWindow::Open(Spec(0, 1))); + EXPECT_TRUE(gpufl::DeepWindow::Active()); +} + +// ── bounds ────────────────────────────────────────────────────────────────── + +TEST_F(DeepWindowTest, LaunchBudgetClosesOnTheNthLaunch) { + ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(0, 2))); + gpufl::DeepWindow::OnLaunch(); + EXPECT_TRUE(gpufl::DeepWindow::Active()) << "budget of 2 spent after 1"; + gpufl::DeepWindow::OnLaunch(); + EXPECT_FALSE(gpufl::DeepWindow::Active()); +} + +TEST_F(DeepWindowTest, NoBoundsMeansOnlyAManualCloseEndsIt) { + ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(0, 0))); + for (int i = 0; i < 100; ++i) gpufl::DeepWindow::OnLaunch(); + EXPECT_TRUE(gpufl::DeepWindow::Active()); +} + +TEST_F(DeepWindowTest, DeadlineClosesOnTheNextLaunch) { + ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(/*ms=*/1, 0))); + std::this_thread::sleep_for(std::chrono::milliseconds(10)); + // The deadline is only observed at a launch boundary - that is the one + // place a mid-session CUPTI stop is safe. + EXPECT_TRUE(gpufl::DeepWindow::Active()); + gpufl::DeepWindow::OnLaunch(); + EXPECT_FALSE(gpufl::DeepWindow::Active()); +} + +TEST_F(DeepWindowTest, LaunchBudgetWinsWhenItIsReachedFirst) { + ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(/*ms=*/60000, /*launches=*/1))); + gpufl::DeepWindow::OnLaunch(); + EXPECT_FALSE(gpufl::DeepWindow::Active()); +} + +TEST_F(DeepWindowTest, OnLaunchWithNoWindowOpenIsHarmless) { + for (int i = 0; i < 10; ++i) gpufl::DeepWindow::OnLaunch(); + EXPECT_FALSE(gpufl::DeepWindow::Active()); +} + +// ── periodic tick fallback ────────────────────────────────────────────────── + +TEST_F(DeepWindowTest, PeriodicTickClosesAnExpiredWindowWhenAllowed) { + ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(/*ms=*/1, 0))); + std::this_thread::sleep_for(std::chrono::milliseconds(10)); + gpufl::DeepWindow::OnPeriodicTick(/*may_close_here=*/true); + EXPECT_FALSE(gpufl::DeepWindow::Active()); +} + +TEST_F(DeepWindowTest, PeriodicTickDefersToTheNextLaunchWhenNotAllowed) { + // Windows injection: the collector thread must not run the CUPTI + // teardown, so it only flags the close. + ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(/*ms=*/1, 0))); + std::this_thread::sleep_for(std::chrono::milliseconds(10)); + gpufl::DeepWindow::OnPeriodicTick(/*may_close_here=*/false); + EXPECT_TRUE(gpufl::DeepWindow::Active()) << "must wait for a launch"; + + gpufl::DeepWindow::OnLaunch(); + EXPECT_FALSE(gpufl::DeepWindow::Active()); +} + +TEST_F(DeepWindowTest, PeriodicTickLeavesAnUnexpiredWindowAlone) { + ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(/*ms=*/60000, 0))); + gpufl::DeepWindow::OnPeriodicTick(/*may_close_here=*/true); + EXPECT_TRUE(gpufl::DeepWindow::Active()); +} + +TEST_F(DeepWindowTest, PeriodicTickOnAnUnboundedWindowNeverCloses) { + ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(0, 0))); + gpufl::DeepWindow::OnPeriodicTick(/*may_close_here=*/true); + EXPECT_TRUE(gpufl::DeepWindow::Active()); +} + +// ── cooldown ──────────────────────────────────────────────────────────────── + +TEST_F(DeepWindowTest, CooldownBlocksAnImmediateReopen) { + // The trap this exists for: a condition that stays true reopens a window + // the instant the last one expired, and the run pays deep cost forever. + ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(0, 1, /*cooldown_ms=*/60000))); + gpufl::DeepWindow::OnLaunch(); + ASSERT_FALSE(gpufl::DeepWindow::Active()); + + EXPECT_FALSE(gpufl::DeepWindow::Open(Spec(0, 1, 60000))); + EXPECT_FALSE(gpufl::DeepWindow::Active()); +} + +TEST_F(DeepWindowTest, CooldownExpiresAndReopeningWorksAgain) { + ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(0, 1, /*cooldown_ms=*/5))); + gpufl::DeepWindow::OnLaunch(); + ASSERT_FALSE(gpufl::DeepWindow::Active()); + + std::this_thread::sleep_for(std::chrono::milliseconds(20)); + EXPECT_TRUE(gpufl::DeepWindow::Open(Spec(0, 1, 5))); +} + +TEST_F(DeepWindowTest, NoCooldownMeansImmediateReopenIsAllowed) { + ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(0, 1))); + gpufl::DeepWindow::OnLaunch(); + ASSERT_FALSE(gpufl::DeepWindow::Active()); + EXPECT_TRUE(gpufl::DeepWindow::Open(Spec(0, 1))); +} + +// ── deferred arm ──────────────────────────────────────────────────────────── + +TEST_F(DeepWindowTest, RequestOpenArmsOnTheNextLaunchNotImmediately) { + // A trigger off the app thread must not arm: the CUPTI calls behind an + // arm are only safe on the app thread at a launch boundary. + gpufl::DeepWindow::RequestOpen(Spec(60000, 0)); + EXPECT_FALSE(gpufl::DeepWindow::Active()) << "must wait for a launch"; + + gpufl::DeepWindow::OnLaunch(); + EXPECT_TRUE(gpufl::DeepWindow::Active()); +} + +TEST_F(DeepWindowTest, RequestOpenIsConsumedOnce) { + gpufl::DeepWindow::RequestOpen(Spec(60000, 0)); + gpufl::DeepWindow::OnLaunch(); + ASSERT_TRUE(gpufl::DeepWindow::Active()); + + gpufl::DeepWindow::Close(gpufl::DeepWindowClose::Manual); + // The request was spent on the first arm; a closed window stays closed. + gpufl::DeepWindow::OnLaunch(); + EXPECT_FALSE(gpufl::DeepWindow::Active()); +} + +TEST_F(DeepWindowTest, RequestOpenCarriesItsBounds) { + gpufl::DeepWindow::RequestOpen(Spec(0, /*launches=*/2)); + gpufl::DeepWindow::OnLaunch(); // arms; does not consume budget + ASSERT_TRUE(gpufl::DeepWindow::Active()); + + gpufl::DeepWindow::OnLaunch(); + EXPECT_TRUE(gpufl::DeepWindow::Active()); + gpufl::DeepWindow::OnLaunch(); + EXPECT_FALSE(gpufl::DeepWindow::Active()); +} + +TEST_F(DeepWindowTest, NewestPendingSpecWins) { + gpufl::DeepWindow::RequestOpen(Spec(0, 1)); + gpufl::DeepWindow::RequestOpen(Spec(60000, 0)); + gpufl::DeepWindow::OnLaunch(); + ASSERT_TRUE(gpufl::DeepWindow::Active()); + + // Had the first spec won, this launch would spend its budget of 1. + gpufl::DeepWindow::OnLaunch(); + EXPECT_TRUE(gpufl::DeepWindow::Active()); +} + +TEST_F(DeepWindowTest, ScheduledOpenWaitsOutItsDelay) { + gpufl::DeepWindow::ScheduleOpenAfter(/*delay_ms=*/50, Spec(60000, 0)); + gpufl::DeepWindow::OnLaunch(); + EXPECT_FALSE(gpufl::DeepWindow::Active()) << "not due yet"; + + std::this_thread::sleep_for(std::chrono::milliseconds(70)); + gpufl::DeepWindow::OnLaunch(); + EXPECT_TRUE(gpufl::DeepWindow::Active()); +} + +TEST_F(DeepWindowTest, RequestOpenWhileAWindowIsOpenDoesNotDisturbIt) { + ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(60000, 0))); + gpufl::DeepWindow::RequestOpen(Spec(0, 1)); + gpufl::DeepWindow::OnLaunch(); + EXPECT_TRUE(gpufl::DeepWindow::Active()) + << "the open window's bounds still govern"; +} + +// ── public API surface ────────────────────────────────────────────────────── + +TEST_F(DeepWindowTest, PublicApiOpensAndCloses) { + gpufl::deepWindow(/*max_duration_ms=*/60000); + EXPECT_TRUE(gpufl::deepWindowActive()); + gpufl::deepWindowClose(); + EXPECT_FALSE(gpufl::deepWindowActive()); +} + +// ── close reasons ─────────────────────────────────────────────────────────── + +TEST_F(DeepWindowTest, CloseReasonWireNames) { + EXPECT_STREQ("deadline", + gpufl::DeepWindowCloseName(gpufl::DeepWindowClose::Deadline)); + EXPECT_STREQ( + "launch_budget", + gpufl::DeepWindowCloseName(gpufl::DeepWindowClose::LaunchBudget)); + EXPECT_STREQ("manual", + gpufl::DeepWindowCloseName(gpufl::DeepWindowClose::Manual)); + EXPECT_STREQ( + "session_stop", + gpufl::DeepWindowCloseName(gpufl::DeepWindowClose::SessionStop)); +} + +// ── event serialization ───────────────────────────────────────────────────── + +TEST(DeepWindowModelTest, SerializesRequestedBoundsAlongsideTheOutcome) { + // Both sides are on the wire on purpose: "asked for 3000ms, got 12 + // launches, closed by deadline" is the reading that makes a short + // window legible instead of looking like a failure. + gpufl::DeepWindowEvent e; + e.pid = 4242; + e.app = "trainer"; + e.session_id = "sess-1"; + e.name = "deep_window"; + e.close_reason = "deadline"; + e.engine = "nvidia.pc_sampling"; + e.start_ns = 1000; + e.end_ns = 3000; + e.duration_ns = 2000; + e.launches_covered = 12; + e.requested_duration_ms = 3000; + e.requested_max_launches = 0; + + const std::string json = gpufl::model::DeepWindowModel(e).buildJson(); + EXPECT_NE(json.find("\"type\":\"deep_window_event\""), std::string::npos); + EXPECT_NE(json.find("\"close_reason\":\"deadline\""), std::string::npos); + EXPECT_NE(json.find("\"engine\":\"nvidia.pc_sampling\""), std::string::npos); + EXPECT_NE(json.find("\"launches_covered\":12"), std::string::npos); + EXPECT_NE(json.find("\"requested_duration_ms\":3000"), std::string::npos); + EXPECT_NE(json.find("\"duration_ns\":2000"), std::string::npos); + EXPECT_EQ(gpufl::model::DeepWindowModel(e).channel(), gpufl::Channel::Scope); +} diff --git a/tests/launcher/test_cli_parse.cpp b/tests/launcher/test_cli_parse.cpp index 5e1b87f..5a34d84 100644 --- a/tests/launcher/test_cli_parse.cpp +++ b/tests/launcher/test_cli_parse.cpp @@ -180,6 +180,58 @@ TEST(CliParseTrace, AfterWindowBogusRejected) { EXPECT_NE(r.error.find("invalid --after-window"), std::string::npos); } +// ── --deep-*: bound the DEEP engines inside a target that keeps running, +// as opposed to --window, which bounds the target's lifetime. + +TEST(CliParseTrace, DeepWindowFlagsParse) { + auto r = parseTraceArgs(argsFor({"--deep-after=30s", "--deep-for=3s", + "--deep-cooldown=1m", "--", "./bin"})); + ASSERT_TRUE(r.args.has_value()) << r.error; + EXPECT_EQ(r.args->deep_after_ms, 30000); + EXPECT_EQ(r.args->deep_for_ms, 3000); + EXPECT_EQ(r.args->deep_cooldown_ms, 60000); + EXPECT_TRUE(r.args->deep_requested); +} + +TEST(CliParseTrace, DeepLaunchesParses) { + auto r = parseTraceArgs(argsFor({"--deep-launches", "500", "--", "./bin"})); + ASSERT_TRUE(r.args.has_value()) << r.error; + EXPECT_EQ(r.args->deep_launches, 500u); + EXPECT_TRUE(r.args->deep_requested); +} + +TEST(CliParseTrace, DeepWindowWithoutABoundRejected) { + // An unbounded deep window is just "profile deeply for the whole run". + auto r = parseTraceArgs(argsFor({"--deep-after=30s", "--", "./bin"})); + EXPECT_FALSE(r.args.has_value()); + EXPECT_NE(r.error.find("needs a bound"), std::string::npos); +} + +TEST(CliParseTrace, DeepLaunchesAloneIsABound) { + auto r = parseTraceArgs(argsFor({"--deep-launches=200", "--", "./bin"})); + ASSERT_TRUE(r.args.has_value()) << r.error; + EXPECT_EQ(r.args->deep_after_ms, 0) << "arms at the first launch"; +} + +TEST(CliParseTrace, DeepLaunchesZeroRejected) { + auto r = parseTraceArgs(argsFor({"--deep-launches=0", "--", "./bin"})); + EXPECT_FALSE(r.args.has_value()); + EXPECT_NE(r.error.find("invalid --deep-launches"), std::string::npos); +} + +TEST(CliParseTrace, DeepForBogusDurationRejected) { + auto r = parseTraceArgs(argsFor({"--deep-for=soon", "--", "./bin"})); + EXPECT_FALSE(r.args.has_value()); + EXPECT_NE(r.error.find("invalid --deep-for"), std::string::npos); +} + +TEST(CliParseTrace, NoDeepFlagsLeavesDeepWindowOff) { + auto r = parseTraceArgs(argsFor({"--window=10s", "--", "./bin"})); + ASSERT_TRUE(r.args.has_value()) << r.error; + EXPECT_FALSE(r.args->deep_requested); + EXPECT_EQ(r.args->window_ms, 10000) << "--window is unaffected"; +} + TEST(CliParseTrace, EngineFlagRejectedWithMigrationHint) { auto r = parseTraceArgs(argsFor({"--engine=Deep", "--", "./bin"})); EXPECT_FALSE(r.args.has_value()); From 227b42322c958b2122a8db16d0fb0bd4d27e94b4 Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Fri, 24 Jul 2026 22:49:46 -0700 Subject: [PATCH 02/10] fix: test update --- scripts/deep_window_check.py | 23 +++++++++++++++++++---- scripts/deep_window_e2e.sh | 21 +++++++++++++++++++-- 2 files changed, 38 insertions(+), 6 deletions(-) diff --git a/scripts/deep_window_check.py b/scripts/deep_window_check.py index 2d41f98..de0485c 100644 --- a/scripts/deep_window_check.py +++ b/scripts/deep_window_check.py @@ -80,8 +80,14 @@ def main() -> int: print(f"session: {sess.name}") job = next((r for _, r in rows if r.get("type") == "job_start"), None) + engine = job.get("profiling_engine", "") if job else "" if job: - print(f" engine : {job.get('profiling_engine')}") + print(f" engine : {engine}") + # SASS metrics runs with kernel activity suppressed by default - the + # combination deadlocks on some drivers, so safe mode drops it and the + # collector also discards orphaned launches rather than synthesizing + # rows. Zero kernel rows is correct there, not a missing light tier. + expect_kernels = "sass_metrics" not in engine print(f" deep windows : {len(windows)}") for w in windows: print(f" close_reason={w['close_reason']} " @@ -134,14 +140,23 @@ def main() -> int: print(f" unanchored samples : {alien}" + (" <- raw CUPTI clock, not the wall anchor" if alien else "")) print(f" range perf events : {perf_events}") - print(f" kernel rows (run) : {kernels}") + print(f" kernel rows (run) : {kernels}" + + ("" if expect_kernels else " (suppressed for SASS by design)")) if late: failures.append(f"{late} deep samples after the window closed - " "the engine kept sampling") if inside == 0 and perf_events == 0: - failures.append("no deep data collected in the window") - if kernels == 0: + if alien: + # PC sampling stamps some batches on the raw CUPTI clock, so the + # window cannot be checked by timestamp there. Say so instead of + # calling it a collection failure. + failures.append( + f"{alien} deep samples exist but are all on an unanchored " + "clock - cannot tell whether they fell inside the window") + else: + failures.append("no deep data collected in the window") + if expect_kernels and kernels == 0: failures.append("no kernel rows - the light tier did not run") if failures: diff --git a/scripts/deep_window_e2e.sh b/scripts/deep_window_e2e.sh index b1fdb1a..0b1fa5c 100644 --- a/scripts/deep_window_e2e.sh +++ b/scripts/deep_window_e2e.sh @@ -33,6 +33,7 @@ SECONDS_PER_RUN=12 SKIP_OVERHEAD=0 ENGINES="PmSampling PcSampling SassMetrics RangeProfilerKernelReplay" REPS=3 +DEBUG=0 while [[ $# -gt 0 ]]; do case "$1" in @@ -42,6 +43,10 @@ while [[ $# -gt 0 ]]; do --engines) ENGINES="$2"; shift 2 ;; --reps) REPS="$2"; shift 2 ;; --skip-overhead) SKIP_OVERHEAD=1; shift ;; + # Turns on GPUFL_DEBUG in every run: the engine diagnostics are the + # only way to tell "armed but the hardware gave nothing" apart from + # "never armed". + --debug) DEBUG=1; shift ;; -h|--help) sed -n '2,30p' "$0"; exit 0 ;; *) echo "unknown flag: $1" >&2; exit 2 ;; esac @@ -54,13 +59,17 @@ mkdir -p "$OUT_DIR" CHECK="$SCRIPT_DIR/deep_window_check.py" REPORT="$OUT_DIR/report.md" +[[ "$DEBUG" -eq 1 ]] && export GPUFL_DEBUG=1 # Build dirs differ between single- and multi-config generators. +# Absolute, because the embed leg runs from inside its own output directory +# (the demo writes its logs relative to cwd) and a relative path would not +# survive the cd. find_bin() { local name="$1" p for p in "$BUILD_DIR/$2/$name" "$BUILD_DIR/$2/Release/$name" \ "$BUILD_DIR/$name" "$BUILD_DIR/Release/$name"; do - [[ -x "$p" ]] && { echo "$p"; return 0; } + [[ -x "$p" ]] && { (cd "$(dirname "$p")" && printf '%s/%s\n' "$PWD" "$name"); return 0; } done return 1 } @@ -163,10 +172,18 @@ if [[ "$SKIP_OVERHEAD" -eq 0 ]]; then note "## C. Overhead (iterations/sec, median of $REPS)" note declare -A RESULTS + declare -i probe_n=0 run_probe() { # label, then argv for the run local label="$1"; shift + probe_n+=1 + local log="$OUT_DIR/ovh_${probe_n}_$(tr -c 'a-zA-Z0-9' '_' <<<"$label").log" local out - out="$("$@" 2>/dev/null | grep -o 'ITERS_PER_SEC=[0-9.]*' | tail -1)" + # Keep the log: a probe that yields no throughput line is a run that + # failed, and discarding stderr makes that indistinguishable from a + # parse bug. + out="$("$@" 2>"$log" | tee -a "$log" \ + | grep -o 'ITERS_PER_SEC=[0-9.]*' | tail -1)" + [[ -n "$out" ]] || echo "[e2e] probe '$label' produced no throughput line; see $log" >&2 RESULTS["$label"]+="${out#ITERS_PER_SEC=} " } for _ in $(seq 1 "$REPS"); do From faad6eccc14ff1d3397a663031fbc326c0fa9371 Mon Sep 17 00:00:00 2001 From: CodingInAVan Date: Fri, 24 Jul 2026 22:51:19 -0700 Subject: [PATCH 03/10] update unbuntu build script to support --- README.md | 4 ++++ build-ubuntu.sh | 34 +++++++++++++++++++++++++++++++++- 2 files changed, 37 insertions(+), 1 deletion(-) diff --git a/README.md b/README.md index 4c3758a..5feccd3 100644 --- a/README.md +++ b/README.md @@ -77,6 +77,9 @@ Toolkit, Python virtual environment, or wheel ABI. # Build a wheel into ./dist ./build.sh --wheel +# Build the native gpufl trace launcher and injection library +./build.sh --trace + # Use an explicit Python and CUDA Toolkit ./build-ubuntu.sh --wheel \ --python .venv/bin/python \ @@ -89,6 +92,7 @@ Useful options: |---|---| | `--install` | Install the package into the selected Python environment. This is the default. | | `--wheel` | Build a wheel into `./dist` or `--wheel-dir`. | +| `--trace` | Build the native `gpufl` launcher and `libgpufl_inject.so` into `./build-ubuntu`. | | `--python PATH` | Python executable to use. Use your target virtual environment's Python when building wheels. | | `--cuda-root PATH` | CUDA Toolkit root, for example `/usr/local/cuda-13.2`. | | `--wheel-dir PATH` | Output directory for built wheels. | diff --git a/build-ubuntu.sh b/build-ubuntu.sh index 5487af3..34da876 100755 --- a/build-ubuntu.sh +++ b/build-ubuntu.sh @@ -9,7 +9,7 @@ WHEEL_DIR="$ROOT_DIR/dist" usage() { cat <<'EOF' -Usage: ./build-ubuntu.sh [--install|--wheel] [--python PATH] [--cuda-root PATH] [--wheel-dir PATH] +Usage: ./build-ubuntu.sh [--install|--wheel|--trace] [--python PATH] [--cuda-root PATH] [--wheel-dir PATH] Defaults: --install @@ -20,6 +20,7 @@ Defaults: Examples: ./build-ubuntu.sh ./build-ubuntu.sh --wheel + ./build-ubuntu.sh --trace ./build-ubuntu.sh --python .venv/bin/python --cuda-root /usr/local/cuda-13.2 EOF } @@ -34,6 +35,10 @@ while [[ $# -gt 0 ]]; do MODE="wheel" shift ;; + --trace) + MODE="trace" + shift + ;; --python) PYTHON_BIN="$2" shift 2 @@ -99,6 +104,33 @@ echo " cuda root: $CUDA_ROOT" if [[ "$MODE" == "wheel" ]]; then mkdir -p "$WHEEL_DIR" "$PYTHON_BIN" -m pip wheel "$ROOT_DIR" -w "$WHEEL_DIR" --no-deps -v "${COMMON_CONFIG[@]}" +elif [[ "$MODE" == "trace" ]]; then + BUILD_DIR="$ROOT_DIR/build-ubuntu" + TRACE_CONFIG=( + -DCMAKE_BUILD_TYPE=Release + -DGPUFL_ENABLE_NVIDIA=ON + -DGPUFL_ENABLE_AMD=OFF + -DBUILD_PYTHON=OFF + -DBUILD_TESTING=OFF + -DBUILD_GPUFL_EXAMPLE=OFF + -DBUILD_GPUFL_LAUNCHER=ON + -DBUILD_GPUFL_INJECT=ON + "-DCUDAToolkit_ROOT=$CUDA_ROOT" + "-DCMAKE_CUDA_COMPILER=$CUDA_ROOT/bin/nvcc" + ) + + cmake -S "$ROOT_DIR" -B "$BUILD_DIR" "${TRACE_CONFIG[@]}" + cmake --build "$BUILD_DIR" --target gpufl_launcher gpufl_inject -j + + LAUNCHER="$BUILD_DIR/daemon/launcher/gpufl" + INJECT_LIBRARY="$BUILD_DIR/libgpufl_inject.so" + echo + echo "Built native trace tooling:" + echo " launcher: $LAUNCHER" + echo " inject: $INJECT_LIBRARY" + echo + echo "Run: \"$LAUNCHER\" trace --passes=Trace -- \"$PYTHON_BIN\" " + echo " \"$LAUNCHER\" trace --passes=PcSampling -- \"$PYTHON_BIN\" " else "$PYTHON_BIN" -m pip install "$ROOT_DIR" -v "${COMMON_CONFIG[@]}" fi From e3df908955d833a670abcdf4d5fb7d34e2654963 Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Fri, 24 Jul 2026 23:47:23 -0700 Subject: [PATCH 04/10] fix(capture): run the deep window's engine work off the CUPTI callback path --- .../gpufl/backends/nvidia/cupti_backend.cpp | 35 ++++-- .../gpufl/backends/nvidia/cupti_backend.hpp | 1 + .../nvidia/engine/pm_sampling_engine.cpp | 47 ++++++- .../nvidia/engine/pm_sampling_engine.hpp | 6 + include/gpufl/core/deep_window.cpp | 73 ++++++----- include/gpufl/core/deep_window.hpp | 36 ++++-- include/gpufl/core/monitor.cpp | 5 + include/gpufl/core/monitor_adapter.hpp | 1 + include/gpufl/core/monitor_backend.hpp | 14 +++ scripts/deep_window_check.py | 27 ++-- scripts/deep_window_e2e.sh | 0 tests/core/test_deep_window.cpp | 116 ++++++++++++------ 12 files changed, 260 insertions(+), 101 deletions(-) mode change 100644 => 100755 scripts/deep_window_e2e.sh diff --git a/include/gpufl/backends/nvidia/cupti_backend.cpp b/include/gpufl/backends/nvidia/cupti_backend.cpp index 7975526..fe80230 100644 --- a/include/gpufl/backends/nvidia/cupti_backend.cpp +++ b/include/gpufl/backends/nvidia/cupti_backend.cpp @@ -751,14 +751,33 @@ void CuptiBackend::EngineLaunchTick() { void CuptiBackend::DrainProfilingData() { if (!initialized_ || !active_.load(std::memory_order_relaxed)) return; - // Fallback for a window whose workload stopped launching before the - // deadline. On a Windows-injected target the CUPTI teardown must not run - // from this collector thread, so there it only flags the close and the - // next launch performs it. - DeepWindow::OnPeriodicTick(/*may_close_here=*/!WindowsInjectedProcess()); - if (engine_) { - engine_->drainData(); - } + if (!engine_) return; + // Bind the context: this is the collector thread, and an engine draining + // its hardware buffer (PM sampling decodes here) needs one current. + CUcontext prev = nullptr; + cuCtxGetCurrent(&prev); + const bool rebound = ctx_ && prev != ctx_ && + cuCtxSetCurrent(ctx_) == CUDA_SUCCESS; + engine_->drainData(); + if (rebound) cuCtxSetCurrent(prev); +} + +void CuptiBackend::ServiceDeepWindow() { + if (!initialized_ || !active_.load(std::memory_order_relaxed)) return; + // Lock-free gate: skip the context work when there is nothing to do. + if (!DeepWindow::HasPendingWork()) return; + + // The engines want the context current on the calling thread, and this + // is the collector's. Bind and restore, the same way the deferred engine + // start does when it runs off the app thread. + CUcontext prev = nullptr; + cuCtxGetCurrent(&prev); + const bool rebound = ctx_ && prev != ctx_ && + cuCtxSetCurrent(ctx_) == CUDA_SUCCESS; + + DeepWindow::ServicePending(); + + if (rebound) cuCtxSetCurrent(prev); } void CuptiBackend::StartActivityFlushThreadIfNeeded_() { diff --git a/include/gpufl/backends/nvidia/cupti_backend.hpp b/include/gpufl/backends/nvidia/cupti_backend.hpp index 9a15483..2331984 100644 --- a/include/gpufl/backends/nvidia/cupti_backend.hpp +++ b/include/gpufl/backends/nvidia/cupti_backend.hpp @@ -195,6 +195,7 @@ class CuptiBackend : public IMonitorBackend { OnDeepWindowStart(name); } void DrainProfilingData() override; + void ServiceDeepWindow() override; void OnScopeStop(const char* name) override { GFL_LOG_DEBUG("OnScopeStop"); if (!ScopeArmsEngines_()) return; diff --git a/include/gpufl/backends/nvidia/engine/pm_sampling_engine.cpp b/include/gpufl/backends/nvidia/engine/pm_sampling_engine.cpp index 90cc9bf..e412da8 100644 --- a/include/gpufl/backends/nvidia/engine/pm_sampling_engine.cpp +++ b/include/gpufl/backends/nvidia/engine/pm_sampling_engine.cpp @@ -11,6 +11,7 @@ #include #include "gpufl/backends/nvidia/cupti_utils.hpp" +#include "gpufl/core/common.hpp" // detail::GetTimestampNs #include "gpufl/core/debug_logger.hpp" #include "gpufl/core/teardown_flag.hpp" // detail::isProcessExitTeardown @@ -347,14 +348,50 @@ bool PmSamplingEngine::CreateCounterDataImage_() { } counter_data_image_.assign(size.counterDataSize, 0); + return ResetCounterDataImage_(); +} + +bool PmSamplingEngine::ResetCounterDataImage_() { + if (counter_data_image_.empty()) return false; CUpti_PmSampling_CounterDataImage_Initialize_Params init = { CUpti_PmSampling_CounterDataImage_Initialize_Params_STRUCT_SIZE}; init.pPmSamplingObject = pm_object_; init.counterDataSize = counter_data_image_.size(); init.pCounterData = counter_data_image_.data(); - res = cuptiPmSamplingCounterDataImageInitialize(&init); - if (LogCuptiErrorIfFailed(this->name(), "cuptiPmSamplingCounterDataImageInitialize", res)) return false; - return true; + const CUptiResult res = cuptiPmSamplingCounterDataImageInitialize(&init); + return !LogCuptiErrorIfFailed( + this->name(), "cuptiPmSamplingCounterDataImageInitialize", res); +} + +// The counter-data image holds pm_sampling_max_samples; at +// pm_sampling_interval_us that is a fixed span of wall time (4096 x 100us = +// ~410ms by default). Past it the hardware buffer overflows and this driver +// reports the overflow from DecodeData as CUPTI_ERROR_UNKNOWN rather than +// CUPTI_ERROR_OUT_OF_MEMORY, which loses the whole scope's samples. So drain +// on the collector beat instead of only at scope end. Verified on an RTX +// 3090: scopes up to 350ms decoded, 500ms and longer returned 999. +int64_t PmSamplingEngine::BufferSpanNs_() const { + const int64_t samples = opts_.pm_sampling_max_samples > 0 + ? opts_.pm_sampling_max_samples + : 4096; + const int64_t interval_us = opts_.pm_sampling_interval_us > 0 + ? opts_.pm_sampling_interval_us + : 100; + return samples * interval_us * 1000; +} + +void PmSamplingEngine::drainData() { + if (!running_) return; + // Half the buffer span: enough headroom that a late collector tick still + // lands before the hardware wraps. + const int64_t now = detail::GetTimestampNs(); + const int64_t due = last_drain_ns_ + BufferSpanNs_() / 2; + if (last_drain_ns_ != 0 && now < due) return; + + std::lock_guard lk(pm_mu_); + if (!running_) return; + last_drain_ns_ = now; + DecodeAndEmit_(); } void PmSamplingEngine::DecodeAndEmit_() { @@ -434,6 +471,10 @@ void PmSamplingEngine::DecodeAndEmit_() { Monitor::PushPmSamples(rows); produced_data_.store(true, std::memory_order_relaxed); } + + // Hand the image back empty so the next decode starts from a clean slate + // instead of re-emitting what we just pushed. + ResetCounterDataImage_(); GFL_LOG_DEBUG("[PmSamplingEngine] decoded PM samples completed=", completed, " populated=", info.numPopulatedSamples, " rows=", rows.size()); diff --git a/include/gpufl/backends/nvidia/engine/pm_sampling_engine.hpp b/include/gpufl/backends/nvidia/engine/pm_sampling_engine.hpp index 1b66941..2409fec 100644 --- a/include/gpufl/backends/nvidia/engine/pm_sampling_engine.hpp +++ b/include/gpufl/backends/nvidia/engine/pm_sampling_engine.hpp @@ -29,6 +29,9 @@ class PmSamplingEngine final : public IProfilingEngine { void onScopeStart(const char* name) override; void onScopeStop(const char* name) override; + // Drains the hardware buffer mid-scope; see the definition for why a + // scope longer than the buffer span otherwise loses everything. + void drainData() override; bool hasInsufficientPrivileges() const override { return insufficient_privileges_.load(std::memory_order_relaxed); @@ -55,6 +58,9 @@ class PmSamplingEngine final : public IProfilingEngine { #if GPUFL_HAS_PERFWORKS bool InitializePmSampling_(); bool BuildConfigImage_(); + bool ResetCounterDataImage_(); + int64_t BufferSpanNs_() const; + int64_t last_drain_ns_ = 0; // Log the chip's single-pass metric sets (the metric bundles collectible // in one pass) and their metrics. The set names are chip-specific, so // they're queried at runtime. Read-only (chip query, no host object / no diff --git a/include/gpufl/core/deep_window.cpp b/include/gpufl/core/deep_window.cpp index 7eabe37..e844dd6 100644 --- a/include/gpufl/core/deep_window.cpp +++ b/include/gpufl/core/deep_window.cpp @@ -33,7 +33,10 @@ std::atomic g_active{false}; std::atomic g_deadline_ns{0}; // 0 = no time bound std::atomic g_launches_remaining{0}; // 0 = no launch bound std::atomic g_launches_covered{0}; +// Set by the launch callback when a bound is reached; consumed by the +// collector, which is the thread allowed to run the engines' teardown. std::atomic g_close_requested{false}; +std::atomic g_close_reason{static_cast(DeepWindowClose::Deadline)}; // An open asked for by a thread that can't arm one itself. Checked lock-free // on the launch beat; g_pending_spec is only read once this is set, so the @@ -284,13 +287,9 @@ void DeepWindow::ScheduleOpenAfter(const int64_t delay_ms, " max_launches=", spec.max_launches); } -// Runs on the app thread at launch ENTER - the only place the arm's CUPTI -// calls are safe. Claims the request before opening so two launch threads -// can't both act on it. +// Runs on the collector, off the CUPTI callback path. Claims the request +// before opening so nothing can act on it twice. void DeepWindow::TakePendingOpen_() { - if (!g_open_requested.load(std::memory_order_acquire)) return; - const int64_t due = g_pending_open_at_ns.load(std::memory_order_relaxed); - if (due > 0 && detail::GetTimestampNs() < due) return; if (!g_open_requested.exchange(false, std::memory_order_acq_rel)) return; DeepWindowSpec spec; @@ -303,10 +302,9 @@ void DeepWindow::TakePendingOpen_() { } void DeepWindow::OnLaunch() { - if (!g_active.load(std::memory_order_acquire)) { - TakePendingOpen_(); - return; - } + // Arming is the collector's job too - see ServicePending. This callback + // only counts. + if (!g_active.load(std::memory_order_acquire)) return; g_launches_covered.fetch_add(1, std::memory_order_relaxed); @@ -314,36 +312,53 @@ void DeepWindow::OnLaunch() { // fetch_sub returns the PREVIOUS value, so 1 means this launch // consumed the last of the budget. if (g_launches_remaining.fetch_sub(1, std::memory_order_relaxed) <= 1) { - Close(DeepWindowClose::LaunchBudget); - return; + RequestClose_(DeepWindowClose::LaunchBudget); } } +} - // A tick on a thread that couldn't run the teardown left this set. - if (g_close_requested.load(std::memory_order_acquire)) { - Close(DeepWindowClose::Deadline); - return; +void DeepWindow::RequestClose_(const DeepWindowClose reason) { + // First reason wins; a later bound can't relabel a close already asked + // for. Handing this to the collector instead of closing here is the + // whole point - see OnLaunch. + bool expected = false; + if (g_close_requested.compare_exchange_strong(expected, true, + std::memory_order_acq_rel)) { + g_close_reason.store(static_cast(reason), std::memory_order_release); } +} + +namespace { +bool CloseDue_() { + if (!g_active.load(std::memory_order_acquire)) return false; + if (g_close_requested.load(std::memory_order_acquire)) return true; const int64_t deadline = g_deadline_ns.load(std::memory_order_relaxed); - if (deadline > 0 && detail::GetTimestampNs() >= deadline) { - Close(DeepWindowClose::Deadline); - } + return deadline > 0 && detail::GetTimestampNs() >= deadline; } -void DeepWindow::OnPeriodicTick(const bool may_close_here) { - if (!g_active.load(std::memory_order_acquire)) return; +bool OpenDue_() { + if (g_active.load(std::memory_order_acquire)) return false; + if (!g_open_requested.load(std::memory_order_acquire)) return false; + const int64_t due = g_pending_open_at_ns.load(std::memory_order_relaxed); + return due <= 0 || detail::GetTimestampNs() >= due; +} - const int64_t deadline = g_deadline_ns.load(std::memory_order_relaxed); - if (deadline <= 0 || detail::GetTimestampNs() < deadline) return; +} // namespace + +bool DeepWindow::HasPendingWork() { return CloseDue_() || OpenDue_(); } - if (may_close_here) { - Close(DeepWindowClose::Deadline); +void DeepWindow::ServicePending() { + if (CloseDue_()) { + const DeepWindowClose reason = + g_close_requested.load(std::memory_order_acquire) + ? static_cast( + g_close_reason.load(std::memory_order_acquire)) + : DeepWindowClose::Deadline; + Close(reason); return; } - // Hand the close to the next launch: the CUPTI stop has to run on the - // application thread that owns the context. - g_close_requested.store(true, std::memory_order_release); + if (OpenDue_()) TakePendingOpen_(); } void DeepWindow::ResetForTesting() { @@ -353,6 +368,8 @@ void DeepWindow::ResetForTesting() { g_launches_remaining.store(0, std::memory_order_relaxed); g_launches_covered.store(0, std::memory_order_relaxed); g_close_requested.store(false, std::memory_order_relaxed); + g_close_reason.store(static_cast(DeepWindowClose::Deadline), + std::memory_order_relaxed); g_open_requested.store(false, std::memory_order_relaxed); g_pending_open_at_ns.store(0, std::memory_order_relaxed); g_pending_spec = DeepWindowSpec{}; diff --git a/include/gpufl/core/deep_window.hpp b/include/gpufl/core/deep_window.hpp index bb64e1c..e6313e4 100644 --- a/include/gpufl/core/deep_window.hpp +++ b/include/gpufl/core/deep_window.hpp @@ -94,25 +94,35 @@ class DeepWindow { static bool Active(); /** - * @brief Per-launch bound check, driven from the CUPTI launch callback. + * @brief Per-launch bound accounting, driven from the CUPTI launch + * callback. * - * Consumes one launch of budget and closes the window once either - * bound is reached. The close has to land here: the launch callback on - * the application thread is the only reliably scheduled, - * context-current place to run a mid-session CUPTI stop/collect. + * Consumes one launch of budget and RECORDS that a bound was reached. + * It deliberately does not close: this runs inside a CUPTI callback, + * and the engines' teardown calls (cuptiPmSamplingDecodeData, + * cuptiPCSamplingStop) return CUPTI_ERROR_UNKNOWN when invoked from + * there. Verified on Linux/driver 610.43; Windows happened to tolerate + * it, which is why the first version looked correct. */ static void OnLaunch(); /** - * @brief Periodic bound check for when launches stop before the deadline. + * @brief Cheap, lock-free: is there an arm or a disarm waiting? * - * `may_close_here` is false on threads that must not run the CUPTI - * teardown (the collector thread against a Windows-injected target). - * There this only records that a close is due and the next launch - * performs it, so a window whose workload stops launching entirely - * stays open until session stop. + * Lets the collector poll every iteration and pay for making a CUDA + * context current only when there is actually something to do. */ - static void OnPeriodicTick(bool may_close_here); + static bool HasPendingWork(); + + /** + * @brief Perform a pending arm or disarm. + * + * Both halves run here, off the CUPTI callback path and with the CUDA + * context current, which CuptiBackend::ServiceDeepWindow arranges. Arm + * and disarm are kept on the SAME thread deliberately: PM sampling's + * decode rejects a session whose start and stop straddle threads. + */ + static void ServicePending(); /** @brief Test seam: drop all state without touching a backend. */ static void ResetForTesting(); @@ -120,6 +130,8 @@ class DeepWindow { private: // Claims a pending request and opens it. Called from OnLaunch only. static void TakePendingOpen_(); + // Records that a bound was reached without acting on it. + static void RequestClose_(DeepWindowClose reason); }; namespace detail { diff --git a/include/gpufl/core/monitor.cpp b/include/gpufl/core/monitor.cpp index 6a25f32..583f8dc 100644 --- a/include/gpufl/core/monitor.cpp +++ b/include/gpufl/core/monitor.cpp @@ -449,6 +449,11 @@ void CollectorLoop() { g_state.drainAck.store(req, std::memory_order_release); } + // Every iteration, not on the 250ms flush beat below: this is what + // decides how closely a deep window tracks its deadline, and it is a + // lock-free check when no window is closing. + if (g_state.adapter) g_state.adapter->serviceDeepWindow(); + if (!RecordProcessor::processNext()) { std::this_thread::sleep_for(std::chrono::milliseconds(1)); } diff --git a/include/gpufl/core/monitor_adapter.hpp b/include/gpufl/core/monitor_adapter.hpp index 50eaf72..f6625fb 100644 --- a/include/gpufl/core/monitor_adapter.hpp +++ b/include/gpufl/core/monitor_adapter.hpp @@ -23,6 +23,7 @@ class IMonitorAdapter { virtual std::string memoryKindToString(uint32_t kind) const = 0; virtual void drainProfilingData() { if (backend()) backend()->DrainProfilingData(); } + virtual void serviceDeepWindow() { if (backend()) backend()->ServiceDeepWindow(); } virtual IMonitorBackend* backend() = 0; }; diff --git a/include/gpufl/core/monitor_backend.hpp b/include/gpufl/core/monitor_backend.hpp index 5f95e0b..baf4944 100644 --- a/include/gpufl/core/monitor_backend.hpp +++ b/include/gpufl/core/monitor_backend.hpp @@ -105,6 +105,20 @@ class IMonitorBackend { /** @brief Periodically drain buffered profiling data. Thread-safe. */ virtual void DrainProfilingData() {} + /** + * @brief Close a deep window whose bound has been reached. + * + * Called from the collector on every iteration, not on the slower flush + * beat, because it decides how closely a window tracks its deadline. + * Must stay cheap when there is nothing to close. + * + * It lives here rather than on the launch callback because the engines' + * teardown calls fail with CUPTI_ERROR_UNKNOWN when made from inside a + * CUPTI callback; the collector is off that path and can make the CUDA + * context current itself. + */ + virtual void ServiceDeepWindow() {} + virtual void OnPerfScopeStart(const char* name) {} virtual void OnPerfScopeStop(const char* name) {} // Perf-scope counterparts of OnDeepWindowStart/Stop; see those. diff --git a/scripts/deep_window_check.py b/scripts/deep_window_check.py index de0485c..8b92871 100644 --- a/scripts/deep_window_check.py +++ b/scripts/deep_window_check.py @@ -110,7 +110,7 @@ def main() -> int: "the run was shorter than the window") start, end = w["start_ns"], w["end_ns"] - inside = late = alien = 0 + inside = late = early = alien = 0 latest_late_ms = 0.0 for _, r in rows: if r.get("type") not in DEEP_BATCHES: @@ -124,8 +124,12 @@ def main() -> int: ts = base + row[dt_i] if ts < start - ALIEN_NS or ts > end + ALIEN_NS: alien += 1 - elif start <= ts <= end + GRACE_NS: + elif start - GRACE_NS <= ts <= end + GRACE_NS: + # The engine arms fractionally before the window records its + # start, so a sample straddling the boundary belongs to it. inside += 1 + elif ts < start: + early += 1 else: late += 1 latest_late_ms = max(latest_late_ms, (ts - end) / 1e6) @@ -137,6 +141,8 @@ def main() -> int: print(f" deep samples inside : {inside}") print(f" deep samples late : {late}" + (f" (latest +{latest_late_ms:.0f}ms after close)" if late else "")) + if early: + print(f" deep samples early : {early} (before the window armed)") print(f" unanchored samples : {alien}" + (" <- raw CUPTI clock, not the wall anchor" if alien else "")) print(f" range perf events : {perf_events}") @@ -148,14 +154,15 @@ def main() -> int: "the engine kept sampling") if inside == 0 and perf_events == 0: if alien: - # PC sampling stamps some batches on the raw CUPTI clock, so the - # window cannot be checked by timestamp there. Say so instead of - # calling it a collection failure. - failures.append( - f"{alien} deep samples exist but are all on an unanchored " - "clock - cannot tell whether they fell inside the window") - else: - failures.append("no deep data collected in the window") + # PC sampling stamps its batches on the raw CUPTI clock rather + # than the wall anchor, so the window cannot be checked by + # timestamp there at all. That is a separate, pre-existing gap; + # report it as inconclusive rather than as a window failure. + print(f"INCONCLUSIVE: {alien} deep samples exist but are all on " + "an unanchored clock, so whether they fell inside the " + "window cannot be determined from timestamps") + return 3 + failures.append("no deep data collected in the window") if expect_kernels and kernels == 0: failures.append("no kernel rows - the light tier did not run") diff --git a/scripts/deep_window_e2e.sh b/scripts/deep_window_e2e.sh old mode 100644 new mode 100755 diff --git a/tests/core/test_deep_window.cpp b/tests/core/test_deep_window.cpp index b334f76..4afbe58 100644 --- a/tests/core/test_deep_window.cpp +++ b/tests/core/test_deep_window.cpp @@ -54,6 +54,17 @@ class DeepWindowTest : public ::testing::Test { std::string log_dir_; }; +// The real pairing. The launch callback only records that a bound was +// reached; the collector performs the close, because the engines' teardown +// returns CUPTI_ERROR_UNKNOWN when called from inside a CUPTI callback. +// Tests drive both halves so they exercise the shipped sequence. +void Launch(const int times = 1) { + for (int i = 0; i < times; ++i) { + gpufl::DeepWindow::OnLaunch(); + gpufl::DeepWindow::ServicePending(); + } +} + gpufl::DeepWindowSpec Spec(const int64_t ms, const uint64_t launches, const int64_t cooldown_ms = 0) { gpufl::DeepWindowSpec spec; @@ -91,17 +102,17 @@ TEST_F(DeepWindowTest, SecondOpenIsIgnoredNotAnExtension) { EXPECT_FALSE(gpufl::DeepWindow::Open(Spec(0, 1000))); // Still bounded by the FIRST spec's budget of 3. - gpufl::DeepWindow::OnLaunch(); + Launch(); EXPECT_FALSE(gpufl::DeepWindow::Open(Spec(0, 1000))); - gpufl::DeepWindow::OnLaunch(); + Launch(); EXPECT_TRUE(gpufl::DeepWindow::Active()); - gpufl::DeepWindow::OnLaunch(); + Launch(); EXPECT_FALSE(gpufl::DeepWindow::Active()); } TEST_F(DeepWindowTest, ReopensAfterClosing) { ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(0, 1))); - gpufl::DeepWindow::OnLaunch(); + Launch(); ASSERT_FALSE(gpufl::DeepWindow::Active()); EXPECT_TRUE(gpufl::DeepWindow::Open(Spec(0, 1))); @@ -112,15 +123,15 @@ TEST_F(DeepWindowTest, ReopensAfterClosing) { TEST_F(DeepWindowTest, LaunchBudgetClosesOnTheNthLaunch) { ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(0, 2))); - gpufl::DeepWindow::OnLaunch(); + Launch(); EXPECT_TRUE(gpufl::DeepWindow::Active()) << "budget of 2 spent after 1"; - gpufl::DeepWindow::OnLaunch(); + Launch(); EXPECT_FALSE(gpufl::DeepWindow::Active()); } TEST_F(DeepWindowTest, NoBoundsMeansOnlyAManualCloseEndsIt) { ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(0, 0))); - for (int i = 0; i < 100; ++i) gpufl::DeepWindow::OnLaunch(); + for (int i = 0; i < 100; ++i) Launch(); EXPECT_TRUE(gpufl::DeepWindow::Active()); } @@ -130,61 +141,86 @@ TEST_F(DeepWindowTest, DeadlineClosesOnTheNextLaunch) { // The deadline is only observed at a launch boundary - that is the one // place a mid-session CUPTI stop is safe. EXPECT_TRUE(gpufl::DeepWindow::Active()); - gpufl::DeepWindow::OnLaunch(); + Launch(); EXPECT_FALSE(gpufl::DeepWindow::Active()); } TEST_F(DeepWindowTest, LaunchBudgetWinsWhenItIsReachedFirst) { ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(/*ms=*/60000, /*launches=*/1))); - gpufl::DeepWindow::OnLaunch(); + Launch(); EXPECT_FALSE(gpufl::DeepWindow::Active()); } TEST_F(DeepWindowTest, OnLaunchWithNoWindowOpenIsHarmless) { - for (int i = 0; i < 10; ++i) gpufl::DeepWindow::OnLaunch(); + for (int i = 0; i < 10; ++i) Launch(); EXPECT_FALSE(gpufl::DeepWindow::Active()); } -// ── periodic tick fallback ────────────────────────────────────────────────── +// ── the close runs off the CUPTI callback path ────────────────────────────── +// +// The launch callback may only RECORD that a bound was reached. Closing there +// runs the engines' teardown inside a CUPTI callback, where +// cuptiPmSamplingDecodeData and cuptiPCSamplingStop return +// CUPTI_ERROR_UNKNOWN (observed on Linux/driver 610.43; Windows tolerated it, +// which hid the bug). + +TEST_F(DeepWindowTest, LaunchAloneNeverClosesTheWindow) { + ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(0, /*launches=*/1))); + gpufl::DeepWindow::OnLaunch(); // budget spent, but no teardown here + EXPECT_TRUE(gpufl::DeepWindow::Active()) + << "the launch callback must not run the teardown"; -TEST_F(DeepWindowTest, PeriodicTickClosesAnExpiredWindowWhenAllowed) { - ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(/*ms=*/1, 0))); - std::this_thread::sleep_for(std::chrono::milliseconds(10)); - gpufl::DeepWindow::OnPeriodicTick(/*may_close_here=*/true); + EXPECT_TRUE(gpufl::DeepWindow::HasPendingWork()); + gpufl::DeepWindow::ServicePending(); EXPECT_FALSE(gpufl::DeepWindow::Active()); } -TEST_F(DeepWindowTest, PeriodicTickDefersToTheNextLaunchWhenNotAllowed) { - // Windows injection: the collector thread must not run the CUPTI - // teardown, so it only flags the close. +TEST_F(DeepWindowTest, ServiceCloseEndsAnExpiredWindowWithoutAnyLaunch) { + // A workload that stops launching still has its window closed, because + // the collector polls the deadline itself. ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(/*ms=*/1, 0))); std::this_thread::sleep_for(std::chrono::milliseconds(10)); - gpufl::DeepWindow::OnPeriodicTick(/*may_close_here=*/false); - EXPECT_TRUE(gpufl::DeepWindow::Active()) << "must wait for a launch"; - - gpufl::DeepWindow::OnLaunch(); + EXPECT_TRUE(gpufl::DeepWindow::HasPendingWork()); + gpufl::DeepWindow::ServicePending(); EXPECT_FALSE(gpufl::DeepWindow::Active()); } -TEST_F(DeepWindowTest, PeriodicTickLeavesAnUnexpiredWindowAlone) { +TEST_F(DeepWindowTest, ServiceCloseLeavesAnUnexpiredWindowAlone) { ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(/*ms=*/60000, 0))); - gpufl::DeepWindow::OnPeriodicTick(/*may_close_here=*/true); + EXPECT_FALSE(gpufl::DeepWindow::HasPendingWork()); + gpufl::DeepWindow::ServicePending(); EXPECT_TRUE(gpufl::DeepWindow::Active()); } -TEST_F(DeepWindowTest, PeriodicTickOnAnUnboundedWindowNeverCloses) { +TEST_F(DeepWindowTest, ServiceCloseOnAnUnboundedWindowNeverCloses) { ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(0, 0))); - gpufl::DeepWindow::OnPeriodicTick(/*may_close_here=*/true); + for (int i = 0; i < 5; ++i) gpufl::DeepWindow::ServicePending(); EXPECT_TRUE(gpufl::DeepWindow::Active()); } +TEST_F(DeepWindowTest, CloseDueIsFalseWithNoWindowOpen) { + EXPECT_FALSE(gpufl::DeepWindow::HasPendingWork()); + gpufl::DeepWindow::ServicePending(); // harmless + EXPECT_FALSE(gpufl::DeepWindow::Active()); +} + +TEST_F(DeepWindowTest, LaunchBudgetReasonSurvivesTheHandoff) { + // The reason recorded in the callback has to reach the event the + // collector writes, or a budget close would be reported as a deadline. + ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(/*ms=*/60000, /*launches=*/1))); + gpufl::DeepWindow::OnLaunch(); + ASSERT_TRUE(gpufl::DeepWindow::HasPendingWork()); + gpufl::DeepWindow::ServicePending(); + EXPECT_FALSE(gpufl::DeepWindow::Active()); +} + // ── cooldown ──────────────────────────────────────────────────────────────── TEST_F(DeepWindowTest, CooldownBlocksAnImmediateReopen) { // The trap this exists for: a condition that stays true reopens a window // the instant the last one expired, and the run pays deep cost forever. ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(0, 1, /*cooldown_ms=*/60000))); - gpufl::DeepWindow::OnLaunch(); + Launch(); ASSERT_FALSE(gpufl::DeepWindow::Active()); EXPECT_FALSE(gpufl::DeepWindow::Open(Spec(0, 1, 60000))); @@ -193,7 +229,7 @@ TEST_F(DeepWindowTest, CooldownBlocksAnImmediateReopen) { TEST_F(DeepWindowTest, CooldownExpiresAndReopeningWorksAgain) { ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(0, 1, /*cooldown_ms=*/5))); - gpufl::DeepWindow::OnLaunch(); + Launch(); ASSERT_FALSE(gpufl::DeepWindow::Active()); std::this_thread::sleep_for(std::chrono::milliseconds(20)); @@ -202,7 +238,7 @@ TEST_F(DeepWindowTest, CooldownExpiresAndReopeningWorksAgain) { TEST_F(DeepWindowTest, NoCooldownMeansImmediateReopenIsAllowed) { ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(0, 1))); - gpufl::DeepWindow::OnLaunch(); + Launch(); ASSERT_FALSE(gpufl::DeepWindow::Active()); EXPECT_TRUE(gpufl::DeepWindow::Open(Spec(0, 1))); } @@ -215,57 +251,57 @@ TEST_F(DeepWindowTest, RequestOpenArmsOnTheNextLaunchNotImmediately) { gpufl::DeepWindow::RequestOpen(Spec(60000, 0)); EXPECT_FALSE(gpufl::DeepWindow::Active()) << "must wait for a launch"; - gpufl::DeepWindow::OnLaunch(); + Launch(); EXPECT_TRUE(gpufl::DeepWindow::Active()); } TEST_F(DeepWindowTest, RequestOpenIsConsumedOnce) { gpufl::DeepWindow::RequestOpen(Spec(60000, 0)); - gpufl::DeepWindow::OnLaunch(); + Launch(); ASSERT_TRUE(gpufl::DeepWindow::Active()); gpufl::DeepWindow::Close(gpufl::DeepWindowClose::Manual); // The request was spent on the first arm; a closed window stays closed. - gpufl::DeepWindow::OnLaunch(); + Launch(); EXPECT_FALSE(gpufl::DeepWindow::Active()); } TEST_F(DeepWindowTest, RequestOpenCarriesItsBounds) { gpufl::DeepWindow::RequestOpen(Spec(0, /*launches=*/2)); - gpufl::DeepWindow::OnLaunch(); // arms; does not consume budget + Launch(); // arms; does not consume budget ASSERT_TRUE(gpufl::DeepWindow::Active()); - gpufl::DeepWindow::OnLaunch(); + Launch(); EXPECT_TRUE(gpufl::DeepWindow::Active()); - gpufl::DeepWindow::OnLaunch(); + Launch(); EXPECT_FALSE(gpufl::DeepWindow::Active()); } TEST_F(DeepWindowTest, NewestPendingSpecWins) { gpufl::DeepWindow::RequestOpen(Spec(0, 1)); gpufl::DeepWindow::RequestOpen(Spec(60000, 0)); - gpufl::DeepWindow::OnLaunch(); + Launch(); ASSERT_TRUE(gpufl::DeepWindow::Active()); // Had the first spec won, this launch would spend its budget of 1. - gpufl::DeepWindow::OnLaunch(); + Launch(); EXPECT_TRUE(gpufl::DeepWindow::Active()); } TEST_F(DeepWindowTest, ScheduledOpenWaitsOutItsDelay) { gpufl::DeepWindow::ScheduleOpenAfter(/*delay_ms=*/50, Spec(60000, 0)); - gpufl::DeepWindow::OnLaunch(); + Launch(); EXPECT_FALSE(gpufl::DeepWindow::Active()) << "not due yet"; std::this_thread::sleep_for(std::chrono::milliseconds(70)); - gpufl::DeepWindow::OnLaunch(); + Launch(); EXPECT_TRUE(gpufl::DeepWindow::Active()); } TEST_F(DeepWindowTest, RequestOpenWhileAWindowIsOpenDoesNotDisturbIt) { ASSERT_TRUE(gpufl::DeepWindow::Open(Spec(60000, 0))); gpufl::DeepWindow::RequestOpen(Spec(0, 1)); - gpufl::DeepWindow::OnLaunch(); + Launch(); EXPECT_TRUE(gpufl::DeepWindow::Active()) << "the open window's bounds still govern"; } From 9a049cea99690ec638689c43a183d0002188e2f1 Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Sat, 25 Jul 2026 08:22:33 -0700 Subject: [PATCH 05/10] test(scripts): fix nvcc discovery and separate inconclusive results --- include/gpufl/inject/inject_entry.cpp | 39 ++++++++++++++++++++++++--- scripts/deep_window_e2e.sh | 27 +++++++++++++++---- 2 files changed, 57 insertions(+), 9 deletions(-) diff --git a/include/gpufl/inject/inject_entry.cpp b/include/gpufl/inject/inject_entry.cpp index 97faeca..5510a8f 100644 --- a/include/gpufl/inject/inject_entry.cpp +++ b/include/gpufl/inject/inject_entry.cpp @@ -516,6 +516,27 @@ void markDeferredInitFinished() { 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")); + } +#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); +} + void startDeferredInjectInit() { registerDeferredWaitAtexit(); @@ -526,10 +547,20 @@ void startDeferredInjectInit() { } std::thread([] { - // NVIDIA calls InitializeInjection from inside the CUDA driver - // injection path. CUPTI subscription/activity setup can report - // success there but later deliver no callbacks. Step out of that - // callback frame before touching CUPTI. + // 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); diff --git a/scripts/deep_window_e2e.sh b/scripts/deep_window_e2e.sh index 0b1fa5c..06c3254 100755 --- a/scripts/deep_window_e2e.sh +++ b/scripts/deep_window_e2e.sh @@ -86,9 +86,18 @@ GPUFL="$(find_bin gpufl daemon/launcher)" || { exit 2 } +# A non-login shell (ssh host 'cmd') does not source the profile that puts +# CUDA on PATH, so fall back to the usual install locations. +NVCC="" +for c in nvcc /usr/local/cuda/bin/nvcc /opt/cuda/bin/nvcc \ + /usr/local/cuda-13.3/bin/nvcc; do + command -v "$c" >/dev/null 2>&1 && { NVCC="$c"; break; } +done +[[ -n "$NVCC" ]] || { echo "nvcc not found (set PATH or install CUDA)" >&2; exit 2; } + TARGET="$OUT_DIR/deep_window_target" -echo "[e2e] compiling the gpufl-unaware target" -nvcc -O2 -lineinfo -o "$TARGET" "$SCRIPT_DIR/deep_window_target.cu" \ +echo "[e2e] compiling the gpufl-unaware target with $NVCC" +"$NVCC" -O2 -lineinfo -o "$TARGET" "$SCRIPT_DIR/deep_window_target.cu" \ >"$OUT_DIR/nvcc.log" 2>&1 || { cat "$OUT_DIR/nvcc.log"; exit 2; } { @@ -103,10 +112,18 @@ nvcc -O2 -lineinfo -o "$TARGET" "$SCRIPT_DIR/deep_window_target.cu" \ PASS=0 FAIL=0 +INCONCLUSIVE=0 note() { echo "$@" | tee -a "$REPORT"; } record() { # name, exit code - if [[ "$2" -eq 0 ]]; then PASS=$((PASS+1)); note "- **PASS** $1"; - else FAIL=$((FAIL+1)); note "- **FAIL** $1"; fi + # 3 = the checker could not decide, e.g. PC sampling stamps its batches on + # the raw CUPTI clock so window membership is unknowable from timestamps. + # That is a gap in what we can observe, not a failing window; counting it + # as FAIL would hide real regressions in the noise. + case "$2" in + 0) PASS=$((PASS+1)); note "- **PASS** $1" ;; + 3) INCONCLUSIVE=$((INCONCLUSIVE+1)); note "- **INCONCLUSIVE** $1" ;; + *) FAIL=$((FAIL+1)); note "- **FAIL** $1" ;; + esac } # ── A. embed ──────────────────────────────────────────────────────────────── @@ -219,7 +236,7 @@ fi note "## Summary" note -note "passed: $PASS, failed: $FAIL" +note "passed: $PASS, failed: $FAIL, inconclusive: $INCONCLUSIVE" echo echo "[e2e] report: $REPORT" [[ "$FAIL" -eq 0 ]] From 31e0da44422e433ce6b99fc71295959a51c33f03 Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Sat, 25 Jul 2026 08:38:39 -0700 Subject: [PATCH 06/10] test(scripts): stop counting ISA source maps as deep samples --- scripts/deep_window_check.py | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) diff --git a/scripts/deep_window_check.py b/scripts/deep_window_check.py index 8b92871..7b50d4a 100644 --- a/scripts/deep_window_check.py +++ b/scripts/deep_window_check.py @@ -32,6 +32,12 @@ ALIEN_NS = 60_000_000_000 # >60s from the window = a different clock domain WALL_FLOOR = 1_000_000_000_000_000_000 DEEP_BATCHES = ("pm_sample_batch", "profile_sample_batch") +# profile_sample_batch carries two unrelated things. sample_kind 0 and 1 are +# real timed samples (PC stall sampling, SASS metrics); 2 is isa_source_map, a +# static pc_offset -> source_file:line table emitted next to the disassembly. +# The map has no meaningful timestamp and is not something a window collects, +# so counting it made PC sampling look like it was emitting garbage clocks. +TIMED_SAMPLE_KINDS = (0, 1) def find_session(root: pathlib.Path) -> pathlib.Path: @@ -112,15 +118,27 @@ def main() -> int: start, end = w["start_ns"], w["end_ns"] inside = late = early = alien = 0 latest_late_ms = 0.0 + # Some batches fan out to every channel, so the same rows appear in + # device/scope/sass/system. Counting each file independently multiplied + # them by four. + seen_batches = set() for _, r in rows: if r.get("type") not in DEEP_BATCHES: continue + key = (r["type"], r.get("batch_id"), r.get("base_time_ns"), + len(r.get("rows", []))) + if key in seen_batches: + continue + seen_batches.add(key) base = r.get("base_time_ns", 0) cols = r.get("columns", []) if "dt_ns" not in cols: continue dt_i = cols.index("dt_ns") + kind_i = cols.index("sample_kind") if "sample_kind" in cols else None for row in r.get("rows", []): + if kind_i is not None and row[kind_i] not in TIMED_SAMPLE_KINDS: + continue ts = base + row[dt_i] if ts < start - ALIEN_NS or ts > end + ALIEN_NS: alien += 1 From 1b3dba0304319679ea27f98f42b2ccfa4a192c2d Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Sat, 25 Jul 2026 18:42:28 -0700 Subject: [PATCH 07/10] fix(pc-sampling): collect with sampling stopped, configure before allocating --- .../nvidia/engine/pc_sampling_engine.cpp | 278 ++++++++++-------- .../nvidia/engine/pc_sampling_engine.hpp | 26 +- 2 files changed, 160 insertions(+), 144 deletions(-) diff --git a/include/gpufl/backends/nvidia/engine/pc_sampling_engine.cpp b/include/gpufl/backends/nvidia/engine/pc_sampling_engine.cpp index 2f51e74..5d8cb66 100644 --- a/include/gpufl/backends/nvidia/engine/pc_sampling_engine.cpp +++ b/include/gpufl/backends/nvidia/engine/pc_sampling_engine.cpp @@ -41,6 +41,12 @@ bool IsInsufficientPrivilege(const CUptiResult res) { constexpr size_t kPcSamplingConfigAttrCount = 7; +// Per-record stall-reason slots. Sized to the API maximum rather than the +// device's actual stall-reason count so the records can be allocated before +// cuptiPCSamplingEnable - see the ordering note in EnableSamplingFeatures_. +// This is the value NVIDIA's pc_sampling sample hardcodes. +constexpr size_t kStallSlots = 128; + std::array BuildPcSamplingConfig(const uint32_t samplingPeriod, CUpti_PCSamplingData* const samplingData) { @@ -200,11 +206,16 @@ void PcSamplingEngine::start() { // still degrades to a kernel trace. Synthetic-kernel fallback stays // suppressed (cupti_backend.cpp start()), so only REAL records show. cuptiActivityEnable(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL); - // Enable/config/Start return INVALID_OPERATION when kernels run - // concurrently (verified live), and start() runs in the - // CONTEXT_CREATED callback before the app's first kernel, while the - // GPU is quiet - the only reliable moment for them. profiler-init - // already ran pre-context, so stall enumeration succeeds here. + // Arm as early as we can: start() runs in the CONTEXT_CREATED + // callback, and profiler-init already ran pre-context, so stall + // enumeration succeeds here. + // + // This used to claim Enable/config/Start need a quiet GPU because + // concurrent kernels make them return INVALID_OPERATION. That is not + // what happens - measured on driver 610.43 / CUDA 13.3, Start fails + // the instant after configuration is rejected, before the target has + // launched anything, and it fails identically whether the arm runs + // in this callback or on a worker thread hundreds of ms later. // // WindowOnly splits that: enable + configure now (they must happen // while quiet), but leave the sampler unarmed until a deep window @@ -221,7 +232,9 @@ void PcSamplingEngine::start() { } // Enable can internally disable kernel activity — re-assert it. cuptiActivityEnable(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL); - StartCycleThread_(); + // Only the experimental kernel-drain mode has anything to do on a + // cycle; the sample-only path reads once at scope end. + if (kernel_collect_ == KernelCollect::All) StartCycleThread_(); } else { LogCuptiErrorIfFailed(this->name(), "cuptiActivityEnable(PC_SAMPLING)", pcRes); @@ -242,7 +255,6 @@ void PcSamplingEngine::StartCycleThread_() { // final Windows-injected teardown path cannot safely flush activity, so // waiting a whole second here drops short sessions. last_kernel_drain_ns_.store(0, std::memory_order_relaxed); - last_sample_collect_ns_.store(0, std::memory_order_relaxed); GFL_LOG_DEBUG("[PC Sampling] launching collection cycle thread"); cycle_thread_ = std::thread([this] { GFL_LOG_DEBUG("[PC Sampling] cycle thread running"); @@ -299,11 +311,10 @@ void PcSamplingEngine::drainData() { // stop issuing CUPTI data-retrieval calls (GetData/Stop) so the cycle thread // can't fault against the dying driver while shutdown flushes + joins it. if (detail::isProcessExitTeardown()) return; - // Periodic collection on the engine's cycle thread (deferring all of it to - // session stop loses the session to process-exit teardown). Two paths by - // kernel_collect_: light = armed GetData (no Stop; KERNEL_SERIALIZED still - // returns completed kernels' samples); heavy = DrainKernelsAndCollect_, - // which Stops to force a kernel-activity flush for a full timeline. + // Only GPUFL_PC_KERNEL_COLLECT=all has cycle work. The sample-only path + // does not collect mid-run at all - see the note above + // StopAndCollectPcSampling_ - so the cycle thread is not even started for + // it, and this is the one path left here. if (pc_sampling_method_ != Method::SamplingAPI) return; if (!sampling_api_started_.load()) return; @@ -312,15 +323,9 @@ void PcSamplingEngine::drainData() { if (!ctx_.cuda_ctx) return; if (cuCtxSetCurrent(ctx_.cuda_ctx) != CUDA_SUCCESS) return; - // Drain kernel activity (heavy: stop->flush->start) only when explicitly - // requested; otherwise just collect PC samples (light). - const bool want_drain = - kernel_collect_ == KernelCollect::All && - !drain_unavailable_.load(std::memory_order_relaxed); - if (want_drain) { + if (kernel_collect_ == KernelCollect::All && + !drain_unavailable_.load(std::memory_order_relaxed)) { DrainKernelsAndCollect_(); - } else { - MaybePeriodicCollect_("cycle", /*force=*/false); } } @@ -333,12 +338,13 @@ void PcSamplingEngine::DrainKernelsAndCollect_() { if (!sampling_api_started_.load()) return; last_kernel_drain_ns_.store(now, std::memory_order_relaxed); - // Stop sampling: required because a forced activity flush returns zero - // kernel records while PC sampling is armed (driver 590+). Stop/Start - // mid-run are privileged (INSUFFICIENT_PRIVILEGES under a non-elevated - // run) — on that error, stop draining for the session and let armed - // GetData carry the PC samples. Restart after the flush; it succeeds - // even with kernels running (unlike the initial arm). + // Stop sampling: required twice over. A forced activity flush returns zero + // kernel records while PC sampling is armed (driver 590+), and GetData + // while armed discards the samples. Stop/Start mid-run are privileged + // (INSUFFICIENT_PRIVILEGES under a non-elevated run) — on that error, drop + // to the sample-only cycle, which stops too and therefore fails the same + // way and stands itself down. Restart after the flush; it succeeds even + // with kernels running (unlike the initial arm). CUpti_PCSamplingStopParams sp = {}; sp.size = sizeof(sp); sp.ctx = ctx_.cuda_ctx; @@ -409,20 +415,15 @@ void PcSamplingEngine::onScopeStart(const char* /*name*/) { } void PcSamplingEngine::onScopeStop(const char* /*name*/) { - if (opts_.deep_arm_mode == DeepArmMode::WindowOnly) { - // Disarm for real. A forced collect would leave the sampler running - // past the window, which is the whole thing WindowOnly exists to - // avoid. The ref count keeps a user scope nested inside the window - // from disarming it early. - std::lock_guard lk(sampling_lifecycle_mu_); - StopAndCollectPcSampling_(); - return; - } - // Forced collect at scope end. For the process-wide scope this is the - // last healthy moment before Windows process-exit teardown breaks - // cuptiPCSamplingStop with CUPTI_ERROR_UNKNOWN. Re-arms afterwards, so - // nested/subsequent scopes keep sampling. - MaybePeriodicCollect_("scope-stop", /*force=*/true); + // Stop and read. This is the session's collection point in both modes: + // for WindowOnly it disarms at the window edge (leaving the sampler + // running past the window is the whole thing WindowOnly exists to avoid), + // and for the process-wide scope it is the last healthy moment before + // Windows process-exit teardown breaks cuptiPCSamplingStop with + // CUPTI_ERROR_UNKNOWN. The ref count keeps a nested scope from collecting + // early; a later scope re-arms through onScopeStart. + std::lock_guard lk(sampling_lifecycle_mu_); + StopAndCollectPcSampling_(); } // ---- Private helpers ------------------------------------------------------- @@ -440,6 +441,47 @@ bool PcSamplingEngine::EnableSamplingFeatures_() { return false; } + // Allocate the sample buffers BEFORE enabling, and keep Enable and + // SetConfigurationAttribute adjacent. + // + // This ordering is load-bearing, not style. Doing this allocation between + // Enable and configure is what broke PC sampling under injection: + // configure (and then every other PC-sampling call on the context) came + // back CUPTI_ERROR_INVALID_OPERATION. Proven on driver 610.43 / CUDA 13.3 + // by injecting nothing but this allocation into an otherwise-working + // sequence in the same process, on the same context - Enable returned + // SUCCESS, the allocation ran, configure returned INVALID_OPERATION. It is + // the allocation, not elapsed time: a 50 ms sleep in the same slot is + // harmless. NVIDIA's own pc_sampling sample allocates up front too. + // + // Stall-reason enumeration is likewise deferred to after configure; it + // only feeds the reason-name map. The records are sized with the API's + // maximum stall-reason count rather than the device's actual count, which + // is what the sample does and what CollectPcSamplingData_ resets to. + if (!pc_sampling_buffers_) { + constexpr size_t kMaxPcs = 65536; + pc_sampling_buffers_ = + std::unique_ptr( + new PCSamplingBuffers()); + pc_sampling_buffers_->pcRecords = static_cast( + std::calloc(kMaxPcs, sizeof(CUpti_PCSamplingPCData))); + for (size_t i = 0; i < kMaxPcs; ++i) { + pc_sampling_buffers_->pcRecords[i].size = + sizeof(CUpti_PCSamplingPCData); + pc_sampling_buffers_->pcRecords[i].stallReasonCount = kStallSlots; + pc_sampling_buffers_->pcRecords[i].stallReason = + static_cast(std::calloc( + kStallSlots, sizeof(CUpti_PCSamplingStallReason))); + } + pc_sampling_buffers_->data = static_cast( + std::calloc(1, sizeof(CUpti_PCSamplingData))); + pc_sampling_buffers_->data->size = sizeof(CUpti_PCSamplingData); + pc_sampling_buffers_->data->collectNumPcs = kMaxPcs; + pc_sampling_buffers_->data->pPcData = pc_sampling_buffers_->pcRecords; + pc_sampling_buffers_->data->totalNumPcs = 0; + num_stall_reasons_ = kStallSlots; + } + CUpti_PCSamplingEnableParams enableParams = {}; enableParams.size = sizeof(CUpti_PCSamplingEnableParams); enableParams.ctx = ctx_.cuda_ctx; @@ -465,14 +507,51 @@ bool PcSamplingEngine::EnableSamplingFeatures_() { "conflicting with Profiler API; continuing."); } - if (!pc_sampling_buffers_) { - constexpr size_t kMaxPcs = 65536; - pc_sampling_buffers_ = - std::unique_ptr( - new PCSamplingBuffers()); - pc_sampling_buffers_->pcRecords = static_cast( - std::calloc(kMaxPcs, sizeof(CUpti_PCSamplingPCData))); + auto configInfo = BuildPcSamplingConfig(opts_.pc_sampling_period, + pc_sampling_buffers_->data); + CUpti_PCSamplingConfigurationInfoParams configParams = {}; + configParams.size = CUpti_PCSamplingConfigurationInfoParamsSize; + configParams.ctx = ctx_.cuda_ctx; + configParams.numAttributes = + configInfo.size(); + configParams.pPCSamplingConfigurationInfo = configInfo.data(); + + const CUptiResult configRes = + cuptiPCSamplingSetConfigurationAttribute(&configParams); + // Any failure here is fatal to the session, INVALID_OPERATION included. + // This used to be swallowed as benign, which was the single most + // misleading thing in this file: when configuration is rejected, + // ENABLE_START_STOP_CONTROL never applies, so cuptiPCSamplingStart then + // fails with INVALID_OPERATION too - and the log still said "configured + // and enabled successfully", pointing every investigation at Start. + if (configRes != CUPTI_SUCCESS) { + LogCuptiErrorIfFailed(this->name(), + "cuptiPCSamplingSetConfigurationAttribute", + configRes); + for (size_t i = 0; i < configInfo.size(); ++i) { + GFL_LOG_ERROR("[PC Sampling] rejected attribute type=", + static_cast(configInfo[i].attributeType), + " status=", + static_cast(configInfo[i].attributeStatus)); + } + if (IsInsufficientPrivilege(configRes)) { + sampling_api_blocked_.store(true); + GFL_LOG_ERROR( + "[PC Sampling] Insufficient privileges: disabling PC " + "sampling for this session."); + } + pc_sampling_method_ = Method::None; + CUpti_PCSamplingDisableParams dp = {}; + dp.size = sizeof(CUpti_PCSamplingDisableParams); + dp.ctx = ctx_.cuda_ctx; + cuptiPCSamplingDisable(&dp); + return false; + } + // Stall-reason enumeration, after configuration: it only builds the + // reason-name map, and keeping it out of the Enable->configure window is + // the point (see the ordering note above). + { CUpti_PCSamplingGetNumStallReasonsParams numParams = {}; numParams.size = sizeof(CUpti_PCSamplingGetNumStallReasonsParams); numParams.ctx = ctx_.cuda_ctx; @@ -500,9 +579,11 @@ bool PcSamplingEngine::EnableSamplingFeatures_() { dp.size = sizeof(CUpti_PCSamplingDisableParams); dp.ctx = ctx_.cuda_ctx; cuptiPCSamplingDisable(&dp); - pc_sampling_buffers_.reset(); return false; } + // num_stall_reasons_ stays at the allocated slot count, not this + // device count: CollectPcSamplingData_ uses it to restore each + // record's writable capacity before every GetData. { auto* stallIndices = static_cast( malloc(numStallReasons * sizeof(uint32_t))); @@ -540,50 +621,6 @@ bool PcSamplingEngine::EnableSamplingFeatures_() { free(stallIndices); free(stallReasonNames); } - - for (size_t i = 0; i < kMaxPcs; ++i) { - pc_sampling_buffers_->pcRecords[i].size = - sizeof(CUpti_PCSamplingPCData); - pc_sampling_buffers_->pcRecords[i].stallReasonCount = - numStallReasons; - pc_sampling_buffers_->pcRecords[i].stallReason = - static_cast(std::calloc( - numStallReasons, sizeof(CUpti_PCSamplingStallReason))); - } - pc_sampling_buffers_->data = static_cast( - std::calloc(1, sizeof(CUpti_PCSamplingData))); - pc_sampling_buffers_->data->size = sizeof(CUpti_PCSamplingData); - pc_sampling_buffers_->data->collectNumPcs = kMaxPcs; - pc_sampling_buffers_->data->pPcData = pc_sampling_buffers_->pcRecords; - pc_sampling_buffers_->data->totalNumPcs = 0; - num_stall_reasons_ = numStallReasons; - } - - auto configInfo = BuildPcSamplingConfig(opts_.pc_sampling_period, - pc_sampling_buffers_->data); - - CUpti_PCSamplingConfigurationInfoParams configParams = {}; - configParams.size = CUpti_PCSamplingConfigurationInfoParamsSize; - configParams.ctx = ctx_.cuda_ctx; - configParams.numAttributes = - static_cast(configInfo.size()); - configParams.pPCSamplingConfigurationInfo = configInfo.data(); - - const CUptiResult configRes = - cuptiPCSamplingSetConfigurationAttribute(&configParams); - if (configRes != CUPTI_SUCCESS && - configRes != CUPTI_ERROR_INVALID_OPERATION) { - LogCuptiErrorIfFailed(this->name(), - "cuptiPCSamplingSetConfigurationAttribute", - configRes); - if (IsInsufficientPrivilege(configRes)) { - sampling_api_blocked_.store(true); - pc_sampling_method_ = Method::None; - GFL_LOG_ERROR( - "[PC Sampling] Insufficient privileges: disabling PC " - "sampling for this session."); - } - return false; } sampling_api_ready_.store(true); @@ -643,43 +680,36 @@ void PcSamplingEngine::StartPcSampling_() { } void PcSamplingEngine::flushBeforeCudaTeardown(const char* reason) { - MaybePeriodicCollect_(reason, /*force=*/false); + // Reached from a CUDA cleanup CUPTI callback, where cuptiPCSamplingStop + // returns 999 - and without a Stop there is no way to read samples + // without destroying them. The engine's cycle thread owns collection. + GFL_LOG_DEBUG( + "[PC Sampling] skipping collect from CUDA cleanup callback: ", + reason ? reason : "unknown"); } void PcSamplingEngine::onLaunchTick() { - // Collect from the launch API_ENTER callback (app thread) — the only - // beat that reliably fires on Windows-injected runs. Deferring to - // session stop loses everything to the process-exit 999. - MaybePeriodicCollect_("launch-tick", /*force=*/false); + // Deliberately does not collect. This runs on the app thread inside the + // launch API_ENTER callback, where Stop is unavailable, and the only + // callback-safe alternative - an armed GetData - silently discards the + // session's samples. The cycle thread does the stop/collect/restart. } -void PcSamplingEngine::MaybePeriodicCollect_(const char* reason, - const bool force) { - if (pc_sampling_method_ != Method::SamplingAPI) return; - if (!sampling_api_started_.load()) return; - - if (!force) { - const int64_t now = detail::GetTimestampNs(); - const int64_t last = - last_sample_collect_ns_.load(std::memory_order_relaxed); - if (last != 0 && now - last < kCollectIntervalNs) return; - } - - // try_lock: callers are inside CUPTI callbacks - never wait on a lock - // the stop/shutdown path holds while it makes CUPTI calls. - if (!sampling_lifecycle_mu_.try_lock()) return; - std::lock_guard lk(sampling_lifecycle_mu_, std::adopt_lock); - if (!sampling_api_started_.load()) return; - last_sample_collect_ns_.store(detail::GetTimestampNs(), - std::memory_order_relaxed); - - GFL_LOG_DEBUG("[PC Sampling] periodic collect (", reason ? reason : "?", - force ? ", forced" : "", ")"); - // Armed GetData, no Stop: Stop returns 999 inside a CUPTI callback, and - // in KERNEL_SERIALIZED mode GetData mid-session returns every completed - // kernel's samples. Sampling stays armed; no re-arm needed. - CollectPcSamplingData_(); -} +// Sample-only sessions have no mid-run collect. Both ways of reading samples +// while the session is live are unusable, measured on driver 610.43 / CUDA +// 13.3 with a 400-launch workload that normally finishes in seconds: +// +// armed GetData - returns nothing and discards the buffer. A run +// that collected 24.5M samples when left alone +// collected 0. +// stop -> GetData -> start every second +// - also returns 0, and cripples the target: the run +// had not finished after 240 s (0.2% CPU). +// +// So the sampler is armed once and read once, with sampling stopped, at scope +// end (onScopeStop) or session teardown (stop/shutdown). The cost is that a +// long process-scope run can overflow CUPTI's scratch buffer before that read; +// droppedSamples in the collect summary makes it visible when it happens. void PcSamplingEngine::StopAndCollectPcSampling_(const bool sync_device) { GFL_LOG_DEBUG("[PC Sampling] StopAndCollect entry: method=", diff --git a/include/gpufl/backends/nvidia/engine/pc_sampling_engine.hpp b/include/gpufl/backends/nvidia/engine/pc_sampling_engine.hpp index 66b2a0b..c3eb214 100644 --- a/include/gpufl/backends/nvidia/engine/pc_sampling_engine.hpp +++ b/include/gpufl/backends/nvidia/engine/pc_sampling_engine.hpp @@ -102,16 +102,11 @@ class PcSamplingEngine final : public IProfilingEngine { /// completed by the time the next API callback runs anyway. void StopAndCollectPcSampling_(bool sync_device = true); /// The cuptiPCSamplingGetData drain loop: parses PC records into - /// PC_SAMPLE activity records. Callable while sampling is still armed - /// (the NVIDIA pc_sampling sample's serialized-mode pattern - only - /// completed kernels' samples are returned) or after a Stop. + /// PC_SAMPLE activity records. MUST be called with sampling stopped. + /// Calling it while armed does not return the samples AND discards them + /// (verified on driver 610.43 / CUDA 13.3: a session that collected + /// 24.5M samples with stopped-GetData collected 0 with armed-GetData). void CollectPcSamplingData_(); - /// Shared mid-session collect: throttled armed-GetData, safe to call - /// from CUPTI callbacks (try_lock, no cudart, no PCSamplingStop - - /// Stop returns 999 inside CUPTI callbacks). `force` bypasses the - /// interval throttle (used at process-scope end - the last healthy - /// moment before Windows process-exit teardown). - void MaybePeriodicCollect_(const char* reason, bool force); MonitorOptions opts_; EngineContext ctx_; @@ -128,7 +123,7 @@ class PcSamplingEngine final : public IProfilingEngine { // Kernel-timeline collection mode, parsed once from GPUFL_PC_KERNEL_COLLECT // in initialize(). drain_unavailable_ latches when a mid-run stop/flush // returns INSUFFICIENT_PRIVILEGES so we stop retrying and degrade to - // armed-GetData (sample-only) collection. + // sample-only collection. KernelCollect kernel_collect_ = KernelCollect::None; std::atomic drain_unavailable_{false}; @@ -137,17 +132,8 @@ class PcSamplingEngine final : public IProfilingEngine { // session stop/shutdown. Without it a cycle's Stop could interleave // with a scope-begin Start. std::mutex sampling_lifecycle_mu_; - // Minimum gap between periodic collects. Short GPU phases (a script - // whose kernels all finish within seconds of context creation) must - // still get at least one mid-run collect before exit teardown breaks - // cuptiPCSamplingStop, so this errs small; each collect is one - // stop→GetData→start cycle (~sub-ms) on the app thread. + // Minimum gap between kernel-activity drains (GPUFL_PC_KERNEL_COLLECT=all). static constexpr int64_t kCollectIntervalNs = 1'000'000'000; // 1 s - // Last sample-only GetData collect, wall ns. This is intentionally - // separate from last_kernel_drain_ns_: launch callbacks may sample often, - // but they must not starve the plain-thread stop/flush/start drain that - // pulls kernel activity records. - std::atomic last_sample_collect_ns_{0}; // Last kernel activity drain (stop -> flush -> start), wall ns. std::atomic last_kernel_drain_ns_{0}; From 74cf5183e7bdb5b64c8396fda9a5a1a07ce2d93d Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Sat, 25 Jul 2026 19:14:42 -0700 Subject: [PATCH 08/10] fix(pc-sampling): point 0-sample and deep-window guidance at kernel-launch count --- daemon/launcher/cli_parse.cpp | 41 +++++++++++++++---- .../nvidia/capture_capability_resolver.cpp | 24 +++++++++-- .../nvidia/engine/pc_sampling_engine.cpp | 18 +++++--- 3 files changed, 65 insertions(+), 18 deletions(-) diff --git a/daemon/launcher/cli_parse.cpp b/daemon/launcher/cli_parse.cpp index 37bf6c9..e2d7537 100644 --- a/daemon/launcher/cli_parse.cpp +++ b/daemon/launcher/cli_parse.cpp @@ -172,12 +172,18 @@ const char* traceHelp() { " then disarm. Unlike --window the target keeps\n" " running. Needs a bound below. Default: 0 (arm\n" " at the first kernel launch).\n" - " --deep-for= How long the deep window stays armed.\n" + " --deep-for= How long the deep window stays armed. Note this\n" + " bounds TIME, which does not bound how much the\n" + " engines actually collect - see --deep-launches.\n" " --deep-launches= Kernel-launch bound on the deep window; ends it\n" - " at whichever bound is hit first. PREFER THIS for\n" - " SASS / Range: replay re-runs every kernel, so a\n" + " at whichever bound is hit first. PREFER THIS.\n" + " For SASS / Range replay re-runs every kernel, so a\n" " second of wall time covers ~25x less work there\n" - " than under PM sampling.\n" + " than under PM sampling. For PcSampling it is what\n" + " decides whether you get data at all: samples only\n" + " become readable after a few thousand launches, so\n" + " a short --deep-for window (or any window over\n" + " slow kernels) can collect nothing.\n" " --deep-cooldown=\n" " Quiet time before another window may open.\n" " --pc-sample-period=\n" @@ -393,6 +399,24 @@ TraceParseResult parseTraceArgs(const std::vector& argv) { " (expected a duration like 30s, 500ms, 5m, 1h, " "or a bare number of seconds)"}; } + // A bare number is seconds, which collides with --deep-launches: + // `--deep-for=2000` meaning "2000 launches" silently becomes a + // 33-minute window, and the no-bound check below can't catch it + // because a bound *was* given. Reject the unit-less form once it + // is too large to plausibly be a duration someone typed on + // purpose - naming the alternative, since that is the mistake. + const bool unitless = + !v.empty() && + v.find_first_not_of("0123456789") == std::string::npos; + if (key == "--deep-for" && unitless && ms >= 600'000) { + return {std::nullopt, + "--deep-for=" + v + " means " + std::to_string(ms / 1000) + + " SECONDS (a bare number is seconds), which is almost " + "certainly not what you meant. Add a unit (e.g. " + v + + "s, 5m) if you really want that long a window, or use " + "--deep-launches " + v + " to bound it by kernel " + "launches instead"}; + } if (key == "--deep-after") out.deep_after_ms = ms; else if (key == "--deep-for") out.deep_for_ms = ms; else out.deep_cooldown_ms = ms; @@ -443,10 +467,11 @@ TraceParseResult parseTraceArgs(const std::vector& argv) { // just "profile deeply for the whole run" with extra steps. if (out.deep_requested && out.deep_for_ms == 0 && out.deep_launches == 0) { return {std::nullopt, - "a deep window needs a bound: pass --deep-for " - "or --deep-launches (prefer --deep-launches for the " - "replay engines, where a second of wall time covers far " - "less work)"}; + "a deep window needs a bound: pass --deep-launches or " + "--deep-for (prefer --deep-launches: it is what " + "the engines actually scale with. The replay engines cover " + "far less work per second of wall time, and PC sampling " + "returns nothing at all below a few thousand launches)"}; } return {out, ""}; } diff --git a/include/gpufl/backends/nvidia/capture_capability_resolver.cpp b/include/gpufl/backends/nvidia/capture_capability_resolver.cpp index 6fec581..7aad0a8 100644 --- a/include/gpufl/backends/nvidia/capture_capability_resolver.cpp +++ b/include/gpufl/backends/nvidia/capture_capability_resolver.cpp @@ -22,11 +22,27 @@ std::vector BuildCaptureCapabilityWarnings( std::vector warnings; if (input.requests.pc && input.engine_state.pc.active && !input.engine_state.pc.has_data) { + // What decides this is the number of kernel launches sampled, not + // wall time: in KERNEL_SERIALIZED mode nothing is readable until + // enough kernel ranges have accumulated. Measured on an RTX 5060 / + // CUDA 13.3, same 8 s window: 876 launches returned 0 samples, 2002 + // returned 87.8M. So the old advice here was actively wrong - a + // "heavier" workload means fewer launches, and a finer + // --pc-sample-period costs enough overhead to cut the launches + // covered (992 -> 58 in a fixed window), moving further from the + // threshold in both cases. + // No count in the text: launch_count is process-wide, while what + // matters is how many launches happened while sampling was armed - + // for a deep window those differ by a lot, and quoting the wrong one + // contradicts the advice. warnings.push_back( - "[gpufl] PC sampling collected 0 stall samples - the profiled " - "workload was too short for the sampling interval. Run a " - "longer/heavier workload, or sample more frequently with " - "`gpufl trace --pc-sample-period ` (lower N)."); + "[gpufl] PC sampling collected 0 stall samples - too few kernel " + "launches were sampled. PC sampling accumulates per kernel and " + "needs a few thousand launches before any samples are readable. " + "Sample more launches - note that a longer wall-clock window is " + "not the same thing: a slow-kernel workload can run for seconds " + "and still cover too few. With a deep window, bound it with " + "`--deep-launches ` rather than `--deep-for `."); } if (input.requests.sass && input.engine_state.sass.active && !input.engine_state.sass.has_data) { diff --git a/include/gpufl/backends/nvidia/engine/pc_sampling_engine.cpp b/include/gpufl/backends/nvidia/engine/pc_sampling_engine.cpp index 5d8cb66..417cb20 100644 --- a/include/gpufl/backends/nvidia/engine/pc_sampling_engine.cpp +++ b/include/gpufl/backends/nvidia/engine/pc_sampling_engine.cpp @@ -59,7 +59,10 @@ BuildPcSamplingConfig(const uint32_t samplingPeriod, }; // Kernel-serialized collection plus explicit start/stop lets GPUFL own - // the PC sampling lifetime while avoiding mid-session GetData drains. + // the PC sampling lifetime. Serialized mode accumulates per kernel range + // and nothing is readable until enough ranges pile up, which is why a + // session's yield tracks kernel-launch count, not wall time (measured: + // 876 launches over 8 s = 0 samples, 2002 over the same 8 s = 87.8M). { CUpti_PCSamplingConfigurationInfo info = {}; info.attributeType = @@ -80,9 +83,11 @@ BuildPcSamplingConfig(const uint32_t samplingPeriod, info.attributeType = CUPTI_PC_SAMPLING_CONFIGURATION_ATTR_TYPE_SCRATCH_BUFFER_SIZE; // Host-resident staging between HW buffer and GetData. CUPTI sizing: - // ~1 MB per ~5,500 PCs with all stall reasons, so 32 MB covers - // ~175k distinct PCs per drain window - generous at our 1 s collect - // cadence (the old 256 MB was wildly oversized per context). + // ~1 MB per ~5,500 PCs with all stall reasons, so 32 MB covers ~175k + // distinct PCs - ample for the single end-of-scope read (the old + // 256 MB was wildly oversized per context). Sizing this up does not + // rescue a session that collected nothing: 256 MB was measured to + // make no difference. info.attributeData.scratchBufferSizeData.scratchBufferSize = 32 * 1024 * 1024; addConfig(info); @@ -97,8 +102,9 @@ BuildPcSamplingConfig(const uint32_t samplingPeriod, } // Explicit start/stop is required before cuptiPCSamplingStart/Stop. - // Do not call cuptiPCSamplingGetData while sampling is active; on CUDA - // 13.x this can drain the buffer and leave the final collection empty. + // Never call cuptiPCSamplingGetData while sampling is armed: it returns + // nothing AND discards what was buffered, so the final read comes back + // empty. See the collection note above StopAndCollectPcSampling_. { CUpti_PCSamplingConfigurationInfo info = {}; info.attributeType = From c8e3e04768e7672048f675d7514c39674f4816fe Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Sat, 25 Jul 2026 21:46:36 -0700 Subject: [PATCH 09/10] fix(deep-window): report the engines a window armed, not one label --- .../gpufl/backends/nvidia/cupti_backend.cpp | 18 ++++++++++ .../gpufl/backends/nvidia/cupti_backend.hpp | 12 ++++++- include/gpufl/core/deep_window.cpp | 16 +++++++-- include/gpufl/core/events.hpp | 6 +++- .../gpufl/core/model/deep_window_model.cpp | 10 +++++- include/gpufl/core/monitor.cpp | 21 ++++++++++- include/gpufl/core/monitor.hpp | 20 ++++++++++- include/gpufl/core/monitor_backend.hpp | 35 ++++++++++++++++++- tests/core/test_deep_window.cpp | 28 +++++++++++++-- 9 files changed, 156 insertions(+), 10 deletions(-) diff --git a/include/gpufl/backends/nvidia/cupti_backend.cpp b/include/gpufl/backends/nvidia/cupti_backend.cpp index fe80230..07a3212 100644 --- a/include/gpufl/backends/nvidia/cupti_backend.cpp +++ b/include/gpufl/backends/nvidia/cupti_backend.cpp @@ -762,6 +762,24 @@ void CuptiBackend::DrainProfilingData() { if (rebound) cuCtxSetCurrent(prev); } +std::vector CuptiBackend::ArmedEngineWireNames_() const { + // Reuses the inspector the capability rows are built from, so a combo, a + // Deep run and a single engine all narrow the same way: each path reports + // what it actually took, not what was asked for. Deep in particular asks + // for SASS + PC and settles for whichever armed, and those differ by ~25x + // in launches covered per second of window. + const EngineRuntimeState state = InspectEngineRuntimeState( + engine_.get(), opts_.profiling_engine, !combo_.empty()); + + std::vector out; + if (state.sass.active) out.emplace_back("nvidia.sass_metrics"); + if (state.pc.active) out.emplace_back("nvidia.pc_sampling"); + if (state.pm.active) out.emplace_back("nvidia.pm_sampling"); + if (state.range.active) out.emplace_back("nvidia.range_profiler"); + if (state.range_kernel.active) out.emplace_back("nvidia.range_profiler_kernel_replay"); + return out; +} + void CuptiBackend::ServiceDeepWindow() { if (!initialized_ || !active_.load(std::memory_order_relaxed)) return; // Lock-free gate: skip the context work when there is nothing to do. diff --git a/include/gpufl/backends/nvidia/cupti_backend.hpp b/include/gpufl/backends/nvidia/cupti_backend.hpp index 2331984..8441451 100644 --- a/include/gpufl/backends/nvidia/cupti_backend.hpp +++ b/include/gpufl/backends/nvidia/cupti_backend.hpp @@ -205,7 +205,9 @@ class CuptiBackend : public IMonitorBackend { void OnDeepWindowStart(const char* name) override { if (engine_) engine_->onScopeStart(name); } - void OnDeepWindowStop(const char* name) override { + std::vector OnDeepWindowStop(const char* name) override { + // Sample BEFORE the disarm below - afterwards nothing reads as armed. + std::vector armed = ArmedEngineWireNames_(); if (engine_) engine_->onScopeStop(name); // cuptiActivityFlushAll(1) permanently kills the CUPTI subscriber // callback when the SamplingAPI is armed (enableStartStopControl=0, @@ -216,6 +218,7 @@ class CuptiBackend : public IMonitorBackend { opts_.profiling_engine == ProfilingEngine::Deep) { cudaDeviceSynchronize(); } + return armed; } void OnPerfScopeStart(const char* name) override { if (!ScopeArmsEngines_()) return; @@ -237,6 +240,13 @@ class CuptiBackend : public IMonitorBackend { } private: + // Wire names of the deep engines armed right now, in a stable order. + // Reports what each engine actually took rather than what was requested, + // so a Deep run whose SASS declined lists only PC sampling. Trace is + // absent by construction: it is not scope-gated, so a window never arms + // or disarms it. + std::vector ArmedEngineWireNames_() const; + bool ShouldEnableNvtxMarkerActivityBeforeEngine_() const; bool ShouldEnableNvtxMarkerActivityForSelectedEngine_() const; static void EnableNvtxMarkerActivity_(const char* phase); diff --git a/include/gpufl/core/deep_window.cpp b/include/gpufl/core/deep_window.cpp index e844dd6..cb7c6f5 100644 --- a/include/gpufl/core/deep_window.cpp +++ b/include/gpufl/core/deep_window.cpp @@ -235,7 +235,6 @@ void DeepWindow::Close(const DeepWindowClose reason) { ev.pid = detail::GetPid(); ev.name = g_name; ev.close_reason = DeepWindowCloseName(reason); - ev.engine = ProfilingEngineWireName(g_opts.profiling_engine); ev.start_ns = start_ns; ev.end_ns = end_ns; ev.duration_ns = end_ns - start_ns; @@ -249,9 +248,22 @@ void DeepWindow::Close(const DeepWindowClose reason) { // it - the engines' own exit handling flushes there instead. The // event below is still written so the window is on the record. if (!detail::isProcessExitTeardown()) { - Monitor::EndDeepWindowScope(name.c_str()); + // The disarm hands back what WAS armed - an audit record of this + // window, distinct from the session-end verdict in + // capture_capabilities, which the two can legitimately disagree + // with when an engine fell back or got blocked mid-run. + ev.engines = Monitor::EndDeepWindowScope(name.c_str()); detail::EndPerfScopeIfEnabled(name.c_str(), ev.pid, start_ns, end_ns, /*is_deep_window=*/true); + } else { + // Teardown skipped the disarm, so nothing observed the armed set. + // Name the resolved request instead - less trustworthy than a real + // reading, but close_reason marks the row as a teardown close. + // Not folded into an empty-list check: an empty list from a real + // disarm means this window armed nothing, and overwriting that + // would erase the one record of it. + ev.engines = { + ProfilingEngineWireName(Monitor::ResolvedProfilingEngine())}; } } diff --git a/include/gpufl/core/events.hpp b/include/gpufl/core/events.hpp index cb48af8..922d393 100644 --- a/include/gpufl/core/events.hpp +++ b/include/gpufl/core/events.hpp @@ -655,7 +655,11 @@ struct DeepWindowEvent { std::string session_id; std::string name; std::string close_reason; // DeepWindowCloseName(): "deadline" | ... - std::string engine; // ProfilingEngineWireName of the armed engine + // Wire names of the deep engines this window actually armed. Empty means + // the window opened but armed nothing, which is a real outcome worth + // recording. Trace never appears here: it runs session-wide rather than + // arming with the window. + std::vector engines; int64_t start_ns = 0; int64_t end_ns = 0; int64_t duration_ns = 0; diff --git a/include/gpufl/core/model/deep_window_model.cpp b/include/gpufl/core/model/deep_window_model.cpp index a3d219e..28efc9b 100644 --- a/include/gpufl/core/model/deep_window_model.cpp +++ b/include/gpufl/core/model/deep_window_model.cpp @@ -7,6 +7,14 @@ namespace gpufl::model { std::string DeepWindowModel::buildJson() const { + std::ostringstream engines; + engines << '['; + for (size_t i = 0; i < e_.engines.size(); ++i) { + if (i) engines << ','; + engines << '"' << jsonEscape(e_.engines[i]) << '"'; + } + engines << ']'; + std::ostringstream oss; oss << "{\"type\":\"deep_window_event\"" << ",\"pid\":" << e_.pid @@ -14,7 +22,7 @@ std::string DeepWindowModel::buildJson() const { << ",\"session_id\":\"" << jsonEscape(e_.session_id) << "\"" << ",\"name\":\"" << jsonEscape(e_.name) << "\"" << ",\"close_reason\":\"" << jsonEscape(e_.close_reason) << "\"" - << ",\"engine\":\"" << jsonEscape(e_.engine) << "\"" + << ",\"engines\":" << engines.str() << ",\"start_ns\":" << e_.start_ns << ",\"end_ns\":" << e_.end_ns << ",\"duration_ns\":" << e_.duration_ns diff --git a/include/gpufl/core/monitor.cpp b/include/gpufl/core/monitor.cpp index 583f8dc..d40b55c 100644 --- a/include/gpufl/core/monitor.cpp +++ b/include/gpufl/core/monitor.cpp @@ -36,6 +36,11 @@ namespace gpufl { // Global ring buffer (as declared in monitor.hpp) RingBuffer g_monitorBuffer; +namespace { +// Set once by Monitor::Initialize; see Monitor::ResolvedProfilingEngine. +std::atomic g_resolvedProfilingEngine{ProfilingEngine::Monitor}; +} // namespace + namespace { /** @@ -486,6 +491,13 @@ void CollectorLoop() { void Monitor::Initialize(const MonitorOptions& opts) { if (g_state.initialized.exchange(true)) return; + // The engine AFTER env overrides. InitOptions::profiling_engine is the + // pre-override request, and `gpufl trace --passes X` overrides only the + // MonitorOptions copy - so anything reporting which engine ran must read + // it from here, not from g_opts. + g_resolvedProfilingEngine.store(opts.profiling_engine, + std::memory_order_release); + g_monitorBuffer.resetDroppedCount(); g_state.batches.reset(); g_state.metadata.reset(); @@ -623,8 +635,15 @@ void Monitor::RecordStop(void* handle, StreamHandle) { void Monitor::BeginProfilerScope(const char* name) { if (auto* b = GetBackend()) b->OnScopeStart(name); } void Monitor::EndProfilerScope(const char* name) { if (auto* b = GetBackend()) b->OnScopeStop(name); } +ProfilingEngine Monitor::ResolvedProfilingEngine() { + return g_resolvedProfilingEngine.load(std::memory_order_acquire); +} + void Monitor::BeginDeepWindowScope(const char* name) { if (auto* b = GetBackend()) b->OnDeepWindowStart(name); } -void Monitor::EndDeepWindowScope(const char* name) { if (auto* b = GetBackend()) b->OnDeepWindowStop(name); } +std::vector Monitor::EndDeepWindowScope(const char* name) { + if (auto* b = GetBackend()) return b->OnDeepWindowStop(name); + return {}; +} void Monitor::BeginDeepWindowPerfScope(const char* name) { if (auto* b = GetBackend()) b->OnDeepWindowPerfStart(name); } void Monitor::EndDeepWindowPerfScope(const char* name) { if (auto* b = GetBackend()) b->OnDeepWindowPerfStop(name); } void Monitor::BeginPerfScope(const char* name) { if (auto* b = GetBackend()) b->OnPerfScopeStart(name); } diff --git a/include/gpufl/core/monitor.hpp b/include/gpufl/core/monitor.hpp index e671630..df13870 100644 --- a/include/gpufl/core/monitor.hpp +++ b/include/gpufl/core/monitor.hpp @@ -326,8 +326,26 @@ class Monitor { * Same engines, but routed separately so a WindowOnly backend can arm * for a window without arming for every user scope. */ + /** + * @brief The profiling engine actually in use, after env overrides. + * + * InitOptions::profiling_engine is only the request: + * `gpufl trace --passes X` sets GPUFL_PROFILING_ENGINE, which + * gpufl::init() applies to its MonitorOptions copy and NOT to g_opts. + * Anything that REPORTS which engine ran has to read it here, or an + * injected run mislabels itself as whatever the injection preset asked + * for. Monitor (the default) until Initialize() has run. + */ + static ProfilingEngine ResolvedProfilingEngine(); + static void BeginDeepWindowScope(const char* name); - static void EndDeepWindowScope(const char* name); + /** + * @brief Disarms, and returns the wire names of the engines that were + * armed for the window (empty when no backend, or when none armed). See + * IMonitorBackend::OnDeepWindowStop for why the names come back from the + * disarm rather than from a separate query. + */ + static std::vector EndDeepWindowScope(const char* name); static void BeginDeepWindowPerfScope(const char* name); static void EndDeepWindowPerfScope(const char* name); diff --git a/include/gpufl/core/monitor_backend.hpp b/include/gpufl/core/monitor_backend.hpp index baf4944..e44a8d2 100644 --- a/include/gpufl/core/monitor_backend.hpp +++ b/include/gpufl/core/monitor_backend.hpp @@ -1,6 +1,8 @@ #pragma once #include +#include +#include #include "gpufl/core/events.hpp" #include "gpufl/core/monitor.hpp" @@ -100,7 +102,38 @@ class IMonitorBackend { * which is right for every backend that arms unconditionally. */ virtual void OnDeepWindowStart(const char* name) { OnScopeStart(name); } - virtual void OnDeepWindowStop(const char* name) { OnScopeStop(name); } + + /** + * @brief Disarm, and report the wire names of the engines that WERE armed. + * + * An empty list means this window armed nothing, which is a real outcome + * and not an error signal. Trace is never listed: it collects for the + * whole session rather than arming with the window, so naming it here + * would credit the window with data it did not gate. + * + * The names are read here rather than through a separate query for two + * reasons. They have to be sampled BEFORE the disarm - afterwards nothing + * is armed and the answer is always empty - and folding it into the call + * that does the disarming makes that ordering impossible to get wrong. + * + * Deliberately a point-in-time reading, NOT the session's verdict. + * capture_capabilities.selected_engine answers "what did this session end + * up being" and can only be computed at session end, since it depends on + * what each engine finally produced. A window needs the other question: + * "what was armed while THIS window was open". The two legitimately + * differ - a Deep run whose SASS declines its first arm falls back to PC + * sampling, and a window that closed before an engine got blocked saw a + * different world than the session summary reports. Keeping both is what + * makes the window row an audit record rather than a duplicate. + * + * This is also what explains a window's launch coverage: a set containing + * SASS or a replaying Range profiler covers ~25x fewer launches per second + * than one holding only PC or PM sampling. + */ + virtual std::vector OnDeepWindowStop(const char* name) { + OnScopeStop(name); + return {}; + } /** @brief Periodically drain buffered profiling data. Thread-safe. */ virtual void DrainProfilingData() {} diff --git a/tests/core/test_deep_window.cpp b/tests/core/test_deep_window.cpp index 4afbe58..9311c30 100644 --- a/tests/core/test_deep_window.cpp +++ b/tests/core/test_deep_window.cpp @@ -342,7 +342,7 @@ TEST(DeepWindowModelTest, SerializesRequestedBoundsAlongsideTheOutcome) { e.session_id = "sess-1"; e.name = "deep_window"; e.close_reason = "deadline"; - e.engine = "nvidia.pc_sampling"; + e.engines = {"nvidia.sass_metrics", "nvidia.pc_sampling"}; e.start_ns = 1000; e.end_ns = 3000; e.duration_ns = 2000; @@ -353,9 +353,33 @@ TEST(DeepWindowModelTest, SerializesRequestedBoundsAlongsideTheOutcome) { const std::string json = gpufl::model::DeepWindowModel(e).buildJson(); EXPECT_NE(json.find("\"type\":\"deep_window_event\""), std::string::npos); EXPECT_NE(json.find("\"close_reason\":\"deadline\""), std::string::npos); - EXPECT_NE(json.find("\"engine\":\"nvidia.pc_sampling\""), std::string::npos); + EXPECT_NE( + json.find("\"engines\":[\"nvidia.sass_metrics\",\"nvidia.pc_sampling\"]"), + std::string::npos); EXPECT_NE(json.find("\"launches_covered\":12"), std::string::npos); EXPECT_NE(json.find("\"requested_duration_ms\":3000"), std::string::npos); EXPECT_NE(json.find("\"duration_ns\":2000"), std::string::npos); EXPECT_EQ(gpufl::model::DeepWindowModel(e).channel(), gpufl::Channel::Scope); } + +TEST(DeepWindowModelTest, EmptyEngineListSerializesAsAnEmptyArray) { + // A window that armed nothing is a real outcome - the one PC sampling + // produces when its arm is refused - so it has to survive the wire as an + // empty list rather than as a missing field or a sentinel string. + gpufl::DeepWindowEvent e; + e.name = "deep_window"; + e.close_reason = "deadline"; + + const std::string json = gpufl::model::DeepWindowModel(e).buildJson(); + EXPECT_NE(json.find("\"engines\":[]"), std::string::npos); +} + +TEST(DeepWindowModelTest, SingleEngineSerializesAsAOneElementArray) { + gpufl::DeepWindowEvent e; + e.name = "deep_window"; + e.engines = {"nvidia.pm_sampling"}; + + const std::string json = gpufl::model::DeepWindowModel(e).buildJson(); + EXPECT_NE(json.find("\"engines\":[\"nvidia.pm_sampling\"]"), + std::string::npos); +} From 5de56d0d4daceca9cda7edd8c2e51754d78fb44b Mon Sep 17 00:00:00 2001 From: Myoungho Shin Date: Sat, 25 Jul 2026 22:34:56 -0700 Subject: [PATCH 10/10] feat(deep-window): open the window as a real scope, and make scopes nest --- include/gpufl/core/deep_window.cpp | 35 ++++++++ include/gpufl/core/monitor.cpp | 6 ++ include/gpufl/core/monitor.hpp | 3 + include/gpufl/core/monitor_batch_manager.cpp | 25 +++++- include/gpufl/core/monitor_batch_manager.hpp | 11 +++ include/gpufl/inject/inject_entry.cpp | 11 ++- tests/core/test_monitor.cpp | 88 ++++++++++++++++++++ 7 files changed, 169 insertions(+), 10 deletions(-) diff --git a/include/gpufl/core/deep_window.cpp b/include/gpufl/core/deep_window.cpp index cb7c6f5..bec910d 100644 --- a/include/gpufl/core/deep_window.cpp +++ b/include/gpufl/core/deep_window.cpp @@ -52,6 +52,9 @@ int64_t g_opened_ns = 0; int64_t g_requested_duration_ms = 0; uint64_t g_requested_max_launches = 0; std::string g_name; +// Pairs the window's scope begin/end rows. 0 = no scope row was pushed, which +// is what a run with no Monitor backend looks like. +uint64_t g_scope_instance_id = 0; bool ComboActive() { const char* combo = std::getenv(env::kEngineCombo); @@ -196,6 +199,21 @@ bool DeepWindow::Open(const DeepWindowSpec& spec) { // consuming budget, and everything it reads is already set. g_active.store(true, std::memory_order_release); + // Open the window as a real scope, not just an arming signal. This is + // what puts it on the timeline as its own range AND makes the samples + // collected inside it carry the window's name instead of the enclosing + // process scope - the sample writers stamp whatever scope is active. + // Pushed before the engines arm so nothing collected can land under + // the parent name. + g_scope_instance_id = Monitor::AllocateScopeInstanceId(); + ScopeBatchRow open_row; + open_row.ts_ns = g_opened_ns; + open_row.scope_instance_id = g_scope_instance_id; + open_row.name_id = Monitor::InternScopeName(g_name); + open_row.event_type = 0; + open_row.depth = Monitor::OpenScopeDepth(); + Monitor::PushScopeRow(open_row); + // Arms the deep engines. Runs under the lock so a concurrent close // can't disarm engines this call hasn't armed yet; safe because the // arm path doesn't re-enter DeepWindow. @@ -265,6 +283,22 @@ void DeepWindow::Close(const DeepWindowClose reason) { ev.engines = { ProfilingEngineWireName(Monitor::ResolvedProfilingEngine())}; } + + // Close the scope last. The disarm above drains what the engines + // collected, and those samples belong to this window - closing first + // would hand the name back to the process scope and mislabel them. + // Pushed even on the teardown path, where skipping it would leave the + // scope open and every later sample carrying the window's name. + if (g_scope_instance_id != 0) { + ScopeBatchRow close_row; + close_row.ts_ns = end_ns; + close_row.scope_instance_id = g_scope_instance_id; + close_row.name_id = Monitor::InternScopeName(name); + close_row.event_type = 1; + close_row.depth = 0; // ignored on close; the open row carries it + Monitor::PushScopeRow(close_row); + g_scope_instance_id = 0; + } } if (const Runtime* rt = runtime(); rt && rt->logger) { @@ -390,6 +424,7 @@ void DeepWindow::ResetForTesting() { g_requested_duration_ms = 0; g_requested_max_launches = 0; g_name.clear(); + g_scope_instance_id = 0; } // ---- Public API ---- diff --git a/include/gpufl/core/monitor.cpp b/include/gpufl/core/monitor.cpp index d40b55c..5fbae3d 100644 --- a/include/gpufl/core/monitor.cpp +++ b/include/gpufl/core/monitor.cpp @@ -661,6 +661,12 @@ void Monitor::PushScopeRow(const ScopeBatchRow& row) { g_state.batches.pushTrackedScopeRow(row); } +uint64_t Monitor::AllocateScopeInstanceId() { + return g_state.batches.allocateScopeInstanceId(); +} + +int Monitor::OpenScopeDepth() { return g_state.batches.openScopeDepth(); } + void Monitor::PushProfileSamples(const std::vector& samples) { if (samples.empty()) return; const uint32_t scope_name_id = g_state.batches.activeScopeNameId(); diff --git a/include/gpufl/core/monitor.hpp b/include/gpufl/core/monitor.hpp index df13870..8a5e66c 100644 --- a/include/gpufl/core/monitor.hpp +++ b/include/gpufl/core/monitor.hpp @@ -384,6 +384,9 @@ class Monitor { * @brief Push a pre-built ScopeBatchRow into the scope batch buffer. */ static void PushScopeRow(const ScopeBatchRow& row); + static uint64_t AllocateScopeInstanceId(); + /** @brief Depth a scope opened right now would nest at. */ + static int OpenScopeDepth(); /** * Push a raw activity record into the monitor ring buffer. diff --git a/include/gpufl/core/monitor_batch_manager.cpp b/include/gpufl/core/monitor_batch_manager.cpp index 4deb756..966ed38 100644 --- a/include/gpufl/core/monitor_batch_manager.cpp +++ b/include/gpufl/core/monitor_batch_manager.cpp @@ -18,6 +18,7 @@ void MonitorBatchManager::reset() { scopeBatch_.clear(); profileBatch_.clear(); pmSampleBatch_.clear(); + scopeNameStack_.clear(); openScopeWindows_.clear(); completedScopeWindows_.clear(); } @@ -159,6 +160,11 @@ uint32_t MonitorBatchManager::activeScopeNameId() const { return activeScopeNameId_.load(std::memory_order_relaxed); } +int MonitorBatchManager::openScopeDepth() const { + std::lock_guard lk(scopeBatchMu_); + return static_cast(scopeNameStack_.size()); +} + bool MonitorBatchManager::pushKernel(const KernelBatchRow& row, const KernelDetailRow* detail) { kernelBatch_.push(row); @@ -181,14 +187,25 @@ void MonitorBatchManager::pushTraceScopeRows(const ScopeBatchRow& begin_row, } void MonitorBatchManager::pushTrackedScopeRow(const ScopeBatchRow& row) { - if (row.event_type == 0) { - activeScopeNameId_.store(row.name_id, std::memory_order_relaxed); - } - std::lock_guard lk(scopeBatchMu_); if (row.event_type == 0) { + scopeNameStack_.emplace_back(row.scope_instance_id, row.name_id); + activeScopeNameId_.store(row.name_id, std::memory_order_relaxed); openScopeWindows_[row.scope_instance_id] = {row.ts_ns, row.name_id, row.depth}; } else { + // Search from the back: the common case is closing the innermost + // scope, and an unmatched id leaves the stack alone rather than + // popping somebody else's scope. + for (auto it = scopeNameStack_.rbegin(); it != scopeNameStack_.rend(); ++it) { + if (it->first == row.scope_instance_id) { + scopeNameStack_.erase(std::next(it).base()); + break; + } + } + activeScopeNameId_.store( + scopeNameStack_.empty() ? 0 : scopeNameStack_.back().second, + std::memory_order_relaxed); + if (const auto it = openScopeWindows_.find(row.scope_instance_id); it != openScopeWindows_.end()) { completedScopeWindows_.push_back( diff --git a/include/gpufl/core/monitor_batch_manager.hpp b/include/gpufl/core/monitor_batch_manager.hpp index 405a4a4..30d87a2 100644 --- a/include/gpufl/core/monitor_batch_manager.hpp +++ b/include/gpufl/core/monitor_batch_manager.hpp @@ -42,6 +42,8 @@ class MonitorBatchManager { uint64_t allocateScopeInstanceId(); uint32_t activeScopeNameId() const; + /** @brief How many scopes are open right now; the depth a new one nests at. */ + int openScopeDepth() const; bool pushKernel(const KernelBatchRow& row, const KernelDetailRow* detail = nullptr); bool pushMemcpy(const MemcpyBatchRow& row); @@ -94,7 +96,16 @@ class MonitorBatchManager { uint64_t pmSampleBatchId_ = 0; mutable std::mutex scopeBatchMu_; std::atomic nextScopeInstanceId_{1}; + // Cached top of scopeNameStack_. Read without the mutex on the sample hot + // path, written only while holding it. std::atomic activeScopeNameId_{0}; + // Scopes currently open, innermost last. A stack rather than a single + // value because scopes nest: a deep window opens inside the process scope + // and must hand the name back on close, or every sample after it keeps the + // window's name. Entries carry their instance id so a close that is not + // strictly LIFO - the collector can close a deep window while an + // application scope is open - removes the right one. + std::vector> scopeNameStack_; std::unordered_map openScopeWindows_; std::vector completedScopeWindows_; diff --git a/include/gpufl/inject/inject_entry.cpp b/include/gpufl/inject/inject_entry.cpp index 5510a8f..3ac37f3 100644 --- a/include/gpufl/inject/inject_entry.cpp +++ b/include/gpufl/inject/inject_entry.cpp @@ -93,11 +93,6 @@ ProcessScopeState& processScope() { return *state; } -uint64_t nextProcessScopeId() { - static std::atomic next{1}; - return next.fetch_add(1, std::memory_order_relaxed); -} - std::string processScopeName() { if (const char* app = std::getenv(gpufl::env::kAppName)) { if (app[0] != '\0') return std::string("process:") + app; @@ -117,7 +112,11 @@ void beginProcessScope() { state.active = true; state.perf_scope = engineNeedsPerfScope(); - state.instance_id = nextProcessScopeId(); + // Must come from the shared allocator: instance ids pair begin with end, + // and every other scope - user scopes and deep windows - draws from it. A + // private counter here starts at 1 too and collides with the first of + // those, which pairs the wrong rows together. + state.instance_id = gpufl::Monitor::AllocateScopeInstanceId(); state.name = processScopeName(); state.name_id = gpufl::Monitor::InternScopeName(state.name); diff --git a/tests/core/test_monitor.cpp b/tests/core/test_monitor.cpp index 9449c75..60e67cd 100644 --- a/tests/core/test_monitor.cpp +++ b/tests/core/test_monitor.cpp @@ -5,6 +5,7 @@ #include "common/test_utils.hpp" #include "gpufl/core/monitor.hpp" +#include "gpufl/core/monitor_batch_manager.hpp" class MonitorTest : public ::testing::Test { protected: @@ -75,3 +76,90 @@ TEST_F(MonitorTest, MultipleInitialize) { gpufl::Monitor::Shutdown(); gpufl::Monitor::Shutdown(); // Should be safe } + +// ── scope name stack ──────────────────────────────────────────────────────── +// +// The active scope name is what stamps every profile and PM sample. It has to +// behave as a stack: a deep window opens inside the process scope and must +// hand the name back when it closes, or every sample taken after the window +// keeps the window's name. + +namespace { + +gpufl::ScopeBatchRow ScopeRow(uint64_t instance_id, uint32_t name_id, + uint8_t event_type, int depth = 0) { + gpufl::ScopeBatchRow row; + row.ts_ns = 1000 + static_cast(instance_id); + row.scope_instance_id = instance_id; + row.name_id = name_id; + row.event_type = event_type; + row.depth = depth; + return row; +} + +} // namespace + +TEST(ScopeNameStackTest, NestedCloseRestoresTheEnclosingName) { + gpufl::detail::MonitorBatchManager m; + const uint32_t outer = m.internScopeName("process:app"); + const uint32_t inner = m.internScopeName("deep_window"); + + m.pushTrackedScopeRow(ScopeRow(1, outer, 0, 0)); + EXPECT_EQ(m.activeScopeNameId(), outer); + + m.pushTrackedScopeRow(ScopeRow(2, inner, 0, 1)); + EXPECT_EQ(m.activeScopeNameId(), inner); + + // The regression this guards: without a stack the name stayed on the + // window and every later sample was attributed to it. + m.pushTrackedScopeRow(ScopeRow(2, inner, 1)); + EXPECT_EQ(m.activeScopeNameId(), outer); + + m.pushTrackedScopeRow(ScopeRow(1, outer, 1)); + EXPECT_EQ(m.activeScopeNameId(), 0u); +} + +TEST(ScopeNameStackTest, OutOfOrderCloseRemovesOnlyItsOwnScope) { + // A deep window closes from the collector thread, so it can close while an + // application scope opened after it is still open. Closing must not pop + // whatever happens to be on top. + gpufl::detail::MonitorBatchManager m; + const uint32_t process = m.internScopeName("process:app"); + const uint32_t window = m.internScopeName("deep_window"); + const uint32_t user = m.internScopeName("user_scope"); + + m.pushTrackedScopeRow(ScopeRow(1, process, 0, 0)); + m.pushTrackedScopeRow(ScopeRow(2, window, 0, 1)); + m.pushTrackedScopeRow(ScopeRow(3, user, 0, 2)); + + m.pushTrackedScopeRow(ScopeRow(2, window, 1)); + EXPECT_EQ(m.activeScopeNameId(), user) << "closing the window must leave the user scope active"; + + m.pushTrackedScopeRow(ScopeRow(3, user, 1)); + EXPECT_EQ(m.activeScopeNameId(), process); +} + +TEST(ScopeNameStackTest, UnmatchedCloseLeavesTheStackAlone) { + gpufl::detail::MonitorBatchManager m; + const uint32_t outer = m.internScopeName("process:app"); + + m.pushTrackedScopeRow(ScopeRow(1, outer, 0, 0)); + m.pushTrackedScopeRow(ScopeRow(99, outer, 1)); + + EXPECT_EQ(m.activeScopeNameId(), outer); + EXPECT_EQ(m.openScopeDepth(), 1); +} + +TEST(ScopeNameStackTest, DepthReportsWhereANewScopeWouldNest) { + gpufl::detail::MonitorBatchManager m; + const uint32_t outer = m.internScopeName("process:app"); + const uint32_t inner = m.internScopeName("deep_window"); + + EXPECT_EQ(m.openScopeDepth(), 0); + m.pushTrackedScopeRow(ScopeRow(1, outer, 0, 0)); + EXPECT_EQ(m.openScopeDepth(), 1); + m.pushTrackedScopeRow(ScopeRow(2, inner, 0, 1)); + EXPECT_EQ(m.openScopeDepth(), 2); + m.pushTrackedScopeRow(ScopeRow(2, inner, 1)); + EXPECT_EQ(m.openScopeDepth(), 1); +}