SYCL: a start short of VRAM runs slower instead of stopping. Requests at once (--batch N) and a long context no longer end in a refusal on an 8 GB A750: when the expert cache sizing leaves no working cache (or only one sized for a prompt loan that cannot happen), the prompt path's own chunk goes down (to 256), then a reserve that was asked for (to 300 MiB); a context that does not fit in VRAM turns KV streaming on by itself (the largest resident window that fits, 20480 at the least); the person's own --prefill N and --kv-resident N are kept. Only a start with no slot even then is refused, with one plain line saying what is taken and what to lower, and an out-of-memory during the load ends the same way. short_cache.hpp holds the arithmetic, short_cache_test (CPU) the A750 figures from the engine logs

sycl/CMakeLists.txt

@@ -428,6 +428,9 @@ if(STRATA_SYCL_PARITY)
   target_link_libraries(mmvq_bench PRIVATE strata_kernels)
   add_executable(dflash2_ring_test ${CMAKE_CURRENT_SOURCE_DIR}/src/core/dflash2_ring_test.cpp)   # DF2-SYCL: the ring log: rebuilt K/V rows are the original ones
   target_link_libraries(dflash2_ring_test PRIVATE strata_engine strata_prefill strata_spec)
+  add_executable(short_cache_test ${CMAKE_CURRENT_SOURCE_DIR}/src/core/short_cache_test.cpp)   # the sizing of a start short of VRAM (CPU, no model)
+  target_include_directories(short_cache_test PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include)
+  add_test(NAME short_cache_test COMMAND short_cache_test)
   add_executable(verify_refusal_test ${CMAKE_CURRENT_SOURCE_DIR}/src/core/verify_refusal_test.cpp)   # a refused start (no VRAM expert tier) ends cleanly: the verifiers leave the default queue alone
   target_link_libraries(verify_refusal_test PRIVATE strata_engine strata_prefill strata_spec)
   add_test(NAME verify_refusal_test COMMAND verify_refusal_test)   # needs one SYCL GPU

sycl/include/strata/core/short_cache.hpp

@@ -0,0 +1,52 @@
+// sycl/include/strata/core/short_cache.hpp - sizing a start that is short of VRAM (owner rule: a long context and requests at once never shrink the context or
+// stop a start that can run; they run slower).  Pure arithmetic, no device.
+//
+//  * size_short_cache: the expert cache's sizing left no working cache.  The speed knobs go in this order, and each only as far as needed: the prompt path's own
+//    chunk (its buffers are 160 + 0.66 MiB a token: 1182 MiB at 2048), then a reserve that was asked for (down to the small-card reserve).  The person's own chunk
+//    (--prefill N) is left alone.  The result is the largest chunk (then reserve) that leaves `work_slots`, else the smallest chunk and reserve (whatever slots that leaves).
+//  * pick_kv_resident: the largest KV window (65536, 32768, 20480 cells a layer) whose sessions fit; 20480 (the engine's smallest) when none does.
+#pragma once
+
+#include <algorithm>
+#include <cstdint>
+
+namespace strata::core {
+
+struct ShortCache {
+    int64_t chunk = 0;
+    int64_t reserve_mib = 0;
+    int64_t slots = 0;
+};
+
+inline int64_t owned_prompt_mib(int64_t chunk) { return 160 + (chunk * 680) / 1024; }
+
+/// Slots the cache gets when the prompt path allocates its own `chunk`-token buffers.
+inline int64_t slots_with_own_chunk(int64_t free_b, int64_t reserve_mib, int64_t chunk, int64_t draft_head_b, int64_t pipe_first_b, int64_t blob_b, int64_t profile_cap) {
+    int64_t n = (free_b - (((reserve_mib + owned_prompt_mib(chunk)) << 20) + draft_head_b + pipe_first_b)) / blob_b;
+    if (profile_cap > 0) n = std::min(n, profile_cap);
+    return std::max<int64_t>(n, 0);
+}
+
+inline ShortCache size_short_cache(int64_t free_b, int64_t reserve_mib, bool reserve_given, int64_t small_reserve_mib, int64_t chunk, bool chunk_given, int64_t min_chunk,
+                                   int64_t work_slots, int64_t draft_head_b, int64_t pipe_first_b, int64_t blob_b, int64_t profile_cap) {
+    auto slots_at = [&](int64_t c, int64_t r) { return slots_with_own_chunk(free_b, r, c, draft_head_b, pipe_first_b, blob_b, profile_cap); };
+    while (!chunk_given && chunk > min_chunk && slots_at(chunk, reserve_mib) < work_slots) chunk /= 2;
+    if (slots_at(chunk, reserve_mib) < work_slots && reserve_given && reserve_mib > small_reserve_mib) {
+        int64_t r = reserve_mib;
+        while (r > small_reserve_mib && slots_at(chunk, r) < work_slots) r = std::max<int64_t>(small_reserve_mib, r - 50);
+        reserve_mib = r;
+    }
+    return ShortCache{chunk, reserve_mib, slots_at(chunk, reserve_mib)};
+}
+
+/// `need_at(window)` is the bytes the sessions and everything after them need with that KV window (0: the whole context in VRAM); `free_b` what the card has.
+/// Returns 0 when the whole context fits (no streaming), else the largest window that fits, else the smallest the engine takes (20480).
+template <class NeedAt>
+inline int64_t pick_kv_resident(int64_t ctx, int64_t free_b, NeedAt need_at) {
+    if (need_at((int64_t) 0) <= free_b) return 0;
+    for (const int64_t w : {(int64_t) 65536, (int64_t) 32768, (int64_t) 20480})
+        if (w < ctx && need_at(w) <= free_b) return w;
+    return 20480;
+}
+
+}  // namespace strata::core

