Conversation cache: STRATA_PARK_VRAM_MIB=N keeps parked conversations' K/V in VRAM in front of the host copy

With STRATA_PARK_ASYNC, a park copies the K/V device-to-device into a VRAM slot (64 MiB granules, up to N MiB in all,
only while 1.5 GB stays free) that the background DMA reads from; the host image is still filled. A switch back
restores the K/V device-to-device from the slot (~1 ms instead of 27-107 ms over a PCIe Gen2 x4 link) and needs no
join on the DMA. Each K/V layer carries its own join-and-release token, so any order of destruction is safe; the
elastic K/V's growth failure path releases every slot and retries once.
Restore of a 3K conversation 28-47 -> 7-12 ms on a CMP 170HX; outputs identical.

HIP: the runtime calls whose result the cache does not check are cast to (void) (the HIP shims declare hipError_t
[[nodiscard]]), so this commit carries its own HIP guard and builds alone for gfx1100 and gfx906.

Co-Authored-By: Claude Fable 5.1 <[email protected]>
Co-Authored-By: Claude Opus 5.5 <[email protected]>

include/strata/core/conversation_cache.hpp

@@ -62,6 +62,14 @@ struct ConversationKv {
     int format = 0;
     int64_t cells = 0, heads = 0, head_dim = 0, page_size = 0, pooled_rows = 0, idx_dim = 0;
     ConversationBuffer k, v, k_scale, v_scale, pooled;
+    // STRATA_PARK_VRAM_MIB: the tag this layer's parts were recorded under in the park's VRAM slot (0: none).  Set by
+    // conversation_kv_save; a restore that finds the slot copies from it instead of from the buffers above.
+    uint64_t dev_tag = 0;
+    // STRATA_PARK_VRAM_MIB: set while a background copy fills the buffers above from the slot; its deleter waits for
+    // that copy and releases the slot.  This layer's own token (not shared): the layers outlive their image when they
+    // are retained for the next park (ConversationCache::retain), and go in any order.  Last member: it goes before
+    // any buffer.
+    std::shared_ptr<void> pending;
     size_t bytes() const {
         return k.bytes() + v.bytes() + k_scale.bytes() + v_scale.bytes() + pooled.bytes();
     }

include/strata/core/conversation_snapshot.hpp

@@ -90,6 +90,21 @@ bool conversation_park_async_begin(size_t need);
 bool conversation_park_async_launch(std::string& error);
 /// Abandons a begun park that was not launched (the error paths): nothing was queued to the background.
 void conversation_park_async_abandon();
+/// STRATA_PARK_VRAM_MIB=N (with STRATA_PARK_ASYNC): a park also keeps the whole K/V of its image in VRAM, in a slot
+/// of its own taken at begin while the slots' total stays within N MiB (and 1.5 GB of VRAM stays free), so a restore
+/// of that image copies device to device from the slot (conversation_kv_restore / _prefix) instead of host to
+/// device, and does not wait for the background copy.  The slot lives as long as the image and its K/V layers
+/// (SavedConversation::pending, ConversationKv::pending).  Default 0: off.
+/// The tag of the begun park's slot (0: none); kept until the next begin, for the tokens.
+uint64_t conversation_park_slot_tag();
+/// Frees the slot under `tag` (no-op for 0 or a slot already gone); the caller joined the copy that reads it.
+void conversation_park_slot_release(uint64_t tag);
+/// VRAM pressure: joins the background copy, then frees every slot and the stage.  Returns the bytes freed (0, and
+/// nothing is touched, while STRATA_PARK_VRAM_MIB is unset).
+size_t conversation_park_vram_release_all();
+/// conversation_park_join, unless every K/V layer of `image` has its slot (the restore then reads no host buffer
+/// the background copy could still be filling; a layer that falls back to the host joins for itself).
+void conversation_park_join_for(const SavedConversation& image);
 bool conversation_checkpoint_restore(const ConversationCheckpoint& checkpoint, SessionState& session,
                                      const ModelGeometry& g, std::string& error);
 // conversation_checkpoint_save(from) + conversation_checkpoint_restore(to) of `ids` without a host image, for two

src/core/conversation_snapshot.cpp

@@ -115,6 +115,103 @@ bool valid(const QsaState& st, const Layout& l, int64_t upto, std::string& error
     return true;
 }
 
+// STRATA_PARK_VRAM_MIB=N (with STRATA_PARK_ASYNC): a park also keeps the whole K/V of its image in VRAM, in a slot
+// of its own, while the slots' total stays within N MiB; a restore of that image copies from the slot device to
+// device (~1 ms) instead of from the host buffers over PCIe, and needs no wait for the background copy.  A slot is
+// released with the image or its K/V layers (their tokens), or all at once under VRAM pressure
+// (conversation_park_vram_release_all).  Host thread only; the background copy only reads a slot's bytes, and every
+// reuse or free of one comes after a join (conversation_park_async_begin, release_all).
+struct VramKvSlot {
+    void* p = nullptr;
+    size_t cap = 0, used = 0;
+    uint64_t tag = 0;   // 0: free (kept for reuse by size)
+    int dev = -1;
+    struct Part { uint64_t tag; size_t part, offset, bytes; };   // one layer's part (k, v, k_scale, v_scale, pooled)
+    std::vector<Part> parts;
+};
+struct VramKv {
+    std::vector<std::unique_ptr<VramKvSlot>> slots;
+    uint64_t next_tag = 1;   // slot and layer tags, never reused
+};
+VramKv& vram_kv() { static VramKv v; return v; }
+size_t park_vram_budget() {
+    static const size_t b = [] {
+        const char* v = std::getenv("STRATA_PARK_VRAM_MIB");
+        const long long mib = v != nullptr ? std::atoll(v) : 0;
+        return mib > 0 ? (size_t) mib << 20 : (size_t) 0;
+    }();
+    return b;
+}
+bool park_trace() { static const bool on = std::getenv("STRATA_PARK_TRACE") != nullptr; return on; }
+void vram_kv_free(VramKvSlot& s) {
+    int cur = s.dev;
+    (void) cudaGetDevice(&cur);
+    if (cur != s.dev) (void) cudaSetDevice(s.dev);
+    (void) cudaFree(s.p);
+    if (cur != s.dev) (void) cudaSetDevice(cur);
+    (void) cudaGetLastError();
+    s.p = nullptr; s.cap = 0;
+}
+void vram_kv_release(VramKvSlot& s) { s.tag = 0; s.used = 0; s.parts.clear(); }   // free, kept for reuse
+// A slot of `need` bytes on `dev`: the best-fitting free one, else (the free ones freed first) a new one, rounded up
+// to 64 MiB so a growing conversation's next parks reuse it, while the slots' total stays within the budget and
+// 1.5 GB of VRAM stays free beside it.  nullptr, with `why`, otherwise.
+VramKvSlot* vram_kv_take(size_t need, int dev, const char*& why) {
+    VramKv& v = vram_kv();
+    const size_t budget = park_vram_budget();
+    if (budget == 0) { why = "off"; return nullptr; }
+    VramKvSlot* best = nullptr;
+    for (auto& s : v.slots)
+        if (s->tag == 0 && s->dev == dev && s->cap >= need && (best == nullptr || s->cap < best->cap)) best = s.get();
+    if (best == nullptr) {
+        size_t held = 0;
+        for (size_t i = 0; i < v.slots.size();) {
+            if (v.slots[i]->tag != 0) { held += v.slots[i]->cap; ++i; continue; }
+            vram_kv_free(*v.slots[i]);
+            v.slots.erase(v.slots.begin() + (std::ptrdiff_t) i);
+        }
+        const size_t want = (need + ((size_t) 64 << 20) - 1) & ~(((size_t) 64 << 20) - 1);
+        if (held > budget || want > budget - held) { why = "budget"; return nullptr; }
+        size_t fr = 0, tot = 0;
+        if (cudaMemGetInfo(&fr, &tot) != cudaSuccess || fr < want + ((size_t) 3 << 29)) {
+            (void) cudaGetLastError();
+            why = "vram";
+            return nullptr;
+        }
+        void* p = nullptr;
+        if (cudaMalloc(&p, want) != cudaSuccess) { (void) cudaGetLastError(); why = "cudaMalloc"; return nullptr; }
+        auto s = std::make_unique<VramKvSlot>();
+        s->p = p; s->cap = want; s->dev = dev;
+        best = s.get();
+        v.slots.push_back(std::move(s));
+    }
+    best->tag = v.next_tag++;
+    best->used = 0;
+    best->parts.clear();
+    return best;
+}
+// The live slot on the current device holding parts under `tag` (nullptr: none, or tag 0)
+const VramKvSlot* vram_kv_find(uint64_t tag) {
+    if (tag == 0) return nullptr;
+    int dev = 0;
+    if (cudaGetDevice(&dev) != cudaSuccess) { (void) cudaGetLastError(); return nullptr; }
+    for (const auto& s : vram_kv().slots) {
+        if (s->tag == 0 || s->dev != dev) continue;
+        for (const auto& p : s->parts)
+            if (p.tag == tag) return s.get();
+    }
+    return nullptr;
+}
+// Whether the slot holds, under `tag`, at least sizes[i] bytes of every part i that has any; `at` gets the offsets
+bool vram_kv_locate(const VramKvSlot& s, uint64_t tag, const std::array<size_t,5>& sizes, std::array<size_t,5>& at) {
+    std::array<bool,5> have{};
+    for (const auto& p : s.parts)
+        if (p.tag == tag && p.part < at.size() && p.bytes >= sizes[p.part]) { have[p.part] = true; at[p.part] = p.offset; }
+    for (size_t i = 0; i < sizes.size(); ++i)
+        if (sizes[i] > 0 && !have[i]) return false;
+    return true;
+}
+
 struct ParkAsync {
     std::mutex mu;
     std::future<cudaError_t> pending;
@@ -126,6 +223,8 @@ struct ParkAsync {
     cudaStream_t stream = nullptr;
     int dev = -1;
     bool active = false;
+    VramKvSlot* slot = nullptr;   // STRATA_PARK_VRAM_MIB: the begun park's slot (the stage is not used then)
+    uint64_t slot_tag = 0;        // its tag, kept past launch for the image's and its layers' tokens
 };
 ParkAsync& park_async() { static ParkAsync a; return a; }
 constexpr size_t kParkPiece = (size_t) 4 << 20;
@@ -211,12 +310,44 @@ bool conversation_kv_save(ConversationKv& image, const QsaState& st, const Model
     const auto src = pools(st);
     const std::array<size_t,5> sizes = {l.data, l.value_data, l.scales, l.value_scales, l.pooled};
     const std::array<ConversationBuffer*,5> dst = {&image.k, &image.v, &image.k_scale, &image.v_scale, &image.pooled};
+    // STRATA_PARK_VRAM_MIB: this layer's parts go into the park's slot under a fresh tag (0: no slot this park)
+    image.dev_tag = park_async().slot != nullptr ? vram_kv().next_tag++ : 0;
     for (size_t i = 0; i < dst.size(); ++i) {
         // Recopy the partial page and the indexer's moving spare row. Completed
         // pages/rows strictly before the first rewritten token remain identical.
         const size_t keep = i == 4 ? (index ? size_t(unchanged_tokens / strata::kernels::qsa_real_shapes().idx_block) * g.idx_key_dim * 4 : 0)
                                   : l.cells ? (sizes[i] / size_t(l.cells)) * size_t(whole_cells) : 0;
         if (keep > dst[i]->size()) { error = "conversation snapshot: missing reusable prefix"; return false; }
+        {   // STRATA_PARK_VRAM_MIB: with a slot, the whole part goes device to device into it (a restore of this image
+            // reads it from there) and the host buffer gets only the bytes past the kept prefix, read from the slot on
+            // the background thread; a part the slot cannot take is copied synchronously - never through the stage,
+            // whose offsets launch would read from the slot
+            ParkAsync& a = park_async();
+            if (a.slot != nullptr) {
+                VramKvSlot& s = *a.slot;
+                const bool fits = src[i] != nullptr && sizes[i] > 0 && s.used + sizes[i] <= s.cap &&
+                                  (st.kv_mode == 0 || on_device(src[i]));   // (kv_mode 0: the pools are device memory)
+                if (fits) {
+                    const size_t at0 = s.used;
+                    const cudaError_t e = conversation_detail::copy_nonblocking(static_cast<uint8_t*>(s.p) + at0, src[i], sizes[i], true);
+                    if (e != cudaSuccess) { error = std::string("conversation snapshot slot copy: ") + cudaGetErrorString(e); return false; }
+                    s.parts.push_back({image.dev_tag, i, at0, sizes[i]});
+                    s.used += sizes[i];
+                    dst[i]->resize_for_overwrite(sizes[i]);
+                    if (!dst[i]->visit(keep, sizes[i] - keep, [&](uint8_t* p, size_t n, size_t at) {
+                            a.items.push_back({p, n, at0 + at});
+                            return true;
+                        })) return false;
+                } else {
+                    dst[i]->resize(sizes[i]);
+                    if (!dst[i]->visit(keep, sizes[i] - keep, [&](uint8_t* p, size_t n, size_t at) {
+                            return transfer(p, src[i] ? static_cast<const uint8_t*>(src[i]) + at : nullptr, n, error);
+                        })) return false;
+                }
+                if (reused_bytes) *reused_bytes += keep;
+                continue;
+            }
+        }
         // STRATA_PARK_ASYNC: device bytes go to the VRAM stage now and into the buffer on the background thread
         ParkAsync& a = park_async();
         const bool trace = park_trace_on();
@@ -323,10 +454,20 @@ bool conversation_kv_copy(const QsaState& from, const QsaState& to, const ModelG
 
 bool conversation_kv_restore(const ConversationKv& image, const QsaState& st, const ModelGeometry& g,
                              int64_t upto, bool index, std::string& error) {
-    conversation_park_join();
+    // STRATA_PARK_VRAM_MIB: from the image's VRAM slot when it holds every part (no host buffer is read: no join)
+    const VramKvSlot* slot = vram_kv_find(image.dev_tag);
+    if (slot == nullptr) conversation_park_join();
     if (!conversation_kv_validate(image, st, g, upto, index, error)) return false;
     const std::array<const ConversationBuffer*,5> src = {&image.k, &image.v, &image.k_scale, &image.v_scale, &image.pooled};
     const auto dst = pools(st);
+    std::array<size_t,5> sizes{}, off{};
+    for (size_t i = 0; i < src.size(); ++i) sizes[i] = src[i]->size();
+    if (slot != nullptr && !vram_kv_locate(*slot, image.dev_tag, sizes, off)) { slot = nullptr; conversation_park_join(); }
+    if (slot != nullptr) {
+        for (size_t i = 0; i < src.size(); ++i)
+            if (!transfer(dst[i], static_cast<const uint8_t*>(slot->p) + off[i], sizes[i], error)) return false;
+        return refresh_residency(st, g, upto, error);
+    }
     for (size_t i = 0; i < src.size(); ++i)
         if (!src[i]->visit(0, src[i]->size(), [&](const uint8_t* p, size_t n, size_t at) {
                 return transfer(static_cast<uint8_t*>(dst[i]) + at, p, n, error);
@@ -336,7 +477,8 @@ bool conversation_kv_restore(const ConversationKv& image, const QsaState& st, co
 
 bool conversation_kv_restore_prefix(const ConversationKv& image, const QsaState& st, const ModelGeometry& g,
                                     int64_t captured, int64_t upto, bool index, std::string& error) {
-    conversation_park_join();
+    const VramKvSlot* slot = vram_kv_find(image.dev_tag);   // (STRATA_PARK_VRAM_MIB, as in conversation_kv_restore)
+    if (slot == nullptr) conversation_park_join();
     if (upto < 0 || upto > captured) {
         error = "conversation snapshot: prefix beyond the captured K/V";
         return false;
@@ -349,6 +491,14 @@ bool conversation_kv_restore_prefix(const ConversationKv& image, const QsaState&
     for (size_t i = 0; i < src.size(); ++i)
         if (sizes[i] > src[i]->size()) { error = "conversation snapshot: prefix beyond the K/V payload"; return false; }
     const auto dst = pools(st);
+    // STRATA_PARK_VRAM_MIB: the prefix is the leading bytes of each part in the slot too
+    std::array<size_t,5> off{};
+    if (slot != nullptr && !vram_kv_locate(*slot, image.dev_tag, sizes, off)) { slot = nullptr; conversation_park_join(); }
+    if (slot != nullptr) {
+        for (size_t i = 0; i < src.size(); ++i)
+            if (!transfer(dst[i], static_cast<const uint8_t*>(slot->p) + off[i], sizes[i], error)) return false;
+        return refresh_residency(st, g, upto, error);
+    }
     for (size_t i = 0; i < src.size(); ++i)
         if (!src[i]->visit(0, sizes[i], [&](const uint8_t* p, size_t n, size_t at) {
                 return transfer(static_cast<uint8_t*>(dst[i]) + at, p, n, error);
@@ -449,10 +599,23 @@ bool conversation_park_async_begin(size_t need) {
     a.active = false;
     a.items.clear();
     a.used = 0;
+    a.slot = nullptr;
+    a.slot_tag = 0;
     if (!park_async_on() || need == 0) return false;
     int dev = 0;
     if (cudaGetDevice(&dev) != cudaSuccess) { (void) cudaGetLastError(); return false; }
-    if (a.stage != nullptr && (a.cap < need || a.dev != dev)) {
+    // STRATA_PARK_VRAM_MIB: a slot of the whole K/V instead of the stage, while the budget allows - taken before the
+    // stage would be re-allocated, so a growing conversation's park pays no cudaFree (a device sync) when it gets one
+    const char* why = "";
+    a.slot = vram_kv_take(need, dev, why);
+    a.slot_tag = a.slot != nullptr ? a.slot->tag : 0;
+    if (park_trace() && park_vram_budget() != 0) {
+        if (a.slot != nullptr)
+            std::fprintf(stderr, "strata parktrace: kv slot %zu MB tag %llu\n", need >> 20, (unsigned long long) a.slot_tag);
+        else
+            std::fprintf(stderr, "strata parktrace: kv slot: none (%s)\n", why);
+    }
+    if (a.stage != nullptr && ((a.slot == nullptr && a.cap < need) || a.dev != dev)) {
         const int cur = dev;
         (void) cudaSetDevice(a.dev); (void) cudaFree(a.stage); (void) cudaSetDevice(cur);
         a.stage = nullptr; a.cap = 0;
@@ -472,8 +635,13 @@ bool conversation_park_async_begin(size_t need) {
         if (a.stream) (void) cudaStreamDestroy(a.stream);
         a.stream = nullptr; a.host[0] = a.host[1] = nullptr;
         (void) cudaGetLastError();
+        if (a.slot != nullptr) { vram_kv_release(*a.slot); a.slot = nullptr; a.slot_tag = 0; }   // (STRATA_PARK_VRAM_MIB)
         return false;
     }
+    if (a.slot != nullptr) {   // STRATA_PARK_VRAM_MIB: the slot stands in for the stage this park
+        a.active = true;
+        return true;
+    }
     if (a.stage == nullptr) {
         // rounded up to 64 MiB: the next parks of a growing conversation reuse it instead of re-allocating (a cudaFree
         // synchronizes the device, and the checkpoint DMA is in flight at this point)
@@ -499,6 +667,7 @@ void conversation_park_async_abandon() {
     cudaStream_t s = nullptr;
     if (conversation_detail::private_stream(s) == cudaSuccess) (void) cudaStreamSynchronize(s);   // nothing reads the stage after
     (void) cudaGetLastError();
+    if (a.slot != nullptr) { vram_kv_release(*a.slot); a.slot = nullptr; a.slot_tag = 0; }   // (STRATA_PARK_VRAM_MIB)
 }
 
 bool conversation_park_async_launch(std::string& error) {
@@ -510,13 +679,16 @@ bool conversation_park_async_launch(std::string& error) {
     if (e == cudaSuccess) e = cudaStreamSynchronize(s);   // the device-to-device copies into the stage
     if (e != cudaSuccess) {
         a.items.clear();
+        if (a.slot != nullptr) { vram_kv_release(*a.slot); a.slot = nullptr; a.slot_tag = 0; }   // (STRATA_PARK_VRAM_MIB)
         error = std::string("conversation snapshot stage: ") + cudaGetErrorString(e);
         return false;
     }
+    // STRATA_PARK_VRAM_MIB: the host copies read the slot; it stays until the image and its K/V layers go (their tokens)
+    const uint8_t* stage = a.slot != nullptr ? static_cast<const uint8_t*>(a.slot->p) : a.stage;
+    a.slot = nullptr;
     if (a.items.empty()) return true;
     std::vector<ParkAsync::Item> items;
     items.swap(a.items);
-    const uint8_t* stage = a.stage;
     void* host[2] = {a.host[0], a.host[1]};
     cudaStream_t st = a.stream;
     const int dev = a.dev;
@@ -554,4 +726,45 @@ bool conversation_park_async_launch(std::string& error) {
     return true;
 }
 
+uint64_t conversation_park_slot_tag() { return park_async().slot_tag; }
+
+void conversation_park_slot_release(uint64_t tag) {
+    if (tag == 0) return;
+    for (auto& s : vram_kv().slots)
+        if (s->tag == tag) { vram_kv_release(*s); return; }
+}
+
+size_t conversation_park_vram_release_all() {
+    if (park_vram_budget() == 0) return 0;   // off: nothing was taken, and the stage stays as it was
+    conversation_park_join();   // the background copy may be reading a slot or the stage
+    ParkAsync& a = park_async();
+    VramKv& v = vram_kv();
+    size_t freed = 0;
+    for (size_t i = 0; i < v.slots.size();) {
+        if (a.active && v.slots[i].get() == a.slot) { ++i; continue; }   // a park in progress keeps its own
+        freed += v.slots[i]->cap;
+        vram_kv_free(*v.slots[i]);
+        v.slots.erase(v.slots.begin() + (std::ptrdiff_t) i);
+    }
+    if (a.stage != nullptr && !a.active) {
+        int cur = a.dev;
+        (void) cudaGetDevice(&cur);
+        if (cur != a.dev) (void) cudaSetDevice(a.dev);
+        (void) cudaFree(a.stage);
+        if (cur != a.dev) (void) cudaSetDevice(cur);
+        (void) cudaGetLastError();
+        freed += a.cap;
+        a.stage = nullptr; a.cap = 0;
+    }
+    return freed;
+}
+
+void conversation_park_join_for(const SavedConversation& image) {
+    size_t hits = 0;
+    for (const auto& k : image.kv) hits += vram_kv_find(k.dev_tag) != nullptr ? 1 : 0;
+    if (park_trace() && park_vram_budget() != 0)
+        std::fprintf(stderr, "strata parktrace: kv restore: %zu of %zu layers from the slot\n", hits, image.kv.size());
+    if (hits < image.kv.size()) conversation_park_join();
+}
+
 } // namespace strata::core

src/core/conversation_state.cpp

@@ -552,6 +552,9 @@ bool conversation_snapshot_save(SavedConversation& image, const ConversationView
     captured.kv = std::move(reuse.kv);
     const size_t layers = owned_qsa(ss);
     captured.kv.resize(layers + (draft ? 1 : 0));
+    // STRATA_PARK_VRAM_MIB: the retained layers' tokens go now (joined above, so no wait): held past the launch below,
+    // their replacement would join the copy just launched
+    for (auto& k : captured.kv) k.pending.reset();
     bool live_staged = false;
     const bool live_saved = checkpoint_save(captured.live, ss, g, error, live_staged);
     // STRATA_CKPT_ASYNC: the running state lands in live.gdn after the save returns; the image waits for that copy
@@ -596,7 +599,22 @@ bool conversation_snapshot_save(SavedConversation& image, const ConversationView
                                        std::max<int64_t>(0, unchanged - 1), reused_bytes)) return false;
     abandon.on = false;
     if (staged && !conversation_park_async_launch(error)) return false;
-    if (staged) captured.pending = std::shared_ptr<void>(new char, [](void* p) { conversation_park_join(); delete (char*) p; });
+    if (staged) {
+        // STRATA_PARK_VRAM_MIB: the image's slot goes with the image, after the copy that reads it; each K/V layer
+        // carries a token of its own too (the layers outlive the image when they are retained for the next park, and
+        // the copy fills their buffers).  Without a slot: the image's join as before, no layer tokens.
+        const uint64_t slot = conversation_park_slot_tag();
+        auto token = [slot] {
+            return std::shared_ptr<void>(new char, [slot](void* p) {
+                conversation_park_join();
+                conversation_park_slot_release(slot);
+                delete (char*) p;
+            });
+        };
+        captured.pending = token();
+        if (slot != 0)
+            for (auto& k : captured.kv) k.pending = token();
+    }
     image = std::move(captured);
     return true;
 }
@@ -675,7 +693,7 @@ ConversationRestore conversation_snapshot_restore(const SavedConversation& image
                                                    const ModelGeometry& g, const QsaState* draft, std::string& error) {
     conversation_checkpoint_join(image.live.gdn.data());
     for (const auto& c : image.checkpoints) conversation_checkpoint_join(c.gdn.data());
-    conversation_park_join();
+    conversation_park_join_for(image);   // (STRATA_PARK_VRAM_MIB: no wait when the K/V comes from its slot)
     if (!conversation_snapshot_validate(image, ss, g, draft, error)) return ConversationRestore::invalid;
     if (!sync(error)) return ConversationRestore::transfer_failed;
     const int64_t upto = (int64_t) image.live.ids.size();
@@ -712,7 +730,7 @@ ConversationRestore conversation_snapshot_restore(const SavedConversation& image
                                                    const ModelGeometry& g, const QsaState& draft, std::string& error) {
     conversation_checkpoint_join(image.live.gdn.data());
     for (const auto& c : image.checkpoints) conversation_checkpoint_join(c.gdn.data());
-    conversation_park_join();
+    conversation_park_join_for(image);   // (STRATA_PARK_VRAM_MIB: no wait when the K/V comes from its slot)
     return conversation_snapshot_restore(image, ss, g, &draft, error);
 }
 ConversationRestore conversation_snapshot_restore_prefix(const SavedConversation& image, int64_t upto,
@@ -720,7 +738,7 @@ ConversationRestore conversation_snapshot_restore_prefix(const SavedConversation
                                                           const QsaState& draft, std::string& error) {
     conversation_checkpoint_join(image.live.gdn.data());
     for (const auto& c : image.checkpoints) conversation_checkpoint_join(c.gdn.data());
-    conversation_park_join();
+    conversation_park_join_for(image);   // (STRATA_PARK_VRAM_MIB: no wait when the K/V comes from its slot)
     return conversation_snapshot_restore_prefix(image, upto, ss, g, &draft, error);
 }
 ConversationRestore conversation_snapshot_restore_prefix(const SavedConversation& image, int64_t upto,
@@ -728,7 +746,7 @@ ConversationRestore conversation_snapshot_restore_prefix(const SavedConversation
                                                           const QsaState* draft, std::string& error) {
     conversation_checkpoint_join(image.live.gdn.data());
     for (const auto& c : image.checkpoints) conversation_checkpoint_join(c.gdn.data());
-    conversation_park_join();
+    conversation_park_join_for(image);   // (STRATA_PARK_VRAM_MIB: no wait when the K/V comes from its slot)
     // a layer split's image (stage images) is restored stage by stage, never in part
     if (!image.stage_images.empty()) {
         fail(error, "a layer split's image restored as a single session");

src/program/generate.cpp

@@ -7869,12 +7869,18 @@ int main(int argc, char** argv) {
             return false;
         }
         int64_t fresh = 0;
-        const bool ok = strata::core::qsa_kv_elastic_grow(target, [&]() -> strata::core::VmmChunk {
+        auto grow_take = [&]() -> strata::core::VmmChunk {
             if (kvg.spare.empty()) { ++fresh; return 0; }   // the cache is at its floor: new memory
             const strata::core::VmmChunk h = kvg.spare.back();
             kvg.spare.pop_back();
             return h;
-        });
+        };
+        bool ok = strata::core::qsa_kv_elastic_grow(target, grow_take);
+        // STRATA_PARK_VRAM_MIB: the parked conversations' VRAM copies of their K/V give way to the K/V itself
+        if (!ok && strata::core::conversation_park_vram_release_all() > 0) {
+            (void) cudaGetLastError();
+            ok = strata::core::qsa_kv_elastic_grow(target, grow_take);
+        }
         kvg.fresh += fresh;
         kvg.cells = strata::core::qsa_kv_elastic_cells();
         ++kvg.grows;
@@ -12762,7 +12768,10 @@ int main(int argc, char** argv) {
                 }
             }
             // the elastic K/V: a parked conversation is restored whole, and it may be longer than this prompt
-            if (incoming && !kvg_ensure((int64_t) incoming->live.ids.size() + 256, kv_quiesce)) {
+            // (STRATA_PARK_VRAM_MIB: the parked K/V copies in VRAM go first, then one more try)
+            if (incoming && !kvg_ensure((int64_t) incoming->live.ids.size() + 256, kv_quiesce) &&
+                (strata::core::conversation_park_vram_release_all() == 0 ||
+                 !kvg_ensure((int64_t) incoming->live.ids.size() + 256, kv_quiesce))) {
                 std::printf("ERR the K/V cannot grow to the parked conversation: no VRAM is left\n");
                 return 1;
             }