Spill gated to discrete GPUs; KV window floor 4096 walked by fit_start; output head to pinned host memory as the last step

The draft-layer drop, the 16-slot prompt-path loan, the own-buffer fallback and the head-to-host move run only on a discrete GPU (strata::spill_allowed:
not Metal, not an integrated GPU such as Strix Halo), with a CHECK for the gate in arch_defaults_test. qsa_kv_resident_min is 4096 (20480 on a Metal build;
STRATA_KV_RESIDENT_MIN overrides): windows of 4096 and 8192 gave the same tokens as 20480 with the cache fixed. fit_start walks the window 20480 / 8192 / 4096 only as
far as the card needs. When every step leaves the cache no slot, the output head (Q6_K, 0.41 GiB) moves to mapped pinned host memory (about 3x slower decode).
Measured on the 8 GB mimic at 128K: DFlash2 alone 391 MiB free, 63-70 tok/s, 32K prompt 19.8 s; MTP asked: layer dropped, 555 MiB free, 61-71 tok/s, 32K prompt 13.4 s;
2 requests 277 MiB free, 46 tok/s per user; 4 requests 171 MiB free, 15 tok/s per user. 12 GB cells unchanged.

include/strata/core/native_head.hpp

@@ -21,6 +21,10 @@ public:
 
     bool load(const std::vector<std::string>& shards, int64_t n_in, int64_t n_out, std::string& err);
     void unload();
+    /// Moves the loaded weights from VRAM to mapped pinned host memory (a card too short for it: every logits row then reads them over PCIe).  False: not loaded, already
+    /// on the host, or a repacked layout that stays in VRAM.
+    bool to_host(std::string& err);
+    bool on_host() const { return host_ != nullptr; }
     bool run(const float* mixed, float* logits, void* stream, std::string& err) const;
     uint64_t weight_bytes() const { return bytes_; }
     bool loaded() const { return weights_ != nullptr; }
@@ -34,6 +38,7 @@ public:
 
 private:
     void* weights_ = nullptr;
+    void* host_ = nullptr;   // the pinned host buffer weights_ aliases when the head did not fit in VRAM
     void* scratch_ = nullptr;
     uint64_t bytes_ = 0;
     int n_in_ = 0, n_out_ = 0;

include/strata/core/spill_gate.hpp

@@ -0,0 +1,10 @@
+#pragma once
+// Which cards may be moved off their GPU memory when a start is short of it (the draft layer dropped before it loads, the output head in pinned host memory, a prompt-path
+// loan that leaves 16 cache slots, the KV window walked below 20480).  Only a discrete GPU: its VRAM is a separate budget and the host's RAM is where the spilled work goes.  A
+// Mac (Metal) and an integrated GPU (Strix Halo gfx1151, any APU) have ONE pool: there the "free VRAM" figure is the OS's available RAM, a low number means something else,
+// and moving a buffer to host memory frees nothing.  Pure, so the gate has a test that needs no GPU (src/core/arch_defaults_test.cpp).
+namespace strata {
+
+inline constexpr bool spill_allowed(bool metal_build, bool integrated_gpu) { return !metal_build && !integrated_gpu; }
+
+}  // namespace strata

src/core/arch_defaults_test.cpp

@@ -3,6 +3,7 @@
 #include "strata/core/gfx12_platform_defaults.hpp"
 #include "strata/core/sm120_series.hpp"
 #include "strata/core/defaults_026.hpp"
+#include "strata/core/spill_gate.hpp"
 
 #include <cstdio>
 #include <cstdlib>
@@ -598,5 +599,12 @@ int main() {
         CHECK(ple_batch_resolve(nullptr, rt130, false, 12, 0));
         CHECK(ple_batch_resolve(nullptr, rt128, false, 8, 6));
     }
+    {   // the spill gate: only a discrete GPU may have its start moved off VRAM (draft layer dropped, head on the host, loan of 16, KV window below 20480)
+        using strata::spill_allowed;
+        CHECK(spill_allowed(false, false));     // CUDA / HIP discrete
+        CHECK(!spill_allowed(true, false));     // Metal: one pool, "free VRAM" is the OS's available RAM
+        CHECK(!spill_allowed(false, true));     // an integrated GPU (Strix Halo gfx1151, APUs): the same
+        CHECK(!spill_allowed(true, true));
+    }
     return fails ? 1 : 0;
 }