sycl/src/core/short_cache_test.cpp

@@ -0,0 +1,43 @@
+// sycl/src/core/short_cache_test.cpp - the sizing of a start short of VRAM (strata/core/short_cache.hpp), on the 8 GB Arc A750's figures from the engine's logs.
+#include "strata/core/short_cache.hpp"
+
+#include <cstdio>
+
+using namespace strata::core;
+static int g_fail = 0;
+static void check(bool ok, const char* what) { std::printf("  %-110s %s\n", what, ok ? "ok" : "FAIL"); if (!ok) ++g_fail; }
+
+int main() {
+    const int64_t MiB = 1 << 20;
+    const int64_t blob = 2330000 + 0;   // the pack's largest blob (2.33 MB)
+    const int64_t head = 73 * MiB;
+    // --batch 2 + MTP, 8192 tokens, 1024 MiB reserve given: the engine's auto sizing left 0 slots (2.46 GiB free, chunk 2048)
+    const int64_t free8k = (int64_t) (2.46 * 1024) * MiB;
+    check(slots_with_own_chunk(free8k, 1024, 2048, head, 0, blob, 24576) == 0, "8K, reserve 1024, chunk 2048: the sizing left no slot (the refused start)");
+    ShortCache a = size_short_cache(free8k, 1024, true, 300, 2048, false, 256, 128, head, 0, blob, 24576);
+    check(a.chunk == 1024 && a.reserve_mib == 1024 && a.slots >= 128, "... the chunk goes 2048 -> 1024, the typed reserve stays, a working cache is left");
+    // the same with 128K sessions of 0.73 GiB (KV window 32768): 0.92 GiB free
+    const int64_t free128 = (int64_t) (0.92 * 1024) * MiB;
+    ShortCache b = size_short_cache(free128, 300, true, 300, 1024, false, 256, 128, head, 0, blob, 24576);
+    check(b.chunk == 256 && b.reserve_mib == 300 && b.slots > 0 && b.slots < 128, "128K, window 32768: the smallest chunk, a cache below 128 slots but a start");
+    // a typed reserve is lowered only after the chunk is as small as it goes
+    ShortCache c = size_short_cache((int64_t) (1.36 * 1024) * MiB, 1024, true, 300, 8192, false, 256, 128, head, 0, blob, 24576);
+    check(c.chunk == 256 && c.reserve_mib < 1024 && c.reserve_mib >= 300 && c.slots >= 128, "128K, window 20480, reserve 1024: chunk 256 first, then the reserve comes down only as far as needed");
+    // the person's own chunk is not touched
+    ShortCache d = size_short_cache(free8k, 1024, true, 300, 2048, true, 256, 128, head, 0, blob, 24576);
+    check(d.chunk == 2048, "a chunk the person asked for stays");
+    ShortCache e = size_short_cache(free8k, 1024, false, 300, 2048, true, 256, 128, head, 0, blob, 24576);
+    check(e.reserve_mib == 1024, "a reserve that was not asked for is not touched here (the engine's own reserve rule runs before)");
+    // impossible: nothing left even at the smallest chunk and reserve
+    ShortCache f = size_short_cache(100 * MiB, 300, true, 300, 2048, false, 256, 128, head, 0, blob, 24576);
+    check(f.slots == 0 && f.chunk == 256, "a card with 100 MiB free: no slot at the smallest chunk - the one plain refusal");
+    // the KV window
+    auto need = [&](int64_t w) { return (w == 0 ? 3 * 1930LL : w == 65536 ? 3 * 1050LL : w == 32768 ? 3 * 750LL : 3 * 585LL) * MiB + 3000 * MiB; };
+    check(pick_kv_resident(131072, 12000 * MiB, need) == 0, "a card that holds the whole context keeps it in VRAM (no streaming)");
+    check(pick_kv_resident(131072, 8500 * MiB, need) == 65536, "the largest window that fits is taken");
+    check(pick_kv_resident(131072, 6100 * MiB, need) == 32768, "... the next one down");
+    check(pick_kv_resident(131072, 4000 * MiB, need) == 20480, "none fits: the engine's smallest window");
+    check(pick_kv_resident(32768, 4000 * MiB, need) == 20480 && pick_kv_resident(20480, 100 * MiB, need) == 20480, "a window is never larger than the context asked for");
+    std::printf("%s\n", g_fail ? "SHORT-CACHE: FAIL" : "SHORT-CACHE: ok");
+    return g_fail ? 1 : 0;
+}

