STRATA_MTP_ROOM_MIN — a long prompt takes the MTP drafter's 512 experts (decode-only) as ring

docs/DETAILS.md

@@ -1957,6 +1957,20 @@ the pairs in which the switched arm won). They are here so you can try them on y
   `tools/batch_test.py --n 16` prompts (solo and batch) were identical too, and to those of a build without the
   option. The uniform cache (no profile) holds no resident expert when a native pack's prompt is read (serve needs the
   profile, and a one-shot run fills the cache only while it decodes), so `expert_cache_slot_guard_test` covers it.
+- **`STRATA_MTP_ROOM_MIN=TOKENS` (`--serve`, `--mtp`): the MTP drafter's 512 routed experts are the prompt path's ring
+  while a request that reads TOKENS prompt tokens or more is read** (0: every request). TOKENS counts what the request
+  reads, not the whole prompt: a turn that reuses a cached conversation prefix and reads a few hundred tokens does not
+  take the room. The experts are 708 MB that only the decode reads - `record_rest` is their only reader; a prompt's
+  passes run the draft layer's dense tensors and its K/V - so the prompt path's ring is pointed at those bytes (nothing
+  is freed, nothing is allocated: the block never moves, so the decode's captured graphs stay valid) and the decode's
+  own bytes are back before the first generated token. Measured on the Q8_0 4-way rig: the ring really did take the
+  675 MiB and the decode did not move, but the deeper ring bought no time at 3,080 or 14,859 tokens (it was not what
+  the prompt waited on). On that rig on 2026-10-19 the read threshold was measured: continuations that read 61, 7
+  and 7 tokens skipped the room, and their first token came 0.19-0.24 s sooner (the room's lend and reclaim).
+  Reported by merbanan on an 8 GB RTX 2060 SUPER, where the ring was the limit: +17% / +17% /
+  +23% prompt tokens/s at 4K / 16K / 32K, decode within noise, logits bitwise equal. The same shape the Maya checkout
+  implemented for its DFlash drafter and then for its MTP config's NextN block (local commits `7df369c`, `5ccf9a4`).
+  Details, how to measure it and what it costs: [MTP_PROMPT_ROOM.md](MTP_PROMPT_ROOM.md).
 
 ## Thread placement: `--aux-cpus` (Linux, off by default)
 

include/strata/core/mtp.hpp

@@ -71,6 +71,25 @@ class MtpDrafter {
     int32_t top2(int j) const { return j >= 0 && j < (int) top2_.size() ? top2_[(size_t) j] : -1; }
     int max_t() const { return max_t_; }
     uint64_t vram_bytes() const { return vram_; }
+
+    /// THE PROMPT ROOM (STRATA_MTP_ROOM_MIN).  The layer's 512 routed experts (708 MB) are read by the DECODE alone:
+    /// `record_rest` is their only reader, and every prompt pass - the capture's K/V (record_front) and the prompt
+    /// path's Prefill::draft_kv - touches the dense tensors and the K/V state instead.  So a long prompt may take
+    /// their bytes: the prompt path's ring is pointed AT this block (Prefill::set_ring_room_base), and every expert
+    /// blob it stages overwrites an expert.  `experts_lend` copies the block into pinned host memory first (taken
+    /// once, so the `experts_reclaim` after the prompt is a DMA and not another read of experts.bin); `reclaim`
+    /// copies it back IN PLACE.
+    /// THE BLOCK IS NEVER FREED, NEVER REALLOCATED AND NEVER MOVES: the round and step graphs were captured with its
+    /// address as the MoE kernel's argument, and lending it in place is what keeps them valid - nothing is dropped
+    /// and no round pays a re-capture.
+    /// While the experts are lent no decode pass may run: `draft`, `draft_first`, `chain_launch` and `record_rest`
+    /// refuse (with `err`), rather than read the prompt's staged bytes as experts.
+    uint64_t experts_bytes() const;    ///< 0: this drafter does not own the experts (shared weights: --batch-mtp)
+    void* experts_base() const { return experts_; }   ///< the bytes the room's slots are carved from
+    bool experts_resident() const { return experts_ != nullptr && !experts_lent_; }
+    bool experts_lend(std::string& err);      ///< snapshot to pinned host (the block stays in place)
+    bool experts_reclaim(std::string& err);   ///< the snapshot back into the block, at the address it never left
+
     /// The draft layer's K/V state (read-only: --serve's STRATA_STATE_HASH check hashes it)
     const QsaState& kv_state() const { return st_; }
     /// KV streaming: refill the ring of the drafter's window from its host copy for a sequence that continues at
@@ -169,6 +188,13 @@ class MtpDrafter {
 
 private:
     bool record_forward(int T, int step_row0, cudaStream_t cs, std::string& err);
+    /// False (with `err`) while the experts are lent to a prompt: the prompt path's ring is pointed at their bytes
+    /// (STRATA_MTP_ROOM_MIN), so a decode pass would read the prompt's staged blobs as experts.  Every decode entry
+    /// checks this before it launches anything.
+    bool decode_ready(std::string& err) const {
+        if (experts_lent_) { err = "mtp: the 512 experts are lent to a prompt (STRATA_MTP_ROOM_MIN)"; return false; }
+        return true;
+    }
     /// The layer's front for T rows at step rows [row0, +T): the embedding, the fc projections, the attention
     /// hyper-connection read (R_, inj_, mixed_) and the K/V appended.
     bool record_front(int T, int row0, cudaStream_t cs, std::string& err);
@@ -246,6 +272,11 @@ class MtpDrafter {
     std::vector<Tensor> tensors_;
     uint8_t* dense_ = nullptr;
     uint8_t* experts_ = nullptr;
+    // the prompt room: the experts' pinned host snapshot, its size, and whether a prompt has taken their contents
+    // (the block itself stays where it is - see experts_lend)
+    void* experts_host_ = nullptr;
+    uint64_t experts_host_bytes_ = 0;
+    bool experts_lent_ = false;
     // Slot drafters borrow immutable weights and head; each still owns its state and scratch.
     bool owns_weights_ = true;
     bool owns_draft_head_ = true;

include/strata/prefill/prefill.hpp

@@ -124,6 +124,33 @@ class Prefill {
     /// fits it keeps 0.1.39's ring).  0 slots = none.  A layer split's set_ring_override and STRATA_PREFILL_RING win.
     static void set_ring_budget(int slots, int64_t small_max);
 
+    /// THE PROMPT ROOM (STRATA_MTP_ROOM_MIN): `slots` ring slots ON TOP of the rule above, POINTED AT another device
+    /// user's bytes - the MTP drafter's 512 routed experts, which only the decode reads (MtpDrafter::experts_lend).
+    /// A slot at or past the rule's ring is `base + (i - rule) * max_blob`: nothing is allocated for it, nothing is
+    /// owned and nothing spills, so the room costs no VRAM of its own and `relayout` keeps working.  The layout holds
+    /// them for the instance's life; ONE PROMPT takes them by `set_ring_room_active(true)`, which moves the ring's
+    /// count (the rule's <-> the rule's + the room's) and nothing else.
+    /// Never past ring_cap() (the ring's events are sized for that), never past what `set_ring_room_base` says the
+    /// lender has, and only where the ring is a big one (a chunk at or above stream_all_min(): below it the ring is
+    /// STAGE slots of a different mechanism).  Set before `init`/`relayout`; 0 (the default): the rule's ring exactly,
+    /// as before.
+    void set_ring_room(int slots) { ring_room_ = slots > 0 ? slots : 0; }
+    int ring_room() const { return ring_room_; }
+
+    /// The memory the room's slots are carved from: `bytes` of another device user's, given to this path's ring for
+    /// the length of one prompt (the drafter's expert block).  Null (the default): no room, whatever `slots` says.
+    void set_ring_room_base(void* base, uint64_t bytes) { ring_room_base_ = base; ring_room_bytes_ = bytes; }
+
+    /// While true the room is THIS prompt's: `ring()` (the streamed ring's slot count) counts it too, and an
+    /// `init`/`relayout` after this re-carves with it.  False (the default): the rule's ring, so a slot of the room
+    /// is never staged into and the lender's bytes are left alone.
+    void set_ring_room_active(bool on);
+    bool ring_room_active() const { return ring_room_active_; }
+
+    /// True when every stream of this path has finished.  The room's close needs it: what a prompt staged into those
+    /// slots is written over by the lender's own bytes, and a prompt pass still reading them would take the DMA.
+    bool idle(std::string& err);
+
     /// Device bytes `init` needs for a chunk of `chunk` tokens (what a borrowed region must hold).  `src`: init gets an
     /// ExpertSource (carve picks the fused MoE layout only then).
     static uint64_t bytes_needed(const core::ModelGeometry& g, const core::SessionState& ss, int64_t chunk,
@@ -241,6 +268,10 @@ class Prefill {
     static uint64_t bytes_needed_impl(const core::ModelGeometry& g, const core::SessionState& ss, int64_t chunk,
                                       bool owned_pages, bool src);
     kernels::cpu::ExpertPool* cpu_pool_ = nullptr;   ///< set_cpu_pool
+    int ring_room_ = 0;                              ///< set_ring_room (this instance's, and it survives `reset`)
+    void* ring_room_base_ = nullptr;                 ///< set_ring_room_base: the bytes the room's slots are
+    uint64_t ring_room_bytes_ = 0;
+    bool ring_room_active_ = false;                  ///< the room is this prompt's (survives `reset` with it)
     // Stage-1 pipeline: intermediate stages return after handing their chunk to
     // the direct successor. The public run() drains the chain once at prompt end.
     bool run_impl(const int64_t* tokens, int64_t n, int64_t pos0, std::string& err);

src/core/mtp.cpp

@@ -155,6 +155,7 @@ MtpDrafter::~MtpDrafter() {
     if (sh_join_) cudaEventDestroy(sh_join_);
     if (owns_weights_ && dense_) cudaFree(dense_);
     if (owns_weights_ && experts_) cudaFree(experts_);
+    if (experts_host_) cudaFreeHost(experts_host_);   // the prompt room's host snapshot (the device copy never left)
     if (state_arena_) cudaFree(state_arena_);
     if (arena_) cudaFree(arena_);
     if (head_logits_) cudaFree(head_logits_);
@@ -537,6 +538,54 @@ void MtpDrafter::record_top2(int j) {
     top2_[(size_t) j] = b2 < 0 ? -1 : dhead_ != nullptr ? dvocab_host_[(size_t) b2] : (int32_t) b2;
 }
 
+// ---- the prompt room (see the header): the 512 experts leave the device while a prompt is read, and come back
+uint64_t MtpDrafter::experts_bytes() const {
+    if (!owns_weights_ || g_ == nullptr || experts_ == nullptr) return 0;   // a shared drafter's experts are the owner's
+    return (uint64_t) g_->n_expert * strata::kernels::cpu::BLOB;
+}
+
+bool MtpDrafter::experts_lend(std::string& err) {
+    const uint64_t bytes = experts_bytes();
+    if (bytes == 0) { err = "mtp: the 512 experts are not this drafter's to lend"; return false; }
+    if (experts_lent_) return true;
+    const OnDevice on_device(device_);
+    if (!idle(err)) return false;
+    if (chain_live_) { err = "mtp: a chain is in flight"; return false; }
+    if (experts_host_ == nullptr) {
+        // taken once and kept: the room opens and closes around every long prompt, and a second read of experts.bin
+        // (675 MiB) per prompt is what this avoids
+        if (cudaMallocHost(&experts_host_, (size_t) bytes) != cudaSuccess) {
+            (void) cudaGetLastError();
+            err = "mtp: the 512 experts have no room in pinned host memory (they stay the decode's)";
+            return false;
+        }
+        experts_host_bytes_ = bytes;
+    }
+    // The block stays in VRAM, where it is: the prompt path's room slots ARE these bytes (set_ring_room_base), so
+    // this snapshot is what they are worth again once the prompt has read over them.
+    if (cudaMemcpy(experts_host_, experts_, (size_t) bytes, cudaMemcpyDeviceToHost) != cudaSuccess) {
+        (void) cudaGetLastError();
+        err = "mtp: copying the 512 experts to host memory failed (they stay the decode's)";
+        return false;
+    }
+    experts_lent_ = true;
+    return true;
+}
+
+bool MtpDrafter::experts_reclaim(std::string& err) {
+    if (!experts_lent_) return true;
+    const OnDevice on_device(device_);
+    if (!idle(err)) return false;
+    // IN PLACE: the address is the one the round and step graphs were captured with, so there is nothing to drop
+    if (cudaMemcpy(experts_, experts_host_, (size_t) experts_host_bytes_, cudaMemcpyHostToDevice) != cudaSuccess) {
+        (void) cudaGetLastError();
+        err = "mtp: copying the 512 experts back failed";
+        return false;
+    }
+    experts_lent_ = false;
+    return true;
+}
+
 uint64_t MtpDrafter::bind_bytes(uint64_t head_row_bytes, int64_t n_vocab) const {
     uint64_t bytes = head_logits_ ? 0 : (uint64_t) max_t_ * (uint64_t) n_vocab * sizeof(float);
     if (dhead_ == nullptr && owns_draft_head_ && !full_head_env()) {
@@ -995,6 +1044,7 @@ void MtpDrafter::norm_rope(float* data, const float* gamma, int rows, int cols,
 
 bool MtpDrafter::record_rest(int step_row, cudaStream_t cs, std::string& err) {
     using namespace strata::kernels;
+    if (!decode_ready(err)) return false;   // this is the 512 experts' only reader
     const ModelGeometry& g = *g_;
     SessionState& ss = *ss_;
     const int T = 1;
@@ -1552,6 +1602,7 @@ DraftTiming g_dt;
 
 bool MtpDrafter::draft(int T, const int32_t* tokens, int64_t p, int a, int32_t* drafts, std::string& err,
                        float* probs, float min_p, int* n_drafts) {
+    if (!decode_ready(err)) return false;   // the steps read the 512 experts (record_rest)
 #if defined(STRATA_GPU_METAL)
     {
         if (T < 1 || T > max_t_ || a < 0 || a >= T) { err = "mtp: draft arguments out of range"; return false; }
@@ -1770,6 +1821,7 @@ bool MtpDrafter::prepare_chain(std::string& err) {
 bool MtpDrafter::chain_launch(int T, const int32_t* tokens, int64_t p, int a, const int32_t* force, int n_force,
                               int n_out, int n_early, std::string& err) {
     const OnDevice on_device(device_);
+    if (!decode_ready(err)) return false;   // the chain's step graphs read the 512 experts
     if (chain_live_) { err = "mtp: a chain is already in flight"; return false; }
     if (!force_on_ || ev_chain_ == nullptr) { err = "mtp: chains need prepare_chain"; return false; }
     if (T < 1 || T > max_t_ || a < 0 || a >= T || n_out < 1 || n_out > max_t_ - 1 || n_force < 0 || n_force > n_out ||
@@ -1851,6 +1903,7 @@ int MtpDrafter::chain_poll(std::string& err) {
 
 bool MtpDrafter::draft_first(int T, const float* R_row, int32_t token, int64_t cell, int32_t* drafts, std::string& err,
                              float* probs, float min_p, int* n_drafts) {
+    if (!decode_ready(err)) return false;   // the first round's steps read the 512 experts
 #if defined(STRATA_GPU_METAL)
     {   // the decode window's row 0 is the residual (the window that just ran)
         (void) R_row;

src/prefill/prefill.cpp

@@ -949,6 +949,10 @@ struct Prefill::Impl {
     uint16_t* dq_d[DQ] = {};
     uint8_t* stage_dev[RING_MAX] = {};
     int ring = STAGE;                        // the slots of this layout's ring (ring_slots)
+    // THE PROMPT ROOM's resolved counts for this layout (carve): the rule's ring, the room's first slot (past the
+    // staging buffers this layout holds), and the lender's slots on top of them.  `ring` is `room_at` + `room` while
+    // a roomed prompt runs and `rule` alone otherwise (set_ring_room_active).
+    int ring_rule = 0, ring_room = 0, ring_room_at = 0;
     // this layout's GU, H and Xq are the fused path's (moe_bufs' `fused`), so a layer that runs MMQ or the FP16 path may
     // only write stream_all_min() - 1 of its rows into them; fused_layout() keeps the two in step (#583, #954)
     bool fused_bufs = false;
@@ -1485,7 +1489,25 @@ bool Prefill::carve(size_t T, void* alloc) {
     }
     m.ring = ring_slots(T);
     const int nstage = stage_slots(T);   // the staged walk may use STAGE_GRP of them
-    if (o.base == nullptr && m.ring > 0) {
+    m.ring_rule = m.ring;    // the rule's ring: inside the region, because the loan was sized for it
+    m.ring_room = 0;         // THE PROMPT ROOM's slots on top of it (set_ring_room), the lender's own bytes
+    m.ring_room_at = nstage; // the room starts past the layout's staging buffers (stage_slots(T) of them)
+    // THE PROMPT ROOM (set_ring_room + set_ring_room_base): another device user's bytes - the MTP drafter's 512
+    // routed experts, which only the decode reads - pointed at as so many more ring slots for one prompt.  The slots
+    // ARE those bytes: nothing is allocated, nothing is owned, nothing spills, and the last layout of this path
+    // keeps them, so a prompt takes them by moving `m.ring` alone (set_ring_room_active) and the lender's own bytes
+    // are never left out of VRAM.  Never past ring_cap() (`copied`/`used` exist for i < ring_cap(), and the arrays
+    // are RING_MAX) and never past what the lender said it has (one slot is one max_blob).
+    if (ring_room_ > 0 && m.ring > STAGE && ring_room_base_ != nullptr) {
+        const int64_t cap = std::min<int64_t>((int64_t) ring_cap() - m.ring_room_at,
+                                              (int64_t) (ring_room_bytes_ / (uint64_t) MAXBLOB()));
+        const int64_t room = std::min<int64_t>((int64_t) ring_room_, cap);
+        if (room > 0) {
+            m.ring_room = (int) room;
+            if (ring_room_active_) m.ring = m.ring_room_at + (int) room;   // the room is this prompt's: the ring counts it
+        }
+    }
+    if (o.base == nullptr && nstage > 0) {
         // OWNED buffers: the ring in ONE allocation.  384 separate 2.7 MiB cudaMallocs each round up to a 2 MiB page
         // (~1.3 MiB a slot, ~0.5 GiB in all) that no count ever saw.  A borrowed region keeps its per-slot layout (and
         // so its price, `bytes_needed`: the loans' slot counts do not move).
@@ -1502,6 +1524,13 @@ bool Prefill::carve(size_t T, void* alloc) {
             m.used_of[i] = i;
         }
     }
+    if (m.ring_room > 0) {   // THE PROMPT ROOM's slots: the lender's bytes, at the lender's address, taken from nobody
+        for (int i = m.ring_room_at; i < m.ring_room_at + m.ring_room; ++i) {
+            m.stage_dev[i] = (uint8_t*) ring_room_base_ + (size_t) (i - m.ring_room_at) * (size_t) MAXBLOB();
+            m.stage_live[i] = false;                        // the lender's bytes are live experts, not a staged blob
+            m.used_of[i] = i;
+        }
+    }
     m.ple_emb = o.take<float>(T * N, ok);
     m.ple_norm = o.take<float>((size_t) strata::kernels::NG_HC_DIM, ok);
     take_stage(o, ss, s, m.stage, ok);
@@ -1547,6 +1576,25 @@ bool Prefill::relayout(int64_t chunk, void* borrow, uint64_t borrow_bytes, std::
 
 int64_t Prefill::chunk() const { return impl_->T; }
 
+void Prefill::set_ring_room_active(bool on) {
+    ring_room_active_ = on;
+    Impl& m = *impl_;
+    // the slots' own pointers do not move: the room is taken and given back by the ring's count alone
+    if (m.ring_room > 0) m.ring = on ? m.ring_room_at + m.ring_room : m.ring_rule;
+}
+
+bool Prefill::idle(std::string& err) {
+    Impl& m = *impl_;
+    if ((m.cs && cudaStreamSynchronize(m.cs) != cudaSuccess) ||
+        (m.copy && cudaStreamSynchronize(m.copy) != cudaSuccess) ||
+        (m.kv_copy && cudaStreamSynchronize(m.kv_copy) != cudaSuccess)) {
+        (void) cudaGetLastError();
+        err = "prefill: the prompt path's streams failed";
+        return false;
+    }
+    return true;
+}
+
 bool Prefill::draft_kv(core::MtpDrafter& mtp, const float* R_rows, const int32_t* next_tokens, int64_t n, int64_t cell0,
                        std::string& err, const char** why) {
     Impl& m = *impl_;