From ac00ff512ced50f60683545a2e592ddaf9dc5d30 Mon Sep 17 00:00:00 2001 From: oceanio Date: Sun, 18 Oct 2026 08:39:55 +0800 Subject: [PATCH 01/11] metal: the session's QSA states carry no K/V pools or pooled indexer keys (the decode window keeps its own) The Metal decode window (metal/src/engine/window.cpp) keeps its own int8 K/V and pooled indexer keys per QSA layer, sized by the same max cells, and never reads or writes the session's. On Metal a session buffer is one shared MTLBuffer in the residency set, resident in full whether touched or not, so those unused pools cost about 24.6 KB per context token. Under STRATA_GPU_METAL, kv_plan gives the session's QSA states no K/V slots and no pooled rows; their pointers stay null, and a conversation snapshot refuses such a state (n_slots == 0) instead of reading it. M5 Max, Qwen3.8-Flash-Next IQ3_S resident, phys_footprint 5 s after the engine is ready: 16.99 -> 10.61 GiB at --max-context 262144, 9.27 -> 7.68 GiB at 65536. Greedy tokens unchanged. Co-Authored-By: Claude Opus 5.5 --- src/core/layer.cpp | 19 +++++++++++++++++++ 1 file changed, 19 insertions(+) diff --git a/src/core/layer.cpp b/src/core/layer.cpp index 273b2293ac3..846ffc85e56 100644 --- a/src/core/layer.cpp +++ b/src/core/layer.cpp @@ -740,6 +740,17 @@ KvPlan kv_plan(const QsaShapes& s, int64_t max_cells, int64_t ring_cells) { p.pages = (max_cells + s.page_size - 1) / s.page_size; p.slots = p.pages; p.pooled_rows = max_cells / s.idx_block + 2; +#if defined(STRATA_GPU_METAL) + // THE METAL DECODE WINDOW KEEPS ITS OWN K/V AND INDEXER per QSA layer (metal/src/engine/window.cpp), sized by + // the same max cells; the session's pools are never read or written there. On Metal a session buffer is one + // shared MTLBuffer in the residency set, resident in full whether touched or not (measured: 6.0 GiB of K/V pools at + // 262K), so its QSA states carry no K/V pools and no pooled indexer keys. `n_slots == 0` makes every + // conversation-snapshot copy refuse the state (conversation_snapshot.cpp `valid`) instead of reading them. + (void) ring_cells; + p.slots = 0; + p.pooled_rows = 0; + return p; +#endif if (g_kv_resident <= 0 || ring_cells < 0) return p; // ring_cells < 0: always fully resident // A/B only: STRATA_KV_RING_OFF keeps the drafter fully resident, STRATA_KV_MAIN_OFF the main layers static const bool ring_off = std::getenv("STRATA_KV_RING_OFF") != nullptr; @@ -811,6 +822,11 @@ uint64_t qsa_state_init(const ModelGeometry& g, int64_t max_cells, void* base, Q const uint64_t rows = (uint64_t) p.slots * s.n_head_kv * s.page_size; // VRAM rows: the slots const uint64_t q4_row = strata::kernels::kv_q4_bytes_per_head((int) s.head_dim); st.kv_elastic = -1; +#if defined(STRATA_GPU_METAL) + if (p.slots == 0) { + // no pools (kv_plan): the pool pointers stay null rather than aliasing the page table that follows + } else +#endif if (g_kv_elastic && p.mode == 0) { // the elastic K/V: each array at a chunk boundary of the state's own range, the first cells mapped const uint64_t slot_rows = (uint64_t) s.n_head_kv * s.page_size; @@ -881,6 +897,9 @@ uint64_t qsa_state_init(const ModelGeometry& g, int64_t max_cells, void* base, Q st.idx_dead = c.take((uint64_t) s.idx_dim); st.idx_pooled_rows = p.pooled_rows; st.idx_pooled = c.take((uint64_t) p.pooled_rows * s.idx_dim); +#if defined(STRATA_GPU_METAL) + if (p.pooled_rows == 0) st.idx_pooled = nullptr; // none on Metal (kv_plan), not an alias of block_pos +#endif st.idx_block_pos = c.take(1); if (share_rope != nullptr) { st.cos_tab = share_rope->cos_tab; From 75a73bbf49f0b602b681a462622d117fadcd9fd2 Mon Sep 17 00:00:00 2001 From: oceanio Date: Sun, 18 Oct 2026 08:39:56 +0800 Subject: [PATCH 02/11] docs (Mac): the engine's own memory grows with --max-context; what it was before the session K/V change Co-Authored-By: Claude Opus 5.5 --- docs/MAC.md | 1 + 1 file changed, 1 insertion(+) diff --git a/docs/MAC.md b/docs/MAC.md index 56d16318ca2..cf9ca2d3961 100644 --- a/docs/MAC.md +++ b/docs/MAC.md @@ -46,6 +46,7 @@ Intel engines are unchanged. | Gyro-S (agentionai, one 58.5 GB file) | 29.7 GB (the other 28.8 GB is the n-gram table) | 38 GB | 48 GB; streams below that | | Gyro-M (agentionai, one 92.0 GB file) | 63.2 GB (+ the 28.8 GB table) | 71 GB | 96 GB; streams below that | +- **The engine's own memory grows with the context.** Beside the mapped file, the engine holds its buffers and the decode window's K/V for the whole `--max-context`. Measured on the M5 Max with Qwen3.8-Flash-Next IQ3_S (phys_footprint 5 s after the engine is ready): 7.68 GiB at 64K context and 10.61 GiB at 262K. Before this version the engine also kept an unused second copy of the K/V in the session buffer, resident in full: 9.27 GiB at 64K and 16.99 GiB at 262K. - **The n-gram table.** Each pack also has a 28.8 GB n-gram table (the model's second file). It is read from the SSD a few rows at a time and does not count against the memory above. - **Disk.** The original Q2_0 and IQ3 packs download roughly 66 to 84 GB, plus preparation and draft data. Unsloth packs are much From 0fe9a5928599db0a60d681f71122d309ef04c8d8 Mon Sep 17 00:00:00 2001 From: oceanio Date: Sun, 18 Oct 2026 08:39:56 +0800 Subject: [PATCH 03/11] metal: the decode window's per-conversation state in one SlotState (slot 0 only, the same buffers) The GDN recurrence and conv history, the QSA int8 K/V, scales, pooled indexer keys, tail and dead rows, and the PLE history move from separate Window members into one SlotState, with their sizes as named constants. The window binds one instance, slot_, allocated as before (same sizes, same order), so every kernel reads the buffers it read before. No behaviour change: this is the state a batch slot will copy. Co-Authored-By: Claude Opus 5.5 --- metal/src/engine/prompt_chain.inc | 28 ++--- metal/src/engine/window.cpp | 186 ++++++++++++++++++------------ 2 files changed, 124 insertions(+), 90 deletions(-) diff --git a/metal/src/engine/prompt_chain.inc b/metal/src/engine/prompt_chain.inc index f9b24133c79..ad0b65a3ce9 100644 --- a/metal/src/engine/prompt_chain.inc +++ b/metal/src/engine/prompt_chain.inc @@ -202,13 +202,13 @@ bool Window::pc_chunk(const int64_t* tokens, int nT, int64_t pos0, std::string& gemm(b.mixh, b.Wbe, b.be, nT, 48, N); PcGdnPArgs ga{(uint32_t) nT, 16, 48, 10240, eps_, 0}; D(P_(b.lgdn, "gdn_pre"), - {mtl::Arg::B(b.qkv), mtl::Arg::B(ghist_[hp][(size_t) gi]), mtl::Arg::B(W_, T(bl + "ssm_conv1d.weight").offset), + {mtl::Arg::B(b.qkv), mtl::Arg::B(slot_.ghist[hp][(size_t) gi]), mtl::Arg::B(W_, T(bl + "ssm_conv1d.weight").offset), mtl::Arg::B(b.al), mtl::Arg::B(b.be), mtl::Arg::B(W_, T(bl + "ssm_dt.bias").offset), mtl::Arg::B(W_, T(bl + "ssm_a").offset), mtl::Arg::B(b.gqn), mtl::Arg::B(b.gkn), mtl::Arg::B(b.gvc), - mtl::Arg::B(b.ggb), mtl::Arg::B(b.gkq), mtl::Arg::B(ghist_[1 - hp][(size_t) gi]), mtl::Arg::V(ga)}, + mtl::Arg::B(b.ggb), mtl::Arg::B(b.gkq), mtl::Arg::B(slot_.ghist[1 - hp][(size_t) gi]), mtl::Arg::V(ga)}, 16 + 24, nT, 256); D(P_(b.lgdn, "gdn_scanr_4", mtl::Constants().i(6, 0)), - {mtl::Arg::B(gst_[(size_t) gi]), mtl::Arg::B(b.gqn), mtl::Arg::B(b.gkn), mtl::Arg::B(b.gvc), mtl::Arg::B(b.ggb), + {mtl::Arg::B(slot_.gst[(size_t) gi]), mtl::Arg::B(b.gqn), mtl::Arg::B(b.gkn), mtl::Arg::B(b.gvc), mtl::Arg::B(b.ggb), mtl::Arg::B(b.gkq), mtl::Arg::B(b.go), mtl::Arg::V(ga)}, 48 * 4, 1, 32 * 4); D(P_(b.lgdn, "gdn_post"), {mtl::Arg::B(b.go), mtl::Arg::B(b.z), mtl::Arg::B(W_, T(bl + "ssm_norm.weight").offset), @@ -233,17 +233,17 @@ bool Window::pc_chunk(const int64_t* tokens, int nT, int64_t pos0, std::string& {mtl::Arg::B(b.qfull), mtl::Arg::B(b.kcur), mtl::Arg::B(b.vcur), mtl::Arg::B(b.iqr), mtl::Arg::B(b.ikr), mtl::Arg::B(W_, T(bl + "attn_q_norm.weight").offset), mtl::Arg::B(W_, T(bl + "attn_k_norm.weight").offset), mtl::Arg::B(W_, T(bl + "indexer.q_norm.weight").offset), mtl::Arg::B(W_, T(bl + "indexer.k_norm.weight").offset), - mtl::Arg::B(cs_), mtl::Arg::B(b.qcur), mtl::Arg::B(b.qidx), mtl::Arg::B(kq_[(size_t) qi]), mtl::Arg::B(vq_[(size_t) qi]), - mtl::Arg::B(ksc_[(size_t) qi]), mtl::Arg::B(vsc_[(size_t) qi]), mtl::Arg::B(tail_[(size_t) qi]), - mtl::Arg::B(dead_[(size_t) qi]), mtl::Arg::B(pooled_[(size_t) qi]), mtl::Arg::V(a), mtl::Arg::B(mrope_)}, + mtl::Arg::B(cs_), mtl::Arg::B(b.qcur), mtl::Arg::B(b.qidx), mtl::Arg::B(slot_.kq[(size_t) qi]), mtl::Arg::B(slot_.vq[(size_t) qi]), + mtl::Arg::B(slot_.ksc[(size_t) qi]), mtl::Arg::B(slot_.vsc[(size_t) qi]), mtl::Arg::B(slot_.tail[(size_t) qi]), + mtl::Arg::B(slot_.dead[(size_t) qi]), mtl::Arg::B(slot_.pooled[(size_t) qi]), mtl::Arg::V(a), mtl::Arg::B(mrope_)}, 33, nT, 256); if (pos0 + nT > PC_QB) { // some row's context is past the budget: the indexer selects if (b.k4) - D(P_(b.lqsa, "qsa_score_mm"), {mtl::Arg::B(b.qidx), mtl::Arg::B(pooled_[(size_t) qi]), mtl::Arg::B(dead_[(size_t) qi]), + D(P_(b.lqsa, "qsa_score_mm"), {mtl::Arg::B(b.qidx), mtl::Arg::B(slot_.pooled[(size_t) qi]), mtl::Arg::B(slot_.dead[(size_t) qi]), mtl::Arg::B(b.scores), mtl::Arg::V(a)}, (uint32_t) (((pos0 + nT) / 4 + 1 + 63) / 64), (nT + 15) / 16, 128); else - D(P_(b.lqsa, "qsa_score"), {mtl::Arg::B(b.qidx), mtl::Arg::B(pooled_[(size_t) qi]), mtl::Arg::B(dead_[(size_t) qi]), + D(P_(b.lqsa, "qsa_score"), {mtl::Arg::B(b.qidx), mtl::Arg::B(slot_.pooled[(size_t) qi]), mtl::Arg::B(slot_.dead[(size_t) qi]), mtl::Arg::B(b.scores), mtl::Arg::V(a)}, (uint32_t) (((pos0 + nT) / 4 + 1 + 7) / 8), nT, 256); } @@ -252,12 +252,12 @@ bool Window::pc_chunk(const int64_t* tokens, int nT, int64_t pos0, std::string& const uint32_t nq = (uint32_t) nT * 24 * 256; D(P_(b.lqsa, "qsa_q_f16"), {mtl::Arg::B(b.qcur), mtl::Arg::B(b.qcurh), mtl::Arg::V(nq)}, (nq + 255) / 256, 1, 256); D(P_(b.lqsa, "qsa_attn_mm", mtl::Constants().b(7, true), "qsa_attn_mm#k4"), - {mtl::Arg::B(b.qcurh), mtl::Arg::B(kq_[(size_t) qi]), mtl::Arg::B(vq_[(size_t) qi]), mtl::Arg::B(ksc_[(size_t) qi]), - mtl::Arg::B(vsc_[(size_t) qi]), mtl::Arg::B(b.qids), mtl::Arg::B(b.part), mtl::Arg::V(a)}, + {mtl::Arg::B(b.qcurh), mtl::Arg::B(slot_.kq[(size_t) qi]), mtl::Arg::B(slot_.vq[(size_t) qi]), mtl::Arg::B(slot_.ksc[(size_t) qi]), + mtl::Arg::B(slot_.vsc[(size_t) qi]), mtl::Arg::B(b.qids), mtl::Arg::B(b.part), mtl::Arg::V(a)}, PC_QCHUNKS, 2, 32, nT); } else { - D(P_(b.lqsa, "qsa_attn"), {mtl::Arg::B(b.qcur), mtl::Arg::B(kq_[(size_t) qi]), mtl::Arg::B(vq_[(size_t) qi]), - mtl::Arg::B(ksc_[(size_t) qi]), mtl::Arg::B(vsc_[(size_t) qi]), mtl::Arg::B(b.qids), + D(P_(b.lqsa, "qsa_attn"), {mtl::Arg::B(b.qcur), mtl::Arg::B(slot_.kq[(size_t) qi]), mtl::Arg::B(slot_.vq[(size_t) qi]), + mtl::Arg::B(slot_.ksc[(size_t) qi]), mtl::Arg::B(slot_.vsc[(size_t) qi]), mtl::Arg::B(b.qids), mtl::Arg::B(b.part), mtl::Arg::V(a)}, PC_QCHUNKS, 2, 256, nT); } @@ -361,7 +361,7 @@ bool Window::pc_chunk(const int64_t* tokens, int nT, int64_t pos0, std::string& {mtl::Arg::B(b.keyn), mtl::Arg::B(b.qn), mtl::Arg::B(b.pval), mtl::Arg::B(W_, off("blk.1.ple_norm_conv.weight")), mtl::Arg::B(b.gated), mtl::Arg::B(b.normW), mtl::Arg::V(pa2)}, HC, nT, 256); D(P(lple_, "ple_conv_add", {}, "ple_conv_add"), - {mtl::Arg::B(pleHist_), mtl::Arg::B(b.normW), kw_arg(), mtl::Arg::B(b.gated), + {mtl::Arg::B(slot_.pleHist), mtl::Arg::B(b.normW), kw_arg(), mtl::Arg::B(b.gated), mtl::Arg::B(b.R), mtl::Arg::V(pa2)}, HC * N / 256, nT, 256); } } @@ -378,7 +378,7 @@ bool Window::pc_chunk(const int64_t* tokens, int nT, int64_t pos0, std::string& if (ple_on_) { using namespace strata::kernels; const float* nw = b.normW.as(); - float* hist = pleHist_.as(); + float* hist = slot_.pleHist.as(); for (int c = 0; c < NG_HC_DIM; ++c) { float* col = hist + (size_t) c * NG_HIST; float tmp[NG_HIST * 2]; diff --git a/metal/src/engine/window.cpp b/metal/src/engine/window.cpp index add35c0e9f3..a4874e57434 100644 --- a/metal/src/engine/window.cpp +++ b/metal/src/engine/window.cpp @@ -116,6 +116,55 @@ const float* const* g_embd_rows = nullptr; const int32_t* g_mrope_host = nullptr; const int32_t* g_mrope_dev = nullptr; +// One conversation's state in the decode window: per GDN layer (l % 4 != 3) the recurrence and the +// conv history (two halves: the window's hp_ says which holds the committed one), per QSA layer (l % 4 == 3) the int8 +// K/V and their scales, the indexer's pooled keys, tail block and dead rows, and the PLE history. The K/V and pooled +// keys are sized from the context; alloc can run for more than one instance (batch slots, later). The MTP draft +// layer's caches (mkq_ ...) are the window's, in the same per-cell layout. +struct SlotState { + // ---- the parts' sizes in bytes, in one place (alloc, the checkpoints and reset use them) + static constexpr size_t kGdnState = 4 * 128 * 48 * 128; // per GDN layer: fp32 recurrence, 48 heads x 128 x 128 + static constexpr size_t kConvHist = 4 * 10240 * 3; // per GDN layer, each half: fp32 conv history, 3 taps + // the fp32 PLE history, 9 rows x 10240 + static constexpr size_t kPleHist = 4 * (size_t) strata::kernels::NG_HIST * strata::kernels::NG_HC_DIM; + static constexpr size_t kTail = 4 * 4 * 128; // per QSA layer: the indexer's tail block (4 fp32 keys) + static constexpr size_t kDead = 4 * 128; // per QSA layer: the tail block's fp32 pooled key + static constexpr size_t kKvCell = 2 * 256; // per cell: int8 K (V the same), 2 KV heads x 256 + static constexpr size_t kScaleCell = 2 * 2 * 4; // per cell: the K scales (V the same) + static constexpr size_t kPooledRow = 4 * 128; // per block of kBlockCells cells: the fp32 pooled key + static constexpr int kBlockCells = 4; + static constexpr size_t kGdnLayers = 36, kQsaLayers = 12; + // the pooled-key rows of a window of `cells` (one per block, the moving spare row and one more) + static int pooled_rows(int cells) { return cells / kBlockCells + 2; } + + std::vector gst, ghist[2]; // 36 GDN layers + std::vector kq, vq, ksc, vsc, tail, dead, pooled; // 12 QSA layers + mtl::Buffer pleHist; + int n_cells = 0; + // the window's cell count for a context of max_cells tokens (the K/V rows, the RoPE and M-RoPE tables) + static int cells_for(int64_t max_cells) { return (int) std::max(4096, (max_cells + 64) / 4 * 4 + 64); } + // the buffers for a window of `cells` = cells_for(the slot's context) (zeroed: a new sequence) + void alloc(mtl::Context& cx, int cells) { + n_cells = cells; + pleHist = cx.alloc(kPleHist); + for (int l = 0; l < 48; ++l) { + if (l % 4 != 3) { + gst.push_back(cx.alloc(kGdnState)); + ghist[0].push_back(cx.alloc(kConvHist)); + ghist[1].push_back(cx.alloc(kConvHist)); + } else { + kq.push_back(cx.alloc((size_t) cells * kKvCell)); + vq.push_back(cx.alloc((size_t) cells * kKvCell)); + ksc.push_back(cx.alloc((size_t) cells * kScaleCell)); + vsc.push_back(cx.alloc((size_t) cells * kScaleCell)); + tail.push_back(cx.alloc(kTail)); + dead.push_back(cx.alloc(kDead)); + pooled.push_back(cx.alloc((size_t) pooled_rows(cells) * kPooledRow)); + } + } + } +}; + class Window { public: bool init(const strata::core::ModelGeometry& g, strata::core::SessionState& ss, const std::string& shard1, int max_t, @@ -229,11 +278,11 @@ class Window { // buffers mtl::Buffer R_, bo_, inj_[2], xn_, lo_, qkv_, z_, al_, be_, yg_, logits_, ids_, wts_, plan_, h_, ye_, gsh_, ush_, ysh_, sg_, qfull_, kcur_, vcur_, iqr_, ikr_, qcur_, qidx_, scores_, qids_, part_, cs_, lm_, apv_, api_, tokout_, - emb_, pkey_, pval_, qkvL_, alL_, beL_, zL_, ydump_, keyn_, qn_, gated_, normW_, pleHist_; + emb_, pkey_, pval_, qkvL_, alL_, beL_, zL_, ydump_, keyn_, qn_, gated_, normW_; Act XA_, XF_, XY_, XQ_, XH_, XE_; int cur_inj_ = 0; - // state - std::vector gst_, ghist_[2], kq_, vq_, ksc_, vsc_, tail_, dead_, pooled_; + // state: the conversation's (slot 0, the only one bound; the window's hp_ picks the committed conv history half) + SlotState slot_; int hp_ = 0; // which of the two history buffers holds the committed conv history (all layers) // commit(n) of a window of several tokens: the GDN fold of the kept prefix waits for the next window and runs at // the start of its command buffer (one submit and wait less a round). Nothing else reads the GDN state between @@ -496,7 +545,7 @@ bool Window::init(const strata::core::ModelGeometry& g, strata::core::SessionSta lsmp_ = cx_->compile({"sample.metal"}, "#define NO_GGML_COMMON 1\n"); lple_ = cx_->compile({strata_metal::resource("STRATA_METAL_ENGINE_KERNELS", "engine-msl", STRATA_METAL_ENGINE_KERNELS_DEFAULT) + "/ple.metal"}, "#define NO_GGML_COMMON 1\n"); - const int MAXC = (int) std::max(4096, (max_cells + 64) / 4 * 4 + 64), MB = MAXC / 4 + 2; + const int MAXC = SlotState::cells_for(max_cells), MB = MAXC / 4 + 2; max_cells_ = MAXC; auto A = [&](size_t bytes) { return cx_->alloc(bytes); }; R_ = A(4 * MAXT * HC * N); bo_ = A(4 * MAXT * N); inj_[0] = A(4 * MAXT * HC); inj_[1] = A(4 * MAXT * HC); @@ -515,7 +564,6 @@ bool Window::init(const strata::core::ModelGeometry& g, strata::core::SessionSta zL_ = A(4 * (size_t) 36 * MAXT * 6144); ydump_ = A(4 * MAXT * 6144); keyn_ = A(4 * MAXT * 10240); qn_ = A(4 * MAXT * 10240); gated_ = A(4 * MAXT * 10240); normW_ = A(4 * MAXT * 10240); mrope_ = A(12 * (size_t) (MAXC + 16)); - pleHist_ = A(4 * 9 * 10240); auto act = [&](int n) { Act a; a.n = n; a.b = A(4 * ((size_t) 4 * MAXT * n + (size_t) MAXT * n / 16)); return a; }; XA_ = act(N); XF_ = act(N); XY_ = act(6144); XQ_ = act(6144); XH_ = act(NFF); XE_ = act(N); { // Gyro (native_experts.txt v5 + hadamard.txt): the rotation's signs and the rotated gate/up input @@ -550,21 +598,7 @@ bool Window::init(const strata::core::ModelGeometry& g, strata::core::SessionSta } for (int t = 0; t < MAXT; ++t) for (int c = 0; c < 2051; ++c) qids_.as()[t * 2051 + c] = c; - for (int l = 0; l < 48; ++l) { - if (l % 4 != 3) { - gst_.push_back(A(4 * 128 * 48 * 128)); - ghist_[0].push_back(A(4 * 10240 * 3)); - ghist_[1].push_back(A(4 * 10240 * 3)); - } else { - kq_.push_back(A((size_t) MAXC * 2 * 256)); - vq_.push_back(A((size_t) MAXC * 2 * 256)); - ksc_.push_back(A(2 * (size_t) MAXC * 2 * 4)); - vsc_.push_back(A(2 * (size_t) MAXC * 2 * 4)); - tail_.push_back(A(4 * 4 * 128)); - dead_.push_back(A(4 * 128)); - pooled_.push_back(A(4 * (size_t) MB * 128)); - } - } + slot_.alloc(*cx_, MAXC); // ---- the HC read in Q8_0 (the owner's default on Metal, as STRATA_HC_Q8 on gfx12): the BF16 down / up / inject // matrices quantized once at load (ggml's q8_0 rule, ref_quant), read by gr_down_q8 / gr_up_q8 { @@ -796,12 +830,12 @@ void Window::layer(std::vector& ds, int l, int ntok, int64_t pos0 mtl::Dispatch d; d.p = P(lgdn_, "gdn_step", mtl::Constants().i(2, ntok), "gdn_step#" + std::to_string(ntok)); const int hp = hp_; - d.args = {mtl::Arg::B(gst_[(size_t) gi]), mtl::Arg::B(ghist_[hp][(size_t) gi]), mtl::Arg::B(qkvL_, o_qkv), + d.args = {mtl::Arg::B(slot_.gst[(size_t) gi]), mtl::Arg::B(slot_.ghist[hp][(size_t) gi]), mtl::Arg::B(qkvL_, o_qkv), mtl::Arg::B(alL_, o_ab), mtl::Arg::B(beL_, o_ab), mtl::Arg::B(zL_, o_z), mtl::Arg::B(W_, T(b + "ssm_conv1d.weight").offset), mtl::Arg::B(W_, T(b + "ssm_dt.bias").offset), mtl::Arg::B(W_, T(b + "ssm_a").offset), mtl::Arg::B(W_, T(b + "ssm_norm.weight").offset), mtl::Arg::B(fuse_ ? XY_.b : yg_), mtl::Arg::V(a), - mtl::Arg::B(ghist_[1 - hp][(size_t) gi])}; + mtl::Arg::B(slot_.ghist[1 - hp][(size_t) gi])}; d.grid[0] = 48; d.tg[0] = 512; push(ds, d); @@ -825,9 +859,9 @@ void Window::layer(std::vector& ds, int l, int ntok, int64_t pos0 d.args = {mtl::Arg::B(qfull_), mtl::Arg::B(kcur_), mtl::Arg::B(vcur_), mtl::Arg::B(iqr_), mtl::Arg::B(ikr_), mtl::Arg::B(W_, T(b + "attn_q_norm.weight").offset), mtl::Arg::B(W_, T(b + "attn_k_norm.weight").offset), mtl::Arg::B(W_, T(b + "indexer.q_norm.weight").offset), mtl::Arg::B(W_, T(b + "indexer.k_norm.weight").offset), - mtl::Arg::B(cs_), mtl::Arg::B(qcur_), mtl::Arg::B(qidx_), mtl::Arg::B(kq_[(size_t) qi]), - mtl::Arg::B(vq_[(size_t) qi]), mtl::Arg::B(ksc_[(size_t) qi]), mtl::Arg::B(vsc_[(size_t) qi]), - mtl::Arg::B(tail_[(size_t) qi]), mtl::Arg::B(dead_[(size_t) qi]), mtl::Arg::B(pooled_[(size_t) qi]), + mtl::Arg::B(cs_), mtl::Arg::B(qcur_), mtl::Arg::B(qidx_), mtl::Arg::B(slot_.kq[(size_t) qi]), + mtl::Arg::B(slot_.vq[(size_t) qi]), mtl::Arg::B(slot_.ksc[(size_t) qi]), mtl::Arg::B(slot_.vsc[(size_t) qi]), + mtl::Arg::B(slot_.tail[(size_t) qi]), mtl::Arg::B(slot_.dead[(size_t) qi]), mtl::Arg::B(slot_.pooled[(size_t) qi]), mtl::Arg::V(a), mtl::Arg::B(mrope_)}; d.grid[0] = 33; d.grid[1] = ntok; @@ -836,7 +870,7 @@ void Window::layer(std::vector& ds, int l, int ntok, int64_t pos0 if (rec_) rec_->push_back({(int) ds.size() - 1, 0}); mtl::Dispatch s; s.p = P(lqsa_, "qsa_score", {}, "qsa_score"); - s.args = {mtl::Arg::B(qidx_), mtl::Arg::B(pooled_[(size_t) qi]), mtl::Arg::B(dead_[(size_t) qi]), mtl::Arg::B(scores_), + s.args = {mtl::Arg::B(qidx_), mtl::Arg::B(slot_.pooled[(size_t) qi]), mtl::Arg::B(slot_.dead[(size_t) qi]), mtl::Arg::B(scores_), mtl::Arg::V(a)}; s.grid[0] = (uint32_t) (((pos0 + ntok) / 4 + 1 + 7) / 8); s.grid[1] = ntok; @@ -857,8 +891,8 @@ void Window::layer(std::vector& ds, int l, int ntok, int64_t pos0 } mtl::Dispatch at; at.p = P(lqsa_, "qsa_attn", {}, "qsa_attn"); - at.args = {mtl::Arg::B(qcur_), mtl::Arg::B(kq_[(size_t) qi]), mtl::Arg::B(vq_[(size_t) qi]), mtl::Arg::B(ksc_[(size_t) qi]), - mtl::Arg::B(vsc_[(size_t) qi]), mtl::Arg::B(qids_), mtl::Arg::B(part_), mtl::Arg::V(a)}; + at.args = {mtl::Arg::B(qcur_), mtl::Arg::B(slot_.kq[(size_t) qi]), mtl::Arg::B(slot_.vq[(size_t) qi]), mtl::Arg::B(slot_.ksc[(size_t) qi]), + mtl::Arg::B(slot_.vsc[(size_t) qi]), mtl::Arg::B(qids_), mtl::Arg::B(part_), mtl::Arg::V(a)}; at.grid[0] = 33; at.grid[1] = 2; at.grid[2] = ntok; @@ -1186,7 +1220,7 @@ bool Window::run(int Tn, const int32_t* tokens, int64_t pos0, bool head, int32_t g3.grid[0] = HC; g3.grid[1] = Tn; g3.tg[0] = 256; push(ds, g3); g4.p = P(lple_, "ple_conv_add", {}, "ple_conv_add"); - g4.args = {mtl::Arg::B(pleHist_), mtl::Arg::B(normW_), kw_arg(), + g4.args = {mtl::Arg::B(slot_.pleHist), mtl::Arg::B(normW_), kw_arg(), mtl::Arg::B(gated_), mtl::Arg::B(R_), mtl::Arg::V(pa)}; g4.grid[0] = HC * N / 256; g4.grid[1] = Tn; g4.tg[0] = 256; push(ds, g4); @@ -1486,11 +1520,11 @@ void Window::fold(std::vector& ds, int n_keep, int hp) { GdnArgs a{16, 48, 10240, 1u, eps_, 0}; mtl::Dispatch d; d.p = P(lgdn_, "gdn_step", mtl::Constants().i(2, n_keep), "gdn_step#" + std::to_string(n_keep)); - d.args = {mtl::Arg::B(gst_[(size_t) gi]), mtl::Arg::B(ghist_[hp][(size_t) gi]), mtl::Arg::B(qkvL_, o_qkv), + d.args = {mtl::Arg::B(slot_.gst[(size_t) gi]), mtl::Arg::B(slot_.ghist[hp][(size_t) gi]), mtl::Arg::B(qkvL_, o_qkv), mtl::Arg::B(alL_, o_ab), mtl::Arg::B(beL_, o_ab), mtl::Arg::B(zL_, o_z), mtl::Arg::B(W_, T(b + "ssm_conv1d.weight").offset), mtl::Arg::B(W_, T(b + "ssm_dt.bias").offset), mtl::Arg::B(W_, T(b + "ssm_a").offset), mtl::Arg::B(W_, T(b + "ssm_norm.weight").offset), - mtl::Arg::B(ydump_), mtl::Arg::V(a), mtl::Arg::B(ghist_[1 - hp][(size_t) gi])}; + mtl::Arg::B(ydump_), mtl::Arg::V(a), mtl::Arg::B(slot_.ghist[1 - hp][(size_t) gi])}; d.grid[0] = 48; d.tg[0] = 512; nb_ = true; // the layers' folds are independent of each other @@ -1508,18 +1542,18 @@ bool Window::checkpoint_save(strata::core::ConversationCheckpoint& c, std::strin cx_->run_graph(ds); pend_fold_ = 0; } - const size_t gs = 4 * 128 * 48 * 128, hs = 4 * 10240 * 3, ts = 4 * 4 * 128, dsz = 4 * 128; - c.gdn.resize(gst_.size() * (gs + hs)); - for (size_t i = 0; i < gst_.size(); ++i) { - std::memcpy(c.gdn.data() + i * (gs + hs), gst_[i].ptr, gs); - std::memcpy(c.gdn.data() + i * (gs + hs) + gs, ghist_[hp_][i].ptr, hs); + const size_t gs = SlotState::kGdnState, hs = SlotState::kConvHist, ts = SlotState::kTail, dsz = SlotState::kDead; + c.gdn.resize(slot_.gst.size() * (gs + hs)); + for (size_t i = 0; i < slot_.gst.size(); ++i) { + std::memcpy(c.gdn.data() + i * (gs + hs), slot_.gst[i].ptr, gs); + std::memcpy(c.gdn.data() + i * (gs + hs) + gs, slot_.ghist[hp_][i].ptr, hs); } - c.ple.assign((const uint8_t*) pleHist_.ptr, (const uint8_t*) pleHist_.ptr + 4 * 9 * 10240); - c.tails.resize(tail_.size() * ts); - c.dead.resize(dead_.size() * dsz); - for (size_t j = 0; j < tail_.size(); ++j) { - std::memcpy(c.tails.data() + j * ts, tail_[j].ptr, ts); - std::memcpy(c.dead.data() + j * dsz, dead_[j].ptr, dsz); + c.ple.assign((const uint8_t*) slot_.pleHist.ptr, (const uint8_t*) slot_.pleHist.ptr + SlotState::kPleHist); + c.tails.resize(slot_.tail.size() * ts); + c.dead.resize(slot_.dead.size() * dsz); + for (size_t j = 0; j < slot_.tail.size(); ++j) { + std::memcpy(c.tails.data() + j * ts, slot_.tail[j].ptr, ts); + std::memcpy(c.dead.data() + j * dsz, slot_.dead[j].ptr, dsz); } c.block_pos.clear(); return true; @@ -1529,39 +1563,39 @@ void Window::reset() { pend_fold_ = 0; last_T_ = 0; auto zero = [](const mtl::Buffer& b, size_t n) { std::memset(b.ptr, 0, n); }; - for (size_t i = 0; i < gst_.size(); ++i) { - zero(gst_[i], 4 * 128 * 48 * 128); - zero(ghist_[0][i], 4 * 10240 * 3); - zero(ghist_[1][i], 4 * 10240 * 3); + for (size_t i = 0; i < slot_.gst.size(); ++i) { + zero(slot_.gst[i], SlotState::kGdnState); + zero(slot_.ghist[0][i], SlotState::kConvHist); + zero(slot_.ghist[1][i], SlotState::kConvHist); } - zero(pleHist_, 4 * 9 * 10240); - for (size_t j = 0; j < tail_.size(); ++j) { - zero(tail_[j], 4 * 4 * 128); - zero(dead_[j], 4 * 128); + zero(slot_.pleHist, SlotState::kPleHist); + for (size_t j = 0; j < slot_.tail.size(); ++j) { + zero(slot_.tail[j], SlotState::kTail); + zero(slot_.dead[j], SlotState::kDead); } for (size_t j = 0; j < mtail_.size(); ++j) { - zero(mtail_[j], 4 * 4 * 128); - zero(mdead_[j], 4 * 128); + zero(mtail_[j], SlotState::kTail); + zero(mdead_[j], SlotState::kDead); } } bool Window::checkpoint_restore(const strata::core::ConversationCheckpoint& c, std::string& err) { - const size_t gs = 4 * 128 * 48 * 128, hs = 4 * 10240 * 3, ts = 4 * 4 * 128, dsz = 4 * 128; - if (c.gdn.size() != gst_.size() * (gs + hs) || c.ple.size() != (size_t) 4 * 9 * 10240 || - c.tails.size() != tail_.size() * ts || c.dead.size() != dead_.size() * dsz) { + const size_t gs = SlotState::kGdnState, hs = SlotState::kConvHist, ts = SlotState::kTail, dsz = SlotState::kDead; + if (c.gdn.size() != slot_.gst.size() * (gs + hs) || c.ple.size() != SlotState::kPleHist || + c.tails.size() != slot_.tail.size() * ts || c.dead.size() != slot_.dead.size() * dsz) { err = "metal window: the checkpoint is not this engine's (sizes differ)"; return false; } pend_fold_ = 0; // the restored state replaces whatever was pending last_T_ = 0; - for (size_t i = 0; i < gst_.size(); ++i) { - std::memcpy(gst_[i].ptr, c.gdn.data() + i * (gs + hs), gs); - std::memcpy(ghist_[hp_][i].ptr, c.gdn.data() + i * (gs + hs) + gs, hs); + for (size_t i = 0; i < slot_.gst.size(); ++i) { + std::memcpy(slot_.gst[i].ptr, c.gdn.data() + i * (gs + hs), gs); + std::memcpy(slot_.ghist[hp_][i].ptr, c.gdn.data() + i * (gs + hs) + gs, hs); } - std::memcpy(pleHist_.ptr, c.ple.data(), c.ple.size()); - for (size_t j = 0; j < tail_.size(); ++j) { - std::memcpy(tail_[j].ptr, c.tails.data() + j * ts, ts); - std::memcpy(dead_[j].ptr, c.dead.data() + j * dsz, dsz); + std::memcpy(slot_.pleHist.ptr, c.ple.data(), c.ple.size()); + for (size_t j = 0; j < slot_.tail.size(); ++j) { + std::memcpy(slot_.tail[j].ptr, c.tails.data() + j * ts, ts); + std::memcpy(slot_.dead[j].ptr, c.dead.data() + j * dsz, dsz); } return true; } @@ -1600,7 +1634,7 @@ bool Window::commit(int n_keep, std::string& err) { if (ple_on_) { const float* norm = normW_.as() + (size_t) t * NG_HC_DIM; for (int c = 0; c < NG_HC_DIM; ++c) { - float* col = pleHist_.as() + (size_t) c * NG_HIST; + float* col = slot_.pleHist.as() + (size_t) c * NG_HIST; for (int r = 0; r + 1 < NG_HIST; ++r) col[r] = col[r + 1]; col[NG_HIST - 1] = norm[c]; } @@ -2354,12 +2388,12 @@ bool Window::prompt_chunk(const int32_t* tokens, int Tn, int64_t pos0, std::stri gemm(p.mixh, p.Wbe, p.be, T, 48, N); PgGdnP ga{(uint32_t) T, 16, 48, 10240, eps_, 0}; D(PP(p.lgdn, "gdn_pre"), - {mtl::Arg::B(p.qkv), mtl::Arg::B(ghist_[hp_][(size_t) gi]), mtl::Arg::B(W_, TN(b + "ssm_conv1d.weight").offset), + {mtl::Arg::B(p.qkv), mtl::Arg::B(slot_.ghist[hp_][(size_t) gi]), mtl::Arg::B(W_, TN(b + "ssm_conv1d.weight").offset), mtl::Arg::B(p.al), mtl::Arg::B(p.be), mtl::Arg::B(W_, TN(b + "ssm_dt.bias").offset), mtl::Arg::B(W_, TN(b + "ssm_a").offset), mtl::Arg::B(p.gqn), mtl::Arg::B(p.gkn), mtl::Arg::B(p.gvc), mtl::Arg::B(p.ggb), mtl::Arg::B(p.gkq), - mtl::Arg::B(ghist_[1 - hp_][(size_t) gi]), mtl::Arg::V(ga)}, 16 + 24, T, 256); + mtl::Arg::B(slot_.ghist[1 - hp_][(size_t) gi]), mtl::Arg::V(ga)}, 16 + 24, T, 256); D(PP(p.lgdn, "gdn_scanr_4", mtl::Constants().i(6, 0), "gdn_scanr_4"), - {mtl::Arg::B(gst_[(size_t) gi]), mtl::Arg::B(p.gqn), mtl::Arg::B(p.gkn), mtl::Arg::B(p.gvc), mtl::Arg::B(p.ggb), + {mtl::Arg::B(slot_.gst[(size_t) gi]), mtl::Arg::B(p.gqn), mtl::Arg::B(p.gkn), mtl::Arg::B(p.gvc), mtl::Arg::B(p.ggb), mtl::Arg::B(p.gkq), mtl::Arg::B(p.go), mtl::Arg::V(ga)}, 48 * 4, 1, 32 * 4); D(PP(p.lgdn, "gdn_post"), {mtl::Arg::B(p.go), mtl::Arg::B(p.z), mtl::Arg::B(W_, TN(b + "ssm_norm.weight").offset), mtl::Arg::B(p.ygh), mtl::Arg::B(p.yg), mtl::Arg::V(ga)}, 48, T, 128); @@ -2382,11 +2416,11 @@ bool Window::prompt_chunk(const int32_t* tokens, int Tn, int64_t pos0, std::stri {mtl::Arg::B(p.qfull), mtl::Arg::B(p.kcur), mtl::Arg::B(p.vcur), mtl::Arg::B(p.iqr), mtl::Arg::B(p.ikr), mtl::Arg::B(W_, TN(b + "attn_q_norm.weight").offset), mtl::Arg::B(W_, TN(b + "attn_k_norm.weight").offset), mtl::Arg::B(W_, TN(b + "indexer.q_norm.weight").offset), mtl::Arg::B(W_, TN(b + "indexer.k_norm.weight").offset), - mtl::Arg::B(cs_), mtl::Arg::B(p.qcur), mtl::Arg::B(p.qidx), mtl::Arg::B(kq_[(size_t) qi]), mtl::Arg::B(vq_[(size_t) qi]), - mtl::Arg::B(ksc_[(size_t) qi]), mtl::Arg::B(vsc_[(size_t) qi]), mtl::Arg::B(tail_[(size_t) qi]), - mtl::Arg::B(dead_[(size_t) qi]), mtl::Arg::B(pooled_[(size_t) qi]), mtl::Arg::V(a), mtl::Arg::B(mrope_)}, 33, T, 256); + mtl::Arg::B(cs_), mtl::Arg::B(p.qcur), mtl::Arg::B(p.qidx), mtl::Arg::B(slot_.kq[(size_t) qi]), mtl::Arg::B(slot_.vq[(size_t) qi]), + mtl::Arg::B(slot_.ksc[(size_t) qi]), mtl::Arg::B(slot_.vsc[(size_t) qi]), mtl::Arg::B(slot_.tail[(size_t) qi]), + mtl::Arg::B(slot_.dead[(size_t) qi]), mtl::Arg::B(slot_.pooled[(size_t) qi]), mtl::Arg::V(a), mtl::Arg::B(mrope_)}, 33, T, 256); if (pos0 + T > 2051) - D(PP(p.lqsa, "qsa_score_mm"), {mtl::Arg::B(p.qidx), mtl::Arg::B(pooled_[(size_t) qi]), mtl::Arg::B(dead_[(size_t) qi]), + D(PP(p.lqsa, "qsa_score_mm"), {mtl::Arg::B(p.qidx), mtl::Arg::B(slot_.pooled[(size_t) qi]), mtl::Arg::B(slot_.dead[(size_t) qi]), mtl::Arg::B(p.scores), mtl::Arg::V(a)}, (uint32_t) (((pos0 + T) / 4 + 1 + 63) / 64), (T + 15) / 16, 128); D(PP(p.lqsa, "qsa_topk"), {mtl::Arg::B(p.scores), mtl::Arg::B(p.qids), mtl::Arg::V(a)}, T, 1, 1024); @@ -2394,16 +2428,16 @@ bool Window::prompt_chunk(const int32_t* tokens, int Tn, int64_t pos0, std::stri const uint32_t nq = (uint32_t) T * 24 * 256; D(PP(p.lqsa, "qsa_q_f16"), {mtl::Arg::B(p.qcur), mtl::Arg::B(p.qcurh), mtl::Arg::V(nq)}, (nq + 255) / 256, 1, 256); D(PP(p.lqsa, "qsa_flash", mtl::Constants().b(7, p.k4), "qsa_flash"), - {mtl::Arg::B(p.qcurh), mtl::Arg::B(kq_[(size_t) qi]), mtl::Arg::B(vq_[(size_t) qi]), mtl::Arg::B(ksc_[(size_t) qi]), - mtl::Arg::B(vsc_[(size_t) qi]), mtl::Arg::B(p.qids), mtl::Arg::B(p.qfull), mtl::Arg::B(p.xqh), mtl::Arg::V(a)}, + {mtl::Arg::B(p.qcurh), mtl::Arg::B(slot_.kq[(size_t) qi]), mtl::Arg::B(slot_.vq[(size_t) qi]), mtl::Arg::B(slot_.ksc[(size_t) qi]), + mtl::Arg::B(slot_.vsc[(size_t) qi]), mtl::Arg::B(p.qids), mtl::Arg::B(p.qfull), mtl::Arg::B(p.xqh), mtl::Arg::V(a)}, 2, T, 128); } else { // v1 attention + merge in slices of QS rows (the partials of a slice only) for (int t0 = 0; t0 < T; t0 += p.QS) { const int ns = std::min(p.QS, T - t0); QsaArgs as{(uint32_t) (pos0 + t0), 0, (uint32_t) max_cells_, (uint32_t) ns, eps_}; - D(PP(p.lqsa, "qsa_attn"), {mtl::Arg::B(p.qcur, (size_t) t0 * 24 * 256 * 4), mtl::Arg::B(kq_[(size_t) qi]), - mtl::Arg::B(vq_[(size_t) qi]), mtl::Arg::B(ksc_[(size_t) qi]), mtl::Arg::B(vsc_[(size_t) qi]), + D(PP(p.lqsa, "qsa_attn"), {mtl::Arg::B(p.qcur, (size_t) t0 * 24 * 256 * 4), mtl::Arg::B(slot_.kq[(size_t) qi]), + mtl::Arg::B(slot_.vq[(size_t) qi]), mtl::Arg::B(slot_.ksc[(size_t) qi]), mtl::Arg::B(slot_.vsc[(size_t) qi]), mtl::Arg::B(p.qids, (size_t) t0 * 2051 * 4), mtl::Arg::B(p.part), mtl::Arg::V(as)}, 33, 2, 256, (uint32_t) ns); D(PP(p.lqsa, "qsa_merge"), {mtl::Arg::B(p.part), mtl::Arg::B(p.qfull, (size_t) t0 * 12288 * 4), @@ -2533,7 +2567,7 @@ bool Window::prompt_chunk(const int32_t* tokens, int Tn, int64_t pos0, std::stri {mtl::Arg::B(p.keyn), mtl::Arg::B(p.qn), mtl::Arg::B(p.pval), mtl::Arg::B(W_, off("blk.1.ple_norm_conv.weight")), mtl::Arg::B(p.gated), mtl::Arg::B(p.normW), mtl::Arg::V(pl)}, HC, T, 256); D(P(lple_, "ple_conv_add", {}, "ple_conv_add"), - {mtl::Arg::B(pleHist_), mtl::Arg::B(p.normW), kw_arg(), mtl::Arg::B(p.gated), + {mtl::Arg::B(slot_.pleHist), mtl::Arg::B(p.normW), kw_arg(), mtl::Arg::B(p.gated), mtl::Arg::B(p.R), mtl::Arg::V(pl)}, HC * N / 256, T, 256); } } @@ -2561,7 +2595,7 @@ bool Window::prompt_chunk(const int32_t* tokens, int Tn, int64_t pos0, std::stri hp_ = 1 - hp_; if (ple_on_) { using namespace strata::kernels; - float* hist = pleHist_.as(); + float* hist = slot_.pleHist.as(); const float* nw = p.normW.as(); for (int c = 0; c < NG_HC_DIM; ++c) { float col[NG_HIST]; From 946a1c99689781dfc1cf07f35676900cdfd920a6 Mon Sep 17 00:00:00 2001 From: oceanio Date: Sun, 18 Oct 2026 08:39:57 +0800 Subject: [PATCH 04/11] metal: one reading of the free memory on macOS (host_memory_now, sample_arena_mem), used by the streaming arena host_memory.cpp reads hw.memsize and host_statistics64 once, with arena_budget.h's arena_avail rule (free and purgeable pages and the clean file cache, not other programs' inactive anonymous memory). sample_arena_mem fills ArenaMem from it. The streaming arena now reads its numbers through these two; the values are the same as its own code read, except the fallback when the VM statistics cannot be read (RAM less the footprint, before: RAM less the Metal buffers). The batch slots use the same reading next. Co-Authored-By: Claude Opus 5.5 --- metal/src/engine/host_memory.cpp | 28 ++++++++++++++ metal/src/engine/host_memory.h | 15 ++++++++ metal/src/engine/window.cpp | 63 ++++++++++++++++++-------------- 3 files changed, 78 insertions(+), 28 deletions(-) create mode 100644 metal/src/engine/host_memory.cpp create mode 100644 metal/src/engine/host_memory.h diff --git a/metal/src/engine/host_memory.cpp b/metal/src/engine/host_memory.cpp new file mode 100644 index 00000000000..fda8240dea9 --- /dev/null +++ b/metal/src/engine/host_memory.cpp @@ -0,0 +1,28 @@ +// metal/src/engine/host_memory.cpp - see host_memory.h. +#include "host_memory.h" + +#include "arena_budget.h" + +#include +#include +#include + +namespace strata::metal_engine { + +bool host_memory_now(double& available, double& phys) { + available = 0; + phys = 0; + uint64_t total = 0; + size_t len = sizeof total; + if (sysctlbyname("hw.memsize", &total, &len, nullptr, 0) != 0 || total == 0) return false; + vm_statistics64_data_t vs{}; + mach_msg_type_number_t cnt = HOST_VM_INFO64_COUNT; + if (host_statistics64(mach_host_self(), HOST_VM_INFO64, (host_info64_t) &vs, &cnt) != KERN_SUCCESS) return false; + const double pg = (double) vm_page_size; + phys = (double) total; + available = arena_avail(vs.free_count * pg, vs.purgeable_count * pg, vs.inactive_count * pg, vs.speculative_count * pg, + vs.external_page_count * pg); + return true; +} + +} // namespace strata::metal_engine diff --git a/metal/src/engine/host_memory.h b/metal/src/engine/host_memory.h new file mode 100644 index 00000000000..6dab45dd7d0 --- /dev/null +++ b/metal/src/engine/host_memory.h @@ -0,0 +1,15 @@ +// metal/src/engine/host_memory.h - the memory macOS can give the engine now, read one way for every user: the +// streaming arena (window.cpp). +// +// The rule is arena_budget.h's arena_avail: free and purgeable pages and the clean file cache (inactive and speculative +// pages, at most as many as are file-backed). Other programs' inactive anonymous memory is left out: macOS frees it +// only by compressing or swapping it. +#pragma once + +namespace strata::metal_engine { + +/// `available`: bytes macOS can give now without paging anything out (the rule above); `phys`: hw.memsize. False +/// (both 0) when the VM statistics or the RAM size cannot be read. +bool host_memory_now(double& available, double& phys); + +} // namespace strata::metal_engine diff --git a/metal/src/engine/window.cpp b/metal/src/engine/window.cpp index a4874e57434..7196b2190ce 100644 --- a/metal/src/engine/window.cpp +++ b/metal/src/engine/window.cpp @@ -16,6 +16,7 @@ #include "resource_path.h" #include "ref_quant.h" #include "arena_budget.h" +#include "host_memory.h" #include "stream.h" #include "../../tests/coop_probe.h" @@ -1276,6 +1277,34 @@ ProcMem proc_mem() { } return r; } +// The engine's memory as arena_budget.h's rules read it, for the streaming arena: the RAM, +// the Metal working set, the Metal buffers (less `model_bytes`, a resident pack's wrapped file pages), what the process +// holds, macOS's file-cache floor, what macOS can give now (host_memory_now) and the overrides from the environment. +ArenaMem sample_arena_mem(const mtl::Context& cx, double model_bytes) { + ArenaMem am; + double avail = 0, phys = 0; + const bool vm_ok = host_memory_now(avail, phys); + if (!vm_ok) { + uint64_t p = 0; + size_t len = sizeof p; + sysctlbyname("hw.memsize", &p, &len, nullptr, 0); + phys = (double) p; + } + am.phys = phys; + am.working_set = (double) cx.recommended_working_set(); + am.allocated = std::max(0.0, (double) cx.allocated() - model_bytes); + am.proc_footprint = proc_mem().footprint; // everything the process holds, not only the Metal buffers + { // the file cache macOS keeps before it compresses anonymous memory (a 4-byte count of pages, read into a zeroed 8) + unsigned long long pages = 0; + size_t pl = sizeof pages; + if (sysctlbyname("vm.vm_page_filecache_min", &pages, &pl, nullptr, 0) == 0) am.filecache_min = (double) pages * (double) vm_page_size; + } + am.available = vm_ok ? avail : am.phys - footprint_base(am) - model_bytes; + if (const char* e = std::getenv("STRATA_METAL_KEEP_FREE_GB")) am.keep_gb = std::atof(e); + if (const char* e = std::getenv("STRATA_METAL_ARENA_GB")) am.arena_gb = std::atof(e); + if (const char* e = std::getenv("STRATA_METAL_LEAVE_SHARE")) am.leave_share = std::atof(e); // for measuring; the default is kLeaveShare + return am; +} void stream_atexit() { if (g_stream_window) g_stream_window->stream_report(stderr); } @@ -1286,9 +1315,6 @@ void stream_atexit() { // froze with an 89 GB arena and 7 GB left beside a 50 GB n-gram table read through the file cache; arena_budget.h); // STRATA_METAL_ARENA_GB sets it. bool Window::stream_arena(std::string& err) { - uint64_t phys = 0; - size_t len = sizeof phys; - sysctlbyname("hw.memsize", &phys, &len, nullptr, 0); const double gb = 1e9; if (!(std::getenv("STRATA_METAL_KEEP_HOST_WEIGHTS") && std::string(std::getenv("STRATA_METAL_KEEP_HOST_WEIGHTS")) == "1")) { const uint64_t g = strata::gpu::purge_device_allocations(32u << 20); @@ -1296,35 +1322,16 @@ bool Window::stream_arena(std::string& err) { "Metal window)\n", g / 1e9); } const double used = (double) cx_->allocated(); - // what macOS gives now without paging anything out (arena_avail: free, purgeable, the clean file cache; not other - // programs' inactive memory, freed only by compressing or swapping it) + ArenaMem am = sample_arena_mem(*cx_, 0.0); // arena_budget.h: the rule, testable + const double phys = am.phys; + const ProcMem pm0{am.proc_footprint, proc_mem().rss}; + // what macOS gives now without paging anything out (host_memory.h), asked again before the arena is filled auto avail_now = [&]() -> double { - vm_statistics64_data_t vs; - mach_msg_type_number_t cnt = HOST_VM_INFO64_COUNT; - if (host_statistics64(mach_host_self(), HOST_VM_INFO64, (host_info64_t) &vs, &cnt) == KERN_SUCCESS) { - const double pg = (double) vm_page_size; - return arena_avail(vs.free_count * pg, vs.purgeable_count * pg, vs.inactive_count * pg, vs.speculative_count * pg, - vs.external_page_count * pg); - } - return (double) phys - (double) cx_->allocated(); + double a = 0, p = 0; + return host_memory_now(a, p) ? a : phys - footprint_base(am); }; - ArenaMem am; // arena_budget.h: the rule, testable - am.phys = (double) phys; - am.working_set = (double) cx_->recommended_working_set(); - am.allocated = used; - const ProcMem pm0 = proc_mem(); // everything the process holds, not only the Metal buffers (the old 67% rule missed it) - am.proc_footprint = pm0.footprint; - { // the file cache macOS keeps before it compresses anonymous memory (a 4-byte count of pages, read into a zeroed 8) - unsigned long long pages = 0; - size_t pl = sizeof pages; - if (sysctlbyname("vm.vm_page_filecache_min", &pages, &pl, nullptr, 0) == 0) am.filecache_min = (double) pages * (double) vm_page_size; - } - am.available = avail_now(); if (ss_ && ss_->ple.table && ss_->ple.table->mode() == strata::kernels::PleIo::Mmap) am.ple_mapped = (double) ss_->ple.table->rows() * 192.0; // a row is at most 192 bytes (Q8_0 + its scale) - if (const char* e = std::getenv("STRATA_METAL_KEEP_FREE_GB")) am.keep_gb = std::atof(e); - if (const char* e = std::getenv("STRATA_METAL_ARENA_GB")) am.arena_gb = std::atof(e); - if (const char* e = std::getenv("STRATA_METAL_LEAVE_SHARE")) am.leave_share = std::atof(e); // for measuring; the default is kLeaveShare const double keep = arena_keep(am); const double avail = am.available; double budget = arena_budget(am); From 5bae944a07eebd964bee6a3aa5429ad9089ea055 Mon Sep 17 00:00:00 2001 From: oceanio Date: Sun, 18 Oct 2026 08:40:00 +0800 Subject: [PATCH 05/11] memory guard: a free-memory reading on macOS, so parking and session files are not refused on the Metal engine (#2502) sample_host_memory() had readings for Windows and Linux only, so on a Mac the guard could not read the RAM and refused every park and session save / restore. Under STRATA_GPU_METAL it now takes the Metal engine's own reading (host_memory_now: free and purgeable pages and the clean file cache, hw.memsize as the total), the one its streaming arena uses. Other programs' inactive anonymous memory is not counted: macOS frees it only by compressing or swapping it. No cgroups on macOS. Windows and Linux compile the file as before. The PLE residency policy (0.2.7) reads the same figure. It stays off on the Metal engine, as in a build that cannot read the RAM: it was measured on the PC, and with this reading it would lock the 26.8 GiB PLE table of Qwen IQ3_S in RAM (the engine's footprint at 262K went from 10.6 to 37.4 GiB). Co-Authored-By: Claude Opus 5.5 --- metal/src/engine/host_memory.h | 3 ++- src/core/conversation_memory.cpp | 12 ++++++++++++ src/program/generate.cpp | 7 +++++++ 3 files changed, 21 insertions(+), 1 deletion(-) diff --git a/metal/src/engine/host_memory.h b/metal/src/engine/host_memory.h index 6dab45dd7d0..795e757b57c 100644 --- a/metal/src/engine/host_memory.h +++ b/metal/src/engine/host_memory.h @@ -1,5 +1,6 @@ // metal/src/engine/host_memory.h - the memory macOS can give the engine now, read one way for every user: the -// streaming arena (window.cpp). +// streaming arena (window.cpp) and the memory guard of parking and session files +// (src/core/conversation_memory.cpp, sample_host_memory). // // The rule is arena_budget.h's arena_avail: free and purgeable pages and the clean file cache (inactive and speculative // pages, at most as many as are file-backed). Other programs' inactive anonymous memory is left out: macOS frees it diff --git a/src/core/conversation_memory.cpp b/src/core/conversation_memory.cpp index 78a8aa0a065..d05380b2ba5 100644 --- a/src/core/conversation_memory.cpp +++ b/src/core/conversation_memory.cpp @@ -17,6 +17,10 @@ #define NOMINMAX #endif #include +#elif defined(STRATA_GPU_METAL) +namespace strata::metal_engine { // metal/src/engine/host_memory.h: the engine's one reading of the free memory +bool host_memory_now(double& available, double& phys); +} // namespace strata::metal_engine #endif namespace strata::core { @@ -286,6 +290,14 @@ std::optional sample_host_memory() { return combine(std::min(status.ullAvailPhys, status.ullAvailPageFile), status.ullTotalPhys, CgroupScan{}, memory_limit_bytes()); #elif defined(__linux__) return sample_host_memory(MemoryProbePaths{}, memory_limit_bytes()); +#elif defined(STRATA_GPU_METAL) + // macOS (the Metal engine): no /proc/meminfo. The engine's own reading (host_memory.h), the rule its streaming + // arena uses: free and purgeable pages and the clean file cache, not other programs' inactive anonymous memory; + // hw.memsize is the total. No cgroups. + double avail = 0, phys = 0; + if (!strata::metal_engine::host_memory_now(avail, phys) || phys <= 0) return {}; + const uint64_t total = (uint64_t) phys; + return combine(std::min((uint64_t) std::max(0.0, avail), total), total, CgroupScan{}, memory_limit_bytes()); #else return {}; #endif diff --git a/src/program/generate.cpp b/src/program/generate.cpp index 824bd19dbe6..4331542ec45 100644 --- a/src/program/generate.cpp +++ b/src/program/generate.cpp @@ -4671,7 +4671,14 @@ int main(int argc, char** argv) { if (pio.mode == strata::kernels::PleIo::Direct && !o.ple_io_explicit) { const strata::kernels::PleIoOptions direct_pio = pio; const char* auto_ram = std::getenv("STRATA_PLE_AUTO_RAM"); +#if defined(STRATA_GPU_METAL) + // (the Mac: the macOS RAM reading is there for the memory guard (parking, session files); this policy was + // measured on the PC, so on the Metal engine it stays off, as in a build that cannot read the RAM) + (void) auto_ram; + const bool allowed = false; +#else const bool allowed = strata::core::ple_auto_ram_allowed(auto_ram, strata::qwen_defaults_026()); +#endif const uint64_t need = ple_table.table_bytes(); const auto available = strata::core::available_host_bytes(); uint64_t experts_ram = 0; // every routed expert's blob, the arena the engine copies into RAM at start From 96c688923d4acb3a51695ddf21eb81197adeed36 Mon Sep 17 00:00:00 2001 From: oceanio Date: Sun, 18 Oct 2026 08:40:00 +0800 Subject: [PATCH 06/11] metal: parking and session files carry the decode window's own K/V and running state (#2502) On Metal a conversation snapshot read the session's QSA K/V, which the decode window does not use (and which the session no longer has), and the checkpoint check expected the CUDA running-state layout, so parking was refused ("invalid checkpoint running-state payload") and a session file could not be written. Under STRATA_GPU_METAL: - conversation_snapshot.cpp: a snapshot's K/V layer is the window's own, cells [0, upto) of the solo conversation's QSA layer or of the MTP draft layer, in the window's layout (format key 0x4D570001, page size 1), copied with memcpy (unified memory, the window idle). - conversation_state.cpp: checkpoint and session-read sizes come from the window's running state (window_checkpoint_sizes); the checkpoint check follows Window::checkpoint_save's layout. - window.cpp: SlotState::parts / layer_cells / cells_of (the state as byte ranges in a fixed order), Window::kv_layer and checkpoint_sizes over them; the draft layer's cache in the same per-cell layout. - generate.cpp: the session fingerprint's model input for the drafter is the MTP GGUF the window loads, and its backend "metal-window". A snapshot costs about 15.4 KB per token plus 118 MB per running state (the live one and each checkpoint kept). Co-Authored-By: Claude Opus 5.5 --- metal/src/engine/window.cpp | 101 +++++++++++++++++++++++++++-- src/core/conversation_snapshot.cpp | 39 +++++++++++ src/core/conversation_state.cpp | 59 +++++++++++++++++ src/program/generate.cpp | 14 +++- 4 files changed, 207 insertions(+), 6 deletions(-) diff --git a/metal/src/engine/window.cpp b/metal/src/engine/window.cpp index 7196b2190ce..05c8b11372d 100644 --- a/metal/src/engine/window.cpp +++ b/metal/src/engine/window.cpp @@ -123,7 +123,8 @@ const int32_t* g_mrope_dev = nullptr; // keys are sized from the context; alloc can run for more than one instance (batch slots, later). The MTP draft // layer's caches (mkq_ ...) are the window's, in the same per-cell layout. struct SlotState { - // ---- the parts' sizes in bytes, in one place (alloc, the checkpoints and reset use them) + // ---- the parts' sizes in bytes, in one place (alloc, parts, the checkpoints, the drafter's caches and the + // session files' layout key use them) static constexpr size_t kGdnState = 4 * 128 * 48 * 128; // per GDN layer: fp32 recurrence, 48 heads x 128 x 128 static constexpr size_t kConvHist = 4 * 10240 * 3; // per GDN layer, each half: fp32 conv history, 3 taps // the fp32 PLE history, 9 rows x 10240 @@ -164,6 +165,39 @@ struct SlotState { } } } + // The conversation's window state for its first `upto` tokens as byte ranges, fn(void* ptr, size_t bytes), in a + // fixed order (two states of one window walk the same sizes): the running state (per GDN layer the recurrence and + // the committed conv-history half `hp`, the PLE history, per QSA layer the indexer tail and dead rows), then per QSA + // layer the positional part of cells [0, upto) (the int8 K/V, their scales, the pooled keys of the blocks up to + // there). The cells past `upto` are rewritten before anything reads them. A host image (parking, session + // files) is the ranges one after the other. (The PLE token window, ple_prev, is the session's: the caller's.) + template void parts(int hp, int64_t upto, F&& fn) const { + for (size_t i = 0; i < gst.size(); ++i) { + fn(gst[i].ptr, kGdnState); + fn(ghist[hp][i].ptr, kConvHist); + } + fn(pleHist.ptr, kPleHist); + for (size_t j = 0; j < tail.size(); ++j) { + fn(tail[j].ptr, kTail); + fn(dead[j].ptr, kDead); + } + for (size_t j = 0; j < kq.size(); ++j) layer_cells(j, upto, fn); + } + // QSA layer j's part of cells [0, upto), in parts()' order: the int8 K, V, their fp16 scales (cell-major, so the + // first `upto` cells are the leading bytes), the pooled keys of the blocks up to there (with the moving spare row) + template void layer_cells(size_t j, int64_t upto, F&& fn) const { + cells_of(kq[j], vq[j], ksc[j], vsc[j], pooled[j], n_cells, upto, fn); + } + // the same walk over any QSA cache of this layout (the MTP draft layer's too) + template static void cells_of(const mtl::Buffer& k, const mtl::Buffer& v, const mtl::Buffer& ks, + const mtl::Buffer& vs, const mtl::Buffer& pool, int cells, int64_t upto, + F&& fn) { + fn(k.ptr, (size_t) upto * kKvCell); + fn(v.ptr, (size_t) upto * kKvCell); + fn(ks.ptr, (size_t) upto * kScaleCell); + fn(vs.ptr, (size_t) upto * kScaleCell); + fn(pool.ptr, (size_t) std::min(pooled_rows(cells), upto / kBlockCells + 1) * kPooledRow); + } }; class Window { @@ -183,6 +217,14 @@ class Window { bool checkpoint_restore(const strata::core::ConversationCheckpoint& c, std::string& err); void reset(); // a new sequence (session_zero): the running state zeroed void set_max_t(int t) { max_t_ = std::min(t, MAXT); } + // ---- conversation snapshots (parking, session files): the solo conversation's K/V, layer by + // layer, as the cells [0, upto) of its 5 parts (k, v, k_scale, v_scale, pooled; SlotState::layer_cells). A layer + // is a QSA state of the solo session (`index`, by its address: QSA layer j of slot_) or the MTP draft layer. + bool kv_layer(const strata::core::QsaState* st, bool index, int64_t upto, void* ptr[5], size_t bytes[5], + std::string& err) const; + // the running state's sizes in a checkpoint (checkpoint_save): gdn, ple, tails, dead + void checkpoint_sizes(size_t z[4]) const; + int64_t cells() const { return max_cells_; } // ---- the MTP draft layer (mtp.hpp's math on mac-kern's kernels; the MTP GGUF from tools/mtp_pack.py) bool mtp_load(const std::string& gguf, int max_t, std::string& err); bool mtp_prefill(const int32_t* next_tokens, int64_t n, int64_t cell0, std::string& err); @@ -1653,6 +1695,36 @@ bool Window::commit(int n_keep, std::string& err) { } +bool Window::kv_layer(const strata::core::QsaState* st, bool index, int64_t upto, void* ptr[5], size_t bytes[5], + std::string& err) const { + if (upto < 0 || upto > max_cells_) { err = "metal window: the snapshot's K/V runs past the context"; return false; } + int k = 0; + auto take = [&](void* p, size_t n) { ptr[k] = p; bytes[k] = n; ++k; }; + if (!index) { // the draft layer (its K/V is the window's: one drafter, the solo conversation's); none: empty + if (!mtp_on_ || mkq_.empty()) { + for (int i = 0; i < 5; ++i) { ptr[i] = nullptr; bytes[i] = 0; } + return true; + } + SlotState::cells_of(mkq_[0], mvq_[0], mksc_[0], mvsc_[0], mpooled_[0], (int) max_cells_, upto, take); + return true; + } + const strata::core::SessionState& ss = *ss_; + for (int64_t j = 0; ss.qsa_states != nullptr && j < ss.qsa_alloc && (size_t) j < slot_.kq.size(); ++j) + if (st == &ss.qsa_states[ss.qsa_ord0 + j]) { + slot_.layer_cells((size_t) j, upto, take); + return true; + } + err = "metal window: a snapshot's K/V layer that is not the solo conversation's"; + return false; +} + +void Window::checkpoint_sizes(size_t z[4]) const { + z[0] = slot_.gst.size() * (SlotState::kGdnState + SlotState::kConvHist); + z[1] = SlotState::kPleHist; + z[2] = slot_.tail.size() * SlotState::kTail; + z[3] = slot_.dead.size() * SlotState::kDead; +} + // ================================================================================================ the MTP draft layer bool Window::mtp_load(const std::string& gguf_path, int max_t, std::string& err) { (void) max_t; @@ -1708,15 +1780,17 @@ bool Window::mtp_load(const std::string& gguf_path, int max_t, std::string& err) } } auto A = [&](size_t bytes) { return cx_->alloc(bytes); }; - const int MAXC = (int) max_cells_, MB = MAXC / 4 + 2; + const int MAXC = (int) max_cells_; mR_ = A(4 * MAXT * HC * N); memb_ = A(4 * MAXT * N); men_ = A(4 * MAXT * N); me2_ = A(4 * MAXT * N); mhn_ = A(4 * MAXT * HC * N); mh2_ = A(4 * HC * MAXT * N); midx_ = A(4 * MAXT * 640); mbo_ = A(4 * MAXT * N); minj_[0] = A(4 * MAXT * HC); minj_[1] = A(4 * MAXT * HC); mqids_ = A(4 * MAXT * 2051); mlm_ = A(4 * (size_t) n_vocab_); auto act = [&](int n) { Act x; x.n = n; x.b = A(4 * ((size_t) 4 * MAXT * n + (size_t) MAXT * n / 16)); return x; }; XEn_ = act(N); XHn_ = act(N); - mkq_.push_back(A((size_t) MAXC * 2 * 256)); mvq_.push_back(A((size_t) MAXC * 2 * 256)); - mksc_.push_back(A(2 * (size_t) MAXC * 2 * 4)); mvsc_.push_back(A(2 * (size_t) MAXC * 2 * 4)); - mtail_.push_back(A(4 * 4 * 128)); mdead_.push_back(A(4 * 128)); mpooled_.push_back(A(4 * (size_t) MB * 128)); + // the draft layer's cache, in SlotState's per-cell layout (SlotState::cells_of walks it for the snapshots) + mkq_.push_back(A((size_t) MAXC * SlotState::kKvCell)); mvq_.push_back(A((size_t) MAXC * SlotState::kKvCell)); + mksc_.push_back(A((size_t) MAXC * SlotState::kScaleCell)); mvsc_.push_back(A((size_t) MAXC * SlotState::kScaleCell)); + mtail_.push_back(A(SlotState::kTail)); mdead_.push_back(A(SlotState::kDead)); + mpooled_.push_back(A((size_t) SlotState::pooled_rows(MAXC) * SlotState::kPooledRow)); { std::string vp; // the subset setup put beside the draft layer (--draft-vocab: /rt/draft_vocab.bin, as CUDA reads it), @@ -2680,6 +2754,22 @@ bool window_checkpoint_restore(const strata::core::ConversationCheckpoint& c, st if (g_the_window == nullptr || !g_the_window->ready()) { err = "metal window: not up"; return false; } return g_the_window->checkpoint_restore(c, err); } +// conversation snapshots (src/core/conversation_snapshot.cpp / conversation_state.cpp under STRATA_GPU_METAL). +// The layout key of a snapshot's K/V layer from this window ("MW", v1): per cell 2 KV heads x 256 int8 K and V, 2 x 4 +// fp16 scales each, no pages (page size 1, cells = tokens); one fp32 pooled key of 128 per 4 cells (+ the spare row). +// A K/V layer of another format (CUDA's, or a later window layout) is not this window's and is refused. +int window_kv_format() { return 0x4D570001; } +bool window_kv_layer(const strata::core::QsaState* st, bool index, int64_t upto, void* ptr[5], size_t bytes[5], + std::string& err) { + if (g_the_window == nullptr || !g_the_window->ready()) { err = "metal window: not up"; return false; } + return g_the_window->kv_layer(st, index, upto, ptr, bytes, err); +} +bool window_checkpoint_sizes(size_t z[4], int64_t& cells, std::string& err) { + if (g_the_window == nullptr || !g_the_window->ready()) { err = "metal window: not up"; return false; } + g_the_window->checkpoint_sizes(z); + cells = g_the_window->cells(); + return true; +} bool window_commit(void* w, int n_keep, std::string& err) { const double t0 = now_ms(); const bool ok = ((Window*) w)->commit(n_keep, err); @@ -2712,6 +2802,7 @@ bool mtp_load(const std::string& rt_dir, int max_t, std::string& err) { g_mtp_max_t = max_t; return true; } +const std::string& mtp_gguf_path() { return g_mtp_path; } bool mtp_load_pending(Window* w, std::string& err) { return g_mtp_path.empty() || w->mtp_load(g_mtp_path, g_mtp_max_t, err); } diff --git a/src/core/conversation_snapshot.cpp b/src/core/conversation_snapshot.cpp index 0e4c440b060..184551a58c1 100644 --- a/src/core/conversation_snapshot.cpp +++ b/src/core/conversation_snapshot.cpp @@ -14,6 +14,14 @@ #include #include +#if defined(STRATA_GPU_METAL) +namespace strata::metal_engine { // metal/src/engine/window.cpp: the decode window holds the conversation's K/V +int window_kv_format(); +bool window_kv_layer(const strata::core::QsaState* st, bool index, int64_t upto, void* ptr[5], size_t bytes[5], + std::string& err); +} // namespace strata::metal_engine +#endif + namespace strata::core { namespace { struct Layout { @@ -32,6 +40,18 @@ bool valid_extent(const QsaState& st, int64_t upto, std::string& error) { } bool layout(const QsaState& st, const ModelGeometry& g, int64_t upto, bool index, Layout& l, std::string& error) { +#if defined(STRATA_GPU_METAL) + // On Metal the session's QSA states hold no K/V (layer.cpp kv_plan): a snapshot's K/V layer is the decode window's, cells + // [0, upto) of the solo conversation's QSA layer (`index`) or of the MTP draft layer, in the window's own layout + // (its format key; no pages: cells = tokens). Any other state is refused. + void* ptr[5] = {}; + size_t n[5] = {}; + if (!strata::metal_engine::window_kv_layer(&st, index, upto, ptr, n, error)) return false; + if (g.idx_key_dim <= 0) { error = "conversation snapshot: invalid K/V geometry"; return false; } + l = {strata::metal_engine::window_kv_format(), upto, (int64_t) (n[4] / ((size_t) g.idx_key_dim * sizeof(float))), 1, + n[0], n[2], n[4], n[1], n[3]}; + return true; +#endif if (!valid_extent(st, upto, error)) return false; if (st.kv_hybrid && (st.kv_mode == 2 || st.kv_q4 || st.kv_int8)) { // a streamed one reads its host copy error = "conversation snapshot: hybrid K8V4 cannot be a ring and needs distinct format flags"; @@ -78,6 +98,17 @@ bool layout(const QsaState& st, const ModelGeometry& g, int64_t upto, bool index } std::array pools(const QsaState& st, bool resident = false) { +#if defined(STRATA_GPU_METAL) + // the window's buffers (layout() refused any state that is not the window's): a solo session QSA state, else the + // draft layer + (void) resident; + void* p[5] = {}; + size_t n[5] = {}; + std::string e; + if (!strata::metal_engine::window_kv_layer(&st, true, 0, p, n, e)) + strata::metal_engine::window_kv_layer(&st, false, 0, p, n, e); + return {p[0], p[1], p[2], p[3], p[4]}; +#endif const bool host = st.kv_mode != 0 && !resident; if (st.kv_hybrid) return {host ? st.host.k_q : st.k_q, host ? st.host.v_q4 : st.v_q4, host ? st.host.k_scale : st.k_scale, @@ -91,6 +122,10 @@ std::array pools(const QsaState& st, bool resident = false) { } bool valid(const QsaState& st, const Layout& l, int64_t upto, std::string& error) { +#if defined(STRATA_GPU_METAL) + (void) st; (void) l; (void) upto; (void) error; + return true; // the window's K/V: no pages, no streaming, no ring (layout() checked the extent) +#endif if (upto < 0 || upto > st.max_cells || st.n_pages < 0 || l.cells / l.page_size > st.n_pages || st.kv_mode < 0 || st.kv_mode > 2 || st.n_slots <= 0 || (st.kv_mode == 0 && st.n_slots < st.n_pages) || @@ -258,6 +293,10 @@ bool park_trace_on() { bool transfer(void* dst, const void* src, size_t n, std::string& error) { if (!n) return true; if (!src || !dst) { error = "conversation snapshot: missing state buffer"; return false; } +#if defined(STRATA_GPU_METAL) + std::memcpy(dst, src, n); // unified memory: the window's buffers are host-visible, and the window is idle here + return true; +#endif // Default handles both device allocations and device-mapped host pool aliases. const cudaError_t e = conversation_detail::copy_nonblocking(dst, src, n); if (e == cudaSuccess) return true; diff --git a/src/core/conversation_state.cpp b/src/core/conversation_state.cpp index b8ce970162d..6b5b3fec3af 100644 --- a/src/core/conversation_state.cpp +++ b/src/core/conversation_state.cpp @@ -12,6 +12,12 @@ #include #include +#if defined(STRATA_GPU_METAL) +namespace strata::metal_engine { // metal/src/engine/window.cpp: the running state's sizes in a checkpoint +bool window_checkpoint_sizes(size_t z[4], int64_t& cells, std::string& err); +} // namespace strata::metal_engine +#endif + namespace strata::core { namespace { using conversation_detail::add; @@ -59,8 +65,32 @@ bool image_keys(const std::vector& images, size_t tokens) return true; } +#if defined(STRATA_GPU_METAL) +// On Metal a checkpoint holds the decode window's running state (Window::checkpoint_save): the whole GDN state, the +// PLE history and per QSA layer the indexer tail and dead rows, no block positions. `z` gets those sizes (tail and +// dead per QSA layer, as the CUDA sizes are); `cells` is the window's. +bool metal_state_sizes(const SessionState& ss, ConversationStateSizes& z, int64_t& cells, std::string& error) { + size_t w[4] = {}; + if (!strata::metal_engine::window_checkpoint_sizes(w, cells, error)) return false; + const size_t layers = owned_qsa(ss); + if (!layers || w[2] % layers || w[3] % layers) return fail(error, "the decode window's QSA layers are not the session's"); + z = {}; + z.gdn = w[0]; z.ple = w[1]; z.tail = w[2] / layers; z.dead = w[3] / layers; z.block_pos = 0; + return true; +} +#endif + bool checkpoint_targets(const SessionState& ss, const ModelGeometry& g, size_t tokens, ConversationStateSizes& z, std::string& error) { +#if defined(STRATA_GPU_METAL) + // the running state and the K/V are the decode window's (the session's QSA states hold neither: layer.cpp kv_plan) + (void) g; + int64_t cells = 0; + if (!metal_state_sizes(ss, z, cells, error)) return false; + if (ss.max_cells < 0 || tokens > (uint64_t) ss.max_cells || tokens > (uint64_t) cells) + return fail(error, "invalid session running-state targets"); + return true; +#endif if (!conversation_session_sizes(g, ss, z, error)) return false; if (ss.max_cells < 0 || tokens > (uint64_t) ss.max_cells || (z.gdn && !ss.gdn_state) || (owned_qsa(ss) && !ss.qsa_states)) return fail(error, "invalid session running-state targets"); @@ -151,6 +181,12 @@ bool conversation_checkpoint_validate(const ConversationCheckpoint& c, const Ses ConversationStateSizes z; if (!checkpoint_targets(ss, g, c.ids.size(), z, error)) return false; const size_t layers = owned_qsa(ss); +#if defined(STRATA_GPU_METAL) + if (c.gdn.size() != z.gdn || c.ple.size() != z.ple || c.tails.size() != layers * z.tail || + c.dead.size() != layers * z.dead || !c.block_pos.empty() || !image_keys(c.imgs, c.ids.size())) + return fail(error, "invalid checkpoint running-state payload"); + return true; +#endif if (c.gdn.size() != z.gdn || c.ple.size() != (ss.ple_hist ? z.ple : 0) || c.tails.size() != layers * z.tail || c.dead.size() != layers * z.dead || c.block_pos.size() != layers * z.block_pos || !image_keys(c.imgs, c.ids.size())) @@ -646,6 +682,29 @@ bool conversation_session_read_limits(SessionReadLimits& limits, const SessionSt const QsaState& draft, uint64_t max_tokens, uint64_t max_checkpoints, std::string& error) { ConversationStateSizes z; +#if defined(STRATA_GPU_METAL) + { // the decode window's running state and K/V: every K/V layer, the draft's too, up to the session's cells + int64_t window_cells = 0; + if (!metal_state_sizes(ss, z, window_cells, error)) return false; + if (ss.max_cells < 0) return fail(error, "invalid session cells"); + const uint64_t tokens = std::min({max_tokens, (uint64_t) ss.max_cells, (uint64_t) window_cells}); + const size_t layers = owned_qsa(ss); + if (!layers || !ss.qsa_states) return fail(error, "invalid session running-state targets"); + SessionReadLimits l = limits; // keeps the caller's admit / progress / max_file_bytes + l.max_tokens = tokens; + l.max_checkpoints = max_checkpoints; + l.max_kv_layers = layers + 1; + l.geometry = geometry_key(g); + l.layer_range = std::make_pair(ss.layer_lo, ss.layer_hi); + l.max_state_bytes = {z.gdn, z.ple, layers * z.tail, layers * z.dead, 0}; + l.max_kv_bytes.assign(layers + 1, {}); + for (size_t j = 0; j < layers; ++j) + if (!conversation_kv_part_sizes(owned(ss, j), g, (int64_t) tokens, true, l.max_kv_bytes[j], error)) return false; + if (!conversation_kv_part_sizes(draft, g, (int64_t) tokens, false, l.max_kv_bytes.back(), error)) return false; + limits = std::move(l); + return true; + } +#endif if (!conversation_session_sizes(g, ss, z, error)) return false; if (ss.max_cells < 0) return fail(error, "invalid session cells"); const uint64_t tokens = std::min(max_tokens, (uint64_t) ss.max_cells); diff --git a/src/program/generate.cpp b/src/program/generate.cpp index 4331542ec45..d839f1c0137 100644 --- a/src/program/generate.cpp +++ b/src/program/generate.cpp @@ -127,7 +127,11 @@ #include #if defined(STRATA_GPU_METAL) -namespace strata::metal_engine { void set_model_shard(const char* shard1); } // metal/src/engine/window.cpp +namespace strata::metal_engine { // metal/src/engine/window.cpp +void set_model_shard(const char* shard1); +// session files: the MTP GGUF the window loads (a model input of the fingerprint), empty without --mtp +const std::string& mtp_gguf_path(); +} #endif #include @@ -10350,7 +10354,11 @@ int main(int argc, char** argv) { in.embedding = o.embd_gguf; in.ple = o.ple_gguf; in.pack = o.pack; +#if defined(STRATA_GPU_METAL) + in.mtp = o.mtp.empty() ? std::string() : strata::metal_engine::mtp_gguf_path(); // the window's MTP GGUF +#else in.mtp = o.mtp.empty() ? std::string() : o.mtp + "/dense.bin"; +#endif if (srcp != nullptr) in.experts = srcp->model_inputs(); std::vector files = strata::core::session_model_inputs(in); // the pack's other files the loader reads (weights.cpp, expert_layout.cpp) @@ -10361,7 +10369,9 @@ int main(int argc, char** argv) { if (srcp == nullptr) files.push_back({"pack experts.bin", o.pack + "/experts.bin", true}); } if (!o.mtp.empty()) { +#if !defined(STRATA_GPU_METAL) // (Metal: the drafter is the one GGUF above) for (const char* n : {"dense.txt", "experts.bin"}) files.push_back({std::string("mtp ") + n, o.mtp + "/" + n}); +#endif files.push_back({"mtp draft_vocab.bin", o.mtp + "/draft_vocab.bin", true}); } uint64_t fp = 0; @@ -10376,6 +10386,8 @@ int main(int argc, char** argv) { c.engine_version = STRATA_VERSION; #if defined(STRATA_USE_HIP) c.backend = "hip"; +#elif defined(STRATA_GPU_METAL) + c.backend = "metal-window"; // the decode window's running state and K/V (window.cpp) #else c.backend = "cuda"; #endif From 0be11be951413914ba31ecb5ecb6693717ea1e44 Mon Sep 17 00:00:00 2001 From: oceanio Date: Sun, 18 Oct 2026 08:40:01 +0800 Subject: [PATCH 07/11] metal: a session file's fingerprint carries the decode window's K/V and running-state layout; a file from another layout is refused The config fingerprint's backend is "metal-window" plus window_layout_key(): the K/V format key, the bytes per cell of each K/V part, the pooled key per block and the running state's part sizes. Two builds of one version with different window layouts get different fingerprints, so a file from the other one is refused whole (SERR, the engine keeps serving) instead of read into the wrong places. Co-Authored-By: Claude Opus 5.5 --- metal/src/engine/window.cpp | 13 +++++++++++++ src/program/generate.cpp | 6 +++++- 2 files changed, 18 insertions(+), 1 deletion(-) diff --git a/metal/src/engine/window.cpp b/metal/src/engine/window.cpp index 05c8b11372d..6e8e8ef9229 100644 --- a/metal/src/engine/window.cpp +++ b/metal/src/engine/window.cpp @@ -2759,6 +2759,19 @@ bool window_checkpoint_restore(const strata::core::ConversationCheckpoint& c, st // fp16 scales each, no pages (page size 1, cells = tokens); one fp32 pooled key of 128 per 4 cells (+ the spare row). // A K/V layer of another format (CUDA's, or a later window layout) is not this window's and is refused. int window_kv_format() { return 0x4D570001; } +// Session files: the window's layout as text, folded into the file's config fingerprint (generate.cpp): the K/V +// format key, the bytes per cell of each K/V part (SlotState::cells_of; the draft layer's are the same), the pooled +// key per block and the running state's parts (SlotState::parts, Window::checkpoint_save). A build whose layout +// differs gets another fingerprint, so its files are refused whole instead of read into the wrong places. +std::string window_layout_key() { + char b[256]; + std::snprintf(b, sizeof b, "kv %08x cell k%zu v%zu ks%zu vs%zu pooled %zu/%d; gdn %zu x (%zu + %zu), ple %zu, qsa %zu x " + "(%zu + %zu)", (unsigned) window_kv_format(), SlotState::kKvCell, SlotState::kKvCell, + SlotState::kScaleCell, SlotState::kScaleCell, SlotState::kPooledRow, SlotState::kBlockCells, + SlotState::kGdnLayers, SlotState::kGdnState, SlotState::kConvHist, SlotState::kPleHist, + SlotState::kQsaLayers, SlotState::kTail, SlotState::kDead); + return b; +} bool window_kv_layer(const strata::core::QsaState* st, bool index, int64_t upto, void* ptr[5], size_t bytes[5], std::string& err) { if (g_the_window == nullptr || !g_the_window->ready()) { err = "metal window: not up"; return false; } diff --git a/src/program/generate.cpp b/src/program/generate.cpp index d839f1c0137..c47920809dd 100644 --- a/src/program/generate.cpp +++ b/src/program/generate.cpp @@ -131,6 +131,8 @@ namespace strata::metal_engine { // metal/src/engine/window.cpp void set_model_shard(const char* shard1); // session files: the MTP GGUF the window loads (a model input of the fingerprint), empty without --mtp const std::string& mtp_gguf_path(); +// session files: the window's K/V and running-state layout (part of the config fingerprint) +std::string window_layout_key(); } #endif @@ -10387,7 +10389,9 @@ int main(int argc, char** argv) { #if defined(STRATA_USE_HIP) c.backend = "hip"; #elif defined(STRATA_GPU_METAL) - c.backend = "metal-window"; // the decode window's running state and K/V (window.cpp) + // the decode window's running state and K/V, and their layout (sizes per cell and per part): a build + // with another layout writes files this one refuses (the fingerprint differs) + c.backend = "metal-window " + strata::metal_engine::window_layout_key(); #else c.backend = "cuda"; #endif From 725200d55f62ce59c79e6b0346ff0fa9e1077c10 Mon Sep 17 00:00:00 2001 From: oceanio Date: Sun, 18 Oct 2026 08:40:01 +0800 Subject: [PATCH 08/11] tools: session_restart_test.py and metal_snapshot_test.py; parking_test.py reads the engine's own log - session_restart_test.py: a session file saved by one engine and restored in a new one decodes the same follow-up as the conversation that never left (--restore-exe / --expect-refused for a file from another build). - metal_snapshot_test.py: on Metal the conversation cache starts and parks, SAVE writes a file, RESTORE of a non-session file answers SERR and keeps the engine in step, RESTORE of the saved file restores it. - parking_test.py: the engine log under tempfile.gettempdir() (on macOS /tmp is not the temp folder, and it read a stale file there). Co-Authored-By: Claude Opus 5.5 --- tools/metal_snapshot_test.py | 110 ++++++++++++++++++++++++++++++++++ tools/parking_test.py | 2 +- tools/session_restart_test.py | 100 +++++++++++++++++++++++++++++++ 3 files changed, 211 insertions(+), 1 deletion(-) create mode 100644 tools/metal_snapshot_test.py create mode 100644 tools/session_restart_test.py diff --git a/tools/metal_snapshot_test.py b/tools/metal_snapshot_test.py new file mode 100644 index 00000000000..e267e77cecd --- /dev/null +++ b/tools/metal_snapshot_test.py @@ -0,0 +1,110 @@ +#!/usr/bin/env python3 +"""Metal snapshots test: on Metal the conversation snapshots (parking, session files) carry the +decode window's own K/V, so the conversation cache and SAVE / RESTORE are on, as on CUDA. Through the --serve +protocol, one engine start (before, the Metal engine never parked and SAVE failed): + + 1. started with --conversation-cache-mib: no start-up line turns the cache off on Metal, and INFO reports + conversation_cache_mib=1024; + 2. a second conversation parks the first (`conversation cache: parked` in the log, no `skip parking`); + 3. SAVE answers `SAVED ` and writes the file; + 4. RESTORE of a file that is not a session file answers `SERR invalid 0 ...` and leaves the engine in step; + 5. RESTORE of the saved file answers `RESTORED ...` with the saved token count; + 6. the engine is still in step: the next GEN decodes. + + python3 tools/metal_snapshot_test.py --exe build-metal/strata --config strata-.json [--extra "..."] + +Exit 0 when every check passes, 1 otherwise. Exactness (the restored follow-up decodes the live tokens) is +tools/parking_test.py's and tools/session_restart_test.py's. +""" +import argparse, json, os, subprocess, sys, tempfile +from pathlib import Path + +sys.path.insert(0, str(Path(__file__).resolve().parent)) +from batch_test import tokenizer # noqa: E402 + + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("--exe", required=True) + ap.add_argument("--config", required=True) + ap.add_argument("--extra", default="--max-context 8192", help="more engine arguments in one string") + a = ap.parse_args() + cfg = json.loads(Path(a.config).read_text()) + tok = tokenizer(cfg["tokenizer"]) + args = [a.exe, "--serve", *cfg["args"], *a.extra.split(), "--conversation-cache-mib", "1024"] + log_path = os.path.join(tempfile.gettempdir(), "metal_snapshot_test_engine.log") + with open(log_path, "w") as log: + p = subprocess.Popen(args, cwd=cfg.get("cwd"), stdin=subprocess.PIPE, stdout=subprocess.PIPE, stderr=log, + text=True, bufsize=1, env=dict(os.environ, **(cfg.get("env") or {}))) + head = [] + for line in p.stdout: + head.append(line.rstrip("\n")) + if line.startswith("READY"): + break + else: + raise SystemExit(f"the engine ended before READY - see {log_path}") + + def send(cmd): + p.stdin.write(cmd + "\n") + p.stdin.flush() + + def answer(prefixes): + for line in p.stdout: + line = line.rstrip("\n") + if line.startswith(prefixes): + return line + return "(the engine ended)" + + def gen(text, n=16): + ids = tok.encode(f"<|im_start|>user\n{text}<|im_end|>\n<|im_start|>assistant\n\n\n\n\n", + parse_special=True) + send(f"GEN {n} " + ",".join(map(str, ids))) + return answer(("DONE", "ERR")) + + fails = [] + + def check(name, ok, seen): + print(f"{'PASS' if ok else 'FAIL'} {name}: {seen}", flush=True) + if not ok: + fails.append(name) + + info = next((l for l in head if l.startswith("INFO")), "") + mib = next((f.split("=", 1)[1] for f in info.split() if f.startswith("conversation_cache_mib=")), "?") + log_text = Path(log_path).read_text() + off = [l for l in log_text.splitlines() if "conversation cache" in l and "Metal" in l] + check("no start-up line turning the conversation cache off on Metal", not off, off[:1] or "none") + check("INFO conversation_cache_mib=1024", mib == "1024", f"conversation_cache_mib={mib}") + + first = gen("Name three primary colors.") + gen("Write one sentence about the sea.") # a second conversation parks the first + log_text = Path(log_path).read_text() + parked = [l for l in log_text.splitlines() if "conversation cache: parked" in l] + skipped = [l for l in log_text.splitlines() if "skip parking" in l] + check("the second conversation parked the first", bool(parked) and not skipped, (skipped or parked)[:1] or "none") + with tempfile.TemporaryDirectory() as d: + path = os.path.join(d, "s.strata-session") + send(f"SAVE {path}") + saved = answer(("SAVED", "SERR", "ERR", "FATAL")) + check("SAVE answers SAVED", saved.startswith("SAVED "), saved) + check("SAVE wrote the file", os.path.exists(path), "exists" if os.path.exists(path) else "no file") + bad = os.path.join(d, "bad.strata-session") + Path(bad).write_bytes(b"not a session") + send(f"RESTORE {bad}") + refused = answer(("RESTORED", "SERR", "ERR", "FATAL")) + check("RESTORE of a non-session file answers SERR invalid 0", refused.startswith("SERR invalid 0 "), refused) + send(f"RESTORE {path}") + restored = answer(("RESTORED", "SERR", "ERR", "FATAL")) + n_saved = saved.split()[1] if saved.startswith("SAVED ") else "?" + check("RESTORE of the saved file answers RESTORED with its tokens", + restored.startswith("RESTORED ") and restored.split()[1] == n_saved, restored) + after = gen("Name three primary colors.") + check("the engine is still in step (GEN after the restores)", after.startswith("DONE"), after) + send("QUIT") + p.wait(timeout=180) + print(f"first request: {first}\nengine log: {log_path}") + print("ALL PASS" if not fails else f"{len(fails)} FAILED") + return 1 if fails else 0 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/tools/parking_test.py b/tools/parking_test.py index 29995825340..a5ba53b3658 100644 --- a/tools/parking_test.py +++ b/tools/parking_test.py @@ -52,7 +52,7 @@ def run(a, cfg, tok, with_b): gen(eng, out, pb, 8) eng.send("QUIT") eng.p.wait(timeout=180) - log = open("/tmp/batch_test_engine.log").read() + log = open(eng.log_path).read() # the engine's log (tempfile.gettempdir(): not /tmp on macOS) restored = re.findall(r"restored \d+ tokens[^\n]*", log) if with_b and "SNAPSHOT_VERIFY draft=" not in log: print("no SNAPSHOT_VERIFY line: the draft ring was not read back", flush=True) diff --git a/tools/session_restart_test.py b/tools/session_restart_test.py new file mode 100644 index 00000000000..500cef9dd86 --- /dev/null +++ b/tools/session_restart_test.py @@ -0,0 +1,100 @@ +#!/usr/bin/env python3 +"""Session file across an engine restart: a follow-up to conversation A must decode the same tokens whether A's state +was still live (first engine: A, SAVE, the follow-up) or came back from the file in a new engine process (second +engine: RESTORE, the follow-up), and the restored engine must resume from the file (reused > 0), not read A again. +Greedy, through the --serve protocol (SAVE / RESTORE ). + + python3 tools/session_restart_test.py --exe "$PWD/build-metal/strata" --config strata-.json \ + [--repeat 60] [--extra "--max-context 65536"] + +Prints the file size, the save and restore times and the follow-up's prompt reuse. Exit 0 when the follow-ups are +identical and the second engine reused the restored state, 1 otherwise. + +--restore-exe ENGINE --expect-refused: the file is written by --exe and restored by another build whose session +layout differs (on a Mac: the decode window's K/V and running-state layout, part of the file's config fingerprint). +Then the restore must be refused cleanly (SERR, the engine keeps serving) and the follow-up is read from its prompt +(reused 0; its tokens may differ from the live follow-up's, as any re-read's can on a Mac). +""" +import argparse, json, os, sys, tempfile, time +from pathlib import Path + +sys.path.insert(0, str(Path(__file__).resolve().parent)) +from batch_test import Engine, tokenizer # noqa: E402 +from parking_test import DOC, chat, gen # noqa: E402 + + +def answer(out, prefixes): + for line in out: + if line.startswith(prefixes): + return line + return "(the engine ended)" + + +def follow_up(tok, pa, ans): + end = tok.encode("<|im_end|>\n", parse_special=True) + return pa + ans + ([] if ans and ans[-1] in end else end) + \ + tok.encode(chat("Now give three keywords for that text, with one sentence each."), parse_special=True) + + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("--exe", required=True) + ap.add_argument("--config", required=True) + ap.add_argument("--repeat", type=int, default=60, help="copies of the paragraph in conversation A") + ap.add_argument("--max-new", type=int, default=120) + ap.add_argument("--extra", default="") + ap.add_argument("--keep", default="", help="save the session file here (kept) instead of a temporary folder") + ap.add_argument("--restore-exe", default="", help="the engine that restores the file (default: --exe)") + ap.add_argument("--expect-refused", action="store_true", + help="the restoring engine must refuse the file (SERR) and still answer the follow-up") + a = ap.parse_args() + cfg = json.loads(Path(a.config).read_text()) + tok = tokenizer(cfg["tokenizer"]) + env = {"STRATA_IQ_MT_MIN": "1"} + pa = tok.encode(chat(DOC * a.repeat + "\nSummarize this text in five sentences."), parse_special=True) + tmp = tempfile.TemporaryDirectory() + path = a.keep or os.path.join(tmp.name, "a.strata-session") + + eng = Engine(a.exe, cfg, 0, env, a.extra.split()) # 1: A, SAVE, the live follow-up + out = eng.lines() + ans, done_a = gen(eng, out, pa, a.max_new) + eng.send(f"SAVE {path}") + saved = answer(out, ("SAVED", "SERR", "ERR", "FATAL")) + follow = follow_up(tok, pa, ans) + live, done_live = gen(eng, out, follow, a.max_new) + eng.send("QUIT") + eng.p.wait(timeout=180) + + restored, got, done_got = "(not run)", [], "" + if saved.startswith("SAVED"): + eng = Engine(os.path.abspath(a.restore_exe or a.exe), cfg, 0, env, a.extra.split()) # 2: a new process, RESTORE + out = eng.lines() + t0 = time.time() + eng.send(f"RESTORE {path}") + restored = answer(out, ("RESTORED", "SERR", "ERR", "FATAL")) + print(f"restore round trip {time.time() - t0:.2f} s", flush=True) + got, done_got = gen(eng, out, follow, a.max_new) + eng.send("QUIT") + eng.p.wait(timeout=180) + + size = os.path.getsize(path) if os.path.exists(path) else 0 + f = done_got.split() + reused = int(f[8]) if len(f) > 8 and f[8].isdigit() else 0 + fl = done_live.split() + print(f"conversation A: {len(pa)} prompt + {len(ans)} generated tokens; follow-up {len(follow)} tokens") + print(f"save: {saved}\nrestore: {restored}\nfile: {size} bytes ({size / max(1, len(pa) + len(ans)):.0f} B/token)") + print(f"live follow-up reused {fl[8] if len(fl) > 8 else '?'}; restored follow-up reused {reused}") + first = next((k for k in range(min(len(live), len(got))) if live[k] != got[k]), None) + same = bool(live) and live == got + print("follow-up tokens:", "IDENTICAL" if same else f"DIFFER at {first}", f"({len(live)} / {len(got)})") + print(" ", repr(tok.decode(got)[:200])) + if a.expect_refused: # refused cleanly: an SERR line, the engine still serves the follow-up (from its prompt) + ok = saved.startswith("SAVED") and restored.startswith("SERR") and len(got) > 0 and reused == 0 + else: + ok = saved.startswith("SAVED") and restored.startswith("RESTORED") and same and reused > 0 + print("PASS" if ok else "FAIL") + return 0 if ok else 1 + + +if __name__ == "__main__": + sys.exit(main()) From 18a20cd353d5d6c87a8234301d7b3210c121e99b Mon Sep 17 00:00:00 2001 From: oceanio Date: Sun, 18 Oct 2026 08:40:02 +0800 Subject: [PATCH 09/11] docs: the conversation cache and session files on the Mac (MAC.md, DETAILS.md), measured at 29K tokens MAC.md: parking and session files work on the Mac and carry the decode window's own K/V; sizes and times for a 29,404-token conversation on the M5 Max with Qwen3.8-Flash-Next IQ3_S. DETAILS.md: the memory guard's macOS reading, the Metal backend and window layout in the session fingerprint, and that session files have run on macOS/Metal. Co-Authored-By: Claude Opus 5.5 --- docs/DETAILS.md | 7 ++++--- docs/MAC.md | 1 + 2 files changed, 5 insertions(+), 3 deletions(-) diff --git a/docs/DETAILS.md b/docs/DETAILS.md index e2565af6f98..47adfc2bfcb 100644 --- a/docs/DETAILS.md +++ b/docs/DETAILS.md @@ -693,7 +693,7 @@ system prompt with a billing header that changes on every request; the server pi | `--conversation-cache-min-free-mib N` | RAM floor, default 2560. Parking and session files are skipped when less than the snapshot plus the floor is available. | | `--memory-limit-mib N` (or `STRATA_MEMORY_LIMIT_MIB`) | Total RAM this engine's container may use, for when a container cannot see its own limit. Config `args`: `"--memory-limit-mib", "102400"`. `0` = none. | -The RAM figure the floor is checked against is the smallest of `MemAvailable`, the room under each cgroup limit visible to +The RAM figure the floor is checked against is the smallest of `MemAvailable` (macOS: the Metal engine's one reading of free memory, also used for its streaming arena - free and purgeable pages and the clean file cache from `host_statistics64`, not other programs' inactive anonymous memory), the room under each cgroup limit visible to the engine (v2 `memory.max` and `memory.high`, v1 `memory.limit_in_bytes`, less clean inactive file cache) and the `--memory-limit-mib` cap. The startup line `strata generate: memory guard: limit X GiB (source: flag|cgroup|meminfo) ...` names the source. It covers parking and session files only, not the startup sizing of the arena, the resident budget or the Windows commit check. @@ -729,7 +729,8 @@ curl -X POST "http://127.0.0.1:8080/slots/0?action=restore" -H "Content-Type: ap - Clients still send their messages and images; the file holds engine state and token identity, not a chat export. A follow-up that continues the restored conversation reuses its state; a short tail may be read again. Only the deepest checkpoint is saved. - A file is bound to the model files (sampled by size and first and last MiB, not by path) and to the settings that change what - the cells mean (engine version, CUDA or HIP, `--kv`, `STRATA_KV_ROT`, `--kv-resident`, `--max-context`, `--mtp-window`, the + the cells mean (engine version, CUDA, HIP or Metal - on Metal also the decode window's layout, so two builds of one version with + different layouts are told apart -, `--kv`, `STRATA_KV_ROT`, `--kv-resident`, `--max-context`, `--mtp-window`, the rope configuration, a loaded control vector and the arithmetic switches). Sampling, seeds, draft tuning and the expert tier are not in it. Do not change model files while their sessions are kept. The hash detects accidental corruption and does not authenticate a file: restore only files this engine wrote. @@ -739,7 +740,7 @@ curl -X POST "http://127.0.0.1:8080/slots/0?action=restore" -H "Content-Type: ap file had replaced the old one. Requests must be `application/json` (else `415`) and come from no browser page, Strata's own page or a trusted origin. - Not supported with `--layer-split`, `--peer-device`, `--batch` or `--prompt-cache 0`. The engine has been run on Linux and CUDA - only for this feature. Exact token replay across restarts is not guaranteed: expert residency and CPU/GPU rounding can + and on macOS/Metal ([MAC.md](MAC.md)) for this feature. Exact token replay across restarts is not guaranteed: expert residency and CPU/GPU rounding can change later output. ### The web page (Chat, Images, Monitor, About) diff --git a/docs/MAC.md b/docs/MAC.md index cf9ca2d3961..05e638853d4 100644 --- a/docs/MAC.md +++ b/docs/MAC.md @@ -224,6 +224,7 @@ from each other: top-1 agreement 95.4% against 95.8%, mean KL 0.0207 against 0.0 - **Context:** measured up to 32K (needles, KL). Setup recommends 32K below 96 GB and 64K from 96 GB. Up to the model's 262K can be set, but it is not measured. - **One request at a time.** `--parallel` (batch slots) is not ported. Several clients queue. +- **Conversation cache and session files** (`--conversation-cache-mib`, `/slots/0?action=save|restore`) work as on the PC; their snapshots carry the Mac decode window's own K/V. Measured on the M5 Max (128 GB) with Qwen3.8-Flash-Next IQ3_S at 262K context, a 28,785-token conversation: parked as 798 MB (with its checkpoints) in 49 ms and back in 41 ms; as a session file 679 MB, saved in 0.16 s and restored in a new engine in 0.10 s. Both times the follow-up decoded the same 95 tokens as the conversation that never left. Size: about 15.4 KB per token plus 118 MB per running state (the live one and each checkpoint kept with it). A session file's fingerprint names the Mac engine and its decode window's layout (the bytes per cell of each K/V part and the sizes of the running state), so one from the PC engine does not restore on the Mac (not tried), and a Mac build whose window layout differs refuses the file instead of misreading it (checked with a file from an earlier build of this change: refused, the engine kept serving). - **Not ported:** the control vector (`--control-vector-scaled`) and the PC tuning (`--calibrate`). - **Streaming** is measured at 32 GB only. - **Families:** the original Qwen3.8-Flash-Next, Unsloth's UD-Q4/Q5/Q6_K_XL and agentionai's Gyro-S are measured. Setup repacks From cdf1da9805ba31e4962e5e87a8644b51c3433104 Mon Sep 17 00:00:00 2001 From: oceanio Date: Tue, 27 Oct 2026 08:39:54 +0800 Subject: [PATCH 10/11] tools: metal_snapshot_test.py passes its own --conversation-cache-min-tokens after the config's arguments The engine takes a repeated flag's last value, so a config with a higher admission floor (for example --conversation-cache-min-tokens 2048) no longer keeps the test's 32-token conversation from parking and failing check 2. Co-Authored-By: Claude Opus 5.5 --- tools/metal_snapshot_test.py | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/tools/metal_snapshot_test.py b/tools/metal_snapshot_test.py index e267e77cecd..b8810b0072a 100644 --- a/tools/metal_snapshot_test.py +++ b/tools/metal_snapshot_test.py @@ -4,7 +4,8 @@ protocol, one engine start (before, the Metal engine never parked and SAVE failed): 1. started with --conversation-cache-mib: no start-up line turns the cache off on Metal, and INFO reports - conversation_cache_mib=1024; + conversation_cache_mib=1024 (the test's own --conversation-cache-min-tokens comes after the config's, so a + config's admission floor does not keep its short conversations from parking); 2. a second conversation parks the first (`conversation cache: parked` in the log, no `skip parking`); 3. SAVE answers `SAVED ` and writes the file; 4. RESTORE of a file that is not a session file answers `SERR invalid 0 ...` and leaves the engine in step; @@ -31,7 +32,8 @@ def main(): a = ap.parse_args() cfg = json.loads(Path(a.config).read_text()) tok = tokenizer(cfg["tokenizer"]) - args = [a.exe, "--serve", *cfg["args"], *a.extra.split(), "--conversation-cache-mib", "1024"] + args = [a.exe, "--serve", *cfg["args"], *a.extra.split(), "--conversation-cache-mib", "1024", + "--conversation-cache-min-tokens", "16"] # the last value wins: the config's floor may be above 32 tokens log_path = os.path.join(tempfile.gettempdir(), "metal_snapshot_test_engine.log") with open(log_path, "w") as log: p = subprocess.Popen(args, cwd=cfg.get("cwd"), stdin=subprocess.PIPE, stdout=subprocess.PIPE, stderr=log, From 450e96239fd5cb316faeaa3f97bab0faca59be6b Mon Sep 17 00:00:00 2001 From: oceanio Date: Wed, 28 Oct 2026 12:16:59 +0800 Subject: [PATCH 11/11] tools: batch_test.py's Engine starts the engine with the config's env, as serve does parking_test.py and session_restart_test.py start the engine through this class, which passed only the shell's environment and the test's own keys. A config's engine settings (the STRATA_METAL_* tuning, for one) were left out unless they were exported first, so the harness tested an untuned engine. The config's env now comes over the shell's, as serve/server.py applies it, and the test's own keys still win. Reported on #3541. Co-Authored-By: Claude Opus 5.5 --- tools/batch_test.py | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/tools/batch_test.py b/tools/batch_test.py index 78060b3dc3f..a87b5e05d08 100644 --- a/tools/batch_test.py +++ b/tools/batch_test.py @@ -52,7 +52,9 @@ def __init__(self, exe, cfg, batch, extra_env, extra_args): if batch: args += ["--batch", str(batch)] args += extra_args - env = dict(os.environ, **extra_env) + env = dict(os.environ) + env.update({str(k): str(v) for k, v in (cfg.get("env") or {}).items()}) # the config's engine settings, as serve passes them + env.update(extra_env) # the test's own keys win env["LD_LIBRARY_PATH"] = ":".join(cfg.get("lib_dirs", []) + [env.get("LD_LIBRARY_PATH", "")]) if os.name == "nt": env["PATH"] = os.pathsep.join(cfg.get("lib_dirs", []) + [env.get("PATH", "")])