src/core/layer.cpp

@@ -768,7 +768,21 @@ uint64_t kv_pool_bytes(const QsaShapes& s, int64_t pages, bool hybrid, bool int8
 
 void qsa_set_kv_resident(int64_t cells) { g_kv_resident = cells > 0 ? cells : 0; }
 int64_t qsa_kv_resident() { return g_kv_resident; }
-int64_t qsa_kv_resident_min() { return 20480; }
+// The least KV window a streamed session keeps in VRAM (cells): 4096 (was 20480; a window of 4096 and of 8192 gave the same tokens as 20480 with the cache fixed and route tail
+// skip off, 2 requests at 128K, 20K- and 8K-token prompts; setup still writes 32768, fit_start walks 20480 / 8192 / 4096 only on a card that is short).  STRATA_KV_RESIDENT_MIN=N overrides (the
+// selection reads at most idx_top_k + idx_block - 1 cells per row, so a window of a few thousand cells holds it; the pool must hold the blocks of every row of a verify window).
+int64_t qsa_kv_resident_min() {
+    static const int64_t v = [] {
+        const char* e = std::getenv("STRATA_KV_RESIDENT_MIN");
+        const long long n = e != nullptr ? std::atoll(e) : 0;
+        #if defined(STRATA_GPU_METAL)
+        return n >= 1024 ? (int64_t) n : (int64_t) 20480;   // a Mac keeps the old floor: its pool is the host's RAM (strata/core/spill_gate.hpp)
+#else
+        return n >= 1024 ? (int64_t) n : (int64_t) 4096;
+#endif
+    }();
+    return v;
+}
 uint64_t qsa_kv_host_bytes() { return g_kv_host_bytes; }
 
 uint64_t qsa_state_bytes(const ModelGeometry& g, int64_t max_cells, bool with_rope, int64_t ring_cells) {

src/core/native_head.cpp

@@ -1,4 +1,5 @@
 #include "strata/core/native_head.hpp"
+#include "strata/core/spill_gate.hpp"
 #include "strata/artifact/gguf_reader.hpp"
 #include "strata/kernels/iq_kernels.hpp"
 #include "strata/kernels/native_mmvq.hpp"
@@ -21,9 +22,11 @@ void NativeHead::unload() {
     if (weights_) strata::kernels::native_q6_k_unpack(weights_);
     if (weights_) strata::kernels::native_q8_0_sm60_release(weights_);
     if (scratch_) cudaFree(scratch_);
-    if (weights_) cudaFree(weights_);
+    if (host_) cudaFreeHost(host_);   // the mapped-host fallback (weights_ is its device alias)
+    else if (weights_) cudaFree(weights_);
     scratch_ = nullptr;
     weights_ = nullptr;
+    host_ = nullptr;
 }
 
 bool NativeHead::load(const std::vector<std::string>& shards, int64_t n_in, int64_t n_out, std::string& err) {
@@ -54,19 +57,46 @@ bool NativeHead::load(const std::vector<std::string>& shards, int64_t n_in, int6
         }
         void* weights = nullptr;
         void* scratch = nullptr;
-        cudaError_t status = cudaMalloc(&weights, bytes);
+        void* host = nullptr;
+        // SHORT OF VRAM: the head in mapped pinned host memory.  The head (about 0.4 GiB) is the one buffer a card that holds its sessions and weights but not this can
+        // spill whole; every logits row then reads it over PCIe (about 18 ms per row at 28 GB/s), so decode is slower, but the start works.  Taken only when the
+        // allocation in VRAM fails (or STRATA_HEAD_HOST=1 asks for it); the GEMV kernels read the same pointer either way.  Q5_K and the plain MMVQ types: the
+        // repacked Q6_K / SM60 layouts rewrite the weights in place and stay in VRAM.
+        const bool ask_host = std::getenv("STRATA_HEAD_HOST") != nullptr;
+        cudaError_t status = ask_host ? cudaErrorMemoryAllocation : cudaMalloc(&weights, bytes);
+        bool discrete = false;
+        { int d = 0; cudaDeviceProp pr{}; if (cudaGetDevice(&d) == cudaSuccess && cudaGetDeviceProperties(&pr, d) == cudaSuccess) discrete = strata::spill_allowed(false, pr.integrated != 0); else (void) cudaGetLastError(); }
+        if (status != cudaSuccess && discrete && std::getenv("STRATA_HEAD_HOST_OFF") == nullptr && !strata::kernels::native_q6_k_packed_enabled()) {
+            (void) cudaGetLastError();
+            void* hp = nullptr;
+            void* dp = nullptr;
+            if (cudaHostAlloc(&hp, bytes, cudaHostAllocMapped | cudaHostAllocPortable) == cudaSuccess && cudaHostGetDevicePointer(&dp, hp, 0) == cudaSuccess) {
+                std::memcpy(hp, gguf.tensor_data(*tensor), bytes);
+                host = hp;
+                weights = dp;
+                status = cudaSuccess;
+                std::fprintf(stderr, "strata generate: short of GPU memory: the output head (%llu MiB) stays in pinned host memory and is read over PCIe for every logits row: "
+                                     "decode is slower%s", (unsigned long long) (bytes >> 20), "\n");
+            } else {
+                if (hp) cudaFreeHost(hp);
+                (void) cudaGetLastError();
+            }
+        }
         if (status == cudaSuccess)
             status = cudaMalloc(&scratch, strata::kernels::native_q8_1_bytes((int) n_in, 1));
         strata::platform::advise_willneed(gguf.tensor_data(*tensor), bytes);
-        if (status == cudaSuccess)
+        if (status == cudaSuccess && host == nullptr)
             status = cudaMemcpy(weights, gguf.tensor_data(*tensor), bytes, cudaMemcpyHostToDevice);
         if (status != cudaSuccess) {
+            if (host) cudaFreeHost(host);
+            weights = host ? nullptr : weights;
             if (scratch) cudaFree(scratch);
             if (weights) cudaFree(weights);
             err = std::string("native head upload: ") + cudaGetErrorString(status);
             return false;
         }
         weights_ = weights;
+        host_ = host;
         { int d = -1; if (cudaGetDevice(&d) == cudaSuccess) device_ = d; }
         scratch_ = scratch;
         bytes_ = bytes;
@@ -84,6 +114,28 @@ bool NativeHead::load(const std::vector<std::string>& shards, int64_t n_in, int6
     }
 }
 
+bool NativeHead::to_host(std::string& err) {
+    if (weights_ == nullptr || host_ != nullptr || (type_ == 14 && strata::kernels::native_q6_k_packed_enabled()) || (type_ == 8 && std::getenv("STRATA_Q8_SM60") != nullptr)) return false;   // (repacked in place: those stay in VRAM)
+    { int d = 0; cudaDeviceProp pr{}; if (cudaGetDevice(&d) != cudaSuccess || cudaGetDeviceProperties(&pr, d) != cudaSuccess || !strata::spill_allowed(false, pr.integrated != 0)) { (void) cudaGetLastError(); return false; } }   // (Q6_K / Q8_0 may be repacked in place: those stay)
+    void* hp = nullptr;
+    void* dp = nullptr;
+    if (cudaHostAlloc(&hp, bytes_, cudaHostAllocMapped | cudaHostAllocPortable) != cudaSuccess) {
+        (void) cudaGetLastError();
+        err = "native head: no pinned host memory for the head";
+        return false;
+    }
+    if (cudaHostGetDevicePointer(&dp, hp, 0) != cudaSuccess || cudaMemcpy(hp, weights_, bytes_, cudaMemcpyDeviceToHost) != cudaSuccess) {
+        (void) cudaGetLastError();
+        cudaFreeHost(hp);
+        err = "native head: the head could not be moved to host memory";
+        return false;
+    }
+    cudaFree(weights_);
+    weights_ = dp;
+    host_ = hp;
+    return true;
+}
+
 bool NativeHead::run(const float* mixed, float* logits, void* stream, std::string& err) const {
     if (!loaded() || !mixed || !logits || !stream) {
         err = "native head requires loaded weights, device buffers and an explicit stream";

src/program/generate.cpp

@@ -74,6 +74,7 @@
 #include "strata/kernels/kv_q4.hpp"
 #include "strata/kernels/qsa.hpp"
 #include "strata/core/native_head.hpp"
+#include "strata/core/spill_gate.hpp"
 #include "strata/core/verify.hpp"
 #include "strata/core/mtp.hpp"
 #include "strata/core/gated_proposal.hpp"
@@ -267,6 +268,21 @@ bool under_wddm() {
 // How many cache slots a prompt-path loan leaves behind (128 on a card with room; a card too short for that lends against 16, see the cache sizing).
 static int64_t g_lend_keep = 128;
 
+// May a start short of VRAM be moved off the GPU (strata/core/spill_gate.hpp)?  A discrete card only: not Metal, not an integrated GPU (Strix Halo and the other APUs share the host's RAM).
+static bool spill_discrete_gpu() {
+#if defined(STRATA_GPU_METAL)
+    return strata::spill_allowed(true, false);
+#else
+    static const bool ok = [] {
+        int d = 0;
+        cudaDeviceProp pr{};
+        if (cudaGetDevice(&d) != cudaSuccess || cudaGetDeviceProperties(&pr, d) != cudaSuccess) { (void) cudaGetLastError(); return false; }
+        return strata::spill_allowed(false, pr.integrated != 0);
+    }();
+    return ok;
+#endif
+}
+
 std::string vram_free_note() {
     (void) cudaGetLastError();
     size_t free_b = 0, total_b = 0;
@@ -6707,7 +6723,7 @@ int main(int argc, char** argv) {
     // the sessions) the start used to stop at "no VRAM is left for the expert cache" or run out of memory in the head upload.  Without the layer
     // the model still answers (the verify window drafts from the prompt alone): slower decode, a start that works.  Not with DFlash2 (its own
     // block decides), a split, or STRATA_MTP_KEEP=1.
-    if (!o.mtp.empty() && !use_df2 && !multi_gpu && stages.empty() && std::getenv("STRATA_MTP_KEEP") == nullptr) {
+    if (!o.mtp.empty() && !use_df2 && !multi_gpu && stages.empty() && std::getenv("STRATA_MTP_KEEP") == nullptr && spill_discrete_gpu()) {
         std::error_code fec1, fec2;
         const uint64_t layer_b = (uint64_t) std::filesystem::file_size(std::filesystem::path(o.mtp) / "dense.bin", fec1) +
                                  (uint64_t) std::filesystem::file_size(std::filesystem::path(o.mtp) / "experts.bin", fec2);
@@ -6715,7 +6731,7 @@ int main(int argc, char** argv) {
         const uint64_t head_b = (!o.native_head_gguf.empty() && wo_mtp != nullptr)
                                     ? (uint64_t) wo_mtp->ne1 * (uint64_t) (g.n_embd / 256) * 176ull : 0;   // Q5_K: 176 bytes per 256 values
         const uint64_t free_now = (uint64_t) strata::core::device_free_bytes();
-        const uint64_t need = layer_b + head_b + (220ull << 20) + (1250ull << 20);   // + the cache the loan needs (214 slots, 543 MiB) + what is allocated after it and the drafter beyond DFlash2 (measured: DFlash2 alone ends 391 MiB free, MTP 41 at the same cache)
+        const uint64_t need = layer_b + head_b + (220ull << 20) + (1550ull << 20);   // + the cache the loan needs (214 slots, 543 MiB) + what is allocated after it and the drafter beyond DFlash2 (measured: DFlash2 alone ends 391 MiB free, MTP 41 at the same cache)
         if (!fec1 && !fec2 && free_now > 0 && free_now < need) {
             std::fprintf(stderr, "strata generate: short of GPU memory: %.2f GiB is free, and the draft layer (%.2f GiB), the head (%.2f GiB) and a working cache "
                                  "need about %.2f GiB: this start runs without the draft layer (decode is slower; STRATA_MTP_KEEP=1 keeps it)\n",
@@ -7586,7 +7602,7 @@ int main(int argc, char** argv) {
         // free and stopped the first long prompt ("mmq: out of memory"), or gave the prompt path its own 600 MiB of buffers on top and stopped with "do not
         // fit".  The loan may leave 16 slots instead: the cache is smaller, prompts are read in smaller chunks and slower, the reserve stays, and the start works.
         if (o.prefill_chunk > 0 && pf_borrow && o.prefill_auto && o.expert_cache < min_slots && !o.vram_reserve_given && !multi_gpu && stages.empty() &&
-            std::getenv("STRATA_SPILL_OFF") == nullptr) {
+            std::getenv("STRATA_SPILL_OFF") == nullptr && spill_discrete_gpu()) {
             constexpr int64_t kTightKeep = 16;
             const int64_t tight_min = ((int64_t) strata::prefill::Prefill::bytes_needed(g, ss, 256) + blob - 1) / blob + kTightKeep;
             if (tight_min < min_slots) {
@@ -7633,7 +7649,7 @@ int main(int argc, char** argv) {
         // here instead, at the largest chunk that still leaves a cache (1024, 512, 256 tokens), and give the cache the rest: slower (a smaller
         // cache, smaller prompt chunks), never a refused start.  Not with a split, a cache or chunk given on the command line, or STRATA_SPILL_OFF=1.
         if (auto_cache && pf_borrow && o.prefill_auto && !o.vram_reserve_given && o.expert_cache < min_slots && !multi_gpu && stages.empty() && !profile.empty() &&
-            std::getenv("STRATA_SPILL_OFF") == nullptr) {
+            std::getenv("STRATA_SPILL_OFF") == nullptr && spill_discrete_gpu()) {
             constexpr int64_t kSpillFloorMib = 100;   // the verify arena, the head check and the cuBLAS workspace beside the late allocations
             int64_t pick_chunk = 0, pick_slots = 0, pick_own = 0, pick_late = 0;
             for (const int64_t c : {(int64_t) 1024, (int64_t) 512, (int64_t) 256}) {
@@ -7657,6 +7673,35 @@ int main(int argc, char** argv) {
                 reserve_adapted = true;
             }
         }
+        // SHORT OF VRAM: THE OUTPUT HEAD ON THE HOST, LAST.  When every step above left the cache no slot (4 requests at 128K on an 8 GB card: the sessions and the weights
+        // leave 0.63 GiB after the head, and the cache needs 543 MiB beside a reserve), the head (0.41 GiB) moves to mapped pinned host memory and every logits row reads it
+        // over PCIe: decode is about 3x slower (12.7 against 36.1 tok/s per user, 2 requests on a 12 GB mimic), but the start works.  Only a card that would otherwise
+        // refuse; STRATA_HEAD_HOST_OFF=1 keeps the refusal.
+        if (o.expert_cache == 0 && native_head.loaded() && !native_head.on_host() && !o.vram_reserve_given && !multi_gpu && stages.empty() &&
+            std::getenv("STRATA_HEAD_HOST_OFF") == nullptr && std::getenv("STRATA_SPILL_OFF") == nullptr && spill_discrete_gpu()) {
+            std::string herr;
+            if (native_head.to_host(herr)) {
+                free_b = strata::core::device_free_bytes();
+                strata::core::AutoCacheSlots a2 = strata::core::auto_cache_slots((int64_t) free_b, reserve, blob, min_slots, strata::qwen_defaults_026() ? 0 : strata::core::kLateAllocMib);
+                int64_t s4 = a2.slots;
+                if (!profile.empty()) s4 = std::min<int64_t>(s4, (int64_t) profile.size());
+                if (s4 < min_slots && o.vram_reserve_mib > kSmallReserveMib) {   // the reserve lowered for the smaller cache, as above
+                    const int64_t fit_mib = ((int64_t) free_b - mtp_bind - pipe_first - min_slots * blob) / (1 << 20) - prefill_mib;
+                    if (fit_mib >= kSmallReserveMib) {
+                        const int r = (int) std::min<int64_t>(fit_mib, o.vram_reserve_mib);
+                        s4 = ((int64_t) free_b - ((((int64_t) r + prefill_mib) << 20) + mtp_bind + pipe_first)) / blob;
+                        if (!profile.empty()) s4 = std::min<int64_t>(s4, (int64_t) profile.size());
+                        o.vram_reserve_mib = r;
+                        reserve_adapted = true;
+                    }
+                }
+                o.expert_cache = (int) std::max<int64_t>(s4, 0);
+                std::fprintf(stderr, "strata generate: short of GPU memory: the cache had no room, so the output head (%llu MiB) moves to pinned host memory and is read over PCIe for "
+                                     "every logits row: decode is slower; the cache is now %d slots\n", (unsigned long long) (native_head.weight_bytes() >> 20), o.expert_cache);
+            } else {
+                std::fprintf(stderr, "strata generate: short of GPU memory: the output head stays in VRAM (%s)\n", herr.empty() ? "type or build not movable" : herr.c_str());
+            }
+        }
         if (o.expert_cache == 0) {
             // what is short, and what makes room: the numbers a small card picks from
             const int64_t at_reserve = o.vram_reserve_given ? o.vram_reserve_mib

strata_control/vram_fit.py

@@ -84,7 +84,8 @@ DRAFTER_DELTA_MIB = {"mtp": 0, "dflash2": -119, "dflash2+mtp": 502, "none": -835
 KNOBS = {"cuda": ("kv_resident",), "intel": ("chunk", "reserve", "kv_resident"), "other": ("chunk", "reserve")}
 NOTE_SLOW = "short of GPU memory: more is read from RAM and the SSD, expect it to be slower"
 SLOT_RUN_MIN = 1                                       # the engine's minimum: generate.cpp refuses a start whose auto sizing leaves the cache no slot
-KV_RESIDENT_MIN = 20480                                # qsa_kv_resident_min(): a smaller --kv-resident is raised to this
+KV_RESIDENT_MIN = 4096                                 # qsa_kv_resident_min(): a smaller --kv-resident is raised to this (was 20480; 4096 and 8192 gave the same tokens, RESULTS-DF2MEM.md section 18)
+KV_RESIDENT_STEPS = (20480, 8192, 4096)                # the windows a short start walks down, mildest first (a session at 128K: 0.58 / 0.44 / 0.39 GiB)
 STREAM_EXTRA_BYTES = 1400                              # per cell of context a streamed session still keeps in VRAM (rope table, indexer rows, page table); measured: the CUDA engine's slot session at 128K with 32768 resident is 0.73 GiB (RESULTS-DF2MEM.md section 17), the 32K session 0.56 GiB
 SLOTS_MAX = 8                                    # the engine's batch window (serve/server.py PARALLEL_MAX)
 
@@ -214,10 +215,11 @@ def plan(model: str, vram_gb: float, ctx: int, kv: str = "int8", reserve_mib: in
 
     for knob in knobs:
         if knob == "kv_resident":
-            # the window is only ever made smaller, to the engine's minimum; a start with no window and a context beyond it gets one
-            if kv_fixed or ctx <= KV_RESIDENT_MIN or (kvr and kvr <= KV_RESIDENT_MIN) or not measured(model, slots, vendor, drafter):
+            # the window is only ever made smaller, down the steps to the engine's minimum; a start with no window and a context beyond it gets one
+            cur = kvr if kvr else ctx
+            steps = [w for w in KV_RESIDENT_STEPS if w < cur and w < ctx]
+            if kv_fixed or not steps or not measured(model, slots, vendor, drafter):
                 continue                 # (a start priced without a measured base does not count the window, so moving it would only be a claim)
-            steps = [KV_RESIDENT_MIN]
         elif knob == "chunk":
             steps = [] if prefill_fixed else [c for c in CHUNK_STEPS if c < chunk_of(chunk)]
         else:

tests/terminal/test_vram_fit.py

@@ -164,7 +164,7 @@ class SlotSessions(unittest.TestCase):
         self.assertAlmostEqual(full, 1929, delta=15)
         self.assertAlmostEqual(streamed, 748, delta=15)                  # measured 0.73 GiB on the CUDA engine (DF2-MEM): the 32K pool (573 MiB with its fixed part) plus what scales with the whole context
         self.assertEqual(V.session_mib(8192, "int8", 32768), V.session_mib(8192))     # a resident window larger than the context changes nothing
-        self.assertAlmostEqual(V.session_mib(131072, "int8", 4096), V.session_mib(131072, "int8", V.KV_RESIDENT_MIN), delta=0.01)   # the engine raises a small window to 20480
+        self.assertAlmostEqual(V.session_mib(131072, "int8", 4096), V.session_mib(131072, "int8", V.KV_RESIDENT_MIN), delta=0.01)   # the engine raises a smaller window to its minimum (4096)
 
     def test_without_slots_nothing_changes(self):
         a = V.slack_mib(FN, 8.0, 8192, "int8", 700, "auto", "intel")
@@ -306,10 +306,13 @@ class CudaSpill(unittest.TestCase):
         self.assertAlmostEqual(sl(7800, 2), 850 - 835 - 598 - 700 - 218 - 256, delta=60)     # 0.83 GiB after the sessions, then the MTP layer and the head: the OOM
 
     def test_an_8gb_card_with_dflash2_alone_gets_the_small_kv_window_and_keeps_128k(self):
-        for mib in (8192, 8050):
+        for mib, window in ((8192, "20480"), (8050, None)):
             args, par, note = self.start(mib)
             self.assertEqual(self.flag(args, "--max-context"), "131072", mib)
-            self.assertEqual(self.flag(args, "--kv-resident"), "20480", mib)            # measured: starts with 105 MiB free and decodes at 72 tok/s; 32768 does not start
+            if window:
+                self.assertEqual(self.flag(args, "--kv-resident"), window, mib)         # measured: starts and decodes at 69-73 tok/s; 32768 does not start
+            else:
+                self.assertIn(self.flag(args, "--kv-resident"), ("20480", "8192", "4096"), mib)   # a card a little smaller walks further down
             self.assertEqual(self.flag(args, "--prefill"), "auto")                       # the engine steps its own chunk down
             self.assertEqual(self.flag(args, "--vram-reserve-mib"), "700")               # a lower reserve starves the DFlash2 prompt staging (measured)
             self.assertIsNone(par)
@@ -320,9 +323,23 @@ class CudaSpill(unittest.TestCase):
             args, par, note = self.start(8192, n, self.MTP)
             self.assertEqual(self.flag(args, "--max-context"), "131072", n)
             self.assertIsNone(par, n)                                                     # the number of requests is the person's
-            self.assertEqual(self.flag(args, "--kv-resident"), "20480", n)
+            self.assertEqual(self.flag(args, "--kv-resident"), "4096", n)               # measured: 2 requests start at 4096 (277 MiB free, 46 tok/s per user) and 8192; 32768 and 20480 do not
             self.assertIn(V.NOTE_SLOW, note)
 
+    def test_the_window_is_walked_down_only_as_far_as_the_card_needs(self):
+        win = lambda mib, n, drafter: self.flag(self.start(mib, n, drafter)[0], "--kv-resident")
+        self.assertEqual(win(8192, None, self.DF2), "20480")       # DFlash2 alone fits at 20480
+        self.assertEqual(win(8192, None, self.MTP), "8192")        # the MTP layer needs a smaller one
+        self.assertEqual(win(10240, 2, self.MTP), "20480")
+        self.assertEqual(win(10240, 4, self.MTP), "4096")
+        self.assertEqual(win(12288, 8, self.MTP), "8192")
+        self.assertEqual(win(12288, 4, self.MTP), "32768")         # nothing moves where it fits
+
+    def test_the_window_floor_and_steps_match_the_engine(self):
+        self.assertEqual(V.KV_RESIDENT_MIN, 4096)                  # qsa_kv_resident_min()
+        self.assertEqual(V.KV_RESIDENT_STEPS, (20480, 8192, 4096))
+        self.assertTrue(all(w >= V.KV_RESIDENT_MIN for w in V.KV_RESIDENT_STEPS))
+
     def test_12_16_and_24gb_keep_everything_for_1_2_and_4_requests(self):
         for mib in (12288, 12282, 16384, 24576):
             for n, drafter in ((None, self.DF2), (None, self.MTP), (2, self.MTP), (4, self.MTP)):