From e555d0893b6f2e0712d73748b10a1f5d0055e8fe Mon Sep 17 00:00:00 2001 From: ghazni101 Date: Thu, 3 Sep 2026 22:21:55 +0400 Subject: [PATCH 1/2] feat(config): make the memory preflight OOM guard configurable MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The preflight OOM guard (kv_slots::preflight_alloc: deployment-target ceiling + MemAvailable headroom, plus the CLI bench-sweep headroom check) assumed GPU memory always comes from system RAM. That is true on unified-memory APUs (Strix Halo, no swap), where an overshoot is a global OOM that kills the desktop — and is exactly why the guard exists. On a discrete-GPU dev box an overshoot is a plain failed hipMalloc, and reviewer feedback points out agents running multiple daemons/serves there get refused loads the hardware would survive. Add a typed, process-scoped schema key memory.oom_guard (default true, env compat HIPFIRE_OOM_GUARD) so the operator can opt out: hipfire config set memory.oom_guard false # or: HIPFIRE_OOM_GUARD=0 Both production guard sites honor the knob at the single preflight_alloc choke point and in preflight_headroom_for_model; a disabled guard prints a one-line stderr note so a skipped refusal is visible in logs. Default behavior is unchanged: the guard stays on unless explicitly disabled, and scripts/run-bounded.sh remains the hard cgroup backstop either way. --- AGENTS.md | 1 + crates/hipfire-cli/src/main.rs | 9 +++++ crates/hipfire-config/src/lib.rs | 60 +++++++++++++++++++++++++++++ crates/rdna-compute/src/kv_slots.rs | 23 +++++++++++ docs/CONFIG.md | 1 + docs/env-vars.md | 1 + 6 files changed, 95 insertions(+) diff --git a/AGENTS.md b/AGENTS.md index 3e0288765..dd71870a7 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -602,6 +602,7 @@ against the A3B MoE DFlash perfmaxx line. | `HIPFIRE_PROMPT_HEAT_LIMIT` | Max rows in heat dump | 64 | | `HIPFIRE_KV_MODE` | Override kv_cache config | (config) | | `HIPFIRE_ATTN_FLASH` | Override flash_mode config | (config) | +| `HIPFIRE_OOM_GUARD` | Memory preflight OOM guard (`kv_slots::preflight_alloc`, SlotPool arena, bench-sweep headroom check). Protects unified-memory APUs (Strix Halo) where an overshoot is a global OOM; opt out on discrete-GPU dev boxes | ON (`memory.oom_guard`) | |`HIPFIRE_DFLASH_DRAFT`|Force a specific draft path. Empty string = explicit opt-out|(filename auto-match alongside target)| |`HIPFIRE_DFLASH_CTX_CAP`|Max rows for draft context-indexed structures (target_hidden, draft K/V caches, hidden ring). Bounds draft-side VRAM on large-`max_seq` serve loads; over-cap requests fall back to AR (identical output, slower). `0` = uncapped legacy.|8192| |`HIPFIRE_DFLASH_WINDOW`|Windowed draft context (NInfer pattern): SWA over the last W rows on draft layers 0..n-2 + full-attention last layer reaching min(physical_cap, 4W). Draft VRAM pins at W regardless of `max_seq`; past-W requests degrade τ instead of falling back to AR. Refused with CASK eviction. `0`/unset = Legacy (cap + AR fallback).|0 (off)| diff --git a/crates/hipfire-cli/src/main.rs b/crates/hipfire-cli/src/main.rs index 86764ac22..66e20abfc 100644 --- a/crates/hipfire-cli/src/main.rs +++ b/crates/hipfire-cli/src/main.rs @@ -3877,7 +3877,16 @@ fn bench_concurrency_command(paths: &Paths, args: &BenchArgs, spec: &str) -> Res /// a leaked first model would show up: if the slots engine did not actually /// release its weights, `MemAvailable` is still depressed here and this stops /// the sweep instead of taking the box down. +/// +/// `memory.oom_guard=false` (`HIPFIRE_OOM_GUARD=0`) opts out: on a +/// discrete-GPU box an overshoot is a plain failed hipMalloc, not a desktop +/// kill, and a sweep that wants to probe past the headroom is the operator's +/// call. fn preflight_headroom_for_model(paths: &Paths, model: &str) -> Result<()> { + if !hipfire_config::oom_guard_enabled() { + eprintln!("memory headroom guard disabled (memory.oom_guard=false)"); + return Ok(()); + } let registry = load_registry(&paths.registry).registry; let Some(path) = find_model_path(paths, ®istry, model) else { return Ok(()); diff --git a/crates/hipfire-config/src/lib.rs b/crates/hipfire-config/src/lib.rs index 1de90bb0b..b949fb36d 100644 --- a/crates/hipfire-config/src/lib.rs +++ b/crates/hipfire-config/src/lib.rs @@ -633,6 +633,18 @@ pub static FIELDS: &[ConfigField] = &[ Some("HIPFIRE_KV_ADAPTIVE"), "Runtime VRAM-fit KV precision policy." ), + // Process-scoped: the preflight guards snapshot this once at startup, and + // a mid-serve flip would make the refusal policy depend on which load ran + // last — dishonest for a long-lived daemon. + process_bool_field!( + "memory.oom_guard", + "oom_guard", + Memory, + true, + false, + "HIPFIRE_OOM_GUARD", + "Memory preflight OOM guard. Default on: on unified-memory APUs (Strix Halo) GPU memory is system RAM with no swap, so an overshoot is a global OOM that kills the desktop, not this process. Set false (HIPFIRE_OOM_GUARD=0) on discrete-GPU boxes, where an overshoot is a plain failed hipMalloc." + ), field!( "model.deepseek4_experts_per_token", "deepseek4_experts_per_token", @@ -3226,6 +3238,29 @@ pub fn process_value(name: &str) -> Option { active_or_local_process_config().legacy_value(name) } +/// Resolve the memory preflight OOM guard (`memory.oom_guard`, compat +/// `HIPFIRE_OOM_GUARD`). Default ON: the guard exists because on +/// unified-memory APUs (Strix Halo) GPU allocations come out of system RAM +/// with no swap, so a bad admission takes the desktop down with a global +/// OOM rather than failing one request. Opting out — discrete-GPU boxes +/// where an overshoot is a plain failed `hipMalloc`, or deliberate +/// multi-daemon development setups — is the operator's informed trade. +pub fn oom_guard_enabled() -> bool { + oom_guard_enabled_for(process_value("HIPFIRE_OOM_GUARD").as_deref()) +} + +/// Pure form of [`oom_guard_enabled`] so the off-spellings have tests without +/// pinning the process-wide config snapshot. +fn oom_guard_enabled_for(value: Option<&str>) -> bool { + match value { + Some(value) => !matches!( + value.to_ascii_lowercase().as_str(), + "0" | "false" | "off" | "no" + ), + None => true, + } +} + /// Compatibility-shaped access for experimental code while its public policy /// is being consolidated. Values come exclusively from the process snapshot. pub fn developer_var(name: &str) -> std::result::Result { @@ -4445,6 +4480,31 @@ mod tests { env::temp_dir().join(format!("hipfire-config-{name}-{}", std::process::id())) } + #[test] + fn oom_guard_defaults_on_and_parses_off_spellings() { + // Unset and explicit-on spellings keep the guard up. + assert!(oom_guard_enabled_for(None)); + assert!(oom_guard_enabled_for(Some("1"))); + assert!(oom_guard_enabled_for(Some("true"))); + assert!(oom_guard_enabled_for(Some("ON"))); + // A garbage value must not silently disable a safety guard. + assert!(oom_guard_enabled_for(Some("banana"))); + // The typed bool renders "0"; raw compat spellings also count. + assert!(!oom_guard_enabled_for(Some("0"))); + assert!(!oom_guard_enabled_for(Some("false"))); + assert!(!oom_guard_enabled_for(Some("off"))); + assert!(!oom_guard_enabled_for(Some("OFF"))); + assert!(!oom_guard_enabled_for(Some("no"))); + } + + #[test] + fn oom_guard_schema_field_is_process_scoped_with_env_compat() { + let field = field("memory.oom_guard").expect("oom_guard schema field"); + assert_eq!(field.env_compat, Some("HIPFIRE_OOM_GUARD")); + assert!(matches!(field.default.to_value(), ConfigValue::Bool(true))); + assert!(!field.include_builtin_in_process_config); + } + #[test] fn schema_has_unique_keys_and_legacy_keys() { let mut canonical = std::collections::BTreeSet::new(); diff --git a/crates/rdna-compute/src/kv_slots.rs b/crates/rdna-compute/src/kv_slots.rs index 93aac5646..af2cee2c4 100644 --- a/crates/rdna-compute/src/kv_slots.rs +++ b/crates/rdna-compute/src/kv_slots.rs @@ -258,6 +258,12 @@ pub fn build_tiles(slot_query_counts: &[usize], br: usize) -> (Vec, Vec Option { /// Refuse a planned allocation that would either exceed the deployment target's /// VRAM or leave this box without enough headroom to stay responsive. /// +/// Skipped entirely when `memory.oom_guard=false` (`HIPFIRE_OOM_GUARD=0`): +/// the deployment-target ceiling and the headroom check both assume GPU +/// memory comes from system RAM, which is true on unified-memory APUs and +/// false on a discrete-GPU dev box. The skip prints once to stderr so a +/// disabled guard is visible in a log instead of reading as a pass. +/// /// `planned_bytes` must be the TOTAL the caller is about to hold live at once, /// not a single buffer. Returns `Err` with an actionable message; callers should /// skip the configuration rather than proceed. @@ -293,6 +305,17 @@ pub fn mem_available_bytes() -> Option { /// forbidden by scripts/check-env-docs.py. Harnesses live in `examples/`, /// which is exempt, so they read any override there and pass it in. pub fn preflight_alloc(planned_bytes: u64, budget_bytes: u64, what: &str) -> Result<(), String> { + if !hipfire_config::oom_guard_enabled() { + static SKIP_NOTE: std::sync::Once = std::sync::Once::new(); + SKIP_NOTE.call_once(|| { + eprintln!( + "[kv_slots] memory preflight guard disabled (memory.oom_guard=false); \ + allocations will not be refused before allocate" + ); + }); + return Ok(()); + } + let budget = budget_bytes; let gib = |b: u64| b as f64 / 1073741824.0; diff --git a/docs/CONFIG.md b/docs/CONFIG.md index 043cc8225..829057a97 100644 --- a/docs/CONFIG.md +++ b/docs/CONFIG.md @@ -109,6 +109,7 @@ Stable/default-on and safety controls: | `kernel.rocblas_off` | `false` | Disable rocBLAS dispatch. | | `fusions.force_unfused` | `false` | Force supported projection paths unfused. | | `speculation.dflash_tree` | `false` | Enable DDTree tree-SWOR verification. | +| `memory.oom_guard` | `true` | Memory preflight OOM guard (`kv_slots::preflight_alloc` and the bench-sweep headroom check). Refuses oversized allocations before they are made. Default on because on unified-memory APUs (Strix Halo) GPU memory is system RAM with no swap — an overshoot is a global OOM that kills the desktop. Set `false` (`HIPFIRE_OOM_GUARD=0`) on discrete-GPU boxes, where an overshoot is a plain failed hipMalloc. | The following default-off keys are experimental kernel-route overrides. They are typed booleans, process-scoped, and visible in `hipfire config list` with diff --git a/docs/env-vars.md b/docs/env-vars.md index c2da45af4..e7a4abb05 100644 --- a/docs/env-vars.md +++ b/docs/env-vars.md @@ -1051,4 +1051,5 @@ When adding a user-facing knob: |---|---|---| | `HIPFIRE_ATTN_TILE_SIZE` | `128` | Tile size for the batched attention tile+reduce path. Must be a positive multiple of 32; anything else falls back to 128. Resolved once via `Gpu::attn_tile_size()`. **Raising it is safe; lowering it increases `max_tiles` and therefore the `partials` bytes per query row, which can exceed buffers sized elsewhere against the 128 default.** | | `HIPFIRE_VRAM_BUDGET_BYTES` | 32 GiB | Deployment-target VRAM ceiling used by the SP1 benchmark harnesses' preflight. Read by `examples/`, not by production code. | +| `HIPFIRE_OOM_GUARD` | on | Typed key `memory.oom_guard`. Disables the production memory preflight (`kv_slots::preflight_alloc`, the `SlotPool` arena check, and the CLI bench-sweep headroom check) when set to `0`/`false`/`off`. Default on: on unified-memory APUs (Strix Halo) GPU memory is system RAM with no swap, so an overshoot is a global OOM that kills the desktop; on a discrete GPU an overshoot is a plain failed hipMalloc, so multi-daemon dev boxes may opt out. `scripts/run-bounded.sh` remains the hard backstop either way. | | `HIPFIRE_MEM_CAP` | `24G` | Read by `scripts/run-bounded.sh`, not by the binaries: cgroup `MemoryMax` for a gated run. Exit 137 means the cap fired — shrink the configuration rather than raising it. | From e7e17cf6448d23771e75957291fee95ace21586a Mon Sep 17 00:00:00 2001 From: ghazni101 Date: Fri, 4 Sep 2026 09:33:54 +0400 Subject: [PATCH 2/2] =?UTF-8?q?feat(config):=20oom=5Fguard=20auto=20mode?= =?UTF-8?q?=20=E2=80=94=20enable=20only=20on=20unified-memory=20APUs?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Follow-up to 375f446e3. The guard's default is now the three-state `auto` (bool spellings still parse) instead of always-on, so dGPU dev boxes run unguarded by default and unified-memory boxes stay protected without anyone setting anything: auto: unified-memory APU arch (gfx1035/1036/1103/1150/1151/1152) → guard ON — GPU allocations come out of system RAM discrete-GPU arch (gfx90x/10x/1100-02/1200-01, CDNA) → guard OFF — an overshoot is a failed hipMalloc, not an OOM unrecognized arch → guard ON — fail safe; extend the table when support lands no GPU arch known in this process (CLI side, pre-init) → host swap state decides: swap → OFF (overcommit degrades, does not kill), no/unreadable swap → ON (fail safe) Mechanics: - hipfire-config: `memory.oom_guard` flips to process_auto_bool_field (AutoBool rule, default "auto"); new OomGuardMode + is_unified_memory_arch tables + /proc/meminfo SwapTotal probe + oom_guard_effective(arch) resolver that logs the auto decision once with its reason. Tables live here because rdna-compute cannot depend back on this crate. - rdna-compute: arch_caps records the DETECTED (not HIPFIRE_TARGET_ARCH- overridden) arch at Gpu::init — first init wins; kv_slots::preflight_alloc resolves through it. The refusal checks are split into preflight_checks so unit tests stay deterministic regardless of host config. - hipfire-cli: bench-sweep headroom check resolves with arch=None (host swap signal). Verified: config list shows default `auto`, HIPFIRE_OOM_GUARD=0 and `config set memory.oom_guard auto` both resolve; hipfire-config and kv_slots suites green; all touched crates compile. --- AGENTS.md | 2 +- crates/hipfire-cli/map.md | 4 +- crates/hipfire-cli/src/main.rs | 13 +- crates/hipfire-config/map.md | 6 +- crates/hipfire-config/src/lib.rs | 251 +++++++++++++++++++++++---- crates/rdna-compute/map.md | 12 +- crates/rdna-compute/src/arch_caps.rs | 19 ++ crates/rdna-compute/src/dispatch.rs | 6 + crates/rdna-compute/src/kv_slots.rs | 50 ++++-- crates/rdna-compute/src/slot_pool.rs | 10 +- docs/CONFIG.md | 2 +- docs/env-vars.md | 2 +- 12 files changed, 305 insertions(+), 72 deletions(-) diff --git a/AGENTS.md b/AGENTS.md index dd71870a7..1e2c0fbaa 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -602,7 +602,7 @@ against the A3B MoE DFlash perfmaxx line. | `HIPFIRE_PROMPT_HEAT_LIMIT` | Max rows in heat dump | 64 | | `HIPFIRE_KV_MODE` | Override kv_cache config | (config) | | `HIPFIRE_ATTN_FLASH` | Override flash_mode config | (config) | -| `HIPFIRE_OOM_GUARD` | Memory preflight OOM guard (`kv_slots::preflight_alloc`, SlotPool arena, bench-sweep headroom check). Protects unified-memory APUs (Strix Halo) where an overshoot is a global OOM; opt out on discrete-GPU dev boxes | ON (`memory.oom_guard`) | +| `HIPFIRE_OOM_GUARD` | Memory preflight OOM guard (`kv_slots::preflight_alloc`, SlotPool arena, bench-sweep headroom check). `auto`: on for unified-memory APUs (Strix Halo — overshoot is a global OOM), off for discrete GPUs, swap-decided for GPU-less processes | `auto` (`memory.oom_guard`) | |`HIPFIRE_DFLASH_DRAFT`|Force a specific draft path. Empty string = explicit opt-out|(filename auto-match alongside target)| |`HIPFIRE_DFLASH_CTX_CAP`|Max rows for draft context-indexed structures (target_hidden, draft K/V caches, hidden ring). Bounds draft-side VRAM on large-`max_seq` serve loads; over-cap requests fall back to AR (identical output, slower). `0` = uncapped legacy.|8192| |`HIPFIRE_DFLASH_WINDOW`|Windowed draft context (NInfer pattern): SWA over the last W rows on draft layers 0..n-2 + full-attention last layer reaching min(physical_cap, 4W). Draft VRAM pins at W regardless of `max_seq`; past-W requests degrade τ instead of falling back to AR. Refused with CASK eviction. `0`/unset = Legacy (cap + AR fallback).|0 (off)| diff --git a/crates/hipfire-cli/map.md b/crates/hipfire-cli/map.md index d87967dec..02d7bf1cd 100644 --- a/crates/hipfire-cli/map.md +++ b/crates/hipfire-cli/map.md @@ -23,7 +23,7 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside | File | Lines | Public items | Tests | |---|---:|---:|---:| | [`src/bench_concurrency.rs`](src/bench_concurrency.rs) | 720 | 21 | 9 | -| [`src/main.rs`](src/main.rs) | 9,692 | 0 | 65 | +| [`src/main.rs`](src/main.rs) | 9,702 | 0 | 65 | | [`src/serve/complete.rs`](src/serve/complete.rs) | 6,754 | 0 | 89 | | [`src/serve/http.rs`](src/serve/http.rs) | 1,089 | 0 | 6 | | [`src/serve/metrics.rs`](src/serve/metrics.rs) | 328 | 0 | 5 | @@ -53,6 +53,6 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside ### Totals -- 7 modules · 22,228 lines · 23 public items · 206 tests · 0 examples +- 7 modules · 22,238 lines · 23 public items · 206 tests · 0 examples diff --git a/crates/hipfire-cli/src/main.rs b/crates/hipfire-cli/src/main.rs index 66e20abfc..18f00435b 100644 --- a/crates/hipfire-cli/src/main.rs +++ b/crates/hipfire-cli/src/main.rs @@ -3878,13 +3878,14 @@ fn bench_concurrency_command(paths: &Paths, args: &BenchArgs, spec: &str) -> Res /// release its weights, `MemAvailable` is still depressed here and this stops /// the sweep instead of taking the box down. /// -/// `memory.oom_guard=false` (`HIPFIRE_OOM_GUARD=0`) opts out: on a -/// discrete-GPU box an overshoot is a plain failed hipMalloc, not a desktop -/// kill, and a sweep that wants to probe past the headroom is the operator's -/// call. +/// `memory.oom_guard` (default `auto`) opts out or forces the check on: this +/// process never initializes a GPU, so `auto` falls back to host swap state — +/// with swap an overcommit degrades rather than kills and the check stands +/// down; without swap it stays up. A discrete-GPU box that wants the check +/// anyway pins `memory.oom_guard=true`. fn preflight_headroom_for_model(paths: &Paths, model: &str) -> Result<()> { - if !hipfire_config::oom_guard_enabled() { - eprintln!("memory headroom guard disabled (memory.oom_guard=false)"); + if !hipfire_config::oom_guard_effective(None) { + eprintln!("memory headroom guard inactive (memory.oom_guard); continuing sweep"); return Ok(()); } let registry = load_registry(&paths.registry).registry; diff --git a/crates/hipfire-config/map.md b/crates/hipfire-config/map.md index ba071529c..b8658d5c6 100644 --- a/crates/hipfire-config/map.md +++ b/crates/hipfire-config/map.md @@ -23,13 +23,13 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside | File | Lines | Public items | Tests | |---|---:|---:|---:| | [`src/bin/hipfire-rocm-resolve.rs`](src/bin/hipfire-rocm-resolve.rs) | 105 | 0 | 0 | -| [`src/lib.rs`](src/lib.rs) | 5,159 | 79 | 28 | +| [`src/lib.rs`](src/lib.rs) | 5,402 | 85 | 33 | | [`src/rocm.rs`](src/rocm.rs) | 2,460 | 39 | 37 | ### Public API surface - [`src/bin/hipfire-rocm-resolve.rs`](src/bin/hipfire-rocm-resolve.rs): — -- [`src/lib.rs`](src/lib.rs): `rocm`, `CONFIG_SCHEMA_VERSION`, `ConfigError`, `Result`, `ConfigValue`, `DeviceSelector`, `Deepseek4ComputePlacement`, `Deepseek4CompressorCache`, `kind`, `ConfigCategory`, `ConfigScope`, `DefaultValue`, +67 more +- [`src/lib.rs`](src/lib.rs): `rocm`, `CONFIG_SCHEMA_VERSION`, `ConfigError`, `Result`, `ConfigValue`, `DeviceSelector`, `Deepseek4ComputePlacement`, `Deepseek4CompressorCache`, `kind`, `ConfigCategory`, `ConfigScope`, `DefaultValue`, +73 more - [`src/rocm.rs`](src/rocm.rs): `DEVICE_COMPILERS`, `configured_root`, `has_configured_root`, `configured_compiler`, `has_configured_compiler`, `configured_compiler_from`, `is_strict_rocm`, `strict_from`, `strict_from_str`, `CompilerSource`, `ResolvedToolchain`, `version_for_root`, +27 more ### Dependencies (from `Cargo.toml`) @@ -45,6 +45,6 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside ### Totals -- 3 modules · 7,724 lines · 118 public items · 65 tests · 0 examples +- 3 modules · 7,967 lines · 124 public items · 70 tests · 0 examples diff --git a/crates/hipfire-config/src/lib.rs b/crates/hipfire-config/src/lib.rs index b949fb36d..4596c91ac 100644 --- a/crates/hipfire-config/src/lib.rs +++ b/crates/hipfire-config/src/lib.rs @@ -636,14 +636,13 @@ pub static FIELDS: &[ConfigField] = &[ // Process-scoped: the preflight guards snapshot this once at startup, and // a mid-serve flip would make the refusal policy depend on which load ran // last — dishonest for a long-lived daemon. - process_bool_field!( + process_auto_bool_field!( "memory.oom_guard", "oom_guard", Memory, - true, false, "HIPFIRE_OOM_GUARD", - "Memory preflight OOM guard. Default on: on unified-memory APUs (Strix Halo) GPU memory is system RAM with no swap, so an overshoot is a global OOM that kills the desktop, not this process. Set false (HIPFIRE_OOM_GUARD=0) on discrete-GPU boxes, where an overshoot is a plain failed hipMalloc." + "Memory preflight OOM guard. Default auto: on for unified-memory APU architectures (GPU allocations come out of system RAM, so an overshoot can globally OOM the desktop), off for discrete GPUs, and for GPU-less processes decided by host swap state. Set true to force on, false to force off (HIPFIRE_OOM_GUARD)." ), field!( "model.deepseek4_experts_per_token", @@ -3239,25 +3238,142 @@ pub fn process_value(name: &str) -> Option { } /// Resolve the memory preflight OOM guard (`memory.oom_guard`, compat -/// `HIPFIRE_OOM_GUARD`). Default ON: the guard exists because on -/// unified-memory APUs (Strix Halo) GPU allocations come out of system RAM -/// with no swap, so a bad admission takes the desktop down with a global -/// OOM rather than failing one request. Opting out — discrete-GPU boxes -/// where an overshoot is a plain failed `hipMalloc`, or deliberate -/// multi-daemon development setups — is the operator's informed trade. -pub fn oom_guard_enabled() -> bool { - oom_guard_enabled_for(process_value("HIPFIRE_OOM_GUARD").as_deref()) +/// `HIPFIRE_OOM_GUARD`). The guard exists because on unified-memory APUs +/// (Strix Halo) GPU allocations come out of system RAM with no swap, so a +/// bad admission takes the desktop down with a global OOM rather than +/// failing one request; on a discrete GPU an overshoot is a plain failed +/// `hipMalloc`. Default `auto` resolves per deployment class — see +/// [`oom_guard_effective`]. +#[derive(Clone, Copy, Debug, PartialEq, Eq)] +pub enum OomGuardMode { + /// Decide by deployment class (unified-memory APU vs discrete GPU). + Auto, + /// Always refuse oversized allocations before they are made. + On, + /// Never refuse (the operator's informed trade). + Off, } -/// Pure form of [`oom_guard_enabled`] so the off-spellings have tests without -/// pinning the process-wide config snapshot. -fn oom_guard_enabled_for(value: Option<&str>) -> bool { - match value { - Some(value) => !matches!( - value.to_ascii_lowercase().as_str(), - "0" | "false" | "off" | "no" - ), - None => true, +/// Read the configured mode: `auto` (also unset or unparseable — validated +/// layers should not produce anything else), or an on/off spelling. +fn oom_guard_mode_for(value: Option<&str>) -> OomGuardMode { + match value.map(|v| v.trim().to_ascii_lowercase()) { + Some(v) if v == "0" || v == "false" || v == "off" || v == "no" => OomGuardMode::Off, + Some(v) if v == "1" || v == "true" || v == "on" || v == "yes" => OomGuardMode::On, + _ => OomGuardMode::Auto, + } +} + +/// The configured mode of the memory preflight OOM guard. +pub fn oom_guard_mode() -> OomGuardMode { + oom_guard_mode_for(process_value("HIPFIRE_OOM_GUARD").as_deref()) +} + +/// GPU architectures whose allocations land in system RAM: the GPU has no +/// private VRAM (or only a small carve-out), so model weights and KV eat the +/// same physical memory as the desktop. An overshoot here is a global OOM, +/// not a failed hipMalloc. +pub const UNIFIED_MEMORY_ARCHS: &[&str] = &[ + "gfx1035", "gfx1036", // RDNA2 APU (Van Gogh / Steam Deck class) + "gfx1103", // RDNA3 APU (Phoenix orphan) + "gfx1150", "gfx1151", "gfx1152", // RDNA3.5 APU (Strix Point / Strix Halo) +]; + +/// GPU architectures with private VRAM: allocations that exceed it fail +/// that one allocation instead of the machine. +pub const DISCRETE_MEMORY_ARCHS: &[&str] = &[ + "gfx906", "gfx908", "gfx940", "gfx941", "gfx942", // CDNA (HBM) + "gfx1010", "gfx1011", "gfx1012", // RDNA1 + "gfx1030", "gfx1031", "gfx1032", // RDNA2 dGPU + "gfx1100", "gfx1101", "gfx1102", // RDNA3 dGPU + "gfx1200", "gfx1201", // RDNA4 +]; + +/// Whether `arch` is a unified-memory APU (GPU memory is system RAM). +pub fn is_unified_memory_arch(arch: &str) -> bool { + UNIFIED_MEMORY_ARCHS + .iter() + .any(|known| arch.eq_ignore_ascii_case(known)) +} + +/// Whether `arch` is a recognized discrete-VRAM GPU. +fn is_discrete_memory_arch(arch: &str) -> bool { + DISCRETE_MEMORY_ARCHS + .iter() + .any(|known| arch.eq_ignore_ascii_case(known)) +} + +/// `SwapTotal` (kB) from a /proc/meminfo body; `None` when absent/unreadable. +fn swap_total_kb_from_meminfo(meminfo: &str) -> Option { + for line in meminfo.lines() { + if let Some(rest) = line.strip_prefix("SwapTotal:") { + return rest.split_whitespace().next()?.parse().ok(); + } + } + None +} + +/// Host swap size in kB; `None` when /proc/meminfo cannot be read. +fn host_has_swap() -> Option { + let meminfo = std::fs::read_to_string("/proc/meminfo").ok()?; + Some(swap_total_kb_from_meminfo(&meminfo)? > 0) +} + +/// Pure auto decision, testable without pinning host state. With a known GPU +/// arch the deployment class decides (unified-memory APU → on, discrete → +/// off, unrecognized → on, failing safe). Without one (no GPU has been +/// initialized in this process) the host's own lethality decides: with swap +/// an overcommit degrades instead of killing, so the guard stands down; +/// without (or with unreadable) swap, it stays up. +fn oom_guard_auto_for(arch: Option<&str>, has_swap: Option) -> bool { + match arch { + Some(arch) if is_unified_memory_arch(arch) => true, + Some(arch) if is_discrete_memory_arch(arch) => false, + Some(_) => true, + None => !matches!(has_swap, Some(true)), + } +} + +/// Resolve whether the memory preflight guard should refuse allocations in +/// this process. +/// +/// `arch` is the GPU arch this process initialized (see +/// `rdna_compute::arch_caps::process_gpu_arch`), or `None` when no GPU is +/// (yet) known — e.g. a CLI process that only supervises the daemon. The +/// `auto` decision is logged once to stderr with its reason so a refusal (or +/// a skipped refusal) in a daemon log explains itself. +pub fn oom_guard_effective(arch: Option<&str>) -> bool { + match oom_guard_mode() { + OomGuardMode::On => true, + OomGuardMode::Off => false, + OomGuardMode::Auto => { + static DECISION_NOTE: std::sync::Once = std::sync::Once::new(); + let has_swap = host_has_swap(); + let enabled = oom_guard_auto_for(arch, has_swap); + DECISION_NOTE.call_once(|| { + let why = match (arch, has_swap) { + (Some(a), _) if is_unified_memory_arch(a) => { + format!("{a}: unified-memory APU; GPU allocations come from system RAM") + } + (Some(a), _) if is_discrete_memory_arch(a) => { + format!("{a}: discrete GPU; an overshoot is a failed hipMalloc, not an OOM") + } + (Some(a), _) => format!("{a}: unrecognized arch; failing safe"), + (None, Some(true)) => { + "no GPU arch known; host has swap, so an overcommit degrades rather than kills" + .to_string() + } + (None, _) => { + "no GPU arch known; host has no readable swap; failing safe".to_string() + } + }; + eprintln!( + "[oom_guard] auto: {why} → guard {}", + if enabled { "on" } else { "off" } + ); + }); + enabled + } } } @@ -4481,28 +4597,95 @@ mod tests { } #[test] - fn oom_guard_defaults_on_and_parses_off_spellings() { - // Unset and explicit-on spellings keep the guard up. - assert!(oom_guard_enabled_for(None)); - assert!(oom_guard_enabled_for(Some("1"))); - assert!(oom_guard_enabled_for(Some("true"))); - assert!(oom_guard_enabled_for(Some("ON"))); - // A garbage value must not silently disable a safety guard. - assert!(oom_guard_enabled_for(Some("banana"))); + fn oom_guard_mode_parses_auto_on_off() { + // Unset, "auto", and unparseable values all land on Auto — a garbage + // value must not silently disable a safety guard, nor force it past + // the deployment-class decision. + assert_eq!(oom_guard_mode_for(None), OomGuardMode::Auto); + assert_eq!(oom_guard_mode_for(Some("auto")), OomGuardMode::Auto); + assert_eq!(oom_guard_mode_for(Some("AUTO")), OomGuardMode::Auto); + assert_eq!(oom_guard_mode_for(Some("banana")), OomGuardMode::Auto); + assert_eq!(oom_guard_mode_for(Some("1")), OomGuardMode::On); + assert_eq!(oom_guard_mode_for(Some("true")), OomGuardMode::On); + assert_eq!(oom_guard_mode_for(Some("ON")), OomGuardMode::On); // The typed bool renders "0"; raw compat spellings also count. - assert!(!oom_guard_enabled_for(Some("0"))); - assert!(!oom_guard_enabled_for(Some("false"))); - assert!(!oom_guard_enabled_for(Some("off"))); - assert!(!oom_guard_enabled_for(Some("OFF"))); - assert!(!oom_guard_enabled_for(Some("no"))); + assert_eq!(oom_guard_mode_for(Some("0")), OomGuardMode::Off); + assert_eq!(oom_guard_mode_for(Some("false")), OomGuardMode::Off); + assert_eq!(oom_guard_mode_for(Some("OFF")), OomGuardMode::Off); + assert_eq!(oom_guard_mode_for(Some("no")), OomGuardMode::Off); + } + + #[test] + fn unified_and_discrete_arch_classes_are_disjoint_and_complete() { + // Every APU arch must resolve to unified, every dGPU/CDNA arch to + // not-unified, and the two tables must never overlap. + for arch in UNIFIED_MEMORY_ARCHS { + assert!(is_unified_memory_arch(arch)); + assert!( + !DISCRETE_MEMORY_ARCHS.contains(arch), + "{arch} in both tables" + ); + // Case-insensitive: arch strings arrive from the HIP runtime. + assert!(is_unified_memory_arch(&arch.to_uppercase())); + } + for arch in DISCRETE_MEMORY_ARCHS { + assert!(!is_unified_memory_arch(arch)); + assert!(is_discrete_memory_arch(arch)); + } + assert!(is_unified_memory_arch("gfx1151")); + assert!(!is_unified_memory_arch("gfx1100")); + } + + #[test] + fn oom_guard_auto_decision_matrix() { + // Known unified-memory APU: guard on regardless of host swap — GPU + // allocations land in RAM either way. + assert!(oom_guard_auto_for(Some("gfx1151"), Some(true))); + assert!(oom_guard_auto_for(Some("gfx1151"), Some(false))); + assert!(oom_guard_auto_for(Some("gfx1103"), None)); + // Known discrete GPU: overshoot is a failed hipMalloc; stand down. + assert!(!oom_guard_auto_for(Some("gfx1100"), Some(true))); + assert!(!oom_guard_auto_for(Some("gfx942"), None)); + assert!(!oom_guard_auto_for(Some("gfx1201"), Some(false))); + // Unrecognized arch: fail safe. + assert!(oom_guard_auto_for(Some("gfx9999"), Some(true))); + // No GPU arch in this process: the host's own lethality decides. + assert!(!oom_guard_auto_for(None, Some(true))); + assert!(oom_guard_auto_for(None, Some(false))); + // Unreadable /proc/meminfo: fail safe. + assert!(oom_guard_auto_for(None, None)); + } + + #[test] + fn swap_total_parses_from_meminfo() { + let with_swap = "MemTotal: 130000000 kB\nSwapTotal: 2000000 kB\nSwapFree: 2000000 kB\n"; + assert_eq!(swap_total_kb_from_meminfo(with_swap), Some(2_000_000)); + let no_swap = "MemTotal: 130000000 kB\nSwapTotal: 0 kB\n"; + assert_eq!(swap_total_kb_from_meminfo(no_swap), Some(0)); + assert_eq!(swap_total_kb_from_meminfo("MemTotal: 100 kB\n"), None); } #[test] fn oom_guard_schema_field_is_process_scoped_with_env_compat() { let field = field("memory.oom_guard").expect("oom_guard schema field"); assert_eq!(field.env_compat, Some("HIPFIRE_OOM_GUARD")); - assert!(matches!(field.default.to_value(), ConfigValue::Bool(true))); + // Default is the string "auto": the deployment-class decision, not a + // blanket on/off. + assert!(matches!( + field.default.to_value(), + ConfigValue::String(v) if v == "auto" + )); + assert!(matches!(field.rule, ValueRule::AutoBool)); assert!(!field.include_builtin_in_process_config); + // The AutoBool rule must accept all three spellings end to end. + assert!(field.validate(&ConfigValue::Bool(false)).is_ok()); + assert!(field.validate(&ConfigValue::Bool(true)).is_ok()); + assert!(field + .validate(&ConfigValue::String("auto".to_string())) + .is_ok()); + assert!(field + .validate(&ConfigValue::String("sometimes".to_string())) + .is_err()); } #[test] diff --git a/crates/rdna-compute/map.md b/crates/rdna-compute/map.md index c8b7ea6b3..66d9d9dc8 100644 --- a/crates/rdna-compute/map.md +++ b/crates/rdna-compute/map.md @@ -23,13 +23,13 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside | File | Lines | Public items | Tests | |---|---:|---:|---:| -| [`src/arch_caps.rs`](src/arch_caps.rs) | 712 | 55 | 19 | +| [`src/arch_caps.rs`](src/arch_caps.rs) | 731 | 57 | 19 | | [`src/attention.rs`](src/attention.rs) | 14,946 | 211 | 3 | | [`src/bin/hipfire-kernel-hash.rs`](src/bin/hipfire-kernel-hash.rs) | 142 | 0 | 0 | | [`src/cdna/gfx942.rs`](src/cdna/gfx942.rs) | 578 | 10 | 1 | | [`src/cdna/mod.rs`](src/cdna/mod.rs) | 11 | 1 | 0 | | [`src/compiler.rs`](src/compiler.rs) | 2,266 | 8 | 24 | -| [`src/dispatch.rs`](src/dispatch.rs) | 5,224 | 113 | 14 | +| [`src/dispatch.rs`](src/dispatch.rs) | 5,230 | 113 | 14 | | [`src/embedding.rs`](src/embedding.rs) | 410 | 10 | 0 | | [`src/feature_flags.rs`](src/feature_flags.rs) | 908 | 12 | 6 | | [`src/flash_attn_ck.rs`](src/flash_attn_ck.rs) | 1,775 | 26 | 15 | @@ -39,7 +39,7 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside | [`src/gemv.rs`](src/gemv.rs) | 15,991 | 235 | 0 | | [`src/graph.rs`](src/graph.rs) | 556 | 33 | 0 | | [`src/kernels.rs`](src/kernels.rs) | 8,088 | 1228 | 37 | -| [`src/kv_slots.rs`](src/kv_slots.rs) | 420 | 9 | 10 | +| [`src/kv_slots.rs`](src/kv_slots.rs) | 459 | 9 | 10 | | [`src/lib.rs`](src/lib.rs) | 88 | 26 | 1 | | [`src/moe.rs`](src/moe.rs) | 1,742 | 27 | 0 | | [`src/norm.rs`](src/norm.rs) | 6,170 | 90 | 0 | @@ -52,11 +52,11 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside | [`src/replay.rs`](src/replay.rs) | 9,126 | 78 | 81 | | [`src/sampling.rs`](src/sampling.rs) | 1,765 | 22 | 3 | | [`src/scratch.rs`](src/scratch.rs) | 1,415 | 21 | 0 | -| [`src/slot_pool.rs`](src/slot_pool.rs) | 236 | 11 | 7 | +| [`src/slot_pool.rs`](src/slot_pool.rs) | 244 | 11 | 7 | ### Public API surface -- [`src/arch_caps.rs`](src/arch_caps.rs): `ArchCaps`, `new`, `should_use_mmq`, `is_gfx906`, `is_gfx908`, `is_gfx1010`, `is_gfx1011`, `is_gfx1012`, `is_gfx1030`, `is_gfx1031`, `is_gfx1032`, `is_gfx1100`, +43 more +- [`src/arch_caps.rs`](src/arch_caps.rs): `ArchCaps`, `note_process_gpu_arch`, `process_gpu_arch`, `new`, `should_use_mmq`, `is_gfx906`, `is_gfx908`, `is_gfx1010`, `is_gfx1011`, `is_gfx1012`, `is_gfx1030`, `is_gfx1031`, +45 more - [`src/attention.rs`](src/attention.rs): `attention_q8_0_kv_independent_lds_bytes`, `attention_q8_0_kv_independent_max_lane_capacity`, `q8_flash_tile_size`, `dspark_stage_kv`, `triattn_accumulate`, `attention_f32`, `attention_flash`, `attention_flash_gqa`, `attention_flash_gqa_fused`, `attention_gqa_warp`, `attention_gqa_warp_dv`, `kv_cache_write_hfq4`, +199 more - [`src/bin/hipfire-kernel-hash.rs`](src/bin/hipfire-kernel-hash.rs): — - [`src/cdna/gfx942.rs`](src/cdna/gfx942.rs): `Gfx942Device`, `try_gfx942`, `mq2_lloyd_moe_gate_up_wave64`, `mq2_lloyd_moe_gate_up_wave64x8_candidate`, `mq_rotate_x_wave64_batched`, `mq2_lloyd_moe_down_expanded_wave64`, `mq2_lloyd_moe_down_residual_wave64`, `indexer_top_k_buf_parallel`, `grouped_olora_e8`, `grouped_olora_e8_wave64x4_candidate` @@ -100,6 +100,6 @@ _Generated by `scripts/check-crate-maps.py` from the tree — do not edit inside ### Totals -- 30 modules · 110,994 lines · 2757 public items · 225 tests · 192 examples +- 30 modules · 111,066 lines · 2759 public items · 225 tests · 192 examples diff --git a/crates/rdna-compute/src/arch_caps.rs b/crates/rdna-compute/src/arch_caps.rs index a3bcbc19c..10858c814 100644 --- a/crates/rdna-compute/src/arch_caps.rs +++ b/crates/rdna-compute/src/arch_caps.rs @@ -78,6 +78,25 @@ pub struct ArchCaps { flags: std::sync::Arc, } +// Process-wide GPU arch, recorded at `Gpu::init[_with_device]` so config-time +// policy that runs without a `Gpu` handle in hand (the `kv_slots` memory +// preflight) can still classify the deployment as unified-memory APU vs +// discrete GPU. First init wins: a mixed APU + dGPU process is not a +// supported topology, and re-inits (device swaps in harnesses) must not +// silently change which class the OOM guard applies to. +static PROCESS_GPU_ARCH: std::sync::OnceLock = std::sync::OnceLock::new(); + +/// Record the arch this process initialized its primary GPU with. First call +/// wins; later calls are no-ops. +pub fn note_process_gpu_arch(arch: &str) { + let _ = PROCESS_GPU_ARCH.set(arch.to_string()); +} + +/// The arch recorded at GPU init, if any GPU has been initialized. +pub fn process_gpu_arch() -> Option<&'static str> { + PROCESS_GPU_ARCH.get().map(|s| s.as_str()) +} + impl ArchCaps { pub fn new(arch: &str, flags: std::sync::Arc) -> Self { // Atoms diff --git a/crates/rdna-compute/src/dispatch.rs b/crates/rdna-compute/src/dispatch.rs index f269ce54d..b249ad28e 100644 --- a/crates/rdna-compute/src/dispatch.rs +++ b/crates/rdna-compute/src/dispatch.rs @@ -1112,6 +1112,12 @@ impl Gpu { // `gfx10-1-generic` (covers Navi 10/12/14) without per-arch JIT // cache fragmentation. Empty / unset preserves prior behavior. let detected_arch = hip.get_arch(id).unwrap_or_else(|_| "gfx1010".to_string()); + // Record the DETECTED (not compile-target-overridden) arch for + // config-time policy that runs without a Gpu handle in hand — the + // kv_slots memory preflight resolves its auto mode (unified-memory + // APU vs discrete GPU) from physical topology, which a + // HIPFIRE_TARGET_ARCH override does not change. + crate::arch_caps::note_process_gpu_arch(&detected_arch); let arch = hipfire_config::developer_var("HIPFIRE_TARGET_ARCH") .ok() .filter(|s| !s.is_empty()) diff --git a/crates/rdna-compute/src/kv_slots.rs b/crates/rdna-compute/src/kv_slots.rs index af2cee2c4..cfc202f22 100644 --- a/crates/rdna-compute/src/kv_slots.rs +++ b/crates/rdna-compute/src/kv_slots.rs @@ -259,11 +259,15 @@ pub fn build_tiles(slot_query_counts: &[usize], br: usize) -> (Vec, Vec Option { /// Refuse a planned allocation that would either exceed the deployment target's /// VRAM or leave this box without enough headroom to stay responsive. /// -/// Skipped entirely when `memory.oom_guard=false` (`HIPFIRE_OOM_GUARD=0`): +/// Gated by `memory.oom_guard` (compat `HIPFIRE_OOM_GUARD`), default `auto`: /// the deployment-target ceiling and the headroom check both assume GPU -/// memory comes from system RAM, which is true on unified-memory APUs and -/// false on a discrete-GPU dev box. The skip prints once to stderr so a -/// disabled guard is visible in a log instead of reading as a pass. +/// memory comes from system RAM, so `auto` keeps them on only for +/// unified-memory APU architectures (and, when no GPU arch is known in this +/// process, for hosts without swap). See `hipfire_config::oom_guard_effective`. /// /// `planned_bytes` must be the TOTAL the caller is about to hold live at once, /// not a single buffer. Returns `Err` with an actionable message; callers should @@ -305,17 +309,25 @@ pub fn mem_available_bytes() -> Option { /// forbidden by scripts/check-env-docs.py. Harnesses live in `examples/`, /// which is exempt, so they read any override there and pass it in. pub fn preflight_alloc(planned_bytes: u64, budget_bytes: u64, what: &str) -> Result<(), String> { - if !hipfire_config::oom_guard_enabled() { - static SKIP_NOTE: std::sync::Once = std::sync::Once::new(); - SKIP_NOTE.call_once(|| { + // Gpu::init records the detected arch; until it runs (or in GPU-less + // processes) the resolver falls back to host swap state. + if !hipfire_config::oom_guard_effective(crate::arch_caps::process_gpu_arch()) { + static INACTIVE_NOTE: std::sync::Once = std::sync::Once::new(); + INACTIVE_NOTE.call_once(|| { eprintln!( - "[kv_slots] memory preflight guard disabled (memory.oom_guard=false); \ + "[kv_slots] memory preflight guard inactive (memory.oom_guard); \ allocations will not be refused before allocate" ); }); return Ok(()); } + preflight_checks(planned_bytes, budget_bytes, what) +} +/// The guard's actual checks, with no config gating — deterministic on every +/// machine so the unit tests below assert refusal behavior rather than this +/// box's config. [`preflight_alloc`] is the config-gated production entry. +pub(crate) fn preflight_checks(planned_bytes: u64, budget_bytes: u64, what: &str) -> Result<(), String> { let budget = budget_bytes; let gib = |b: u64| b as f64 / 1073741824.0; @@ -364,15 +376,19 @@ mod tests { #[test] fn preflight_refuses_over_target_budget() { // 64 GiB against the 32 GiB R9700 target: must refuse even though this - // dev box has 125 GiB. - let e = preflight_alloc(64 * 1024 * 1024 * 1024, R9700_VRAM_BYTES, "test").unwrap_err(); + // dev box has 125 GiB. The pure checks, not the config-gated entry — + // this test must refuse identically whether this box is an APU or a + // dGPU. + let e = preflight_checks(64 * 1024 * 1024 * 1024, R9700_VRAM_BYTES, "test").unwrap_err(); assert!(e.contains("deployment target"), "unexpected message: {e}"); } #[test] fn preflight_allows_a_small_allocation() { - // 64 MiB is under budget and under any plausible MemAvailable. - assert!(preflight_alloc(64 * 1024 * 1024, R9700_VRAM_BYTES, "test").is_ok()); + // 64 MiB is under budget and under any plausible MemAvailable. Pure + // checks: this test must pass regardless of this box's oom_guard + // setting. + assert!(preflight_checks(64 * 1024 * 1024, R9700_VRAM_BYTES, "test").is_ok()); } #[test] diff --git a/crates/rdna-compute/src/slot_pool.rs b/crates/rdna-compute/src/slot_pool.rs index c6661a9ee..c78f9b954 100644 --- a/crates/rdna-compute/src/slot_pool.rs +++ b/crates/rdna-compute/src/slot_pool.rs @@ -230,7 +230,15 @@ mod tests { // the budget, so `new` correctly returned Ok and the `unwrap_err` here // panicked. The test's comment said "8.7 TB", off by 1000x; the // refusal it is checking was never actually being exercised. - let e = SlotPool::new(8, 4_000_000, PPB).unwrap_err(); + // + // Calls `preflight_checks` directly (not `SlotPool::new`) because + // `SlotPool::new` goes through the config-gated `preflight_alloc`, + // whose `auto` default stands down on discrete GPUs and GPU-less CI + // runners. The unguarded checks are deterministic on every machine. + let cap = 4_000_000usize.div_ceil(PAGE_TOKENS) * PAGE_TOKENS; + let total = (cap * PPB) as u64 * 8 * 2; + let e = crate::kv_slots::preflight_checks(total, R9700_VRAM_BYTES, "SlotPool arena") + .unwrap_err(); assert!(e.contains("budget") || e.contains("GiB"), "unexpected: {e}"); } } diff --git a/docs/CONFIG.md b/docs/CONFIG.md index 829057a97..4c6efba52 100644 --- a/docs/CONFIG.md +++ b/docs/CONFIG.md @@ -109,7 +109,7 @@ Stable/default-on and safety controls: | `kernel.rocblas_off` | `false` | Disable rocBLAS dispatch. | | `fusions.force_unfused` | `false` | Force supported projection paths unfused. | | `speculation.dflash_tree` | `false` | Enable DDTree tree-SWOR verification. | -| `memory.oom_guard` | `true` | Memory preflight OOM guard (`kv_slots::preflight_alloc` and the bench-sweep headroom check). Refuses oversized allocations before they are made. Default on because on unified-memory APUs (Strix Halo) GPU memory is system RAM with no swap — an overshoot is a global OOM that kills the desktop. Set `false` (`HIPFIRE_OOM_GUARD=0`) on discrete-GPU boxes, where an overshoot is a plain failed hipMalloc. | +| `memory.oom_guard` | `auto` | Memory preflight OOM guard (`kv_slots::preflight_alloc` and the bench-sweep headroom check). Refuses oversized allocations before they are made. `auto` decides by deployment class: on for unified-memory APU archs (gfx1035/1036/1103/1150/1151/1152 — GPU allocations come out of system RAM, so an overshoot is a global OOM that kills the desktop), off for discrete GPUs (an overshoot is a plain failed hipMalloc), and for GPU-less processes by host swap state (no swap → on). `true`/`false` force the decision. | The following default-off keys are experimental kernel-route overrides. They are typed booleans, process-scoped, and visible in `hipfire config list` with diff --git a/docs/env-vars.md b/docs/env-vars.md index e7a4abb05..7d51496ca 100644 --- a/docs/env-vars.md +++ b/docs/env-vars.md @@ -1051,5 +1051,5 @@ When adding a user-facing knob: |---|---|---| | `HIPFIRE_ATTN_TILE_SIZE` | `128` | Tile size for the batched attention tile+reduce path. Must be a positive multiple of 32; anything else falls back to 128. Resolved once via `Gpu::attn_tile_size()`. **Raising it is safe; lowering it increases `max_tiles` and therefore the `partials` bytes per query row, which can exceed buffers sized elsewhere against the 128 default.** | | `HIPFIRE_VRAM_BUDGET_BYTES` | 32 GiB | Deployment-target VRAM ceiling used by the SP1 benchmark harnesses' preflight. Read by `examples/`, not by production code. | -| `HIPFIRE_OOM_GUARD` | on | Typed key `memory.oom_guard`. Disables the production memory preflight (`kv_slots::preflight_alloc`, the `SlotPool` arena check, and the CLI bench-sweep headroom check) when set to `0`/`false`/`off`. Default on: on unified-memory APUs (Strix Halo) GPU memory is system RAM with no swap, so an overshoot is a global OOM that kills the desktop; on a discrete GPU an overshoot is a plain failed hipMalloc, so multi-daemon dev boxes may opt out. `scripts/run-bounded.sh` remains the hard backstop either way. | +| `HIPFIRE_OOM_GUARD` | `auto` | Typed key `memory.oom_guard`. Memory preflight OOM guard (`kv_slots::preflight_alloc`, the `SlotPool` arena check, and the CLI bench-sweep headroom check). `auto` decides by deployment class: on for unified-memory APU archs (gfx1035/1036/1103/1150/1151/1152 — GPU memory is system RAM, an overshoot is a global OOM that kills the desktop), off for recognized discrete GPUs (an overshoot is a plain failed hipMalloc), and for processes with no known GPU arch by host swap state (no swap → on; unreadable → on). `1`/`true`/`0`/`false` force it either way. The auto decision is logged once to stderr with its reason. `scripts/run-bounded.sh` remains the hard backstop. | | `HIPFIRE_MEM_CAP` | `24G` | Read by `scripts/run-bounded.sh`, not by the binaries: cgroup `MemoryMax` for a gated run. Exit 137 means the cap fired — shrink the configuration rather than raising it. |