sycl/src/program/generate.cpp

@@ -33,6 +33,7 @@
 #include "strata/core/device.hpp"
 #include "strata/core/remote_expert_opt.hpp"
 #include "strata/core/expert_cache.hpp"
+#include "strata/core/short_cache.hpp"
 #include "strata/core/conversation_snapshot.hpp"
 #include "strata/core/conversation_file.hpp"
 #include "strata/core/conversation_memory.hpp"
@@ -612,6 +613,8 @@ struct Options {
     int vram_reserve_later_mib = -1;
     /// Plan v0.3 P5: batched prompt processing in chunks of this many tokens (0 = the token path).
     int64_t prefill_chunk = 0;
+    /// `--prefill N` (a number, not auto): the person's chunk, which the cache sizing leaves alone.
+    bool prefill_given = false;
     /// `--prefill auto`: the largest chunk (up to 8192) whose buffers the expert cache can lend.  Every expert a chunk
     /// routes to is streamed once per chunk, so a bigger chunk streams fewer bytes per token (the "ubatch" effect).
     bool prefill_auto = false;
@@ -1944,6 +1947,7 @@ int main(int argc, char **argv) try {
                                      : (o.prefill_auto && env_max != nullptr ? std::atoll(env_max) : 8192);
             o.prefill_auto_max = want_max >= 32768 ? 32768 : want_max >= 16384 ? 16384 : 8192;
             o.prefill_chunk = o.prefill_auto ? o.prefill_auto_max : std::atoll(v.c_str());
+            o.prefill_given = !o.prefill_auto && o.prefill_chunk > 0;
         }
         else if (a == "--no-split-rows") o.no_split_rows = true;
         else if (a == "--no-prefill-borrow") { o.no_prefill_borrow = true; borrow_explicit = true; }
@@ -4097,6 +4101,48 @@ int main(int argc, char **argv) try {
     // stage that runs [lb, le) carves only those layers' GDN rows and QSA pools - before the carve every stage
     // held all 48 layers' state whatever layers it ran, which is the same disease the chunked-QSA-prefill PR
     // fixed in llama.cpp: allocation sized by the whole model instead of the device's own work.
+    // A context the sessions cannot hold in VRAM is not a reason to stop (owner rule: a long context and requests at once never shrink the context, they run
+    // slower): with no --kv-resident given, the K/V streams from RAM, the resident window as large as the card leaves room for.
+    // Needed: every session (the main one and each request slot), what loads after them (the draft layer, the native head, the verify buffers), the reserve,
+    // the prompt path's smallest own buffers (chunk 256) and a cache of 128 slots.
+    if (o.kv_resident <= 0 && o.max_context > 20480 && !multi_gpu && !o.kv_grow && o.expert_cache != 0) {
+        const int64_t kLateLoadMib = (!o.mtp.empty() || o.draft == "dflash2") ? 1500 : 700;   // Arc A750: MTP 835 + head 417 + verify 70 + the rest; DFlash2 alone is about the same
+        const int64_t blob_b = (int64_t) strata::kernels::cpu::expert_layout().max_blob;
+        const int64_t n_sess = 1 + std::max(o.batch, 0);
+        const int64_t reserve_b = (int64_t) (o.vram_reserve_given ? o.vram_reserve_mib : std::min(o.vram_reserve_mib, 300)) << 20;
+        const int64_t fixed_b = (kLateLoadMib << 20) + reserve_b + ((160 + (256 * 680) / 1024) << 20) + 128 * blob_b;
+        size_t free_k = 0, total_k = 0;
+        free_k = strata::core::device_free_bytes(); (void) total_k;
+        auto need_at = [&](int64_t res) -> int64_t {
+            strata::core::qsa_set_kv_resident(res);
+            return n_sess * (int64_t) strata::core::session_bytes(g, o.max_context, K, 0, g.n_layers) + fixed_b;
+        };
+        int64_t pick = strata::core::pick_kv_resident(o.max_context, (int64_t) free_k, need_at);
+        if (pick > 0) {
+            // the pinned RAM the streamed K/V takes (it is a copy of the whole context per session): 1.8 GB a session at 128K int8
+            int64_t host_avail = -1;
+            if (std::FILE* mf = std::fopen("/proc/meminfo", "r")) {
+                char line[256];
+                while (std::fgets(line, sizeof(line), mf)) {
+                    long long kb = 0;
+                    if (std::sscanf(line, "MemAvailable: %lld kB", &kb) == 1) { host_avail = (int64_t) kb << 10; break; }
+                }
+                std::fclose(mf);
+            }
+            const int64_t host_need = n_sess * (int64_t) o.max_context * 13 * 1056 + (6ll << 30);
+            if (host_avail >= 0 && host_avail < host_need) {
+                std::fprintf(stderr, "strata generate: the context does not fit in VRAM and %.1f GiB of free RAM is less than the %.1f GiB KV streaming would take: not turned on\n",
+                             (double) host_avail / 1073741824.0, (double) host_need / 1073741824.0);
+                pick = 0;
+            } else {
+                std::fprintf(stderr, "strata generate: %lld tokens of context in %lld sessions do not fit in the %.2f GiB free: KV streaming on, %lld cells a layer in VRAM "
+                                     "(--kv-resident %lld; the rest of the K/V in RAM, so decode and prompts run slower)\n",
+                             (long long) o.max_context, (long long) n_sess, (double) free_k / 1073741824.0, (long long) pick, (long long) pick);
+            }
+        }
+        o.kv_resident = pick;
+        strata::core::qsa_set_kv_resident(o.kv_resident);
+    }
     {
         const strata::core::OnDevice on0(0);
         const int64_t hi0 = multi_gpu ? split_at[0] : -1;
@@ -4686,7 +4732,7 @@ int main(int argc, char **argv) try {
         // under WDDM an over-subscribed allocation does not fail, it pages to system memory and crawls.
         // (with borrowing - the default with a profile - the prompt path lends cache slots instead; `pf_borrow` is
         // the predicate a local `borrow` was here, hoisted above so both cache-size branches read the same one)
-        const int64_t prefill_mib = owned_prefill_mib();
+        int64_t prefill_mib = owned_prefill_mib();
         // the draft layer's head and logits are allocated when it binds, after this: 0.1.27's CJK subset made them
         // ~110-180 MiB larger, and out of the reserve they left 16 GB cards below the stall line (#199)
         int64_t mtp_bind = (!o.mtp.empty() && native_head.loaded())
@@ -4739,6 +4785,36 @@ int main(int argc, char **argv) try {
                 reserve_adapted = true;
             }
         }
+        // Requests at once and a long context never make a start refuse when it can run slower (owner rule).  Still short of a working cache, the speed knobs
+        // go first, in this order: the prompt path's own chunk (its buffers are 160 + 0.66 MiB a token, 1182 MiB at 2048; --prefill N is the person's and stays),
+        // then a reserve that was asked for (down to kSmallReserveMib, the one a small card runs on).  The cache keeps what is left; only a start that
+        // has no slot even then is refused, below.
+        constexpr int64_t kWorkCacheSlots = 128;   // below this the cache is too small to be worth its name (the prompt path cannot lend, every miss goes to the CPU)
+        constexpr int64_t kMinChunk = 256;
+        // (only a start that would stop: no slot, or - sized for a loan - too few slots to lend the prompt path its chunk, so that it allocates its own on top
+        // of a cache that took the VRAM; a start that runs is left as it was sized)
+        if (o.expert_cache < min_slots && o.prefill_chunk > 0) {
+            const strata::core::ShortCache sc = strata::core::size_short_cache(
+                (int64_t) free_b, o.vram_reserve_mib, o.vram_reserve_given, kSmallReserveMib,
+                pf_borrow ? std::min<int64_t>(o.prefill_chunk, 1024) : o.prefill_chunk,   // (what the loan-less path takes when nothing can be lent)
+                o.prefill_given, kMinChunk, kWorkCacheSlots, mtp_bind, pipe_first, blob, (int64_t) profile.size());
+            const int64_t chunk = sc.chunk, reserve_mib = sc.reserve_mib, n = sc.slots;
+            if (n > o.expert_cache || (pf_borrow && n > 0)) {
+                std::fprintf(stderr, "strata generate: expert cache auto: short of VRAM, so the start runs slower instead of stopping: %s%s%s -> %lld slots\n",
+                             chunk != o.prefill_chunk ? ("prompt chunk " + std::to_string(o.prefill_chunk) + " -> " + std::to_string(chunk) + " tokens").c_str() : "",
+                             chunk != o.prefill_chunk && reserve_mib != o.vram_reserve_mib ? ", " : "",
+                             reserve_mib != o.vram_reserve_mib ? ("reserve " + std::to_string(o.vram_reserve_mib) + " -> " + std::to_string(reserve_mib) + " MiB").c_str() : "",
+                             (long long) n);
+                if (reserve_mib != o.vram_reserve_mib) reserve_adapted = true;
+                if (pf_borrow) o.prefill_auto = false;   // the chunk is the person's-sized now: its own buffers, no loan to size it from
+                o.prefill_chunk = chunk;
+                o.vram_reserve_mib = (int) reserve_mib;
+                o.expert_cache = (int) n;
+                prefill_mib = owned_prefill_mib();
+            } else if (pf_borrow && n == 0) {
+                o.expert_cache = 0;   // sized for a loan that cannot happen, and no slot once the prompt path's own buffers are priced: the refusal below says so
+            }
+        }
         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
@@ -4760,11 +4836,12 @@ int main(int argc, char **argv) try {
                          (long long) min_slots, (long long) ((min_slots * blob) >> 20), (long long) short_mib,
                          (long long) session_mib, (long long) o.max_context, (long long) (mtp_bind >> 20),
                          reserve_tip.c_str());
-            g_no_cache_hint = "the GPU memory is full before the expert cache (it needs at least " + std::to_string(min_slots * blob >> 20) + " MiB; this card is about " +
-                              std::to_string(short_mib) + " MiB short). Taken: " + (o.batch > 1 ? std::to_string(o.batch) + " request slots " + std::to_string(o.batch * session_mib) + " MiB (--batch / parallel), " : std::string()) +
-                              "the context " + std::to_string(session_mib) + " MiB at " + std::to_string(o.max_context) + " tokens (--max-context), the draft head " +
-                              std::to_string(mtp_bind >> 20) + " MiB, the reserve " + std::to_string(at_reserve) + " MiB (--vram-reserve-mib). Lower " +
-                              (o.batch > 1 ? "the number of requests at once (--batch / parallel), " : "") + "--max-context or --vram-reserve-mib";
+            g_no_cache_hint = "this card cannot hold even the smallest start: with prompt chunks of 256 tokens and the reserve at " + std::to_string(at_reserve) +
+                              " MiB, the expert cache still needs " + std::to_string(min_slots * blob >> 20) + " MiB and the card is about " + std::to_string(short_mib) +
+                              " MiB short. Taken: " + (o.batch > 1 ? std::to_string(o.batch) + " request slots " + std::to_string(o.batch * session_mib) + " MiB, " : std::string()) +
+                              "the context " + std::to_string(session_mib) + " MiB at " + std::to_string(o.max_context) + " tokens, the draft head " +
+                              std::to_string(mtp_bind >> 20) + " MiB. Lower " + (o.batch > 1 ? "the number of requests at once (--batch / parallel), " : "") +
+                              "use --kv q4_0 (smaller K/V), turn off the draft layer, or close other programs that use the GPU";
         }
     } else if (multi_gpu && o.expert_cache > 0) {
         // an explicit cache size leaves room for the prompt path's buffers and the reserve, or the first prompt
@@ -13806,6 +13883,13 @@ int main(int argc, char **argv) try {
     return 0;
 }
 catch (sycl::exception const &exc) {
+  const std::string what = exc.what();
+  if (exc.code() == sycl::errc::memory_allocation || what.find("OUT_OF_DEVICE_MEMORY") != std::string::npos) {
+    // the one plain refusal for a start no sizing could fit (the sizing above already shrank the prompt chunk, the reserve and the resident K/V window)
+    std::cerr << "strata: this card does not have the memory for this start: the GPU ran out of memory while the engine was loading (" << what << "). Lower the number of "
+                 "requests at once (--batch / parallel), use --kv q4_0 (a smaller K/V), turn off the draft layer, or close other programs that use the GPU" << std::endl;
+    std::exit(1);
+  }
   std::cerr << exc.what() << "Exception caught at file:" << __FILE__
             << ", line:" << __LINE__ << std::endl;
   std::exit(1);