Merge f-029-par-harness

include/strata/prefill/prefill.hpp

@@ -10,6 +10,8 @@
 // Requires the native weights (`--native`): every quantized projection must carry its GGUF blocks.
 #pragma once
 
+#include <array>
+
 #include "strata/core/expert_cache.hpp"
 #include "strata/core/expert_source.hpp"
 #include "strata/core/layer.hpp"
@@ -265,6 +267,19 @@ public:
     /// Layer-major prefill (STRATA_PREFILL_LAYER_MAJOR=1): true while `on_chunk` reports a sub-chunk before the end of
     /// its super-chunk.  The session state is already at the super-chunk's end then, so no checkpoint may be taken.
     /// On a layer split it is the last stage's (the one that reports the rows), asked through the first.
+    /// One prompt of a multi-prompt read (run_multi): its tokens, the position it starts at and, per stage of a layer split
+    /// (index 0 = the first stage), the session it reads into (nullptr / missing: that stage's own main session).
+    struct Seg {
+        const int64_t* tokens = nullptr;
+        int64_t n = 0;
+        int64_t pos0 = 0;
+        std::vector<core::SessionState*> stage_ss;
+    };
+    /// Several prompts in ONE layer-major pass (the streamed experts of a layer are read once for all of them).  Each
+    /// prompt fills its own sessions; `on_chunk` / `on_stage_chunk` report the FIRST prompt only.  Returns false with `err`
+    /// when the layer-major walk does not apply (then nothing was read; the caller reads the prompts one by one).
+    bool run_multi(const std::vector<Seg>& segs, std::string& err);
+
     bool layer_major_mid() const {
         const Prefill* p = this;
         while (p->next_ != nullptr) p = p->next_;
@@ -282,8 +297,11 @@ private:
     int64_t lm_prepare_chain(int64_t n, std::string& why);
     void lm_release_chain();
     bool run_layer_major(const int64_t* tokens, int64_t n, int64_t pos0, std::string& err);
+    struct LmSub { int seg = 0; int64_t c0 = 0, T = 0, t0 = 0; };   ///< a sub-chunk: prompt, first token, tokens, row in the super-chunk
+    bool lm_run(const std::vector<Seg>& segs, int depth, const std::vector<LmSub>* given_super, std::string& err);
     void lm_release();
     bool lm_mid_ = false;
+    bool lm_force_ = false;        ///< run_multi: the walk even for a short total
     bool lm_by_default_ = false;   ///< the walk is on because nothing was set (a wider VRAM margin, the measured super-chunk size)
     // Stage-1 pipeline: intermediate stages return after handing their chunk to
     // the direct successor. The public run() drains the chain once at prompt end.

include/strata/program/par_counters.hpp

@@ -0,0 +1,93 @@
+// include/strata/program/par_counters.hpp - STRATA_PAR_COUNTERS=1: counters-only parallelism accounting (research, opt-in).
+//
+// Off (the default), every hook is one branch on a const bool and reads no clock.  On, the batch loop (one host thread)
+// adds plain integers and, at the places that already wait (a drain, an admission), two clock reads.  Nothing here
+// synchronizes a device.  Printed to stderr as "strata par: ..." with the batch summary line (when the loop goes idle):
+//   windows      group windows launched, rows carried, rows used by live slots, padded rows, histogram of live slots per window
+//   drains       pipeline drains before a control line, by line kind (BGEN = an admission), groups in flight, ms waited
+//   admissions   prompts admitted into a slot: prompt tokens, ms in the prompt phase, ms of the drain before it
+//   idle         ticks with stage 0 free and no launchable group (no live slot / every group in flight)
+// Stage busy/idle with reason stays in STRATA_P4_TIMING (pipe_probe.hpp); this adds what it lacked.
+#pragma once
+
+#include <chrono>
+#include <cstdint>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <string>
+
+namespace strata::program {
+
+/// Expert-dispatch totals (cumulative in the engine); a span's delta is printed at flush.
+struct ParEx { long long misses = 0, entries = 0, hits = 0, offload = 0, pcie = 0; double run_ms = 0, jobs_ms = 0, vwait_ms = 0, vpool_ms = 0; };
+
+class ParCounters {
+public:
+    using Ex = ParEx;
+    using Clock = std::chrono::steady_clock;
+    ParCounters() { const char* e = std::getenv("STRATA_PAR_COUNTERS"); on_ = e != nullptr && e[0] == '1' && e[1] == '\0'; }
+    bool on() const { return on_; }
+
+    void window(int rows, int live) {
+        if (!on_) return;
+        ++windows_; rows_ += rows; used_ += live;
+        ++hist_[live < 0 ? 0 : (live > kHist - 1 ? kHist - 1 : live)];
+    }
+    void drain(const std::string& line, int groups_in_flight, double ms) {
+        if (!on_) return;
+        const bool adm = line.rfind("BGEN", 0) == 0;
+        if (adm) { ++adm_drains_; adm_drain_ms_ += ms; } else { ++other_drains_; other_drain_ms_ += ms; }
+        if (groups_in_flight > 0) ++drains_busy_;
+        groups_at_drain_ += groups_in_flight;
+    }
+    void admission(long long prompt_tokens, double prompt_ms) {
+        if (!on_) return;
+        ++adm_; adm_tokens_ += prompt_tokens; adm_ms_ += prompt_ms;
+    }
+    void begin(const Ex& e, int layers) { if (on_) { ex0_ = e; layers_ = layers > 0 ? layers : 1; } }
+    void adapt(double ms) { if (on_) { ++adapt_; adapt_ms_ += ms; } }
+    void no_launch(bool any_live) { if (on_) { if (any_live) ++idle_all_inflight_; else ++idle_no_live_; } }
+
+    void flush(double wall_ms, const ParEx& e = ParEx{}) {
+        if (!on_ || (windows_ == 0 && adm_ == 0 && adm_drains_ + other_drains_ == 0)) return;
+        std::fprintf(stderr, "strata par: wall %.0f ms; windows %lld rows %lld used %lld padded %lld (%.1f%% used); live slots/window:",
+                     wall_ms, windows_, rows_, used_, rows_ - used_, rows_ ? 100.0 * used_ / rows_ : 0.0);
+        for (int i = 0; i < kHist; ++i) if (hist_[i]) std::fprintf(stderr, " %d:%lld", i, hist_[i]);
+        std::fprintf(stderr, "; drains: admission %lld (%.0f ms) other %lld (%.0f ms), %lld with groups in flight, avg %.2f groups; "
+                             "admissions %lld, %lld prompt tokens in %.0f ms; adapt rounds %lld %.0f ms; "
+                             "ticks without a launch: no live slot %lld, all groups in flight %lld\n",
+                     adm_drains_, adm_drain_ms_, other_drains_, other_drain_ms_, drains_busy_,
+                     (adm_drains_ + other_drains_) ? (double) groups_at_drain_ / (adm_drains_ + other_drains_) : 0.0, adm_, adm_tokens_, adm_ms_,
+                     adapt_, adapt_ms_, idle_no_live_, idle_all_inflight_);
+        if (windows_ > 0 && (e.misses | e.entries | e.hits | e.offload) != 0) {
+            const double wl = (double) windows_ * layers_;
+            std::fprintf(stderr, "strata par experts: per window-layer: routed entries (CPU) %.2f, distinct CPU-computed experts %.2f, "
+                                 "GPU cache hits %.2f, PCIe/offload entries %.2f, distinct PCIe experts %.2f; ms per window: CPU pool run %.2f, "
+                                 "job build %.2f, verifier wait %.2f, verifier pool %.2f (of %.2f wall)\n",
+                         (e.entries - ex0_.entries) / wl, (e.misses - ex0_.misses) / wl, (e.hits - ex0_.hits) / wl,
+                         (e.offload - ex0_.offload) / wl, (e.pcie - ex0_.pcie) / wl, (e.run_ms - ex0_.run_ms) / windows_,
+                         (e.jobs_ms - ex0_.jobs_ms) / windows_, (e.vwait_ms - ex0_.vwait_ms) / windows_,
+                         (e.vpool_ms - ex0_.vpool_ms) / windows_, wall_ms / windows_);
+        }
+        std::fflush(stderr);
+        reset();
+    }
+
+private:
+    static constexpr int kHist = 18;
+    void reset() {
+        windows_ = rows_ = used_ = adm_drains_ = other_drains_ = drains_busy_ = groups_at_drain_ = adm_ = adm_tokens_ = adapt_ = 0;
+        idle_no_live_ = idle_all_inflight_ = 0;
+        adm_drain_ms_ = other_drain_ms_ = adm_ms_ = adapt_ms_ = 0;
+        std::memset(hist_, 0, sizeof hist_);
+    }
+    bool on_ = false;
+    Ex ex0_{};
+    int layers_ = 1;
+    long long windows_ = 0, rows_ = 0, used_ = 0, adm_drains_ = 0, other_drains_ = 0, drains_busy_ = 0, groups_at_drain_ = 0;
+    long long adm_ = 0, adm_tokens_ = 0, adapt_ = 0, idle_no_live_ = 0, idle_all_inflight_ = 0, hist_[kHist] = {};
+    double adm_drain_ms_ = 0, other_drain_ms_ = 0, adm_ms_ = 0, adapt_ms_ = 0;
+};
+
+}  // namespace strata::program

serve/server.py

@@ -327,6 +327,7 @@ def _timeout_env(name: str, default: float) -> float | None:
 # still working is never ended by this, only reported once (the recommendation: engine_silence_s, or this variable).
 # STRATA_ENGINE_STALL_S overrides it (0 = off).  Needs psutil (setup installs it); without it nothing is ended.
 ENGINE_STALL_S = _timeout_env("STRATA_ENGINE_STALL_S", 90.0)
+PAR_COUNTERS = os.environ.get("STRATA_PAR_COUNTERS") == "1"   # research counters (queue wait for a slot), shown in /v1/status
 # --batch: new requests take the control lines (one prompt read at a time) short prompts first: a prompt of at most
 # STRATA_ADMIT_SHORT_FIRST tokens goes before longer ones (opt-in; unset or 0: whoever gets the lock first), and a
 # longer one that waited STRATA_ADMIT_MAX_WAIT_S (default 30) counts as short.  A chat question then waits for the one
@@ -336,6 +337,13 @@ try:
 except ValueError:
     ADMIT_SHORT_FIRST = 0
 ADMIT_MAX_WAIT_S = _timeout_env("STRATA_ADMIT_MAX_WAIT_S", 30.0) or 0.0
+# STRATA_COALESCE_PREFILL=N (research, PAR-H; needs the engine's BCO / BCOK lines): a request that takes the control lines
+# for its admission also sends the prompts of up to N-1 requests still waiting for them (BCO <slot> <ids>, into free slots
+# it reserves for them); the engine reads all the prompts in ONE layer-major pass, so the experts a layer streams are
+# read once for all of them, and each waiting request then finds its prompt in its slot.  0 / unset: off.
+COALESCE_PREFILL = max(0, int(os.environ.get("STRATA_COALESCE_PREFILL") or 0))
+COALESCE_TOKENS = max(256, int(os.environ.get("STRATA_COALESCE_TOKENS") or 16384))   # prompt tokens read together at most
+COALESCE_WAIT_MS = max(0, int(os.environ.get("STRATA_COALESCE_WAIT_MS") or 0))       # a burst's later requests register within this long
 STALL_CPU_EPS_S = 0.5       # CPU seconds the engine may use in the whole stall window and still count as frozen
 STALL_IO_EPS_B = 1 << 20    # bytes read or written in that window
 
@@ -1193,6 +1201,7 @@ class StrataEngine:
             self.slot_cv = threading.Condition()
             self.waiting = 0                            # requests waiting for the control lines (ctl)
             self.admitting = 0                          # #1603: ... and those that gave the lines back to wait for a slot
+            self.par_counters = {"queued": 0, "queue_wait_s": 0.0}
             self.wait_lens: list[list[int]] = []        # ... their prompt lengths (a long read gives way to short ones)
             self.ctl_epoch = 0                          # how often the control lines were taken
             self.ctl = threading.Lock()                 # one admission or solo request on the control lines at a time
@@ -1575,6 +1584,8 @@ class StrataEngine:
                 f = line.split()
                 self._ctl_result = ("badm", len(f) >= 3 and f[2] == "1")
                 return
+            elif line.startswith("BCOK "):               # STRATA_COALESCE_PREFILL: <slot> <tokens>: the slot holds a waiting prompt
+                pass
             elif line.startswith("YIELDED "):            # the read gave way (BYIELD): <slot> <tokens read>
                 f = line.split()
                 if len(f) >= 3 and f[1].lstrip("-").isdigit() and f[2].isdigit():
@@ -1726,11 +1737,11 @@ class StrataEngine:
 
     YIELDS_MAX = 2   # #656: how often one request's prompt read gives way to a shorter waiting one
 
-    def _take_control(self, cancel, plen: int, after_epoch: int | None = None, new: bool = False):
+    def _take_control(self, cancel, plen: int, after_epoch: int | None = None, new: bool = False, prompt=None):
         """Waits for the control lines (one prompt read at a time), yielding None heartbeats; False when cancelled.
         `after_epoch`: a request whose read gave way lets the requests waiting then go first.
         `new`: the request's first admission (STRATA_ADMIT_SHORT_FIRST orders only these)."""
-        entry = [plen, time.monotonic(), new]           # (length, arrival, ordered: STRATA_ADMIT_SHORT_FIRST)
+        entry = [plen, time.monotonic(), new, prompt, cancel]           # (length, arrival, ordered: STRATA_ADMIT_SHORT_FIRST; the prompt and request, for STRATA_COALESCE_PREFILL)
         with self.slot_cv:
             self.waiting += 1
             self.wait_lens.append(entry)
@@ -1784,6 +1795,51 @@ class StrataEngine:
             self.slot_cv.notify_all()
         return True
 
+    def _release_preslot(self, cancel):
+        """STRATA_COALESCE_PREFILL: the request that was given a slot ahead of its turn left before using it."""
+        slot = getattr(cancel, "_strata_preslot", None)
+        if slot is None:
+            return
+        cancel._strata_preslot = None
+        with self.slot_cv:
+            self.slot_busy[slot] = False
+            self.slot_claims.discard(slot)
+            self.slot_cv.notify_all()
+
+    def _coalesce(self, cancel, lead_prompt):
+        """STRATA_COALESCE_PREFILL (the caller holds the control lines and sends its BGEN right after): reserve free slots
+        for requests still waiting for the control lines and send their prompts as BCO lines; the engine reads them in the
+        same layer-major pass as the BGEN's prompt.  Each follower finds its slot in cancel._strata_preslot."""
+        picked = []
+        total = len(lead_prompt)
+        if COALESCE_WAIT_MS > 0:
+            time.sleep(COALESCE_WAIT_MS / 1000.0)    # the requests of a burst reach the server a few ms apart
+        with self.slot_cv:
+            for e in list(self.wait_lens):
+                if len(picked) >= COALESCE_PREFILL - 1:
+                    break
+                prompt = e[3] if len(e) > 4 else None
+                rq = e[4] if len(e) > 4 else None
+                if prompt is None or rq is None or rq is cancel or getattr(rq, "_strata_preslot", None) is not None:
+                    continue
+                if getattr(rq, "_strata_co", False) or len(prompt) < 32 or total + len(prompt) > COALESCE_TOKENS:
+                    continue
+                slot = self.pick_slot(prompt)
+                if slot is None:
+                    break
+                self.slot_busy[slot] = True
+                self.slot_claims.add(slot)
+                self.slot_held[slot] = []
+                rq._strata_preslot = slot
+                rq._strata_co = True
+                total += len(prompt)
+                picked.append((slot, prompt))
+        for slot, prompt in picked:
+            self._send(f"BCO {slot} {','.join(str(int(t)) for t in prompt)}")
+        if picked:
+            print(f"[strata] coalesce: {len(picked)} waiting prompt(s) ({total - len(lead_prompt)} tokens) read with this "
+                  f"request's {len(lead_prompt)}", flush=True)
+
     def _admit_next(self):
         """STRATA_ADMIT_SHORT_FIRST=<tokens> (the caller holds slot_cv): the new request that takes the control lines
         next - one whose prompt is at most that long (it reads little) before a longer one, in arrival order within
@@ -1839,12 +1895,17 @@ class StrataEngine:
         prompt, left = list(ids), int(max_new)
         reused0 = None      # the first read's reused tokens: a continued leg's count includes this request's own output
         carry = {}          # #1869: the figures of the legs that ended in a hand-back to the solo path (the final DONE adds to them)
-        ok = yield from self._take_control(cancel, len(prompt), new=True)
+        ok = yield from self._take_control(cancel, len(prompt), new=True,
+                                           prompt=prompt if (COALESCE_PREFILL > 1 and not embeddings) else None)
         if not ok:
+            self._release_preslot(cancel)
             return
         holding, born = True, self.gen                  # (the engine it now has the control lines of)
         btrace("ctl acquired")
         slot, gen0, reserved = None, None, None
+        pre = getattr(cancel, "_strata_preslot", None)
+        if pre is not None:                             # STRATA_COALESCE_PREFILL: a request before this one read the prompt into this slot
+            slot, cancel._strata_preslot = pre, None
         phase = "none"            # solo -> (admit -> slot) ; "done" once the engine has finished with this request
         stop_sent = False
         yields, solo_again, refused = 0, 0, 0
@@ -1941,6 +2002,9 @@ class StrataEngine:
                                     if admit_wait:
                                         admit_wait = False
                                         self.admitting = max(0, self.admitting - 1)
+                                        if PAR_COUNTERS:                # STRATA_PAR_COUNTERS=1: time spent waiting for a free slot
+                                            self.par_counters["queued"] += 1
+                                            self.par_counters["queue_wait_s"] += time.monotonic() - admit_since
                                     break
                                 if not admit_wait:          # #1603: counted while it waits for a slot (it holds no
                                     admit_wait, admit_since = True, time.monotonic()   # control lines then)
@@ -1993,6 +2057,8 @@ class StrataEngine:
                         raise QueueWaitExpired("the inference queue wait deadline expired; retry later")
                     self.slot_live[slot] = live
                     self._last_done = None                  # (#879: this admission's DONE decides below)
+                    if COALESCE_PREFILL > 1 and not embeddings:
+                        self._coalesce(cancel, prompt)
                     self._send(f"{head} {','.join(str(int(t)) for t in prompt)}")
                     if deadline is not None:
                         cancel._strata_queue_deadline = None
@@ -4795,6 +4861,8 @@ class Service:
             slots = self.engine.slots_view() if hasattr(self.engine, "slots_view") else []
             # the requests in flight that are not in a slot: the one alone on the solo path, or waiting their turn
             in_slots = sum(1 for x in slots if x["state"] != "idle")
+            if PAR_COUNTERS:
+                live["par_counters"] = dict(self.par_counters)
             live.update(parallel=par, running=running, slots=slots, outside_slots=max(0, running - in_slots),
                         waiting=int(getattr(self.engine, "waiting", 0) or 0),
                         admitting=int(getattr(self.engine, "admitting", 0) or 0))

src/prefill/prefill.cpp

@@ -5890,6 +5890,11 @@ const char* lm_env_text() {   // STRATA_PREFILL_LAYER_MAJOR as set (nullptr: uns
 }
 bool lm_env_explicit() { const char* e = lm_env_text(); return e != nullptr && e[0] != '\0'; }
 
+bool lm_force_env() {   // STRATA_PREFILL_LM_FORCE=1 (research): the layer-major walk even for a prompt under two chunks (the identity reference of a multi-prompt read)
+    static const bool v = [] { const char* e = std::getenv("STRATA_PREFILL_LM_FORCE"); return e != nullptr && e[0] == '1'; }();
+    return v;
+}
+
 bool lm_even_env() {   // even super-chunks (STRATA_PREFILL_LM_EVEN=0: the largest that fits, then the rest)
     static const bool v = [] { const char* e = std::getenv("STRATA_PREFILL_LM_EVEN"); return !(e && e[0] == '0'); }();
     return v;
@@ -6004,13 +6009,13 @@ int64_t Prefill::lm_prepare(int64_t n, std::string& why, int64_t cap) {
     if (m.fused_bufs) { why = "the fused expert layout"; return 0; }
     if (!m.GU || !m.Dm || (n_mmq > 0 && !m.ids_identity)) { why = "the MoE buffers are missing"; return 0; }
     // (an explicit STRATA_PREFILL_LM_TOKENS may be one chunk: the A/B against the default path, experts aside)
-    const int64_t min_s = lm_tokens_env() > 0 ? m.T : 2 * m.T;
+    const int64_t min_s = (lm_force_ || lm_force_env()) ? 1 : (lm_tokens_env() > 0 ? m.T : 2 * m.T);   // (lm_force_: run_multi - several prompts, however short)
     if (n < min_s) { why = "a short prompt"; return 0; }
     const int64_t want = lm_tokens_env() > 0 ? lm_tokens_env() : (lm_by_default_ ? kLmDefaultOnTokens : kLmDefaultTokens);
     int64_t S = std::max<int64_t>(min_s, want / m.T * m.T);
     S = std::min<int64_t>(S, (n + m.T - 1) / m.T * m.T);
     if (cap > 0) {
-        if (cap % m.T != 0 || cap < min_s) { why = "the layer split's stages read different chunks"; return 0; }
+        if ((!lm_force_ && cap % m.T != 0) || cap < min_s) { why = "the layer split's stages read different chunks"; return 0; }
         S = std::min(S, cap);
     }
     const int64_t G = m.T * K;
@@ -6156,15 +6161,57 @@ void Prefill::lm_release_chain() {
 }
 
 bool Prefill::run_layer_major(const int64_t* tokens, int64_t n, int64_t pos0, std::string& err) {
+    Seg one;
+    one.tokens = tokens;
+    one.n = n;
+    one.pos0 = pos0;
+    return lm_run(std::vector<Seg>{one}, 0, nullptr, err);
+}
+
+// Several prompts in ONE layer-major pass (PAR-H, 2026-10-30): the super-chunk is a list of sub-chunks, each from one
+// prompt ("segment") with its own sessions; a layer's attention halves run per sub-chunk against that segment's
+// session (GDN / QSA / PLE state), and the layer's experts are staged ONCE for every sub-chunk of the super-chunk -
+// the streamed experts cost is shared by the prompts.  A single segment with no sessions given is exactly the usual run.
+bool Prefill::run_multi(const std::vector<Seg>& segs, std::string& err) {
+    err.clear();
+    if (segs.empty()) { err = "prefill: no prompt"; return false; }
+    int64_t total = 0;
+    for (const Seg& sg : segs) total += sg.n;
+    std::string why;
+    if (hand_in_ != nullptr || !impl_) { err = "prefill: multi-prompt read declined (chunk hand-off in use)"; return false; }
+    for (Prefill* q = this; q != nullptr; q = q->next_) q->lm_force_ = true;
+    struct ForceOff { Prefill* h; ~ForceOff() { for (Prefill* q = h; q != nullptr; q = q->next_) q->lm_force_ = false; } } force_off{this};
+    if (lm_prepare_chain(total, why) <= 0) { err = "prefill: multi-prompt read declined: " + why; return false; }
+    bool ok = lm_run(segs, 0, nullptr, err);
+    std::string drain_err;
+    if (!drain_pipeline(drain_err) && ok) {
+        ok = false;
+        err = drain_err;
+    }
+    for (Prefill* q = this; q != nullptr; q = q->next_) q->lm_mid_ = false;
+    lm_release_chain();
+    return ok;
+}
+
+bool Prefill::lm_run(const std::vector<Seg>& segs, int depth, const std::vector<LmSub>* given_super, std::string& err) {
     err.clear();
     Impl& m = *impl_;
     const core::OnDevice on_device(m.device);
     const core::ModelGeometry& g = *m.g;
-    core::SessionState& ss = *m.ss;
+    // the session of segment i on THIS stage: the segment's own (stage depth `depth`), else this stage's main session
+    auto ss_of = [&](size_t i) -> core::SessionState* {
+        const Seg& sg = segs[i];
+        return (size_t) depth < sg.stage_ss.size() && sg.stage_ss[(size_t) depth] != nullptr ? sg.stage_ss[(size_t) depth] : m.ss;
+    };
+    core::SessionState* ssp = ss_of(0);
+    const bool multi = segs.size() > 1;
+    int64_t n = 0;                       // all the prompts' tokens
+    for (const Seg& sg : segs) n += sg.n;
+    const int64_t pos0 = segs[0].pos0;   // (only the progress line's)
     const auto t_start = Clock::now();
     const cudaStream_t cs = (cudaStream_t) m.cs;
     Impl::LmBufs& lb = m.lm;
-    const int64_t S = lb.S, L = g.n_layers;
+    const int64_t L = g.n_layers;
     // a layer split's stage (2026-10-24): its layers [LB, LE); a later stage's rows arrive in R_host, written by the
     // stage before it, and a stage before the last writes its final rows into the next stage's R_host, then runs
     // the next stage on that super-chunk on a thread while it reads its own next super-chunk (one hand-off buffer:
@@ -6174,38 +6221,89 @@ bool Prefill::run_layer_major(const int64_t* tokens, int64_t n, int64_t pos0, st
     strata::kernels::QsaShapes s = core::qsa_shapes(g);   // as 0.2.2's run_impl
     const uint64_t gdn_floats = (uint64_t) g.ssm_state_size * g.ssm_v_heads * g.ssm_state_size +
                                 (uint64_t) g.ssm_conv_channels * (g.ssm_d_conv - 1);
-    int32_t prev[2] = {ss.ple_prev[0], ss.ple_prev[1]};
+    std::vector<std::array<int32_t, 2>> prevs(segs.size());
+    for (size_t i = 0; i < segs.size(); ++i) prevs[i] = {ss_of(i)->ple_prev[0], ss_of(i)->ple_prev[1]};
     PfTimer pt;
     // the MMQ row table lives in the borrowed cache slots when the prompt path borrows: written for every prompt
     // (none without an MMQ layer: the FP16 path has no such table)
     if (m.ids_identity) mmq::iota(m.ids_identity, m.T * K, m.cs);
     // ---- the PLE rows, read per chunk on a thread as run_impl does (the same state machine, in chunk order)
-    const bool ple_on = ss.ple.ready();
+    const bool ple_on = m.ss->ple.ready();
     static const bool ple_batch_env = [] {
         const char* v = std::getenv("STRATA_PLE_BATCH");
         return v == nullptr || std::atoi(v) != 0;
     }();
     const bool ple_batch = ple_on && ple_batch_env && strata::kernels::ple_native_postops_enabled() &&
-                           (ss.ple.w.key_bf16 != nullptr || ss.ple.w.key_native_data != nullptr) &&
+                           (m.ss->ple.w.key_bf16 != nullptr || m.ss->ple.w.key_native_data != nullptr) &&
                            m.region_bytes / ((uint64_t) (3 * strata::kernels::NG_HC_DIM + N + 4) * 4 + (uint64_t) N * 2 + 4096) >= 64;
-    const int32_t prev0[2] = {prev[0], prev[1]};
-    auto ple_gather = [&m, &ss, tokens, n, prev0](int64_t c0, int buf, std::string& e) -> bool {
-        const int64_t T = std::min(m.T, n - c0);
-        auto at = [&](int64_t i) { return i < 2 ? prev0[i] : (int32_t) tokens[i - 2]; };
+    // ---- the plan.  A segment is cut into pieces of at most S tokens (the usual super-chunks), each piece into
+    // sub-chunks of m.T tokens from the piece's start; pieces of several segments are packed into one super-chunk while
+    // their rows fit S.  One segment: exactly the usual super-chunks.
+    const int64_t S = m.lm.S;
+    // the sub-chunk size of a multi-prompt read: what every stage of the chain can take (the stages' chunks may differ)
+    int64_t tsub = m.T;
+    if (lm_force_)
+        for (const Prefill* q = this; q != nullptr; q = q->next_) tsub = std::min(tsub, q->impl_->T);
+    std::vector<std::vector<LmSub>> supers_own;
+    const std::vector<std::vector<LmSub>>* supers_p = &supers_own;
+    std::vector<std::vector<LmSub>> one_given;
+    if (given_super != nullptr) {
+        one_given.push_back(*given_super);
+        supers_p = &one_given;
+    } else {
+        std::vector<LmSub> cur;
+        int64_t cur_rows = 0;
+        for (size_t si = 0; si < segs.size(); ++si)
+            for (int64_t p0s = 0; p0s < segs[si].n; p0s += S) {
+                const int64_t piece = std::min(S, segs[si].n - p0s);
+                if (!cur.empty() && cur_rows + piece > S) {
+                    supers_own.push_back(std::move(cur));
+                    cur.clear();
+                    cur_rows = 0;
+                }
+                for (int64_t c0s = 0; c0s < piece; c0s += tsub) {
+                    LmSub sb;
+                    sb.seg = (int) si;
+                    sb.c0 = p0s + c0s;
+                    sb.T = std::min(tsub, piece - c0s);
+                    sb.t0 = cur_rows + c0s;
+                    cur.push_back(sb);
+                }
+                cur_rows += piece;
+            }
+        if (!cur.empty()) supers_own.push_back(std::move(cur));
+    }
+    const std::vector<std::vector<LmSub>>& supers = *supers_p;
+    std::vector<std::pair<int, int>> flat;   // every sub-chunk in order: (super-chunk, sub-chunk)
+    for (size_t si = 0; si < supers.size(); ++si)
+        for (size_t c = 0; c < supers[si].size(); ++c) flat.push_back({(int) si, (int) c});
+    std::vector<int64_t> flat_base(supers.size(), 0);
+    for (size_t si = 1; si < supers.size(); ++si) flat_base[si] = flat_base[si - 1] + (int64_t) supers[si - 1].size();
+    auto sub_at = [&](int64_t gi) -> const LmSub& { return supers[(size_t) flat[(size_t) gi].first][(size_t) flat[(size_t) gi].second]; };
+    const std::array<int32_t, 2> prev_none{-1, -1};
+    (void) prev_none;
+    std::vector<std::array<int32_t, 2>> prev0s = prevs;
+    auto ple_gather = [&m, &segs, &ss_of, &prev0s, &sub_at](int64_t gi, int buf, std::string& e) -> bool {
+        const LmSub& sb = sub_at(gi);
+        const int64_t* tokens = segs[(size_t) sb.seg].tokens;
+        core::SessionState* ssq = ss_of((size_t) sb.seg);
+        const int64_t c0 = sb.c0, T = sb.T;
+        const std::array<int32_t, 2>& prev0 = prev0s[(size_t) sb.seg];
+        auto at = [&](int64_t i) { return i < 2 ? prev0[(size_t) i] : (int32_t) tokens[i - 2]; };
         int32_t pv[2] = {at(c0), at(c0 + 1)};
         for (int64_t t = 0; t < T; ++t) {
             const int32_t tok = (int32_t) tokens[c0 + t];
-            strata::kernels::ngram_rows(&tok, pv, 1, ss.ple.consts,
+            strata::kernels::ngram_rows(&tok, pv, 1, m.ss->ple.consts,
                                         m.ple_rows[buf].data() + t * strata::kernels::PLE_N_HEADS);
             pv[0] = pv[1];
             pv[1] = tok;
         }
-        return ss.ple.table->gather_batch(m.ple_rows[buf].data(), (size_t) T, m.ple_emb_host[buf], e);
+        return m.ss->ple.table->gather_batch(m.ple_rows[buf].data(), (size_t) T, m.ple_emb_host[buf], e);
     };
     std::string ple_next_err;
     std::future<bool> ple_next;
     int ple_buf = 0;
-    auto ple_land = [&](int64_t c0, int64_t T) -> bool {
+    auto ple_land = [&](int64_t gi, int64_t T) -> bool {
         const auto tp = Clock::now();
         if (!ple_next.get()) { err = ple_next_err; return false; }
         // #1996: the pageable fallback (no RAM to pin) goes up through a pinned slot; the pinned buffers as before
@@ -6216,12 +6314,12 @@ bool Prefill::run_layer_major(const int64_t* tokens, int64_t n, int64_t pos0, st
             err = std::string("prefill: the PLE rows' upload failed: ") + cudaGetErrorString(cudaGetLastError());
             return false;
         }
-        if (c0 + m.T < n) {
+        if (gi + 1 < (int64_t) flat.size()) {
             if (cudaEventSynchronize(m.ple_copied[ple_buf ^ 1]) != cudaSuccess) {
                 err = std::string("prefill: the PLE rows' upload failed: ") + cudaGetErrorString(cudaGetLastError());
                 return false;
             }
-            ple_next = std::async(std::launch::async, [&ple_gather, &ple_next_err, c1 = c0 + m.T, b = ple_buf ^ 1] {
+            ple_next = std::async(std::launch::async, [&ple_gather, &ple_next_err, c1 = gi + 1, b = ple_buf ^ 1] {
                 strata::aux_cpus::pin_current_thread();
                 return ple_gather(c1, b, ple_next_err);
             });
@@ -6564,9 +6662,13 @@ bool Prefill::run_layer_major(const int64_t* tokens, int64_t n, int64_t pos0, st
         return true;
     };
 
-    for (int64_t sc0 = 0; sc0 < n; sc0 += S) {
-        const int64_t Sn = std::min(S, n - sc0);
-        const int64_t nc = (Sn + m.T - 1) / m.T;
+    for (size_t sci = 0; sci < supers.size(); ++sci) {
+        const std::vector<LmSub>& subs = supers[sci];
+        int64_t Sn = 0;
+        for (const LmSub& sb : subs) Sn = std::max(Sn, sb.t0 + sb.T);
+        const int64_t nc = (int64_t) subs.size();
+        int64_t last_seg0 = -1;
+        for (int64_t c = 0; c < nc; ++c) if (subs[(size_t) c].seg == 0) last_seg0 = c;
         ++lm_supers;
         lm_mid_ = true;
         for (int64_t l = LB; l <= LE; ++l) {
@@ -6575,8 +6677,8 @@ bool Prefill::run_layer_major(const int64_t* tokens, int64_t n, int64_t pos0, st
             // host mode: sub-chunk c's rows go up on the copy stream into buffer c % 2 (issued one sub-chunk ahead)
             auto r_buf = [&](int64_t c) { return (c & 1) ? lb.R2 : R_chunk; };
             auto issue_up = [&](int64_t c) {
-                const int64_t cc0 = sc0 + c * m.T, TT = std::min(m.T, n - cc0);
-                cudaMemcpyAsync(r_buf(c), lb.R_host + (size_t) (c * m.T) * D, (size_t) TT * D * 4, cudaMemcpyHostToDevice,
+                const int64_t TT = subs[(size_t) c].T;
+                cudaMemcpyAsync(r_buf(c), lb.R_host + (size_t) subs[(size_t) c].t0 * D, (size_t) TT * D * 4, cudaMemcpyHostToDevice,
                                 m.copy);
                 cudaEventRecord(lb.ev_in[c & 1], m.copy);
             };
@@ -6587,12 +6689,17 @@ bool Prefill::run_layer_major(const int64_t* tokens, int64_t n, int64_t pos0, st
             for (int64_t c = 0; c < nc; ++c) {
                 if (should_stop && should_stop()) { err = "cancelled"; return false; }
                 core::progress_beat();
-                core::progress_at("reading the prompt (layer-major): layer", l, pos0 + sc0 + c * m.T);
+                const LmSub& sub = subs[(size_t) c];
+                const Seg& sgm = segs[(size_t) sub.seg];
+                const int64_t* const tokens = sgm.tokens;
+                ssp = ss_of((size_t) sub.seg);
+                core::progress_at("reading the prompt (layer-major): layer", l, sgm.pos0 + sub.c0);
                 qsa_index = qsa_index0;
                 gdn_index = gdn_index0;
-                const int64_t c0 = sc0 + c * m.T;          // in this run's tokens
-                const int64_t T = std::min(m.T, n - c0), p0 = pos0 + c0;
-                const int64_t t0 = c * m.T;                // in the super-chunk
+                const int64_t c0 = sub.c0;                 // in this prompt's tokens
+                const int64_t T = sub.T, p0 = sgm.pos0 + c0;
+                const int64_t t0 = sub.t0;                 // in the super-chunk
+                const int64_t gi = flat_base[sci] + c;     // in the whole run
                 m.R = lb.r_dev ? lb.R + t0 * D : r_buf(c);
                 if (!lb.r_dev) {
                     if (l > 0) {
@@ -6643,11 +6750,14 @@ bool Prefill::run_layer_major(const int64_t* tokens, int64_t n, int64_t pos0, st
                         if (!core::embed_row(*m.wt, g, tokens[c0 + t], m.emb + t * N, cs, err)) return false;
                     gr_broadcast(m.emb, m.R, T, cs);
                     if (ple_on && !ple_next.valid())
-                        ple_next = std::async(std::launch::async, [&ple_gather, &ple_next_err, c0, b = ple_buf] {
+                        ple_next = std::async(std::launch::async, [&ple_gather, &ple_next_err, gi, b = ple_buf] {
                             strata::aux_cpus::pin_current_thread();
-                            return ple_gather(c0, b, ple_next_err);
+                            return ple_gather(gi, b, ple_next_err);
                         });
-                    for (int64_t t = 0; t < T; ++t) { prev[0] = prev[1]; prev[1] = (int32_t) tokens[c0 + t]; }
+                    {
+                        std::array<int32_t, 2>& pvv = prevs[(size_t) sub.seg];
+                        for (int64_t t = 0; t < T; ++t) { pvv[0] = pvv[1]; pvv[1] = (int32_t) tokens[c0 + t]; }
+                    }
                 } else if (l > LB) {   // (a later stage's first layer: the stage before wrote layer LB - 1)
                     // the previous layer's experts: the hyper-connection write of its FFN half
                     gr_write(m.R, acc, m.inj, HC, T, cs);
@@ -6692,19 +6802,21 @@ bool Prefill::run_layer_major(const int64_t* tokens, int64_t n, int64_t pos0, st
                             err = std::string("prefill: ") + cudaGetErrorString(cudaGetLastError());
                             return false;
                         }
-                        lm_mid_ = c + 1 < nc;
-                        if (on_stage_chunk && !lm_mid_ && !on_stage_chunk(p0 + T, err)) return false;
-                        if (on_chunk && !on_chunk(m.R, T, p0, err)) return false;
+                        lm_mid_ = c != last_seg0 && last_seg0 >= 0;
+                        if (sub.seg == 0) {   // (the reports are the first prompt's; the others only fill their sessions)
+                            if (on_stage_chunk && !lm_mid_ && !on_stage_chunk(p0 + T, err)) return false;
+                            if (on_chunk && !on_chunk(m.R, T, p0, err)) return false;
+                        }
                     }
                     continue;
                 }
                 const core::LayerView v(*m.wt, l);
-                if (l == 1 && ple_on && !ple_land(c0, T)) return false;
+                if (l == 1 && ple_on && !ple_land(gi, T)) return false;
                 // ---- the PLE block at layer 1 (run_impl's, batched or token by token)
                 if (l == 1 && ple_on && ple_batch) {
                     pt.mark(kPfPle, cs);
                     const auto tp = Clock::now();
-                    const strata::kernels::PleWeights& pw = ss.ple.w;
+                    const strata::kernels::PleWeights& pw = m.ss->ple.w;
                     constexpr int64_t HD = strata::kernels::NG_HC_DIM;
                     const uint64_t per_token = (uint64_t) (3 * HD + N + 4) * 4 + (uint64_t) N * (bf16x2(m.f16_io) ? 4 : 2) + 4096;
                     const int64_t SB = std::min<int64_t>(T, (int64_t) (m.region_bytes / per_token));
@@ -6732,7 +6844,7 @@ bool Prefill::run_layer_major(const int64_t* tokens, int64_t n, int64_t pos0, st
                         m.gemm.bf16(e16, pw.value_bf16, val, nb, N, N);
                         if (e16_lo) m.gemm.bf16(e16_lo, pw.value_bf16, val, nb, N, N, 0, 1.0f);
                         try {
-                            strata::kernels::native_ple_postops_batch(key, m.R + s0 * D, val, ss.ple.hist, pw, qn, gated,
+                            strata::kernels::native_ple_postops_batch(key, m.R + s0 * D, val, ssp->ple_hist, pw, qn, gated,
                                                                       gate, (int) nb, cs);
                         } catch (const std::exception& e) { err = std::string("prefill PLE: ") + e.what(); return false; }
                     }
@@ -6745,10 +6857,10 @@ bool Prefill::run_layer_major(const int64_t* tokens, int64_t n, int64_t pos0, st
                         po.normalized = m.ple_norm;
                         po.result = m.R + t * D;
                         try {
-                            strata::kernels::ple_block(m.ple_emb + t * N, m.R + t * D, ss.ple.hist, ss.ple.w, po,
-                                                       ss.ple.scratch, cs);
+                            strata::kernels::ple_block(m.ple_emb + t * N, m.R + t * D, ssp->ple_hist, m.ss->ple.w, po,
+                                                       m.ss->ple.scratch, cs);
                         } catch (const std::exception& e) { err = std::string("prefill PLE: ") + e.what(); return false; }
-                        strata::kernels::ple_history_advance(ss.ple.hist, m.ple_norm, cs);
+                        strata::kernels::ple_history_advance(ssp->ple_hist, m.ple_norm, cs);
                     }
                     stats_.ms_ple += ms_since(tp);
                 }
@@ -6806,7 +6918,7 @@ bool Prefill::run_layer_major(const int64_t* tokens, int64_t n, int64_t pos0, st
                                               *wsa = need(v, "ssm_a", err);
                         if (!wqkv || !wg || !wo || !wa || !wb || !wc || !wnm || !wdt || !wsa) return false;
                         pt.mark(kPfGdn, cs);
-                        float* state = ss.gdn_state + (size_t) (gdn_index - ss.gdn_ord0) * gdn_floats;
+                        float* state = ssp->gdn_state + (size_t) (gdn_index - ssp->gdn_ord0) * gdn_floats;
                         float* conv = state + (uint64_t) g.ssm_state_size * g.ssm_v_heads * g.ssm_state_size;
                         if (!native_proj(m.gemm, wqkv, m.mixed_h, m.qkv, T, v.name("attn_qkv.weight"), err)) return false;
                         if (!native_proj(m.gemm, wg, m.mixed_h, m.z, T, v.name("attn_gate.weight"), err)) return false;
@@ -6828,7 +6940,7 @@ bool Prefill::run_layer_major(const int64_t* tokens, int64_t n, int64_t pos0, st
                         ++gdn_index;
                     } else if (half == 0) {
                         // ======================= QSA (run_impl's, without its debug dumps) =======================
-                        const core::QsaState& st = ss.qsa_states[qsa_index];
+                        const core::QsaState& st = ssp->qsa_states[qsa_index];
                         const core::WeightRef *wq = need(v, "attn_q.weight", err), *wk = need(v, "attn_k.weight", err),
                                               *wv = need(v, "attn_v.weight", err), *wo = need(v, "attn_output.weight", err),
                                               *wik = need(v, "indexer.k_proj.weight", err),
@@ -7000,17 +7112,22 @@ bool Prefill::run_layer_major(const int64_t* tokens, int64_t n, int64_t pos0, st
                       cudaGetErrorString(cudaGetLastError());
                 return false;
             }
-            if (on_stage_chunk && !on_stage_chunk(pos0 + sc0 + Sn, err)) return false;
+            if (on_stage_chunk && last_seg0 >= 0 &&
+                !on_stage_chunk(segs[0].pos0 + subs[(size_t) last_seg0].c0 + subs[(size_t) last_seg0].T, err))
+                return false;
             next_err_.clear();
-            next_run_ = std::async(std::launch::async, [this, tokens, sc0, Sn, pos0] {
+            next_run_ = std::async(std::launch::async, [this, segs_c = std::make_shared<std::vector<Seg>>(segs), depth,
+                                                        sp_super = std::make_shared<std::vector<LmSub>>(supers[sci])] {
                 strata::aux_cpus::pin_current_thread();
-                return next_->run_layer_major(tokens + sc0, Sn, pos0 + sc0, next_err_);
+                return next_->lm_run(*segs_c, depth + 1, sp_super.get(), next_err_);
             });
         }
     }
     lm_mid_ = false;
-    ss.ple_prev[0] = prev[0];
-    ss.ple_prev[1] = prev[1];
+    for (size_t i = 0; i < segs.size(); ++i) {
+        ss_of(i)->ple_prev[0] = prevs[i][0];
+        ss_of(i)->ple_prev[1] = prevs[i][1];
+    }
     if (ple_next.valid()) ple_next.wait();
     if (cudaStreamSynchronize(cs) != cudaSuccess) {
         err = std::string("prefill: ") + cudaGetErrorString(cudaGetLastError());

src/program/generate.cpp

@@ -101,6 +101,13 @@
 #include "strata/spec/suffix_drafter.hpp"
 #include "strata/spec/spec_start.hpp"
 #include "strata/program/batch_groups.hpp"
+#include "strata/program/par_counters.hpp"
+#if defined(__linux__)
+#include <execinfo.h>
+#include <fcntl.h>
+#include <signal.h>
+#include <unistd.h>
+#endif
 #include "strata/program/pipe_probe.hpp"
 #include "strata/kernels/cvec.hpp"
 #include "strata/core/progress.hpp"
@@ -235,6 +242,35 @@ std::string vram_free_note() {
     return buf;
 }
 
+#if defined(__linux__)
+// STRATA_CRASH_BT=1 (research): a fatal signal prints the call stack to stderr (the addresses go through addr2line -e <binary>)
+namespace {
+char g_crash_path[512] = "";
+void crash_bt_handler(int sig) {
+    void* frames[48];
+    const int n = backtrace(frames, 48);
+    const char msg[] = "strata: fatal signal, call stack:\n";
+    int fd = 2;
+    if (g_crash_path[0] != '\0') {
+        const int f = open(g_crash_path, O_WRONLY | O_CREAT | O_APPEND, 0644);
+        if (f >= 0) fd = f;
+    }
+    (void) !write(fd, msg, sizeof msg - 1);
+    backtrace_symbols_fd(frames, n, fd);
+    signal(sig, SIG_DFL);
+    raise(sig);
+}
+const bool crash_bt_installed = [] {
+    const char* v = std::getenv("STRATA_CRASH_BT");
+    if (v == nullptr || v[0] == '\0' || v[0] == '0') return false;
+    if (v[0] != '1' || v[1] != '\0') std::snprintf(g_crash_path, sizeof g_crash_path, "%s", v);
+    for (int sg : {SIGFPE, SIGSEGV, SIGBUS, SIGABRT}) signal(sg, crash_bt_handler);
+    return true;
+}();
+}  // namespace
+#endif
+
+
 // STRATA_PIPE_BYIELD_NODRAIN=1 (read once; off by default): with the pipelined groups, a BYIELD that reaches the main
 // loop is dropped without draining the groups first.  A pipelined read never gives way (batch_read_piece), so the
 // BYIELD is for a read that has ended, and the drain (and the refill after it) did nothing for it
@@ -12480,6 +12516,14 @@ int main(int argc, char** argv) {
         const bool batch_telemetry_enabled = batch_telemetry_env != nullptr &&
                                              batch_telemetry_env[0] == '1' && batch_telemetry_env[1] == '\0';
         strata::core::BatchWindowSampler batch_telemetry(batch_telemetry_enabled);
+        strata::program::ParCounters pc;   // STRATA_PAR_COUNTERS=1: counters only (par_counters.hpp); off: a branch
+        auto pc_ex = [&]() {
+            strata::program::ParCounters::Ex e;
+            e.misses = drive.d.multi_misses; e.entries = drive.d.multi_entries; e.hits = drive.d.cache_hits;
+            e.offload = drive.d.offload_entries; e.pcie = drive.d.pcie_experts; e.run_ms = drive.d.ms_run; e.jobs_ms = drive.d.ms_jobs;
+            e.vwait_ms = ver.ms_wait; e.vpool_ms = ver.ms_pool;
+            return e;
+        };
         auto report_batch_window = [&](int processed_rows, int active_rows, bool idle, Clock::time_point started) {
             if (!batch_telemetry_enabled) return;
             const auto sample = batch_telemetry.complete(processed_rows, active_rows, idle, started, Clock::now());
@@ -12491,6 +12535,12 @@ int main(int argc, char** argv) {
         };
         int admit_slot = -1;               ///< the slot the request being read will continue in (BGEN)
         long long admit_max_new = 0;
+        // BCO <slot> <ids> (research, PAR-H; STRATA_COALESCE_PREFILL on the server): prompts of waiting requests, read in
+        // the SAME layer-major pass as the next BGEN's prompt into the named idle slots' sessions (the experts a layer
+        // streams are shared), so their own admissions find the prompt in the slot and read nothing.  BCOK <slot> <tokens>
+        // says a slot holds its prompt; a declined or failed co-read says nothing and costs nothing.
+        struct CoReq { int slot = -1; std::vector<int64_t> ids; };
+        std::vector<CoReq> co_pending, co_set;
         auto batch_on = [&] { for (const BSlot& b : bs) if (b.active) return true; return false; };
         auto try_next_line = [&](std::string& out) -> bool {
             std::lock_guard<std::mutex> lk(in_mu);
@@ -12757,6 +12807,7 @@ int main(int argc, char** argv) {
                     bt_stage_pool0[k] = stage_verifier((int) k).ms_pool;
                 }
                 bt_cpu_ent0 = drive.d.multi_entries; bt_off_ent0 = drive.d.offload_entries;
+                if (pc.on()) pc.begin(pc_ex(), (int) g.n_layers);
                 bt_a_rounds0 = a_rounds; bt_a_swapped0 = a_swapped;
             }
             const Clock::time_point w0 = Clock::now();
@@ -12869,8 +12920,10 @@ int main(int argc, char** argv) {
             }
             strata::core::progress().busy.store(was_busy);
             report_batch_window(S, A, !batch_on(), w0);
+            pc.window(S, A);
             if (!batch_on() && bt_windows > 0) {
                 const double w = (double) bt_windows, wall = msd(bt_start, Clock::now());
+                pc.flush(wall, pc_ex());
                 const double L = (double) g.n_layers;
                 const double e_vram = (double) (drive.d.cache_hits - bt_hits0), e_pcie = (double) (drive.d.offload_entries - bt_off_ent0),
                              e_cpu = (double) (drive.d.multi_entries - bt_cpu_ent0);
@@ -13173,6 +13226,7 @@ int main(int argc, char** argv) {
                 p4.span_start(bt_start);
                 bt_wait0 = ver.ms_wait; bt_pool0 = ver.ms_pool;
                 bt_miss0 = drive.d.multi_misses; bt_hits0 = drive.d.cache_hits; bt_pcie0 = drive.d.pcie_experts;
+                if (pc.on()) pc.begin(pc_ex(), (int) g.n_layers);
             }
             for (int k = 0; k < n_pipe; ++k) {
                 const int gi = stage_group[(size_t) k];
@@ -13239,6 +13293,7 @@ int main(int argc, char** argv) {
                 ++pa_windows;
                 G.inflight = false;
                 report_batch_window(G.S, G.active, !pipe_inflight() && !batch_on(), G.started);
+                pc.window(G.S, G.active);
             }
             // start waiting groups on free stages, the longest waiting first; a new step on stage 0 round-robin
             for (int k = n_pipe - 1; k >= 0; --k) {
@@ -13366,6 +13421,7 @@ int main(int argc, char** argv) {
             if (!pipe_inflight() && !batch_on() && bt_windows > 0) {   // all idle: the timing line, as batch_step
                 const double w = (double) bt_windows, wall = std::chrono::duration<double, std::milli>(Clock::now() - bt_start).count();
                 p4.span_end(wall);
+                pc.flush(wall, pc_ex());
                 char pa_txt[160] = "";
                 if (pipe_adapt_every > 0)
                     std::snprintf(pa_txt, sizeof pa_txt, "; adaptive tier (STRATA_PIPE_ADAPT_EVERY): %lld rounds, %lld experts "
@@ -13406,10 +13462,11 @@ int main(int argc, char** argv) {
                 // (STRATA_PIPE_BYIELD_NODRAIN: not for a BYIELD, which is dropped below)
                 if (piped && line.rfind("BSTOP ", 0) != 0 &&
                     !(pipe_byield_nodrain_env() && line.rfind("BYIELD", 0) == 0)) {
-                    const int in_flight = trace ? (int) std::count_if(pg.begin(), pg.end(), [](const PGroup& G) {
+                    const int in_flight = (trace || pc.on()) ? (int) std::count_if(pg.begin(), pg.end(), [](const PGroup& G) {
                                                       return G.inflight; }) : 0;
                     const auto td = Clock::now();
                     if (!pipe_drain()) return 1;
+                    if (pc.on()) pc.drain(line, in_flight, std::chrono::duration<double, std::milli>(Clock::now() - td).count());
                     if (trace) {   // STRATA_TRACE: what the drain before this line cost (BGEN: an admission's first phase)
                         line_drain_ms = std::chrono::duration<double, std::milli>(Clock::now() - td).count();
                         std::fprintf(stderr, "strata trace: drain for %.*s: %d groups in flight, %.1f ms\n",
@@ -13430,6 +13487,21 @@ int main(int argc, char** argv) {
                 if (b >= 0 && b < (int) bs.size()) bs[(size_t) b].stop = true;
                 continue;
             }
+            if (line.rfind("BCO ", 0) == 0) {   // BCO <slot> <id,id,...>: kept for the next BGEN
+                CoReq cr;
+                char* e1 = nullptr;
+                cr.slot = (int) std::strtol(line.c_str() + 4, &e1, 10);
+                for (const char* q = e1; q != nullptr && *q != '\0';) {
+                    while (*q == ' ' || *q == ',') ++q;
+                    if (*q == '\0') break;
+                    char* e2 = nullptr;
+                    cr.ids.push_back((int64_t) std::strtoll(q, &e2, 10));
+                    if (e2 == q) { cr.ids.clear(); break; }
+                    q = e2;
+                }
+                if (cr.slot >= 0 && cr.slot < (int) bs.size() && cr.ids.size() >= 3) co_pending.push_back(std::move(cr));
+                continue;
+            }
             if (line.rfind("BYIELD", 0) == 0) {   // for a prompt read that has ended meanwhile
                 if (trace && piped)
                     std::fprintf(stderr, "strata trace: BYIELD dropped, %d groups in flight\n",
@@ -13465,6 +13537,8 @@ int main(int argc, char** argv) {
                 admit_slot = (int) b;
                 admit_max_new = mn;
                 adm_drain_ms = line_drain_ms;
+                co_set = std::move(co_pending);
+                co_pending.clear();
                 line = std::string(bimg ? "GENI 1" : "GEN 1") + e2;
             }
             // #477: every --expert-profile-save-every minutes, before the next request (at QUIT: after the loop)
@@ -15061,9 +15135,77 @@ int main(int argc, char** argv) {
                     ConvCheckpoint::PromptPlan& p;
                     ~SegPlanScope() { p = {}; }
                 } seg_plan_scope{seg_plan};
+                // the first run of a BGEN with co-reads: the waiting prompts are read in the same layer-major pass
+                bool co_tried = false;
+                auto co_run = [&](int64_t q, int64_t r, std::string& ce) -> bool {
+                    using PSeg = strata::prefill::Prefill::Seg;
+                    std::vector<PSeg> segs;
+                    PSeg lead;
+                    lead.tokens = ids.data() + q;
+                    lead.n = r - q;
+                    lead.pos0 = q;
+                    segs.push_back(lead);
+                    std::vector<const CoReq*> used;
+                    for (const CoReq& c : co_set) {
+                        if (c.slot < 0 || c.slot >= (int) bs.size() || c.slot == admit_slot || bs[(size_t) c.slot].active ||
+                            bs[(size_t) c.slot].partial)
+                            continue;
+                        bool dup = false;
+                        for (const CoReq* u : used) dup = dup || u->slot == c.slot;
+                        if (dup) continue;
+                        PSeg sg;
+                        sg.tokens = c.ids.data();
+                        sg.n = (int64_t) c.ids.size() - 1;   // (the last token goes through the slot's own first window)
+                        sg.pos0 = 0;
+                        for (size_t k = 0; k < bslot_ss.size(); ++k) sg.stage_ss.push_back(bslot_ss[k][(size_t) c.slot].get());
+                        segs.push_back(sg);
+                        used.push_back(&c);
+                    }
+                    if (segs.size() < 2) return sp.run(lead.tokens, lead.n, lead.pos0, ce);
+                    for (const CoReq* c : used) {
+                        park_slot(c->slot);   // the conversation the slot holds, before this read overwrites it
+                        for (size_t k = 0; k < bslot_ss.size(); ++k) {
+                            const strata::core::OnDevice on_k(k == 0 ? 0 : stages[k - 1]->dev);
+                            strata::core::session_zero(*bslot_ss[k][(size_t) c->slot], g, nullptr, nullptr);
+                            cudaDeviceSynchronize();
+                        }
+                        bs[(size_t) c->slot] = BSlot{};   // (it holds nothing until the read is done)
+                    }
+                    const auto tco = Clock::now();
+                    std::fprintf(stderr, "strata batch: co-read starting: %zu more prompt(s) with this request's %lld tokens\n",
+                                 used.size(), (long long) lead.n);
+                    std::fflush(stderr);
+                    if (!sp.run_multi(segs, ce)) {
+                        if (ce.rfind("prefill: multi-prompt read declined", 0) == 0) {   // nothing was read: the usual run
+                            std::fprintf(stderr, "strata batch: co-read of %zu prompt(s) declined (%s)\n", used.size(), ce.c_str());
+                            ce.clear();
+                            return sp.run(lead.tokens, lead.n, lead.pos0, ce);
+                        }
+                        return false;
+                    }
+                    for (const CoReq* c : used) {
+                        BSlot& sl = bs[(size_t) c->slot];
+                        sl = BSlot{};
+                        sl.ids.assign(c->ids.begin(), c->ids.end() - 1);
+                        sl.cached = true;
+                        sl.cvec = want_cvec;
+                        std::printf("BCOK %d %lld\n", c->slot, (long long) sl.ids.size());
+                    }
+                    std::fflush(stdout);
+                    std::fprintf(stderr, "strata batch: co-read: %zu more prompt(s) read in the same pass as this request's %lld tokens "
+                                         "(%.0f ms)\n", used.size(), (long long) lead.n,
+                                 std::chrono::duration<double, std::milli>(Clock::now() - tco).count());
+                    return true;
+                };
                 auto run = [&](int64_t q, int64_t r) -> bool {
                     if (planned) sp.force_chunk(plan.chunk, plan_single, plan.chunked);
-                    const bool run_ok = sp.run(ids.data() + q, r - q, q, e);
+                    bool run_ok;
+                    if (!co_set.empty() && !co_tried && admit_slot >= 0 && req_imgs.empty()) {
+                        co_tried = true;
+                        run_ok = co_run(q, r, e);
+                    } else {
+                        run_ok = sp.run(ids.data() + q, r - q, q, e);
+                    }
                     if (run_ok && q == 0 && sp.kva_floor() > 0) {   // KVA approximated the start of this prompt: remember which tokens
                         kva_floors.note(ids.data(), sp.kva_floor());
                         if (use_mtp && mtp.cell_floor_enabled()) mtp.set_cell_floor(sp.kva_floor());
@@ -15221,7 +15363,14 @@ int main(int argc, char** argv) {
                     }
                     return true;
                 }();
-                if (piped && pr_conc && req_imgs.empty()) {   // (with no slot decoding it only watches for a BYIELD)
+                // STRATA_PIPED_CONCURRENT_MIN=N (opt-in, with STRATA_PIPED_CONCURRENT=1): only a read of at least N tokens
+                // runs beside the decoding slots on its own thread; a shorter one takes the usual path (drained windows,
+                // no priority streams' cost on the short prompts' TTFT and throughput).  Unset or 0: every read, as before.
+                static const int64_t pr_conc_min = [] {
+                    const char* v = std::getenv("STRATA_PIPED_CONCURRENT_MIN");
+                    return v != nullptr ? (int64_t) std::max(0LL, std::atoll(v)) : (int64_t) 0;
+                }();
+                if (piped && pr_conc && req_imgs.empty() && (b0 - a0) >= pr_conc_min) {   // (with no slot decoding it only watches for a BYIELD)
                     // the drafter's prompt pass captures its graphs and grows its records at first use, and the
                     // prompt path's pinned upload slots grow at their first larger upload: done here, with nothing in
                     // flight (on the reader's thread a capture, an allocation or a free takes the driver's lock and
@@ -17244,6 +17393,7 @@ int main(int argc, char** argv) {
                     const auto tc4 = Clock::now();
                     std::fprintf(stderr, "strata batch: slot %d takes %lld tokens (copied in %.1f ms)\n", admit_slot,
                                  (long long) live.size(), std::chrono::duration<double, std::milli>(tc4 - tc0).count());
+                    pc.admission((long long) live.size(), prompt_ms);
                     if (trace) {   // STRATA_TRACE: the admission's phases (ms), from the drain before BGEN to this copy
                         auto span = [](Clock::time_point a, Clock::time_point b) {
                             return std::chrono::duration<double, std::milli>(b - a).count();

tools/batch_test.py

@@ -88,6 +88,12 @@ def main():
     ap.add_argument("--long-tokens", type=int, default=0,
                     help="the LAST prompt is read from about this many tokens of filler text first: its chunks run "
                          "beside the earlier slots' windows (the prompt path lends experts then)")
+    ap.add_argument("--co", action="store_true",
+                    help="PAR-H: before the first BGEN send BCO lines for the other prompts, so the engine reads every prompt "
+                         "in one layer-major pass (the engine needs STRATA_PREFILL_LM_FORCE=1 for prompts this short to be "
+                         "layer-major in the arm without --co)")
+    ap.add_argument("--seq", action="store_true", help="PAR-H: admit the prompts one at a time, each run to its end (no batch composition)")
+    ap.add_argument("--pad", type=int, default=0, help="PAR-H: every prompt gets about this many tokens of its own filler text first")
     ap.add_argument("--mt-min", default="1", help="STRATA_IQ_MT_MIN for the engine (1: exact; empty: the default)")
     a = ap.parse_args()
     cfg = json.loads(Path(a.config).read_text())
@@ -101,6 +107,13 @@ def main():
             per = max(1, len(tok.encode(para % 100)))
             filler = "".join(para % i for i in range(a.long_tokens // per + 1))
             text = f"<|im_start|>user\n{filler}\n\n{q}<|im_end|>\n<|im_start|>assistant\n<think>\n\n</think>\n\n"
+        if a.pad:
+            para = ("Note %d-%d: the survey team measured the northern ridge, logged the readings of station %d, and sent the "
+                    "figures to the office, where a clerk filed the sheets under reference %d. ")
+            qi = QUESTIONS.index(q)
+            per = max(1, len(tok.encode(para % (qi, 100, 100, 100))))
+            filler = "".join(para % (qi, i, (i * 7 + qi) % 97, i * 13 + qi) for i in range(a.pad // per + 1))
+            text = f"<|im_start|>user\n{filler}\n\n{q}<|im_end|>\n<|im_start|>assistant\n<think>\n\n</think>\n\n"
         prompts.append(tok.encode(text, parse_special=True))
     env = {"STRATA_DECODE_TIMING": "1", **({"STRATA_VERIFY_PROFILE": "1"} if os.environ.get("PROF") else {})}
     if a.mt_min:
@@ -126,6 +139,10 @@ def main():
     done = {}
     t_admit = time.time()
     first_bt = None
+    if a.co:
+        for i, ids in enumerate(prompts):
+            if i >= 1:
+                eng.send(f"BCO {i} " + ",".join(map(str, ids)))
     for i, ids in enumerate(prompts):
         eng.send(" ".join(x for x in ("BGEN", str(i), str(a.max_new), a.keys, ",".join(map(str, ids))) if x))
         for line in out:
@@ -144,8 +161,19 @@ def main():
                 if line.split()[2] == "0":
                     done[i] = line
                 break
+        if a.seq and i not in done:   # one prompt at a time: this slot runs to its end before the next is admitted
+            for line in out:
+                if line.startswith("BT "):
+                    _, s, y = line.split()
+                    got[int(s)].append(int(y))
+                elif line.startswith("BDONE "):
+                    done[int(line.split()[1])] = line
+                    break
+                elif line.startswith("ERR"):
+                    print("batch:", line)
+                    return 1
     t_admitted = time.time()
-    for line in out:
+    for line in (out if len(done) < len(prompts) else []):
         if line.startswith("BT "):
             _, s, y = line.split()
             got[int(s)].append(int(y))

tools/par/both.py

@@ -0,0 +1,16 @@
+#!/usr/bin/env python3
+"""both.py URL OUT.json - armrun --loadcmd: the sustained closed-loop reference (this harness, W2 1k/128, c from BOTH_CONC, 6 rounds per client)
+followed by PAR-A's synchronized-burst loader (env PA_LOAD, 2 measured rounds after a warm round).  Both read in one server start."""
+import os, subprocess, sys
+here = os.path.dirname(os.path.abspath(__file__))
+sys.path.insert(0, os.path.join(here, ".."))
+import par_load
+url, out = sys.argv[1], sys.argv[2]
+conc = os.environ.get("BOTH_CONC", "8,16,32")
+print("### SUSTAINED (par_load.py, closed loop, rounds=6)", flush=True)
+par_load.main([url, "--workload", "w2", "--conc", conc, "--rounds", "6", "--max-tokens", "128", "--warm-requests", "2",
+               "--nonce", "sus", "--out", out + ".sustained.json"])
+print("### BURST (PAR-A par_load.py, synchronized waves)", flush=True)
+sys.stdout.flush()
+subprocess.run([sys.executable, os.environ["PA_LOAD"], url, "--conc", conc, "--rounds", "2", "--prompt-tokens", "1000", "--max-tokens", "128",
+                "--out", out + ".burst.json"])

tools/par/cmp_texts.py

@@ -0,0 +1,19 @@
+#!/usr/bin/env python3
+"""cmp_texts.py A.json B.json - compare the answers (par_load --keep-text) of two ledgers request by request (same nonce)."""
+import json, sys
+a, b = (json.load(open(x)) for x in sys.argv[1:3])
+ra = {(r["workload"], r.get("level"), r["idx"]): r for r in a["requests"]}
+rb = {(r["workload"], r.get("level"), r["idx"]): r for r in b["requests"]}
+same = diff = 0
+for k, r in sorted(ra.items(), key=str):
+    o = rb.get(k)
+    if o is None or "text" not in r or "text" not in o:
+        continue
+    if r["text"] == o["text"]:
+        same += 1
+    else:
+        diff += 1
+        n = next((i for i, (x, y) in enumerate(zip(r["text"], o["text"])) if x != y), min(len(r["text"]), len(o["text"])))
+        print("DIFF", k, "first difference at char", n, "|", repr(r["text"][max(0, n - 20):n + 30]), "vs", repr(o["text"][max(0, n - 20):n + 30]))
+print("identical", same, "different", diff)
+sys.exit(1 if diff else 0)

tools/par/co_bt.sh

@@ -0,0 +1,26 @@
+#!/bin/bash
+# co_bt.sh ROOT OUT [N] - batch_test arms, dumps compared: A/A2 layer-major alone (noise), B layer-major co-read, C the usual chunk-major read
+R=$1; O=$2; N=${3:-4}; mkdir -p $O
+PY=/data/external/niko/t027/venv/bin/python
+export STRATA_ROUTE_TAIL_SKIP=0 STRATA_PREFILL_CPU_SHARE=0 STRATA_CRASH_BT=$O/crash.txt
+CFG=$R/res/co/cog_on-cfg.json
+for arm in A A2 B C; do
+  EXTRA=""; [ $arm = B ] && EXTRA="--co"
+  if [ $arm = C ]; then unset STRATA_PREFILL_LM_FORCE; else export STRATA_PREFILL_LM_FORCE=1; fi
+  STRATA_DUMP_FIRST_LOGITS=$O/lg-$arm BATCH_TEST_LOG=$O/engine-$arm.log $PY $R/h/tools/batch_test.py --exe $R/h/b86/strata --config $CFG --batch $N --n $N --max-new 96 --pad 900 --skip-solo --dump $O/dump-$arm.json $EXTRA > $O/bt-$arm.out 2>&1
+  echo "== $arm rc=$?"; grep -E "batch: [0-9]+ slots|co-read:" $O/bt-$arm.out $O/engine-$arm.log | cut -c1-240 | head -4
+done
+$PY - <<PY
+import json
+d={a:json.load(open("$O/dump-%s.json"%a))["batch"] for a in ("A","A2","B","C")}
+def cmp(x,y):
+    r=[]
+    for p,q in zip(d[x],d[y]):
+        r.append("same" if p==q else "DIFF@%d"%next(k for k in range(min(len(p),len(q))) if p[k]!=q[k]))
+    return r
+for x,y in (("A","A2"),("A","B"),("A","C"),("B","C")):
+    print(x,"vs",y,cmp(x,y))
+PY
+
+$PY $R/h/tools/par/kl_logits.py $O lg C A A2 B 2>&1 | head -8
+$PY $R/h/tools/par/kl_logits.py $O lg A A2 B 2>&1 | head -8

tools/par/e00.py

@@ -0,0 +1,31 @@
+#!/usr/bin/env python3
+"""e00.py OUT.json [CUDA_HOME] - machine capability JSON: nvidia-smi links, e00_probe.cu results, RAM, NUMA, CPU (read-only)."""
+import json, os, platform, subprocess, sys, tempfile
+
+out = sys.argv[1]
+ch = sys.argv[2] if len(sys.argv) > 2 else "/usr/local/cuda"
+here = os.path.dirname(os.path.abspath(__file__))
+w = tempfile.mkdtemp()
+
+
+def sh(c, t=120):
+    try:
+        return subprocess.run(c, shell=True, capture_output=True, text=True, timeout=t).stdout.strip()
+    except Exception as e:  # noqa: BLE001
+        return str(e)
+
+
+subprocess.run([f"{ch}/bin/nvcc", "-O2", "-std=c++17", f"{here}/e00_probe.cu", "-o", f"{w}/p", "-lpthread"], check=True)
+p = json.loads(sh(f"{w}/p", 300))
+p["host"] = platform.node()
+p["nvidia_smi_pcie"] = sh("nvidia-smi --query-gpu=index,name,pci.bus_id,pcie.link.gen.current,pcie.link.gen.max,pcie.link.width.current,pcie.link.width.max,power.limit,clocks.max.sm,display_active,memory.total,memory.used,driver_version --format=csv").splitlines()
+p["nvidia_topo"] = sh("nvidia-smi topo -m").splitlines()
+want = ("Model name", "Socket(s)", "Core(s) per socket", "Thread(s) per core", "NUMA node(s)", "L3 cache")
+p["lscpu"] = {l.split(":", 1)[0].strip(): l.split(":", 1)[1].strip() for l in sh("lscpu").splitlines() if ":" in l and l.split(":", 1)[0].strip() in want}
+p["meminfo_gb"] = {l.split(":")[0]: int(l.split()[1]) // 1048576 for l in open("/proc/meminfo") if l.split(":")[0] in ("MemTotal", "MemAvailable", "SwapTotal")}
+p["dimms"] = sh("sudo -n dmidecode -t memory 2>/dev/null | grep -E 'Speed:|Size: [0-9]' | sort | uniq -c").splitlines()
+p["numa"] = sh("numactl -H 2>/dev/null | head -8").splitlines()
+p["kernel"] = platform.release()
+p["note"] = "read-only probe; IOMMU/ACS/clocks untouched"
+json.dump(p, open(out, "w"), indent=1)
+print("wrote", out)

tools/par/e00_probe.cu

@@ -0,0 +1,94 @@
+// E00 capability probe: peer access, GPU<->GPU / H2D / D2H transfer latency+bandwidth at plan sizes, RAM bandwidth.
+// build: nvcc -O2 -std=c++17 e00_probe.cu -o e00_probe -lpthread ; run: ./e00_probe > e00.json
+#include <cuda_runtime.h>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <chrono>
+#include <thread>
+#include <vector>
+#include <algorithm>
+using Clk = std::chrono::steady_clock;
+static double now_us() { return std::chrono::duration<double, std::micro>(Clk::now().time_since_epoch()).count(); }
+#define CK(x) do { cudaError_t e = (x); if (e != cudaSuccess) { fprintf(stderr, "%s:%d %s\n", __FILE__, __LINE__, cudaGetErrorString(e)); exit(2);} } while (0)
+static const size_t SIZES[] = {4096, 16384, 65536, 262144, 1 << 20, 8 << 20, 64 << 20};
+
+template <class F> double med(F f, int reps) {
+    f(); f();
+    std::vector<double> v;
+    for (int i = 0; i < reps; ++i) { double t = now_us(); f(); v.push_back(now_us() - t); }
+    std::sort(v.begin(), v.end());
+    return v[v.size() / 2];
+}
+int main() {
+    int n = 0; CK(cudaGetDeviceCount(&n));
+    printf("{\n \"gpus\": [");
+    for (int d = 0; d < n; ++d) {
+        cudaDeviceProp p; CK(cudaGetDeviceProperties(&p, d));
+        size_t fr, tot; CK(cudaSetDevice(d)); CK(cudaMemGetInfo(&fr, &tot));
+        printf("%s{\"index\":%d,\"name\":\"%s\",\"sm\":%d,\"sms\":%d,\"free_mib\":%zu,\"total_mib\":%zu,\"pci_bus\":%d}", d ? "," : "", d, p.name, p.major * 10 + p.minor, p.multiProcessorCount, fr >> 20, tot >> 20, p.pciBusID);
+    }
+    printf("],\n \"peer\": [");
+    bool first = true;
+    for (int a = 0; a < n; ++a) for (int b = 0; b < n; ++b) if (a != b) {
+        int can = 0; cudaDeviceGetP2PAttribute(&can, cudaDevP2PAttrAccessSupported, a, b);
+        int nat = 0; cudaDeviceGetP2PAttribute(&nat, cudaDevP2PAttrNativeAtomicSupported, a, b);
+        printf("%s{\"from\":%d,\"to\":%d,\"access\":%d,\"native_atomics\":%d}", first ? "" : ",", a, b, can, nat); first = false;
+    }
+    printf("],\n \"transfers\": [");
+    first = true;
+    auto emit = [&](const char* kind, int a, int b, size_t sz, double us) {
+        double gbs = (double) sz / us / 1e3;
+        printf("%s\n  {\"kind\":\"%s\",\"a\":%d,\"b\":%d,\"bytes\":%zu,\"median_us\":%.2f,\"gb_s\":%.3f}", first ? "" : ",", kind, a, b, sz, us, gbs); first = false;
+    };
+    size_t maxs = 64 << 20;
+    void* h = nullptr; CK(cudaMallocHost(&h, maxs)); memset(h, 1, maxs);
+    std::vector<void*> dbuf(n), dbuf2(n);
+    std::vector<cudaStream_t> st(n);
+    for (int d = 0; d < n; ++d) { CK(cudaSetDevice(d)); CK(cudaMalloc(&dbuf[d], maxs)); CK(cudaMalloc(&dbuf2[d], maxs)); CK(cudaStreamCreateWithFlags(&st[d], cudaStreamNonBlocking)); }
+    for (int d = 0; d < n; ++d) {
+        CK(cudaSetDevice(d));
+        for (size_t sz : SIZES) {
+            emit("h2d_pinned", d, -1, sz, med([&] { CK(cudaMemcpyAsync(dbuf[d], h, sz, cudaMemcpyHostToDevice, st[d])); CK(cudaStreamSynchronize(st[d])); }, 30));
+            emit("d2h_pinned", d, -1, sz, med([&] { CK(cudaMemcpyAsync(h, dbuf[d], sz, cudaMemcpyDeviceToHost, st[d])); CK(cudaStreamSynchronize(st[d])); }, 30));
+        }
+        cudaStream_t s2; CK(cudaStreamCreateWithFlags(&s2, cudaStreamNonBlocking));
+        double us = med([&] { CK(cudaMemcpyAsync(dbuf[d], h, maxs / 2, cudaMemcpyHostToDevice, st[d])); CK(cudaMemcpyAsync((char*) h + maxs / 2, dbuf2[d], maxs / 2, cudaMemcpyDeviceToHost, s2)); CK(cudaStreamSynchronize(st[d])); CK(cudaStreamSynchronize(s2)); }, 10);
+        emit("h2d_d2h_bidir_total", d, -1, maxs, us);
+    }
+    for (int a = 0; a < n; ++a) for (int b = 0; b < n; ++b) if (a != b) {
+        int can = 0; cudaDeviceCanAccessPeer(&can, a, b);
+        if (can) { CK(cudaSetDevice(a)); cudaDeviceEnablePeerAccess(b, 0); cudaGetLastError(); }
+        CK(cudaSetDevice(a));
+        for (size_t sz : SIZES)
+            emit(can ? "p2p_peer_memcpy" : "gpu_to_gpu_memcpy_staged", a, b, sz, med([&] { CK(cudaMemcpyPeerAsync(dbuf2[b], b, dbuf[a], a, sz, st[a])); CK(cudaStreamSynchronize(st[a])); }, 30));
+    }
+    if (n >= 2) {
+        CK(cudaSetDevice(0)); cudaDeviceEnablePeerAccess(1, 0); cudaGetLastError(); CK(cudaSetDevice(1)); cudaDeviceEnablePeerAccess(0, 0); cudaGetLastError();
+        double us = med([&] { CK(cudaMemcpyPeerAsync(dbuf2[1], 1, dbuf[0], 0, maxs, st[0])); CK(cudaMemcpyPeerAsync(dbuf2[0], 0, dbuf[1], 1, maxs, st[1])); CK(cudaStreamSynchronize(st[0])); CK(cudaStreamSynchronize(st[1])); }, 10);
+        emit("p2p_bidir_total_bytes", 0, 1, 2 * maxs, us);
+        void* h2 = nullptr; CK(cudaMallocHost(&h2, maxs));
+        us = med([&] { CK(cudaMemcpyAsync(dbuf[0], h, maxs, cudaMemcpyHostToDevice, st[0])); CK(cudaMemcpyAsync(dbuf[1], h2, maxs, cudaMemcpyHostToDevice, st[1])); CK(cudaStreamSynchronize(st[0])); CK(cudaStreamSynchronize(st[1])); }, 10);
+        emit("h2d_both_gpus_total_bytes", 0, 1, 2 * maxs, us);
+    }
+    printf("\n ],\n");
+    size_t N = (512u << 20) / sizeof(double);
+    std::vector<double> A(N, 1.0), Bv(N, 2.0), C(N, 0.0);
+    printf(" \"ram_triad_gb_s\": {");
+    unsigned hw = std::thread::hardware_concurrency();
+    bool f2 = true;
+    for (unsigned t : {1u, 2u, 4u, 8u, 16u, hw}) {
+        if (t > hw) continue;
+        double best = 1e30;
+        for (int r = 0; r < 3; ++r) {
+            double t0 = now_us();
+            std::vector<std::thread> th;
+            for (unsigned i = 0; i < t; ++i) th.emplace_back([&, i] { size_t lo = N * i / t, hi = N * (i + 1) / t; for (size_t k = lo; k < hi; ++k) C[k] = A[k] + 3.0 * Bv[k]; });
+            for (auto& x : th) x.join();
+            best = std::min(best, now_us() - t0);
+        }
+        printf("%s\"%u\": %.1f", f2 ? "" : ", ", t, 3.0 * N * sizeof(double) / best / 1e3); f2 = false;
+    }
+    printf("},\n \"hw_threads\": %u\n}\n", hw);
+    return 0;
+}

tools/par/host_seq.sh

@@ -0,0 +1,18 @@
+#!/bin/bash
+# host_seq.sh ROOT OUT - 4 prompts ONE AT A TIME (--seq) in the engine: A, A2 (stack off), H (host stack only), same fixed flags, no MTP
+R=$1; O=$2; mkdir -p $O
+PY=/data/external/niko/t027/venv/bin/python
+export STRATA_ROUTE_TAIL_SKIP=0 STRATA_PREFILL_CPU_SHARE=0 STRATA_CRASH_BT=$O/crash.txt
+CFG=$R/res/co/cog_on-cfg.json
+HOST="STRATA_CKPT_POOL=1 STRATA_CKPT_POOL_READY=4 STRATA_CKPT_THREADS=1 STRATA_SLOT_CKPT_MOVE=1 STRATA_SLOT_COPY_DIRECT=1 STRATA_SLOT_COPY_BATCHSYNC=1"
+for arm in A A2 H; do
+  EXTRA=""; [ $arm = H ] && EXTRA="$HOST"
+  env $EXTRA BATCH_TEST_LOG=$O/engine-$arm.log $PY $R/h/tools/batch_test.py --exe $R/h/b86/strata --config $CFG --batch 4 --n 4 --max-new 64 --pad 900 --seq --skip-solo --dump $O/dump-$arm.json > $O/bt-$arm.out 2>&1
+  echo "== $arm rc=$?"; grep -E "batch: [0-9]+ slots" $O/bt-$arm.out | cut -c1-200
+done
+$PY - <<PY
+import json
+d={a:json.load(open("$O/dump-%s.json"%a))["batch"] for a in ("A","A2","H")}
+for x,y in (("A","A2"),("A","H")):
+    print(x,"vs",y,["same" if p==q else "DIFF@%d"%next(k for k in range(min(len(p),len(q))) if p[k]!=q[k]) for p,q in zip(d[x],d[y])])
+PY

tools/par/kl_logits.py

@@ -0,0 +1,20 @@
+#!/usr/bin/env python3
+"""kl_logits.py DIR PREFIX REF ARM... - first-window logits dumps (STRATA_DUMP_FIRST_LOGITS=DIR/PREFIX-<arm>, files .0 .1 ...):
+KL(REF || ARM) per request, mean/median/max, and top-1 agreement (the klcmp.py of phaseA-tests/026q, for dumps of one arm per label)."""
+import sys
+from pathlib import Path
+import numpy as np
+d = Path(sys.argv[1]); pre = sys.argv[2]; ref = sys.argv[3]; arms = sys.argv[4:]
+def ls(x):
+    x = x.astype(np.float64); m = x.max(); return x - m - np.log(np.exp(x - m).sum())
+for arm in arms:
+    kls, top, n = [], 0, 0
+    i = 0
+    while (d / f"{pre}-{ref}.{i}").exists() and (d / f"{pre}-{arm}.{i}").exists():
+        a = np.fromfile(d / f"{pre}-{ref}.{i}", "<f4"); b = np.fromfile(d / f"{pre}-{arm}.{i}", "<f4")
+        la, lb = ls(a), ls(b)
+        kls.append(float((np.exp(la) * (la - lb)).sum())); top += int(a.argmax() == b.argmax()); n += 1; i += 1
+    if not n:
+        print(arm, "no dumps"); continue
+    k = np.array(kls)
+    print(f"{arm:4s} vs {ref}: requests {n}  mean KL {k.mean():.6f}  median {np.median(k):.6f}  max {k.max():.6f}  top-1 same {top}/{n}")

tools/par/my_load.py

@@ -0,0 +1,6 @@
+#!/usr/bin/env python3
+"""my_load.py URL OUT.json - par_load.py (this harness) with PAR-A armrun's workload (W2 1k prompt, 128 out, c=8,16, 2 rounds), as armrun --loadcmd."""
+import os, sys
+sys.path.insert(0, os.path.join(os.path.dirname(os.path.abspath(__file__)), ".."))
+import par_load
+sys.exit(par_load.main([sys.argv[1], "--workload", "w2", "--conc", "8,16", "--rounds", "2", "--max-tokens", "128", "--warm-requests", "2", "--out", sys.argv[2]]))

tools/par/overhead.sh

@@ -0,0 +1,15 @@
+#!/bin/bash
+# overhead.sh ROOT OUT [REPS] - counters overhead: arms M (main binary), H0 (branch binary, counters off), H1 (branch, STRATA_PAR_COUNTERS=1), interleaved
+R=$1; O=$2; REPS=${3:-3}; mkdir -p $O
+PY=python3
+for rep in $(seq 1 $REPS); do
+ for arm in M H0 H1; do
+  case $arm in
+   M)  EXE=$R/main/b86/strata; REPO=$R/main; ENV="";;
+   H0) EXE=$R/h/b86/strata;    REPO=$R/h;    ENV="";;
+   H1) EXE=$R/h/b86/strata;    REPO=$R/h;    ENV="STRATA_PAR_COUNTERS=1";;
+  esac
+  $PY $R/h/tools/par_run.py --box boston --tag ov${rep}_$arm --exe $EXE --repo $REPO --gpus 0,1 --parallel 8 --env "$ENV" --work "w1;w2:4@--max-tokens 128 --rounds 3;w2:8@--max-tokens 128 --rounds 3" --rounds 3 --out $O > $O/ov${rep}_$arm.out 2>&1
+  echo "rep $rep $arm: $(grep -E '^(w1|w2)' $O/ov${rep}_$arm.out | awk '{print $2, $3, $4, $5, $6, $7, $8, $9}' | tr '\n' '|')"
+ done
+done

tools/par/recon.py

@@ -0,0 +1,62 @@
+#!/usr/bin/env python3
+"""recon.py - reconcile PAR-A's boston numbers with this harness (runs on boston, one arm after the other)."""
+import json, os, subprocess, sys, time
+PA = "/data/external/niko/parA"
+I = "/data/external/niko/int029"
+R = "/data/external/niko/par29"
+W = R + "/armw"
+OUT = R + "/res/recon"
+PY = "/data/external/niko/t027/venv/bin/python"
+os.makedirs(W, exist_ok=True)
+for name, tgt in (("b5", PA + "/b5"), ("src5", PA + "/src5"), ("bint", I + "/b-int"), ("srcint", I + "/int")):
+    if not os.path.lexists(W + "/" + name):
+        os.symlink(tgt, W + "/" + name)
+REPL = '[{"gpus":[0]},{"gpus":[1]}]'
+PA_ENV = ["STRATA_QSA_SLOTS=1", "STRATA_BATCH_SELF_COMMIT=1", "STRATA_SLOT_COPY_DIRECT=1", "STRATA_CKPT_POOL=1", "STRATA_CKPT_POOL_READY=4",
+          "STRATA_SLOT_CKPT_MOVE=1", "STRATA_SLOT_COPY_BATCHSYNC=1", "STRATA_REPLICA_SHARED_ARENA=1"]
+
+
+def armrun(tag, build, src, loadcmd=None):
+    out = f"{OUT}/{tag}"
+    cmd = [PY, PA + "/armrun.py", "--out", out, "--gpus", "0,1", "--replicas", REPL, "--parallel", "8", "--build", build, "--src", src,
+           "--conc", "8,16", "--rounds", "2", "--max-tokens", "128", "--port", "8191"]
+    for e in PA_ENV:
+        cmd += ["--env", e]
+    if loadcmd:
+        cmd += ["--loadcmd", loadcmd]
+    env = dict(os.environ, PARA_W=W)
+    t = time.time()
+    subprocess.run(cmd, env=env, capture_output=True, text=True)
+    print(f"== {tag} ({time.time() - t:.0f}s)", flush=True)
+    for f in ("load.txt",):
+        p = f"{out}/{f}"
+        if os.path.exists(p):
+            print("".join(l for l in open(p).read().splitlines(True) if l.strip()[:2] in ("8 ", "16") or "tok/s" in l and l.startswith("c=")), flush=True)
+            if loadcmd:
+                print(open(p).read()[-700:])
+
+
+def mine(tag, env, post=None, exe=None, repo=None, rounds=2):
+    out = f"{OUT}/{tag}"
+    cmd = [PY, R + "/h/tools/par_run.py", "--box", "boston", "--tag", tag, "--exe", exe or I + "/b-int/strata", "--repo", repo or I + "/int",
+           "--gpus", "0,1", "--parallel", "8", "--server-extra", ("--replicas 2" if "A9" not in tag else ""), "--env", ";".join(env), "--rounds", "2",
+           "--work", f"w2:8,16@--nonce rc --max-tokens 128 --rounds {rounds} --warm-requests 2", "--out", out]
+    if post:
+        cmd += ["--post-cmd", post]
+    t = time.time()
+    r = subprocess.run(cmd, capture_output=True, text=True)
+    print(f"== {tag} ({time.time() - t:.0f}s)\n" + "\n".join(l for l in r.stdout.splitlines() if l.startswith(("w2", "  ", "c=", "DONE", "ABORT", " 8", " 16"))), flush=True)
+
+
+THEIR_LOAD = f"{PY} {PA}/par_load.py {{url}} --conc 8,16 --rounds 2 --prompt-tokens 1000 --max-tokens 128 --out /dev/null"
+MY_STACK = PA_ENV + ["STRATA_BATCH_GDN_GROUPS=1", "STRATA_BATCH_ATTN_ROWS=1"]
+if "A7" in sys.argv: mine("A7_myarm_burst_rounds1", PA_ENV)
+if "A8" in sys.argv: mine("A8_myarm_sustained_rounds6", PA_ENV, rounds=6)
+if "A9" in sys.argv: mine("A9_layersplit_b1like_rounds1", ["STRATA_QSA_SLOTS=1"] , rounds=1, exe=I + "/b-base/strata", repo=I + "/base")
+which = [x for x in sys.argv[1:] if x not in ("A7", "A8", "A9")] if len(sys.argv) > 1 else ["A1", "A2", "A3", "A4", "A5", "A6"]
+if "A1" in which: armrun("A1_pa_exact_theirload", "b5", "src5")
+if "A2" in which: armrun("A2_pa_exact_myload", "b5", "src5", R + "/h/tools/par/my_load.py")
+if "A3" in which: armrun("A3_pa_cfg_intbuild_theirload", "bint", "srcint")
+if "A4" in which: armrun("A4_pa_cfg_intbuild_myload", "bint", "srcint", R + "/h/tools/par/my_load.py")
+if "A5" in which: mine("A5_myarm_theirload", MY_STACK, THEIR_LOAD)
+if "A6" in which: mine("A6_myarm_paenv_theirload", PA_ENV, THEIR_LOAD)

tools/par/recon3.py

@@ -0,0 +1,63 @@
+#!/usr/bin/env python3
+"""recon3.py BOX ARMS... - sustained + burst for the honest headline.  Runs PAR-A's armrun.py with this harness's both.py as the loader.
+BOX boston: W=armw (symlinks to the int029 / base builds); BOX r9700: W=/data/par-a (b4 / src4)."""
+import os, subprocess, sys, time
+box = sys.argv[1]
+arms = sys.argv[2:]
+HOST = ["STRATA_QSA_SLOTS=1", "STRATA_BATCH_SELF_COMMIT=1", "STRATA_SLOT_COPY_DIRECT=1", "STRATA_CKPT_POOL=1", "STRATA_CKPT_POOL_READY=4",
+        "STRATA_SLOT_CKPT_MOVE=1", "STRATA_SLOT_COPY_BATCHSYNC=1"]
+CUDA = ["STRATA_BATCH_GDN_GROUPS=1", "STRATA_BATCH_ATTN_ROWS=1"]
+if box == "boston":
+    R = "/data/external/niko/par29"; PA = "/data/external/niko/parA"; I = "/data/external/niko/int029"; W = R + "/armw"
+    PY = "/data/external/niko/t027/venv/bin/python"
+    for name, tgt in (("bbase", I + "/b-base"), ("srcbase", I + "/base"), ("bint", I + "/b-int"), ("srcint", I + "/int")):
+        if not os.path.lexists(W + "/" + name):
+            os.symlink(tgt, W + "/" + name)
+    G2 = ["--gpus", "0,1"]
+    REP = '[{"gpus":[0]},{"gpus":[1]}]'
+    ARMS = {
+        "B1": (["--parallel", "8"], [], "bbase", "srcbase"),
+        "LS": (["--parallel", "8"], HOST + CUDA, "bint", "srcint"),
+        "REP": (["--parallel", "8", "--replicas", REP], HOST + CUDA + ["STRATA_REPLICA_SHARED_ARENA=1", "STRATA_COALESCE_PREFILL=3", "STRATA_COALESCE_WAIT_MS=40"], "bint", "srcint"),
+        "REPCR": (["--parallel", "8", "--replicas", REP, "--arg=--no-prefill-borrow"], HOST + CUDA + ["STRATA_REPLICA_SHARED_ARENA=1", "STRATA_COALESCE_PREFILL=3", "STRATA_COALESCE_WAIT_MS=40", "STRATA_BATCH_CONCURRENT_READ=1"], "bint", "srcint"),
+        "REPCR_NOCO": (["--parallel", "8", "--replicas", REP, "--arg=--no-prefill-borrow"], HOST + CUDA + ["STRATA_REPLICA_SHARED_ARENA=1", "STRATA_BATCH_CONCURRENT_READ=1"], "bint", "srcint"),
+        "REP_NOCO": (["--parallel", "8", "--replicas", REP], HOST + CUDA + ["STRATA_REPLICA_SHARED_ARENA=1"], "bint", "srcint"),
+    }
+    GP = G2; CONC = "8,16,32"
+else:
+    R = "/data/par29h"; PA = "/data/par-a"; W = R + "/armw"; PY = "/data/p4/venv/bin/python"
+    os.makedirs(W, exist_ok=True)
+    for name, tgt in (("b4", PA + "/b4"), ("src4", PA + "/src4"), ("bint", R + "/int/b-int"), ("srcint", R + "/int/src")):
+        if not os.path.lexists(W + "/" + name):
+            os.symlink(tgt, W + "/" + name)
+    REP2 = '[{"gpus":[0]},{"gpus":[1]}]'
+    REP4 = '[{"gpus":[0]},{"gpus":[1]},{"gpus":[2]},{"gpus":[3]}]'
+    CO = ["STRATA_REPLICA_SHARED_ARENA=1", "STRATA_LM_HIP=1", "STRATA_COALESCE_PREFILL=3", "STRATA_COALESCE_WAIT_MS=40"]
+    ARMS = {
+        "B1p8_2c": (["--gpus", "0,1", "--parallel", "8"], [], "b4", "src4"),
+        "B1p16_2c": (["--gpus", "0,1", "--parallel", "16"], [], "b4", "src4"),
+        "REP2": (["--gpus", "0,1", "--parallel", "8", "--replicas", REP2], HOST + CO, "b4", "src4"),
+        "B1_4c": (["--gpus", "0,1,2,3", "--profile4"], [], "b4", "src4"),
+        "REP2CR": (["--gpus", "0,1", "--parallel", "8", "--replicas", REP2, "--arg=--no-prefill-borrow"], HOST + CO + ["STRATA_BATCH_CONCURRENT_READ=1"], "bint", "srcint"),
+        "REP4CR": (["--gpus", "0,1,2,3", "--parallel", "4", "--replicas", REP4, "--arg=--no-prefill-borrow"], HOST + CO + ["STRATA_BATCH_CONCURRENT_READ=1"], "bint", "srcint"),
+        "REP2I": (["--gpus", "0,1", "--parallel", "8", "--replicas", REP2], HOST + CO, "bint", "srcint"),
+        "REP4": (["--gpus", "0,1,2,3", "--parallel", "4", "--replicas", REP4], HOST + CO, "b4", "src4"),
+    }
+    GP = None; CONC = "8,16,32"
+OUT = R + "/res/recon3"
+for tag in arms:
+    flags, env, build, src = ARMS[tag][0], ARMS[tag][1], ARMS[tag][2], ARMS[tag][3]
+    out = f"{OUT}/{box}_{tag}"
+    cmd = [PY, PA + "/armrun.py", "--out", out, "--build", build, "--src", src, "--conc", CONC, "--rounds", "2", "--max-tokens", "128",
+           "--port", "8197", "--loadcmd", R + "/h/tools/par/both.py"] + (GP or []) + flags
+    for e in env:
+        cmd += ["--env", e]
+    e2 = dict(os.environ, PARA_W=W, PA_LOAD=PA + "/par_load.py", BOTH_CONC=CONC)
+    t = time.time()
+    subprocess.run(cmd, env=e2, capture_output=True, text=True)
+    print(f"=== {box} {tag} ({time.time() - t:.0f}s)", flush=True)
+    p = out + "/load.txt"
+    if os.path.exists(p):
+        for l in open(p).read().splitlines():
+            if l.startswith(("###", "w2 c=")) or l[:3].strip().isdigit() and "tok/s" not in l and len(l.split()) > 5:
+                print(l[:170], flush=True)

tools/par/sweep.py

@@ -0,0 +1,47 @@
+#!/usr/bin/env python3
+"""sweep.py - run par_run.py over a list of variants (one server start each) and print a one-line summary per variant.
+
+variants file, one per line, '#' comments:   tag | gpus | parallel | groups | env (K=V;K=V) | engine/server extras | work
+  s8g1 | 0,1 | 8 | 1 | STRATA_BATCH_MTP=1 | | w2:4,8@--max-tokens 128
+Empty fields = unset.  Run on the box:  python3 sweep.py --box boston --root ~/par29 --exe ~/par29/main/b86/strata --out RES variants.txt
+"""
+import argparse, json, os, subprocess, sys
+
+ap = argparse.ArgumentParser()
+ap.add_argument("variants")
+ap.add_argument("--box", required=True)
+ap.add_argument("--root", required=True)
+ap.add_argument("--exe", required=True)
+ap.add_argument("--repo", default=None)
+ap.add_argument("--out", required=True)
+ap.add_argument("--rounds", default="2")
+ap.add_argument("--default-work", default="w1;w2:4,8@--max-tokens 128")
+ap.add_argument("--exe-base", default=None)
+a = ap.parse_args()
+repo = a.repo or a.root + "/main"
+here = os.path.dirname(os.path.abspath(__file__))
+for line in open(a.variants):
+    line = line.strip()
+    if not line or line.startswith("#"):
+        continue
+    f = [x.strip() for x in line.split("|")] + [""] * 8
+    tag, gpus, par, groups, env, extra, work, srv = f[:8]
+    cmd = [sys.executable, f"{here}/../par_run.py", "--box", a.box, "--tag", tag, "--exe", a.exe, "--repo", repo, "--gpus", gpus or "0,1",
+           "--out", a.out, "--rounds", a.rounds, "--work", work or a.default_work]
+    if "nomtp" in tag:
+        cmd += ["--no-mtp"]
+    if par:
+        cmd += ["--parallel", par]
+    if groups:
+        cmd += ["--batch-groups" if False else "--groups", groups]
+    if env:
+        cmd += ["--env", env]
+    if extra:
+        cmd += ["--extra", extra]
+    if srv:
+        cmd += ["--server-extra", srv]
+    r = subprocess.run(cmd, capture_output=True, text=True)
+    out = [l for l in r.stdout.splitlines() if l.startswith(("w1 ", "w2 ", "w4 ", "w6 ", "SERVER", "ABORT", "DONE", "[", "par_load"))]
+    print(f"## {tag}: rc={r.returncode}", flush=True)
+    for l in out:
+        print("   " + l[:200], flush=True)

tools/par/variants/b_best.txt

@@ -0,0 +1,3 @@
+b_ck    | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1 | | w1;w2:4,8@--nonce bw --max-tokens 256 --rounds 3;w4:4@--rounds 1 --max-tokens 128
+b_co    | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=60 | | w1;w2:4,8@--nonce bw --max-tokens 256 --rounds 3;w4:4@--rounds 1 --max-tokens 128
+b_coc   | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=60;STRATA_PIPED_CONCURRENT=1;STRATA_PIPED_CONCURRENT_MIN=4096 | | w1;w2:4,8@--nonce bw --max-tokens 256 --rounds 3;w4:4@--rounds 1 --max-tokens 128

tools/par/variants/cnt_2g.txt

@@ -0,0 +1,4 @@
+cnt_c1  | 0,1 | 4 |   | STRATA_PAR_COUNTERS=1;STRATA_P4_TIMING=1 | | w2:1@--max-tokens 128
+cnt_c4  | 0,1 | 4 |   | STRATA_PAR_COUNTERS=1;STRATA_P4_TIMING=1 | | w2:4@--max-tokens 128
+cnt_c8  | 0,1 | 8 |   | STRATA_PAR_COUNTERS=1;STRATA_P4_TIMING=1 | | w2:8@--max-tokens 128
+cnt_s4  | 0,1 | 4 | 1 | STRATA_PAR_COUNTERS=1 | | w2:4@--max-tokens 128

tools/par/variants/co_id.txt

@@ -0,0 +1,2 @@
+co_ref | 0,1 | 8 |   |  | | w2:4@--nonce idt --keep-text --max-tokens 48 --rounds 1;w2:8@--nonce idt --keep-text --max-tokens 48 --rounds 1
+co_on  | 0,1 | 8 |   | STRATA_COALESCE_PREFILL=8;STRATA_CRASH_BT=/data/external/niko/par29/res/co/crash.txt | | w2:4@--nonce idt --keep-text --max-tokens 48 --rounds 1;w2:8@--nonce idt --keep-text --max-tokens 48 --rounds 1

tools/par/variants/co_idz.txt

@@ -0,0 +1,2 @@
+coz_ref | 0,1 | 8 |   | STRATA_ROUTE_TAIL_SKIP=0;STRATA_PREFILL_CPU_SHARE=0 | --vram-reserve-later-mib 1536 --adapt-every 100000 --pcie-frac 0 --suffix-draft 0 --expert-cache 1500 | w2:4@--nonce idz --keep-text --max-tokens 48 --rounds 1
+coz_on  | 0,1 | 8 |   | STRATA_ROUTE_TAIL_SKIP=0;STRATA_PREFILL_CPU_SHARE=0;STRATA_COALESCE_PREFILL=8;STRATA_CRASH_BT=/data/par29/res/co/crash.txt | --vram-reserve-later-mib 1536 --adapt-every 100000 --pcie-frac 0 --suffix-draft 0 --expert-cache 1500 | w2:4@--nonce idz --keep-text --max-tokens 48 --rounds 1

tools/par/variants/co_lm.txt

@@ -0,0 +1,2 @@
+co_lref | 0,1 | 8 |   | STRATA_PREFILL_LM_FORCE=1 | | w2:4@--nonce idl --keep-text --max-tokens 48 --rounds 2;w2:8@--nonce idl --keep-text --max-tokens 48 --rounds 2
+co_lon  | 0,1 | 8 |   | STRATA_PREFILL_LM_FORCE=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=40;STRATA_CRASH_BT=/data/external/niko/par29/res/co/crash.txt | | w2:4@--nonce idl --keep-text --max-tokens 48 --rounds 2;w2:8@--nonce idl --keep-text --max-tokens 48 --rounds 2

tools/par/variants/co_lm3.txt

@@ -0,0 +1,3 @@
+cog_ref  | 0,1 | 8 |   | STRATA_ROUTE_TAIL_SKIP=0;STRATA_PREFILL_CPU_SHARE=0;STRATA_PREFILL_LM_FORCE=1 | --expert-cache 1500 --adapt-every 100000 --pcie-frac 0 --suffix-draft 0 | w2:4@--nonce idg --keep-text --max-tokens 48 --rounds 2;w2:8@--nonce idg --keep-text --max-tokens 48 --rounds 2
+cog_ref2 | 0,1 | 8 |   | STRATA_ROUTE_TAIL_SKIP=0;STRATA_PREFILL_CPU_SHARE=0;STRATA_PREFILL_LM_FORCE=1 | --expert-cache 1500 --adapt-every 100000 --pcie-frac 0 --suffix-draft 0 | w2:4@--nonce idg --keep-text --max-tokens 48 --rounds 2;w2:8@--nonce idg --keep-text --max-tokens 48 --rounds 2
+cog_on   | 0,1 | 8 |   | STRATA_ROUTE_TAIL_SKIP=0;STRATA_PREFILL_CPU_SHARE=0;STRATA_PREFILL_LM_FORCE=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=40;STRATA_CRASH_BT=/data/external/niko/par29/res/co/crash.txt | --expert-cache 1500 --adapt-every 100000 --pcie-frac 0 --suffix-draft 0 | w2:4@--nonce idg --keep-text --max-tokens 48 --rounds 2;w2:8@--nonce idg --keep-text --max-tokens 48 --rounds 2

tools/par/variants/co_perf.txt

@@ -0,0 +1,2 @@
+co_off | 0,1 | 8 |   |  | | w2:4@--nonce pf --keep-text --max-tokens 128 --rounds 1;w2:8@--nonce pf --keep-text --max-tokens 128 --rounds 1
+co_on  | 0,1 | 8 |   | STRATA_COALESCE_PREFILL=8 | | w2:4@--nonce pf --keep-text --max-tokens 128 --rounds 1;w2:8@--nonce pf --keep-text --max-tokens 128 --rounds 1

tools/par/variants/co_perf2.txt

@@ -0,0 +1,2 @@
+co_off | 0,1 | 8 |   |  | --vram-reserve-later-mib 1536 | w2:4@--nonce pf --keep-text --max-tokens 128 --rounds 1;w2:8@--nonce pf --keep-text --max-tokens 128 --rounds 1
+co_on  | 0,1 | 8 |   | STRATA_COALESCE_PREFILL=8;STRATA_CRASH_BT=/data/par29/res/co/crash.txt | --vram-reserve-later-mib 1536 | w2:4@--nonce pf --keep-text --max-tokens 128 --rounds 1;w2:8@--nonce pf --keep-text --max-tokens 128 --rounds 1

tools/par/variants/co_perf3.txt

@@ -0,0 +1,3 @@
+co_off | 0,1 | 8 |   |  | --vram-reserve-later-mib 1536 | w2:4@--nonce pf --max-tokens 128 --rounds 2;w2:8@--nonce pf --max-tokens 128 --rounds 2
+co_on  | 0,1 | 8 |   | STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=40;STRATA_CRASH_BT=/data/par29/res/co/crash.txt | --vram-reserve-later-mib 1536 | w2:4@--nonce pf --max-tokens 128 --rounds 2;w2:8@--nonce pf --max-tokens 128 --rounds 2
+co_onck | 0,1 | 8 |   | STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=40;STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1 | --vram-reserve-later-mib 1536 | w2:4@--nonce pf --max-tokens 128 --rounds 2;w2:8@--nonce pf --max-tokens 128 --rounds 2

tools/par/variants/co_perf4.txt

@@ -0,0 +1,3 @@
+co_w150 | 0,1 | 8 |   | STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=150;STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1 | --vram-reserve-later-mib 1536 | w2:4@--nonce pf --max-tokens 128 --rounds 2;w2:8@--nonce pf --max-tokens 128 --rounds 2
+co_p16  | 0,1 | 16 | 2 | STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=150;STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1 | --vram-reserve-later-mib 1536 | w2:4@--nonce pf --max-tokens 128 --rounds 2;w2:8@--nonce pf --max-tokens 128 --rounds 2
+ck_only | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1 | --vram-reserve-later-mib 1536 | w2:4@--nonce pf --max-tokens 128 --rounds 2;w2:8@--nonce pf --max-tokens 128 --rounds 2

tools/par/variants/co_perf5.txt

@@ -0,0 +1,3 @@
+off25  | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1 | --vram-reserve-mib 2500 | w2:4@--nonce pf --max-tokens 128 --rounds 2;w2:8@--nonce pf --max-tokens 128 --rounds 2
+on25   | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=100;STRATA_CRASH_BT=/data/par29/res/co/crash.txt | --vram-reserve-mib 2500 | w2:4@--nonce pf --max-tokens 128 --rounds 2;w2:8@--nonce pf --max-tokens 128 --rounds 2
+on40   | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=100;STRATA_CRASH_BT=/data/par29/res/co/crash.txt | --vram-reserve-mib 4000 | w2:4@--nonce pf --max-tokens 128 --rounds 2;w2:8@--nonce pf --max-tokens 128 --rounds 2

tools/par/variants/co_t2.txt

@@ -0,0 +1 @@
+co_t4 | 0,1 | 8 |   | STRATA_COALESCE_PREFILL=8;STRATA_CRASH_BT=/data/external/niko/par29/res/co/crash.txt;STRATA_TRACE=1 | | w2:4@--nonce idt --keep-text --max-tokens 16 --rounds 1

tools/par/variants/p1_2g_a.txt

@@ -0,0 +1,7 @@
+# tag | gpus | parallel | groups | env | extras | work
+ctl8   | 0,1 | 8  |   |  |  | w1;w2:4,8@--max-tokens 128;w4:4@--rounds 1 --max-tokens 128
+ctl4   | 0,1 | 4  |   |  |  | w2:4,8@--max-tokens 128
+g1_8   | 0,1 | 8  | 1 |  |  | w2:4,8@--max-tokens 128
+g1_4   | 0,1 | 4  | 1 |  |  | w2:4,8@--max-tokens 128
+p16g2  | 0,1 | 16 | 2 |  |  | w2:4,8@--max-tokens 128
+mtp_g1_4 | 0,1 | 4 | 1 | STRATA_BATCH_MTP=1 | | w2:4,8@--max-tokens 128

tools/par/variants/p1_2g_b.txt

@@ -0,0 +1,6 @@
+pmtp1_8 | 0,1 | 8 | 4 | STRATA_BATCH_MTP=1;STRATA_PIPED_MTP=1 | | w2:4,8@--max-tokens 128
+pmtp2_8 | 0,1 | 8 | 4 | STRATA_BATCH_MTP=1;STRATA_PIPED_MTP=2 | | w2:4,8@--max-tokens 128
+pmtp3_4 | 0,1 | 4 | 2 | STRATA_BATCH_MTP=1;STRATA_PIPED_MTP=3 | | w2:4,8@--max-tokens 128
+conc8  | 0,1 | 8 |   | STRATA_PIPED_CONCURRENT=1 | | w2:4,8@--max-tokens 128;w4:4@--rounds 1 --max-tokens 128
+piece8 | 0,1 | 8 |   | STRATA_PIPED_READ_PIECE=4096;STRATA_PIPED_DECODE_MIN_MS=50 | | w2:4,8@--max-tokens 128;w4:4@--rounds 1 --max-tokens 128
+adapt8 | 0,1 | 8 |   | STRATA_PIPE_ADAPT_EVERY=2000 | | w2:4,8@--max-tokens 128

tools/par/variants/p1_2g_b1.txt

@@ -0,0 +1,3 @@
+pmtp1_8 | 0,1 | 8 | 4 | STRATA_BATCH_MTP=1;STRATA_PIPED_MTP=1 | | w2:4,8@--max-tokens 128
+pmtp2_8 | 0,1 | 8 | 4 | STRATA_BATCH_MTP=1;STRATA_PIPED_MTP=2 | | w2:4,8@--max-tokens 128
+pmtp3_4 | 0,1 | 4 | 2 | STRATA_BATCH_MTP=1;STRATA_PIPED_MTP=3 | | w2:4,8@--max-tokens 128

tools/par/variants/p1_2g_b2.txt

@@ -0,0 +1,3 @@
+conc8  | 0,1 | 8 |   | STRATA_PIPED_CONCURRENT=1 | | w2:4,8@--max-tokens 128;w4:4@--rounds 1 --max-tokens 128
+piece8 | 0,1 | 8 |   | STRATA_PIPED_READ_PIECE=4096;STRATA_PIPED_DECODE_MIN_MS=50 | | w2:4,8@--max-tokens 128;w4:4@--rounds 1 --max-tokens 128
+adapt8 | 0,1 | 8 |   | STRATA_PIPE_ADAPT_EVERY=2000 | | w2:4,8@--max-tokens 128

tools/par/variants/p2_adm.txt

@@ -0,0 +1,3 @@
+adm_sw  | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1 | | w2:4,8@--max-tokens 128
+adm_nb  | 0,1 | 8 |   |  | --no-prefill-borrow | w2:4,8@--max-tokens 128
+adm_all | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_PIPED_CONCURRENT=1;STRATA_PIPED_CONCURRENT_MIN=4096 | --no-prefill-borrow | w2:4,8@--max-tokens 128;w4:4@--rounds 1 --max-tokens 128

tools/par/variants/p2_conc.txt

@@ -0,0 +1,4 @@
+a_def   | 0,1 | 8 |   |  | | w2:4,8@--max-tokens 128;w4:4@--rounds 1 --max-tokens 128
+b_conc  | 0,1 | 8 |   | STRATA_PIPED_CONCURRENT=1 | | w2:4,8@--max-tokens 128;w4:4@--rounds 1 --max-tokens 128
+c_min4k | 0,1 | 8 |   | STRATA_PIPED_CONCURRENT=1;STRATA_PIPED_CONCURRENT_MIN=4096 | | w2:4,8@--max-tokens 128;w4:4@--rounds 1 --max-tokens 128
+d_min12k | 0,1 | 8 |   | STRATA_PIPED_CONCURRENT=1;STRATA_PIPED_CONCURRENT_MIN=12000 | | w2:4,8@--max-tokens 128;w4:4@--rounds 1 --max-tokens 128

tools/par/variants/p2_pp.txt

@@ -0,0 +1,4 @@
+pp0    | 0,1 | 8 |   |  | | w2:4,8@--max-tokens 128;w3:1,4@--rounds 1
+pp1    | 0,1 | 8 |   |  | --prefill-pipe 1 | w2:4,8@--max-tokens 128;w3:1,4@--rounds 1
+pp1k   | 0,1 | 8 |   |  | --prefill-pipe 1024 | w2:4,8@--max-tokens 128;w3:1,4@--rounds 1
+pp4k   | 0,1 | 8 |   |  | --prefill-pipe 4096 | w2:4,8@--max-tokens 128;w3:1,4@--rounds 1

tools/par/variants/p2z_adm.txt

@@ -0,0 +1,5 @@
+a_def   | 0,1 | 8 |   |  | | w2:4@--max-tokens 128;w4:4@--rounds 1 --max-tokens 128
+adm_sw  | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1 | | w2:4@--max-tokens 128
+adm_nb  | 0,1 | 8 |   |  | --no-prefill-borrow | w2:4@--max-tokens 128
+c_min4k | 0,1 | 8 |   | STRATA_PIPED_CONCURRENT=1;STRATA_PIPED_CONCURRENT_MIN=4096 | | w2:4@--max-tokens 128;w4:4@--rounds 1 --max-tokens 128
+pp1k    | 0,1 | 8 |   |  | --prefill-pipe 1024 | w2:4@--max-tokens 128

tools/par/variants/stack_b1.txt

@@ -0,0 +1 @@
+st_b1   | 0,1 | 8 |   |  | | w1;w2:4,8,16@--nonce st --max-tokens 256 --rounds 2;w2:16,8,4@--nonce st --max-tokens 128 --rounds 2

tools/par/variants/stack_id.txt

@@ -0,0 +1,4 @@
+id_nomtp_A  | 0,1 | 8 |   | STRATA_ROUTE_TAIL_SKIP=0;STRATA_PREFILL_CPU_SHARE=0 | --expert-cache 1500 --adapt-every 100000 --pcie-frac 0 --suffix-draft 0 | w2:4@--nonce idz --keep-text --max-tokens 48 --rounds 1
+id_nomtp_A2 | 0,1 | 8 |   | STRATA_ROUTE_TAIL_SKIP=0;STRATA_PREFILL_CPU_SHARE=0 | --expert-cache 1500 --adapt-every 100000 --pcie-frac 0 --suffix-draft 0 | w2:4@--nonce idz --keep-text --max-tokens 48 --rounds 1
+id_nomtp_S0 | 0,1 | 8 |   | STRATA_ROUTE_TAIL_SKIP=0;STRATA_PREFILL_CPU_SHARE=0;STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_BATCH_GDN_GROUPS=1;STRATA_BATCH_ATTN_ROWS=1;STRATA_BATCH_SELF_COMMIT=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_MIN_REQS=3;STRATA_COALESCE_WAIT_MS=60 | --expert-cache 1500 --adapt-every 100000 --pcie-frac 0 --suffix-draft 0 | w2:4@--nonce idz --keep-text --max-tokens 48 --rounds 1
+id_nomtp_S1 | 0,1 | 8 |   | STRATA_ROUTE_TAIL_SKIP=0;STRATA_PREFILL_CPU_SHARE=0;STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_BATCH_GDN_GROUPS=1;STRATA_BATCH_ATTN_ROWS=1;STRATA_BATCH_SELF_COMMIT=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_MIN_REQS=3;STRATA_COALESCE_WAIT_MS=60;STRATA_REPLICA_SHARED_ARENA=1 | --expert-cache 1500 --adapt-every 100000 --pcie-frac 0 --suffix-draft 0 | w2:4@--nonce idz --keep-text --max-tokens 48 --rounds 1 | --replicas 2

tools/par/variants/stack_id2.txt

@@ -0,0 +1,4 @@
+id_nomtp_host | 0,1 | 8 |   | STRATA_ROUTE_TAIL_SKIP=0;STRATA_PREFILL_CPU_SHARE=0;STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1 | --expert-cache 1500 --adapt-every 100000 --pcie-frac 0 --suffix-draft 0 | w2:4@--nonce idz --keep-text --max-tokens 48 --rounds 1
+id_nomtp_cuda | 0,1 | 8 |   | STRATA_ROUTE_TAIL_SKIP=0;STRATA_PREFILL_CPU_SHARE=0;STRATA_BATCH_GDN_GROUPS=1;STRATA_BATCH_ATTN_ROWS=1;STRATA_BATCH_SELF_COMMIT=1 | --expert-cache 1500 --adapt-every 100000 --pcie-frac 0 --suffix-draft 0 | w2:4@--nonce idz --keep-text --max-tokens 48 --rounds 1
+id_nomtp_co   | 0,1 | 8 |   | STRATA_ROUTE_TAIL_SKIP=0;STRATA_PREFILL_CPU_SHARE=0;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_MIN_REQS=3;STRATA_COALESCE_WAIT_MS=60 | --expert-cache 1500 --adapt-every 100000 --pcie-frac 0 --suffix-draft 0 | w2:4@--nonce idz --keep-text --max-tokens 48 --rounds 1
+id_nomtp_rep  | 0,1 | 8 |   | STRATA_ROUTE_TAIL_SKIP=0;STRATA_PREFILL_CPU_SHARE=0;STRATA_REPLICA_SHARED_ARENA=1 | --expert-cache 1500 --adapt-every 100000 --pcie-frac 0 --suffix-draft 0 | w2:4@--nonce idz --keep-text --max-tokens 48 --rounds 1 | --replicas 2

tools/par/variants/stack_s0.txt

@@ -0,0 +1 @@
+st_ls   | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_BATCH_GDN_GROUPS=1;STRATA_BATCH_ATTN_ROWS=1;STRATA_BATCH_SELF_COMMIT=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_MIN_REQS=3;STRATA_COALESCE_WAIT_MS=60 | | w1;w2:4,8,16@--nonce st --max-tokens 256 --rounds 2;w2:16,8,4@--nonce st --max-tokens 128 --rounds 2

tools/par/variants/stack_s1.txt

@@ -0,0 +1,2 @@
+st_rep  | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_BATCH_GDN_GROUPS=1;STRATA_BATCH_ATTN_ROWS=1;STRATA_BATCH_SELF_COMMIT=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_MIN_REQS=3;STRATA_COALESCE_WAIT_MS=60;STRATA_REPLICA_SHARED_ARENA=1 | | w1;w2:4,8,16@--nonce st --max-tokens 256 --rounds 2;w2:16,8,4@--nonce st --max-tokens 128 --rounds 2 | --replicas 2
+st_repw | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_BATCH_GDN_GROUPS=1;STRATA_BATCH_ATTN_ROWS=1;STRATA_BATCH_SELF_COMMIT=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_MIN_REQS=3;STRATA_COALESCE_WAIT_MS=60;STRATA_REPLICA_SHARED_ARENA=1;STRATA_PF_W16_CACHE_MIB=2048 | | w1;w2:4,8,16@--nonce st --max-tokens 256 --rounds 2;w2:16,8,4@--nonce st --max-tokens 128 --rounds 2 | --replicas 2

tools/par/variants/trace_2g.txt

@@ -0,0 +1,2 @@
+tr_c4 | 0,1 | 8 |   | STRATA_PAR_COUNTERS=1;STRATA_TRACE=1 | | w2:4@--max-tokens 128
+tr_c1 | 0,1 | 8 |   | STRATA_PAR_COUNTERS=1;STRATA_TRACE=1 | | w2:1@--max-tokens 128

tools/par/variants/z_ctx.txt

@@ -0,0 +1,3 @@
+zc16 | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=100 | --vram-reserve-mib 2500 --max-context 16384 | w2:4@--nonce pf --max-tokens 128 --rounds 2;w2:8@--nonce pf --max-tokens 128 --rounds 2
+zc8  | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=100 | --vram-reserve-mib 2500 --max-context 8192 | w2:4@--nonce pf --max-tokens 128 --rounds 2;w2:8@--nonce pf --max-tokens 128 --rounds 2
+zc8p4 | 0,1 | 4 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=100 | --vram-reserve-mib 2500 --max-context 8192 | w2:4@--nonce pf --max-tokens 128 --rounds 2;w2:8@--nonce pf --max-tokens 128 --rounds 2

tools/par/variants/z_final.txt

@@ -0,0 +1,4 @@
+zf_b1_1  | 0,1 | 8 |   |  | | w1;w2:4@--nonce zf --max-tokens 128 --rounds 2;w2:8@--nonce zf --max-tokens 128 --rounds 2
+zf_best_1| 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=100 | --vram-reserve-mib 2500 --max-context 8192 | w1;w2:4@--nonce zf --max-tokens 128 --rounds 2;w2:8@--nonce zf --max-tokens 128 --rounds 2
+zf_b1_2  | 0,1 | 8 |   |  | | w1;w2:4@--nonce zf --max-tokens 128 --rounds 2;w2:8@--nonce zf --max-tokens 128 --rounds 2
+zf_best_2| 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=100 | --vram-reserve-mib 2500 --max-context 8192 | w1;w2:4@--nonce zf --max-tokens 128 --rounds 2;w2:8@--nonce zf --max-tokens 128 --rounds 2

tools/par/variants/z_pcie.txt

@@ -0,0 +1,3 @@
+pf60 | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=100 | --vram-reserve-mib 2500 --pcie-frac 0.6 | w2:4@--nonce pf --max-tokens 128 --rounds 2;w2:8@--nonce pf --max-tokens 128 --rounds 2
+pf85 | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=100 | --vram-reserve-mib 2500 --pcie-frac 0.85 | w2:4@--nonce pf --max-tokens 128 --rounds 2;w2:8@--nonce pf --max-tokens 128 --rounds 2
+pf0  | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=100 | --vram-reserve-mib 2500 --pcie-frac 0 | w2:4@--nonce pf --max-tokens 128 --rounds 2;w2:8@--nonce pf --max-tokens 128 --rounds 2

tools/par/variants/z_pool.txt

@@ -0,0 +1,3 @@
+pw5  | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=100 | --vram-reserve-mib 2500 | w2:4@--nonce pf --max-tokens 128 --rounds 2;w2:8@--nonce pf --max-tokens 128 --rounds 2
+pw8  | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=100 | --vram-reserve-mib 2500 --pool-workers 8 | w2:4@--nonce pf --max-tokens 128 --rounds 2;w2:8@--nonce pf --max-tokens 128 --rounds 2
+pw11 | 0,1 | 8 |   | STRATA_CKPT_POOL=1;STRATA_CKPT_POOL_READY=4;STRATA_CKPT_THREADS=1;STRATA_SLOT_CKPT_MOVE=1;STRATA_SLOT_COPY_DIRECT=1;STRATA_SLOT_COPY_BATCHSYNC=1;STRATA_COALESCE_PREFILL=8;STRATA_COALESCE_WAIT_MS=100 | --vram-reserve-mib 2500 --pool-workers 11 | w2:4@--nonce pf --max-tokens 128 --rounds 2;w2:8@--nonce pf --max-tokens 128 --rounds 2

tools/par_analyze.py

@@ -0,0 +1,126 @@
+#!/usr/bin/env python3
+"""Paired-block analysis for par_load.py ledgers (standard library only).
+
+A pair = one A block and one B block of the same workload/level, run on the same hardware and cache regime and labelled
+with the same --pair id (--arm A|B).  For every (workload, level, metric) it prints the per-pair ratios B/A, their
+geometric mean and a paired bootstrap 95% interval (seeded), how many pairs B won, and the decision label from the
+plan (screening = few pairs: <5% in every pair low priority, >=10% keep; confirmation = 20+ pairs, lower bound vs --target).
+
+    python3 tools/par_analyze.py a1.json b1.json a2.json b2.json ...           (arm/pair read from the ledgers)
+    python3 tools/par_analyze.py runs/*.json --metric agg_tps --target 2.0 --json out.json
+
+Metrics where lower is better (ttft_*, tpot_*, max_gap_*, latency_*) are inverted so ratio > 1 always means B is better.
+`--baseline-arm`/`--cand-arm` choose the arms (default A and B).  A pair missing either arm, an invalid ledger
+(no server usage) or a block with errors is reported and left out.
+"""
+import argparse, glob, json, math, random, sys
+from collections import defaultdict
+
+LOWER_BETTER = ("ttft", "tpot", "max_gap", "latency", "lag", "joiner_ttft", "incumbent_gap", "makespan")
+
+
+def lower_is_better(metric):
+    return metric.startswith(LOWER_BETTER)
+
+
+def geomean(xs):
+    return math.exp(sum(math.log(x) for x in xs) / len(xs))
+
+
+def boot(ratios, n=10000, seed=1):
+    rng = random.Random(seed)
+    k = len(ratios)
+    gm = sorted(geomean([ratios[rng.randrange(k)] for _ in range(k)]) for _ in range(n))
+    return gm[int(0.025 * n)], gm[int(0.975 * n) - 1]
+
+
+def load(paths):
+    rows, skipped = [], []
+    for p in paths:
+        try:
+            L = json.load(open(p))
+        except Exception as e:  # noqa: BLE001
+            skipped.append((p, f"unreadable {e}"))
+            continue
+        if not L.get("valid", True):
+            skipped.append((p, "invalid ledger: " + str(L.get("invalid_reason"))))
+            continue
+        for b in L.get("blocks", []):
+            if b.get("errors"):
+                skipped.append((p, f"block has {b['errors']} errors"))
+                continue
+            rows.append((p, b))
+    return rows, skipped
+
+
+def analyze(rows, metric, a_arm="A", b_arm="B", target=None, min_confirm=20, boot_n=10000):
+    groups = defaultdict(lambda: defaultdict(dict))
+    for p, b in rows:
+        if metric not in b or b[metric] is None or not b.get("pair"):
+            continue
+        key = (b.get("workload"), b.get("level", b.get("offered_rate")), b.get("share"), b.get("round"))
+        groups[key][b["pair"]][b["arm"]] = b[metric]
+    out = []
+    for key, pairs in sorted(groups.items(), key=lambda kv: str(kv[0])):
+        ratios, used = [], []
+        for pid, arms in pairs.items():
+            if a_arm in arms and b_arm in arms and arms[a_arm] and arms[b_arm]:
+                r = arms[b_arm] / arms[a_arm]
+                ratios.append(1 / r if lower_is_better(metric) else r)
+                used.append(pid)
+        if not ratios:
+            continue
+        gm = geomean(ratios)
+        lo, hi = boot(ratios, boot_n) if len(ratios) > 1 else (gm, gm)
+        n = len(ratios)
+        faster = sum(1 for r in ratios if r > 1.0)
+        if n < min_confirm:
+            kind = "screen"
+            if all(r < 1.05 for r in ratios):
+                label = "low priority (<5% in every pair)" if all(r > 0.95 for r in ratios) else "reject/regression in every pair" if all(r < 1.0 for r in ratios) else "low priority"
+            elif gm >= 1.10 and faster == n:
+                label = "keep (>=10%)"
+            elif gm >= 1.05:
+                label = "investigate (5-10%)"
+            else:
+                label = "mixed, investigate"
+        else:
+            kind = "confirm"
+            label = "n/a"
+            if target is not None:
+                label = f"PASS (lower bound {lo:.3f} >= {target})" if lo >= target else f"NOT PROVEN (lower bound {lo:.3f} < {target})"
+        out.append({"workload": key[0], "level": key[1], "share": key[2], "round": key[3], "metric": metric, "pairs": n,
+                    "ratios": [round(r, 4) for r in ratios], "geomean": gm, "ci95": [lo, hi], "pairs_better": faster,
+                    "kind": kind, "decision": label, "mean_A": sum(pairs[p][a_arm] for p in used) / n, "mean_B": sum(pairs[p][b_arm] for p in used) / n})
+    return out
+
+
+def main(argv=None):
+    ap = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter)
+    ap.add_argument("ledgers", nargs="+")
+    ap.add_argument("--metric", action="append", default=None, help="repeatable; default agg_tps ttft_p95 max_gap_p95")
+    ap.add_argument("--baseline-arm", default="A")
+    ap.add_argument("--cand-arm", default="B")
+    ap.add_argument("--target", type=float, default=None, help="confirmation: the lower bound must clear this ratio")
+    ap.add_argument("--min-confirm", type=int, default=20)
+    ap.add_argument("--json", default=None)
+    a = ap.parse_args(argv)
+    paths = [p for g in a.ledgers for p in (glob.glob(g) or [g])]
+    rows, skipped = load(paths)
+    res = []
+    for m in a.metric or ["agg_tps", "ttft_p95", "max_gap_p95"]:
+        res += analyze(rows, m, a.baseline_arm, a.cand_arm, a.target, a.min_confirm)
+    for s in skipped:
+        print("skipped:", *s)
+    print(f"{'workload':8} {'level':>6} {'metric':14} {'pairs':>5} {'geomean':>8} {'95% CI':>17} {'better':>7}  decision")
+    for r in res:
+        lvl = r["level"] if r["share"] is None else f"{r['level']}/s{r['share']}"
+        print(f"{r['workload']:8} {str(lvl):>6} {r['metric']:14} {r['pairs']:>5} {r['geomean']:>8.3f} [{r['ci95'][0]:.3f},{r['ci95'][1]:.3f}] "
+              f"{r['pairs_better']:>3}/{r['pairs']:<3}  {r['kind']}: {r['decision']}")
+    if a.json:
+        json.dump(res, open(a.json, "w"), indent=1)
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())

tools/par_load.py

@@ -0,0 +1,544 @@
+#!/usr/bin/env python3
+"""Shared load harness for the parallelism program (research tool, standard library only).
+
+Wraps the OpenAI-compatible server the way tools/serve_load.py does, and adds what an SLO-capacity harness needs:
+
+  * closed-loop concurrency staircases (fixed work, shared wall window incl. the final drain),
+  * open-loop seeded Poisson arrivals (offered rate independent of completions; late/unfinished requests are failures),
+  * token counts from the server's `usage` (completion_tokens / prompt_tokens), NEVER from SSE chunks; a request with
+    no usage is marked `tokens_source: "sse"` and the run is flagged invalid,
+  * per request: TTFT, TPOT, max gap between content chunks (a chunk can hold several tokens when the engine flushes
+    speculative windows, so "gap" is a chunk gap), client scheduling lag, finish reason, cached prompt tokens,
+  * aggregate tok/s = total committed completion tokens / ONE shared wall-clock window (first submit to last completion),
+  * the workloads W1 solo, W2 short concurrent, W4 long joiner, W6 shared-prefix cohorts,
+  * one JSON ledger per run (args, server status before/after, per-request records, per-level blocks).
+
+    python3 tools/par_load.py http://127.0.0.1:8080 --workload w2 --conc 1,2,4,8 --rounds 3 --out run.json
+    python3 tools/par_load.py URL --workload w1 --prompt-tokens 512 --max-tokens 512 --rounds 3 --out w1.json
+    python3 tools/par_load.py URL --workload w4 --conc 4 --joiner-tokens 32768 --out w4.json
+    python3 tools/par_load.py URL --workload w6 --conc 4 --share 0,0.5,0.9 --out w6.json
+    python3 tools/par_load.py URL --workload w2 --open-rates 0.25,0.5,1 --duration 120 --slo-ttft 10 --out open.json
+
+Labelling for the paired analysis (tools/par_analyze.py): --arm A|B --pair N --label "free text".  The ledger's
+`blocks` list holds one entry per (workload, level) with the headline metrics.  See Desktop 30thOctober results/par/HARNESS.md.
+Prompts are deterministic from --seed and carry a unique nonce, so the prompt cache cannot help unless W6 asks for a shared prefix.
+"""
+import argparse, http.client, json, math, os, platform, random, statistics, sys, threading, time, urllib.parse, urllib.request
+
+HARNESS_VERSION = "1.0"
+WORDS = ("time year people way day man thing woman life child world school state family student group country problem hand "
+         "part place case week company system program question work government number night point home water room mother area "
+         "money story fact month lot right study book eye job word business issue side kind head house service friend father "
+         "power hour game line end member law car city community name president team minute idea kid body information back "
+         "parent face others level office door health person art war history party result change morning reason research girl "
+         "guy moment air teacher force education foot boy age policy music market sense nation plan college interest death "
+         "experience effect use class control care field development role effort rate heart drug show leader light voice wife "
+         "police mind price report decision son view relationship town road arm difference value building action model season "
+         "society tax director position player record paper space ground form event official matter center couple site project "
+         "activity star table need court oil situation cost industry figure street image phone data picture practice piece "
+         "land product doctor wall patient worker news test movie north love support technology step baby computer type "
+         "attention film tree source organization hair window culture chance brother energy period course summer plant").split()
+TOPICS = ["the history of lighthouses", "how a mechanical watch keeps time", "why bridges are built as arches",
+          "the water cycle and droughts", "how bread rises", "the design of a medieval castle",
+          "how a transistor switches", "the life of a honey bee colony"]
+
+
+# ----------------------------------------------------------------------------------------------------------- prompts
+class Prompter:
+    """Deterministic filler text of a target token count.  The words-per-token ratio is calibrated once against the
+    server's own usage.prompt_tokens (a max_tokens=1 probe), so the target counts tokens after the real template."""
+
+    def __init__(self, url, key, model, seed, think, calib=None):
+        self.url, self.key, self.model, self.seed, self.think = url, key, model, seed, think
+        self.tok_per_word = (calib or {}).get("tok_per_word")
+        self.overhead = (calib or {}).get("overhead", 40)
+        self.calib = calib or {}
+
+    def words(self, n, rng):
+        return " ".join(rng.choice(WORDS) for _ in range(n))
+
+    def body_text(self, key, n_words):
+        return self.words(n_words, random.Random(f"{self.seed}/{key}"))
+
+    def calibrate(self):
+        if self.tok_per_word:
+            return
+        n = 1500
+        text = "[calibration] " + self.body_text("calib", n) + "\nSay ok."
+        u = chat_once(self.url, self.key, self.model, text, 1, self.think)
+        pt = u.get("usage", {}).get("prompt_tokens")
+        if not pt:
+            self.tok_per_word, self.overhead = 1.3, 40            # server gave no usage: a guess, the ledger says so
+            self.calib = {"tok_per_word": 1.3, "overhead": 40, "guess": True}
+            return
+        # prompt = template overhead (about 25-40 tokens) + n words; assume the overhead is the same as the tail
+        self.overhead = 40
+        self.tok_per_word = max(0.5, (pt - self.overhead) / n)
+        self.calib = {"tok_per_word": self.tok_per_word, "overhead": self.overhead, "probe_prompt_tokens": pt, "probe_words": n}
+
+    def make(self, key, target_tokens, prefix=None, share=0.0, tail="Now write a long, detailed essay about "):
+        """A prompt of ~target_tokens.  With prefix (a stable string) and share: the first share*target tokens are
+        the prefix's words (the same text for the whole cohort), the rest is unique to `key`."""
+        self.calibrate()
+        nw = max(8, int((target_tokens - self.overhead) / self.tok_per_word))
+        topic = TOPICS[hash_int(key) % len(TOPICS)]
+        nonce = f"[req {key}] "
+        if prefix is not None and share > 0:
+            ns = int(nw * share)
+            pre = self.body_text("prefix/" + prefix, ns)
+            rest = self.body_text(key, nw - ns)
+            return pre + "\n" + nonce + rest + f"\n{tail}{topic}."
+        return nonce + self.body_text(key, nw) + f"\n{tail}{topic}."
+
+
+def hash_int(s):
+    h = 1469598103934665603
+    for ch in s.encode():
+        h = ((h ^ ch) * 1099511628211) & 0xFFFFFFFFFFFFFFFF
+    return h
+
+
+# ------------------------------------------------------------------------------------------------------------ request
+def _conn(url):
+    u = urllib.parse.urlparse(url)
+    cls = http.client.HTTPSConnection if u.scheme == "https" else http.client.HTTPConnection
+    return cls(u.hostname, u.port or (443 if u.scheme == "https" else 80), timeout=1800), u.path.rstrip("/")
+
+
+def _body(model, text, max_tokens, think, stream, extra=None):
+    b = {"model": model, "stream": stream, "max_tokens": max_tokens, "temperature": 0,
+         "messages": [{"role": "user", "content": text}]}
+    if stream:
+        b["stream_options"] = {"include_usage": True}
+    if not think:
+        b["reasoning_effort"] = "none"
+        b["chat_template_kwargs"] = {"enable_thinking": False}
+    if extra:
+        b.update(extra)
+    return b
+
+
+def chat_once(url, key, model, text, max_tokens, think):
+    c, base = _conn(url)
+    h = {"Content-Type": "application/json", **({"Authorization": "Bearer " + key} if key else {})}
+    c.request("POST", base + "/v1/chat/completions", json.dumps(_body(model, text, max_tokens, think, False)), h)
+    r = c.getresponse()
+    d = json.loads(r.read() or b"{}")
+    c.close()
+    return d
+
+
+def stream_request(url, key, model, text, max_tokens, think, sched_t0, extra=None, keep_times=False, keep_text=False):
+    """One streamed chat request.  Times are seconds on the perf_counter clock of this process; the record also
+    carries them relative to the run start (t_rel_*) by the caller."""
+    rec = {"t_sched": sched_t0, "t_send": None, "t_first": None, "t_last": None, "t_end": None, "completion_tokens": None,
+           "prompt_tokens": None, "cached_tokens": 0, "finish": None, "tokens_source": None, "chunks": 0, "max_gap": 0.0,
+           "gaps_ge1s": 0, "err": None, "timings": None, "reasoning_chars": 0, "content_chars": 0}
+    times = [] if keep_times else None
+    texts = [] if keep_text else None
+    rec["t_send"] = time.perf_counter()
+    last = None
+    sse_tokens = 0
+    try:
+        c, base = _conn(url)
+        h = {"Content-Type": "application/json", **({"Authorization": "Bearer " + key} if key else {})}
+        c.request("POST", base + "/v1/chat/completions", json.dumps(_body(model, text, max_tokens, think, True, extra)), h)
+        r = c.getresponse()
+        if r.status != 200:
+            rec["err"] = f"http {r.status}: {r.read(300).decode('utf-8', 'replace')}"
+        else:
+            while True:
+                raw = r.readline()
+                if not raw:
+                    break
+                line = raw.decode("utf-8", "replace").strip()
+                if not line.startswith("data:"):
+                    continue
+                d = line[5:].strip()
+                if d == "[DONE]":
+                    break
+                try:
+                    j = json.loads(d)
+                except ValueError:
+                    continue
+                now = time.perf_counter()
+                if j.get("error"):
+                    rec["err"] = str(j["error"])[:300]
+                for ch in j.get("choices", []):
+                    dl = ch.get("delta", {})
+                    a, b = dl.get("content") or "", dl.get("reasoning_content") or ""
+                    if a or b:
+                        if rec["t_first"] is None:
+                            rec["t_first"] = now
+                        elif last is not None:
+                            g = now - last
+                            if g > rec["max_gap"]:
+                                rec["max_gap"] = g
+                            if g >= 1.0:
+                                rec["gaps_ge1s"] += 1
+                        last = now
+                        rec["t_last"] = now
+                        rec["chunks"] += 1
+                        sse_tokens += 1
+                        rec["content_chars"] += len(a)
+                        if texts is not None:
+                            texts.append(a + b)
+                        rec["reasoning_chars"] += len(b)
+                        if times is not None:
+                            times.append(now)
+                    if ch.get("finish_reason"):
+                        rec["finish"] = ch["finish_reason"]
+                u = j.get("usage")
+                if u:
+                    rec["completion_tokens"] = u.get("completion_tokens")
+                    rec["prompt_tokens"] = u.get("prompt_tokens")
+                    rec["cached_tokens"] = (u.get("prompt_tokens_details") or {}).get("cached_tokens", 0) or 0
+                    rec["tokens_source"] = "usage"
+                if j.get("timings"):
+                    rec["timings"] = j["timings"]
+        c.close()
+    except Exception as e:  # noqa: BLE001
+        rec["err"] = repr(e)[:300]
+    rec["t_end"] = time.perf_counter()
+    if rec["completion_tokens"] is None:
+        rec["completion_tokens"] = sse_tokens
+        rec["tokens_source"] = "sse"                      # NOT a valid token count; the ledger marks the run invalid
+    if times is not None:
+        rec["chunk_times"] = times
+    if texts is not None:
+        rec["text"] = "".join(texts)
+    return rec
+
+
+def finish_record(rec, t0):
+    """Derived per-request metrics, and times relative to the run start."""
+    ttft = rec["t_first"] - rec["t_send"] if rec["t_first"] is not None else None
+    n = rec["completion_tokens"] or 0
+    tpot = (rec["t_last"] - rec["t_first"]) / (n - 1) if (ttft is not None and n > 1 and rec["t_last"] is not None) else None
+    out = {k: rec[k] for k in ("completion_tokens", "prompt_tokens", "cached_tokens", "finish", "tokens_source", "chunks", "max_gap",
+                               "gaps_ge1s", "err", "timings", "reasoning_chars", "content_chars")}
+    out.update({"t_sched": rec["t_sched"] - t0, "t_send": rec["t_send"] - t0, "t_end": rec["t_end"] - t0,
+                "t_first": (rec["t_first"] - t0) if rec["t_first"] is not None else None,
+                "lag": rec["t_send"] - rec["t_sched"], "ttft": ttft, "tpot": tpot, "latency": rec["t_end"] - rec["t_send"]})
+    if "text" in rec:
+        out["text"] = rec["text"]
+    if "chunk_times" in rec:
+        out["chunk_t"] = [round(x - t0, 4) for x in rec["chunk_times"]]
+    return out
+
+
+# --------------------------------------------------------------------------------------------------------- aggregates
+def pct(v, p):
+    v = sorted(x for x in v if x is not None)
+    if not v:
+        return None
+    k = (len(v) - 1) * p
+    lo = int(k)
+    hi = min(lo + 1, len(v) - 1)
+    return v[lo] + (v[hi] - v[lo]) * (k - lo)
+
+
+def aggregate(recs, window_s, slo=None):
+    """Level metrics.  window_s = one shared wall window for all requests (first submit to last completion)."""
+    ok = [r for r in recs if not r["err"]]
+    toks = sum(r["completion_tokens"] or 0 for r in ok)
+    agg = {"requests": len(recs), "errors": len(recs) - len(ok), "window_s": window_s, "tokens": toks,
+           "agg_tps": toks / window_s if window_s > 0 else None,
+           "prompt_tokens": sum(r["prompt_tokens"] or 0 for r in ok),
+           "invalid_token_source": sum(1 for r in recs if r["tokens_source"] != "usage"),
+           "length_stops": sum(1 for r in ok if r["finish"] == "length"), "early_stops": sum(1 for r in ok if r["finish"] == "stop")}
+    for k in ("ttft", "tpot", "max_gap", "latency", "lag"):
+        v = [r[k] for r in ok]
+        agg[k + "_p50"], agg[k + "_p95"], agg[k + "_max"] = pct(v, .5), pct(v, .95), (max(x for x in v if x is not None) if any(x is not None for x in v) else None)
+    per = [(r["completion_tokens"] - 1) / (r["t_end"] - r["t_first"]) for r in ok
+           if r["t_first"] is not None and (r["completion_tokens"] or 0) > 1 and r["t_end"] > r["t_first"]]
+    agg["per_req_tps_p50"] = pct(per, .5)
+    if slo:
+        good = [r for r in ok if (slo.get("ttft") is None or (r["ttft"] is not None and r["ttft"] <= slo["ttft"]))
+                and (slo.get("gap") is None or r["max_gap"] <= slo["gap"])]
+        agg["slo"] = slo
+        agg["slo_pass"] = len(good)
+        agg["goodput_tps"] = sum(r["completion_tokens"] or 0 for r in good) / window_s if window_s > 0 else None
+    return agg
+
+
+# --------------------------------------------------------------------------------------------------------- runners
+class Ctx:
+    def __init__(self, a, prompter):
+        self.a, self.p = a, prompter
+
+
+def closed_loop(ctx, level, n_requests, make_job, label):
+    """level concurrent clients pull jobs 0..n_requests-1 from one queue; the window is first submit -> last end."""
+    a = ctx.a
+    jobs = list(range(n_requests))
+    lock = threading.Lock()
+    recs = [None] * n_requests
+    t0 = time.perf_counter()
+
+    def worker():
+        while True:
+            with lock:
+                if not jobs:
+                    return
+                i = jobs.pop(0)
+            text, mt, extra = make_job(i)
+            recs[i] = stream_request(a.url, a.key, a.model, text, mt, a.think, time.perf_counter(), extra, a.keep_times, a.keep_text)
+
+    th = [threading.Thread(target=worker) for _ in range(level)]
+    for t in th:
+        t.start()
+    for t in th:
+        t.join()
+    t1 = max(r["t_end"] for r in recs)
+    fr = [finish_record(r, t0) for r in recs]
+    return fr, t1 - t0
+
+
+def run_closed(ctx, wl, levels):
+    a = ctx.a
+    blocks, allreq = [], []
+    for c in levels:
+        n = max(c * a.rounds, a.min_requests)
+        tag = f"{a.nonce}-{wl}-c{c}"
+
+        def make_job(i, c=c, tag=tag):
+            return ctx.p.make(f"{tag}-{i}", a.prompt_tokens), a.max_tokens, None
+
+        if a.warm_requests:                                      # warmup outside the measured window
+            for w in range(a.warm_requests):
+                stream_request(a.url, a.key, a.model, ctx.p.make(f"{tag}-warm{w}", a.prompt_tokens), min(a.max_tokens, 64), a.think, time.perf_counter())
+        fr, wall = closed_loop(ctx, c, n, make_job, tag)
+        agg = aggregate(fr, wall)
+        agg.update({"workload": wl, "level": c, "mode": "closed", "prompt_target": a.prompt_tokens, "max_tokens": a.max_tokens})
+        blocks.append(agg)
+        for i, r in enumerate(fr):
+            r.update({"workload": wl, "level": c, "idx": i})
+        allreq += fr
+        print(f"{wl} c={c}: {agg['tokens']} tokens / {wall:.1f}s = {agg['agg_tps']:.1f} tok/s  ttft p50/p95 "
+              f"{fmt(agg['ttft_p50'])}/{fmt(agg['ttft_p95'])}  tpot p50 {fmt(agg['tpot_p50'], 4)}  max_gap p95 {fmt(agg['max_gap_p95'])}"
+              f"  err {agg['errors']}", flush=True)
+    return blocks, allreq
+
+
+def fmt(x, d=2):
+    return "n/a" if x is None else f"{x:.{d}f}"
+
+
+def run_w4(ctx, levels):
+    """W4: c W2-style decoders running, then one long joiner (joiner_tokens) arrives joiner_at seconds after they start."""
+    a = ctx.a
+    blocks, allreq = [], []
+    for c in levels:
+        for rd in range(a.rounds):
+            tag = f"{a.nonce}-w4-c{c}-r{rd}"
+            res = {}
+            t0 = time.perf_counter()
+
+            def dec(i):
+                res[("d", i)] = stream_request(a.url, a.key, a.model, ctx.p.make(f"{tag}-d{i}", a.prompt_tokens), a.decoder_tokens, a.think, time.perf_counter(), None, a.keep_times or True)
+
+            def joiner():
+                time.sleep(a.joiner_at)
+                res[("j", 0)] = stream_request(a.url, a.key, a.model, ctx.p.make(f"{tag}-j", a.joiner_tokens), a.joiner_out, a.think, time.perf_counter(), None, True)
+
+            th = [threading.Thread(target=dec, args=(i,)) for i in range(c)] + [threading.Thread(target=joiner)]
+            for t in th:
+                t.start()
+            for t in th:
+                t.join()
+            wall = max(r["t_end"] for r in res.values()) - t0
+            fr = {k: finish_record(r, t0) for k, r in res.items()}
+            j = fr[("j", 0)]
+            ts_j = j["t_send"]
+            te_j = j["t_first"] if j["t_first"] is not None else j["t_end"]
+            # incumbent gaps between the joiner's submit and its first token (the interference window)
+            gaps = []
+            for (k, i), r in fr.items():
+                if k != "d":
+                    continue
+                ct = [x for x in r.get("chunk_t", [])]
+                for x, y in zip(ct, ct[1:]):
+                    if y >= ts_j and x <= te_j + 1.0:
+                        gaps.append(y - x)
+            agg = aggregate(list(fr.values()), wall)
+            allreq += [dict(r, workload="w4", level=c, round=rd, role=k[0]) for k, r in fr.items()]
+            for r in allreq[-len(fr):]:
+                r.pop("chunk_t", None) if not a.keep_times else None
+            decs = [r for k, r in fr.items() if k[0] == "d"]
+            agg.update({"workload": "w4", "level": c, "mode": "closed", "round": rd, "joiner_tokens": a.joiner_tokens,
+                        "joiner_ttft": j["ttft"], "joiner_prompt_tokens": j["prompt_tokens"],
+                        "incumbent_gap_max_during_join": max(gaps) if gaps else None,
+                        "incumbent_gap_p99_during_join": pct(gaps, .99),
+                        "incumbent_tps_p50": pct([(d["completion_tokens"] - 1) / max(d["t_end"] - d["t_first"], 1e-9) for d in decs if d["t_first"] is not None and d["completion_tokens"] and d["completion_tokens"] > 1], .5),
+                        "makespan_s": wall})
+            blocks.append(agg)
+            print(f"w4 c={c} r={rd}: makespan {wall:.1f}s agg {fmt(agg['agg_tps'], 1)} tok/s, joiner ttft {fmt(j['ttft'])}s ({j['prompt_tokens']} prompt tok), "
+                  f"incumbent max gap during join {fmt(agg['incumbent_gap_max_during_join'])}s", flush=True)
+    return blocks, allreq
+
+
+def run_w6(ctx, levels):
+    a = ctx.a
+    blocks, allreq = [], []
+    for c in levels:
+        for share in [float(x) for x in a.share.split(",")]:
+            tag = f"{a.nonce}-w6-c{c}-s{share}"
+            prefix = f"{tag}"
+            if a.prime and share > 0:                           # prefix-warm: one priming request, not measured
+                stream_request(a.url, a.key, a.model, ctx.p.make(f"{tag}-prime", a.prompt_tokens, prefix, share), 8, a.think, time.perf_counter())
+            n = max(c * a.rounds, a.min_requests)
+
+            def make_job(i, tag=tag, prefix=prefix, share=share):
+                return ctx.p.make(f"{tag}-{i}", a.prompt_tokens, prefix, share), a.max_tokens, None
+
+            fr, wall = closed_loop(ctx, c, n, make_job, tag)
+            agg = aggregate(fr, wall)
+            agg.update({"workload": "w6", "level": c, "share": share, "primed": bool(a.prime), "mode": "closed",
+                        "cached_tokens_mean": statistics.mean(r["cached_tokens"] for r in fr) if fr else None})
+            blocks.append(agg)
+            for i, r in enumerate(fr):
+                r.update({"workload": "w6", "level": c, "share": share, "idx": i})
+            allreq += fr
+            print(f"w6 c={c} share={share}: {fmt(agg['agg_tps'], 1)} tok/s ttft p50 {fmt(agg['ttft_p50'])} cached mean {agg['cached_tokens_mean']:.0f}", flush=True)
+    return blocks, allreq
+
+
+def run_open(ctx, wl, rates):
+    """Open loop: seeded exponential inter-arrival times, independent of completions.  Requests still running at
+    duration + drain_s are failures (kept in the ledger as unfinished)."""
+    a = ctx.a
+    blocks, allreq = [], []
+    for rate in rates:
+        rng = random.Random(f"{a.seed}/open/{rate}")
+        sched, t = [], 0.0
+        while True:
+            t += rng.expovariate(rate)
+            if t >= a.duration:
+                break
+            sched.append(t)
+        tag = f"{a.nonce}-{wl}-open{rate}"
+        t0 = time.perf_counter()
+        res = [None] * len(sched)
+        th = []
+
+        def one(i, at):
+            res[i] = stream_request(a.url, a.key, a.model, ctx.p.make(f"{tag}-{i}", a.prompt_tokens), a.max_tokens, a.think, t0 + at, None, False)
+
+        def launcher():
+            for i, at in enumerate(sched):
+                d = t0 + at - time.perf_counter()
+                if d > 0:
+                    time.sleep(d)
+                t_ = threading.Thread(target=one, args=(i, at), daemon=True)
+                t_.start()
+                th.append(t_)
+
+        lt = threading.Thread(target=launcher)
+        lt.start()
+        lt.join()
+        deadline = t0 + a.duration + a.drain_s
+        for t_ in th:
+            t_.join(max(0.0, deadline - time.perf_counter()))
+        done = [r for r in res if r is not None and r["t_end"] is not None]
+        unfinished = len(sched) - len(done)
+        window = max([r["t_end"] for r in done], default=t0) - t0
+        fr = [finish_record(r, t0) for r in done]
+        slo = {"ttft": a.slo_ttft, "gap": a.slo_gap}
+        # tokens finished inside the arrival window over that window; the drain is reported separately
+        inwin = [r for r in fr if r["t_end"] <= a.duration]
+        agg = aggregate(fr, window, slo)
+        agg["tps_in_window"] = sum(r["completion_tokens"] or 0 for r in inwin if not r["err"]) / a.duration
+        agg.update({"workload": wl, "mode": "open", "offered_rate": rate, "offered_requests": len(sched), "unfinished": unfinished,
+                    "duration": a.duration, "drain_s": a.drain_s, "achieved_rate": len(inwin) / a.duration,
+                    "slo_pass_frac": (agg.get("slo_pass", 0) / len(sched)) if sched else None})
+        blocks.append(agg)
+        allreq += [dict(r, workload=wl, offered_rate=rate, idx=i) for i, r in enumerate(fr)]
+        print(f"{wl} open rate={rate}/s: offered {len(sched)}, done {len(done)}, unfinished {unfinished}, slo pass {agg.get('slo_pass')}, "
+              f"goodput {fmt(agg.get('goodput_tps'), 1)} tok/s, ttft p95 {fmt(agg['ttft_p95'])}, lag max {fmt(agg['lag_max'], 3)}", flush=True)
+    return blocks, allreq
+
+
+# ------------------------------------------------------------------------------------------------------------- main
+def get_json(url):
+    try:
+        with urllib.request.urlopen(url, timeout=10) as r:
+            return json.loads(r.read())
+    except Exception as e:  # noqa: BLE001
+        return {"error": repr(e)[:200]}
+
+
+def main(argv=None):
+    ap = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter)
+    ap.add_argument("url")
+    ap.add_argument("--workload", default="w2", choices=["w1", "w2", "w3", "w4", "w6"])
+    ap.add_argument("--conc", default="1,2,4,8", help="closed-loop concurrency levels (w1 forces 1)")
+    ap.add_argument("--rounds", type=int, default=3, help="requests per client per level (a level runs conc*rounds requests)")
+    ap.add_argument("--min-requests", type=int, default=0)
+    ap.add_argument("--warm-requests", type=int, default=0, help="unmeasured short requests before each level")
+    ap.add_argument("--prompt-tokens", type=int, default=None, help="default: w1 512, w2 1024, w3 8192, w6 2048")
+    ap.add_argument("--max-tokens", type=int, default=None, help="default: w1 512, w2 256, w3 128, w6 128")
+    ap.add_argument("--joiner-tokens", type=int, default=28000)
+    ap.add_argument("--joiner-out", type=int, default=64)
+    ap.add_argument("--joiner-at", type=float, default=6.0, help="seconds after the decoders start")
+    ap.add_argument("--decoder-tokens", type=int, default=512)
+    ap.add_argument("--share", default="0,0.5,0.9", help="w6 shared-prefix fractions")
+    ap.add_argument("--prime", action="store_true", help="w6: one unmeasured request primes the prefix (prefix-warm)")
+    ap.add_argument("--open-rates", default=None, help="open-loop offered rates (req/s), e.g. 0.25,0.5,1")
+    ap.add_argument("--duration", type=float, default=120.0, help="open loop: arrival window (s)")
+    ap.add_argument("--drain-s", type=float, default=120.0, help="open loop: extra time to finish before counting failures")
+    ap.add_argument("--slo-ttft", type=float, default=None)
+    ap.add_argument("--slo-gap", type=float, default=None)
+    ap.add_argument("--seed", default="par29")
+    ap.add_argument("--nonce", default=str(int(time.time())))
+    ap.add_argument("--model", default="strata")
+    ap.add_argument("--key", default=os.environ.get("STRATA_KEY", ""))
+    ap.add_argument("--think", action="store_true")
+    ap.add_argument("--keep-times", action="store_true", help="keep every chunk timestamp in the ledger (large)")
+    ap.add_argument("--keep-text", action="store_true", help="keep every answer's text in the ledger (identity checks)")
+    ap.add_argument("--calib", default=None, help="a ledger/calibration JSON to reuse the words-per-token calibration from")
+    ap.add_argument("--arm", default="")
+    ap.add_argument("--pair", default="")
+    ap.add_argument("--label", default="")
+    ap.add_argument("--settings", default=None, help="file with the effective engine settings/provenance to embed (launcher writes it)")
+    ap.add_argument("--out", required=True)
+    a = ap.parse_args(argv)
+    a.url = a.url.rstrip("/")
+    d = {"w1": (512, 512), "w2": (1024, 256), "w3": (8192, 128), "w4": (1024, 256), "w6": (2048, 128)}[a.workload]
+    a.prompt_tokens = a.prompt_tokens or d[0]
+    a.max_tokens = a.max_tokens or d[1]
+    levels = [1] if a.workload == "w1" else [int(x) for x in a.conc.split(",")]
+    calib = None
+    if a.calib and os.path.exists(a.calib):
+        calib = json.load(open(a.calib)).get("calibration")
+    p = Prompter(a.url, a.key, a.model, a.seed, a.think, calib)
+    ctx = Ctx(a, p)
+    status0 = get_json(a.url + "/v1/status")
+    started = time.time()
+    if a.open_rates:
+        blocks, reqs = run_open(ctx, a.workload, [float(x) for x in a.open_rates.split(",")])
+    elif a.workload == "w4":
+        blocks, reqs = run_w4(ctx, levels)
+    elif a.workload == "w6":
+        blocks, reqs = run_w6(ctx, levels)
+    else:
+        blocks, reqs = run_closed(ctx, a.workload, levels)
+    status1 = get_json(a.url + "/v1/status")
+    settings = None
+    if a.settings and os.path.exists(a.settings):
+        settings = open(a.settings, errors="replace").read()[:20000]
+    for b in blocks:
+        b.update({"arm": a.arm, "pair": a.pair, "label": a.label})
+    invalid = sum(b.get("invalid_token_source", 0) for b in blocks)
+    ledger = {"harness": "par_load", "version": HARNESS_VERSION, "started": started, "host": platform.node(), "argv": sys.argv[1:],
+              "args": {k: v for k, v in vars(a).items() if k != "key"}, "calibration": p.calib, "valid": invalid == 0,
+              "invalid_reason": None if invalid == 0 else f"{invalid} requests without server usage: token counts are SSE chunks",
+              "status_before": status0, "status_after": status1, "settings": settings, "blocks": blocks, "requests": reqs}
+    os.makedirs(os.path.dirname(os.path.abspath(a.out)), exist_ok=True)
+    json.dump(ledger, open(a.out, "w"), indent=1)
+    print("ledger:", a.out, "valid" if invalid == 0 else "INVALID (no usage)")
+    return 0 if invalid == 0 and all(b["errors"] == 0 for b in blocks) else 1
+
+
+if __name__ == "__main__":
+    sys.exit(main())

tools/par_run.py

@@ -0,0 +1,215 @@
+#!/usr/bin/env python3
+"""par_run.py - start serve/server.py with one configuration variant on a test box, run par_load.py workloads, stop.
+
+Research tool for the parallelism program.  One call = one server start = one arm/variant.  Writes next to --out:
+  <tag>-cfg.json, <tag>-engine.log, <tag>-server.out, <tag>-<workload>.json (par_load ledgers), <tag>-run.json
+(effective engine lines, PSS/RSS of the engine tree, GPU memory, ready time, the cmd/env).
+
+    python3 tools/par_run.py --box boston --tag base1g --gpus 0 --work "w1;w2:4,8" --out /path/dir
+    python3 tools/par_run.py --box boston --tag s8g4 --gpus 0,1 --parallel 8 --groups 4 --env "STRATA_PIPED_CONCURRENT=1" \
+        --extra "--foo bar" --work "w2:4,8;w4:4" --out dir
+
+--work items are workload[:conc list] separated by ';', optionally with par_load flags after '@' (e.g. "w4:4@--joiner-tokens 16384").
+--default-config uses what the box's 0.2.8 launcher would write (no parallel key, layer_split auto when 2 GPUs).
+"""
+import argparse, json, os, shlex, signal, subprocess, sys, time, urllib.request
+
+BOXES = {
+    "boston": {
+        "py": "/data/external/niko/t027/venv/bin/python", "libs": ["/usr/local/cuda-13.1/lib64"],
+        "pack": "/data/external/niko/pm/pack/iq3_s", "mtp": "/data/external/niko/pm/mtp/rt",
+        "gguf1": "/data3/models/ISTA-DASLab/Qwen3.8-Flash-Next-GSQ-RCO-GGUF/IQ3_S/Qwen3.8-Flash-Next-GSQ-RCO-IQ3_S-00001-of-00002.gguf",
+        "gguf2": "/data3/models/ISTA-DASLab/Qwen3.8-Flash-Next-GSQ-RCO-GGUF/IQ3_S/Qwen3.8-Flash-Next-GSQ-RCO-IQ3_S-00002-of-00002.gguf",
+        "profile": "{repo}/data/expert-profile.bin",
+    },
+    "z440": {
+        "py": "/data/Strata/.venv/bin/python", "libs": ["/usr/local/cuda-12.9/lib64"],
+        "pack": "/data/strata-models/IQ3_S/pack", "mtp": "/data/Strata-data/mtp/rt",
+        "gguf1": "/data/strata-models/IQ3_S/Qwen3.8-Flash-Next-GSQ-RCO-IQ3_S-00001-of-00002.gguf",
+        "gguf2": "/data/strata-models/IQ3_S/Qwen3.8-Flash-Next-GSQ-RCO-IQ3_S-00002-of-00002.gguf",
+        "profile": "{repo}/data/expert-profile.bin",
+    },
+}
+
+
+def pss_tree(pid, exe=None):
+    """Sum of PSS (MB) and RSS of a process and all its descendants (Linux, via /proc ppid)."""
+    kids = {}
+    for d in os.listdir("/proc"):
+        if d.isdigit():
+            try:
+                st = open(f"/proc/{d}/stat").read()
+                ppid = int(st[st.rindex(")") + 2:].split()[1])
+                kids.setdefault(ppid, []).append(int(d))
+            except (OSError, ValueError):
+                pass
+    seen, todo = set(), [pid]
+    for d in os.listdir("/proc"):          # engines are started detached: also every process running our --exe
+        if d.isdigit() and exe:
+            try:
+                if exe in open(f"/proc/{d}/cmdline").read().replace("\0", " "):
+                    todo.append(int(d))
+            except OSError:
+                pass
+    while todo:
+        p = todo.pop()
+        if p not in seen:
+            seen.add(p)
+            todo += kids.get(p, [])
+    pss = rss = 0
+    for p in seen:
+        try:
+            for line in open(f"/proc/{p}/smaps_rollup"):
+                if line.startswith("Pss:"):
+                    pss += int(line.split()[1])
+                elif line.startswith("Rss:"):
+                    rss += int(line.split()[1])
+        except OSError:
+            pass
+    return {"pss_mb": pss // 1024, "rss_mb": rss // 1024, "procs": len(seen)}
+
+
+def gpu_mem():
+    try:
+        o = subprocess.run(["nvidia-smi", "--query-gpu=index,memory.used,memory.total", "--format=csv,noheader,nounits"],
+                           capture_output=True, text=True, timeout=20).stdout
+        return [[int(x) for x in l.split(",")] for l in o.strip().splitlines()]
+    except Exception:  # noqa: BLE001
+        return None
+
+
+def main():
+    ap = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter)
+    ap.add_argument("--box", required=True, choices=list(BOXES))
+    ap.add_argument("--tag", required=True)
+    ap.add_argument("--out", required=True, help="output directory")
+    ap.add_argument("--exe", required=True)
+    ap.add_argument("--repo", required=True, help="source tree (serve/server.py, data/expert-profile.bin)")
+    ap.add_argument("--gpus", default="0,1")
+    ap.add_argument("--parallel", type=int, default=None)
+    ap.add_argument("--groups", default=None)
+    ap.add_argument("--replicas", type=int, default=None)
+    ap.add_argument("--split", default="auto")
+    ap.add_argument("--ctx", type=int, default=32768)
+    ap.add_argument("--prefill", default="auto")
+    ap.add_argument("--no-mtp", action="store_true")
+    ap.add_argument("--mtp-depth", type=int, default=4, help="--spec value (the default in setup configs is 4)")
+    ap.add_argument("--env", default="", help="K=V;K=V for the engine/server environment")
+    ap.add_argument("--extra", default="", help="extra engine args (one string)")
+    ap.add_argument("--server-extra", default="", help="extra server.py args")
+    ap.add_argument("--port", type=int, default=18930)
+    ap.add_argument("--work", default="w2:4,8")
+    ap.add_argument("--rounds", type=int, default=3)
+    ap.add_argument("--cap", type=int, default=1500, help="engine ready timeout (s)")
+    ap.add_argument("--min-free-gb", type=int, default=14)
+    ap.add_argument("--post-cmd", default=None, help="a shell command run after the work items; {url} is the server URL; its output goes to <tag>-post.txt")
+    a = ap.parse_args()
+    B = BOXES[a.box]
+    os.makedirs(a.out, exist_ok=True)
+    # memory guard: never start with too little RAM (the replicas and the 2-GPU engine each pin tens of GB)
+    avail = 0
+    for line in open("/proc/meminfo"):
+        if line.startswith("MemAvailable:"):
+            avail = int(line.split()[1]) // 1048576
+    if avail < a.min_free_gb:
+        print(f"ABORT: only {avail} GB RAM available (< {a.min_free_gb})")
+        return 3
+    log = os.path.join(a.out, f"{a.tag}-engine.log")
+    args = ["--pack", B["pack"], "--native", B["gguf1"], "--ple-gguf", B["gguf2"],
+            "--expert-profile", B["profile"].format(repo=a.repo), "--expert-cache", "auto", "--prefill", a.prefill,
+            "--spec", str(a.mtp_depth), "--spec-min-p", "0.5", "--layer-split", a.split, "--max-context", str(a.ctx), "--kv", "int8"]
+    if not a.no_mtp:
+        args += ["--mtp", B["mtp"]]
+    args += shlex.split(a.extra)
+    gpus = [int(x) for x in a.gpus.split(",")]
+    cfg = {"exe": a.exe, "args": args, "cwd": a.repo, "tokenizer": B["pack"] + "/tokenizer", "model_name": "qwen3.8-flash-next-iq3_s",
+           "log": log, "lib_dirs": B["libs"], "port": a.port, "gpu": gpus, "gpus_asked": True, "engine_ready_s": a.cap}
+    if len(gpus) == 1:
+        cfg["args"] = [x for i, x in enumerate(args) if x != "--layer-split" and (i == 0 or args[i - 1] != "--layer-split")]
+    if a.replicas:
+        cfg["replicas"] = a.replicas
+    cpath = os.path.join(a.out, f"{a.tag}-cfg.json")
+    json.dump(cfg, open(cpath, "w"), indent=1)
+    env = {k: v for k, v in os.environ.items() if not k.startswith("STRATA_")}
+    env["LD_LIBRARY_PATH"] = ":".join(B["libs"] + [env.get("LD_LIBRARY_PATH", "")])
+    for kv in a.env.split(";"):
+        if kv.strip():
+            k, v = kv.split("=", 1)
+            env[k.strip()] = v.strip()
+    cmd = [B["py"], f"{a.repo}/serve/server.py", "--engine", "strata", "--config", cpath, "--port", str(a.port)]
+    if a.parallel is not None:
+        cmd += ["--parallel", str(a.parallel)]
+    if a.groups:
+        cmd += ["--batch-groups", str(a.groups)]
+    cmd += shlex.split(a.server_extra)
+    srvout = open(os.path.join(a.out, f"{a.tag}-server.out"), "w")
+    t0 = time.time()
+    srv = subprocess.Popen(cmd, env=env, stdout=srvout, stderr=subprocess.STDOUT, start_new_session=True, cwd=a.repo)
+    ready = None
+    while time.time() - t0 < a.cap:
+        if srv.poll() is not None:
+            print("SERVER DIED", flush=True)
+            break
+        try:
+            urllib.request.urlopen(f"http://127.0.0.1:{a.port}/v1/models", timeout=5).read()
+            st = json.loads(urllib.request.urlopen(f"http://127.0.0.1:{a.port}/v1/status", timeout=5).read())
+            if (st.get("state") or "ready") in ("ready", "idle", "serving") or st.get("ready"):
+                ready = time.time() - t0
+                break
+        except Exception:  # noqa: BLE001
+            pass
+        time.sleep(3)
+    print(f"[{a.tag}] ready after {ready}s", flush=True)
+    rec = {"tag": a.tag, "box": a.box, "cmd": cmd, "env": a.env, "args": args, "ready_s": ready, "work": {}, "mem": {}}
+    url = f"http://127.0.0.1:{a.port}"
+    if ready is not None:
+        time.sleep(2)
+        rec["mem"]["after_load"] = pss_tree(srv.pid, a.exe)
+        rec["gpu_after_load"] = gpu_mem()
+        here = os.path.dirname(os.path.abspath(__file__))
+        # a short unmeasured warmup so that lazy code objects/JIT and the first expert-cache fills are out of the way
+        subprocess.run([sys.executable, f"{here}/par_load.py", url, "--workload", "w2", "--conc", "2", "--rounds", "1", "--max-tokens", "48",
+                        "--nonce", "warm", "--out", os.path.join(a.out, f"{a.tag}-warm.json")], capture_output=True)
+        for item in a.work.split(";"):
+            if not item.strip():
+                continue
+            main_part, _, extra = item.partition("@")
+            wl, _, conc = main_part.partition(":")
+            wl = wl.strip()
+            outj = os.path.join(a.out, f"{a.tag}-{wl}{conc.replace(',', '_')}.json")
+            c2 = [sys.executable, f"{here}/par_load.py", url, "--workload", wl, "--rounds", str(a.rounds), "--arm", a.tag, "--pair", "0",
+                  "--label", a.tag, "--settings", log, "--out", outj]
+            if conc:
+                c2 += ["--conc", conc]
+            c2 += shlex.split(extra)
+            r = subprocess.run(c2, capture_output=True, text=True)
+            print(r.stdout[-1500:], flush=True)
+            if r.returncode:
+                print("par_load rc", r.returncode, r.stderr[-300:], flush=True)
+            rec["work"][item] = outj
+            rec["mem"][item] = pss_tree(srv.pid, a.exe)
+        if a.post_cmd:
+            r = subprocess.run(a.post_cmd.format(url=url), shell=True, capture_output=True, text=True)
+            open(os.path.join(a.out, f"{a.tag}-post.txt"), "w").write(r.stdout + r.stderr)
+            print(r.stdout[-1500:], flush=True)
+        rec["gpu_after_work"] = gpu_mem()
+    try:
+        os.killpg(srv.pid, signal.SIGTERM)
+        srv.wait(120)
+    except Exception:  # noqa: BLE001
+        try:
+            os.killpg(srv.pid, signal.SIGKILL)
+        except Exception:  # noqa: BLE001
+            pass
+    text = open(log, errors="replace").read() if os.path.exists(log) else ""
+    keys = ("batch_slots", "batch_groups", "expert cache", "slot sessions", "groups of", "pipelined", "VRAM free", "layer split", "layer-split",
+            "strata p4", "strata par", "mtp", "adapt")
+    rec["engine_lines"] = [l[:220] for l in text.splitlines() if any(k in l for k in keys)][:60]
+    rec["errors_in_log"] = len([1 for l in text.splitlines() if any(k in l for k in ("illegal instruction", "Traceback", " ERR ", "segfault", "Segmentation", "out of memory"))])
+    json.dump(rec, open(os.path.join(a.out, f"{a.tag}-run.json"), "w"), indent=1)
+    print("DONE", a.tag, "errors_in_log", rec["errors_in_log"])
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())

tools/test_par_load.py

@@ -0,0 +1,69 @@
+"""Self-test of tools/par_load.py and par_analyze.py against an in-process fake SSE server."""
+import json, os, sys, tempfile, threading, time, unittest
+from http.server import BaseHTTPRequestHandler, ThreadingHTTPServer
+sys.path.insert(0, os.path.dirname(__file__))
+import par_load, par_analyze
+
+
+class H(BaseHTTPRequestHandler):
+    def log_message(self, *a): pass
+    def do_GET(self):
+        self.send_response(200); self.end_headers(); self.wfile.write(b'{"ok":1}')
+    def do_POST(self):
+        n = int(self.headers["Content-Length"]); b = json.loads(self.rfile.read(n))
+        text = b["messages"][0]["content"]
+        pt = len(text.split()) + 10
+        if not b["stream"]:
+            self.send_response(200); self.end_headers()
+            self.wfile.write(json.dumps({"usage": {"prompt_tokens": pt, "completion_tokens": 1}}).encode()); return
+        self.send_response(200); self.send_header("Content-Type", "text/event-stream"); self.end_headers()
+        k = min(b["max_tokens"], 12)
+        for i in range(k):
+            time.sleep(0.01)
+            self.wfile.write(("data: " + json.dumps({"choices": [{"delta": {"content": "w "}}]}) + "\n\n").encode()); self.wfile.flush()
+        # two tokens per chunk at the end would make SSE counts wrong: usage is the truth
+        self.wfile.write(("data: " + json.dumps({"choices": [{"delta": {}, "finish_reason": "length"}], "usage": {"prompt_tokens": pt, "completion_tokens": k + 3}}) + "\n\n").encode())
+        self.wfile.write(b"data: [DONE]\n\n"); self.wfile.flush()
+
+
+class T(unittest.TestCase):
+    @classmethod
+    def setUpClass(cls):
+        cls.s = ThreadingHTTPServer(("127.0.0.1", 0), H)
+        threading.Thread(target=cls.s.serve_forever, daemon=True).start()
+        cls.url = f"http://127.0.0.1:{cls.s.server_address[1]}"
+    @classmethod
+    def tearDownClass(cls): cls.s.shutdown()
+
+    def run_(self, *args):
+        out = tempfile.mktemp(suffix=".json")
+        rc = par_load.main([self.url, "--out", out, *args])
+        return rc, json.load(open(out))
+
+    def test_w2_usage_tokens(self):
+        rc, L = self.run_("--workload", "w2", "--conc", "1,2", "--rounds", "2", "--max-tokens", "12", "--prompt-tokens", "200")
+        self.assertEqual(rc, 0); self.assertTrue(L["valid"])
+        b = L["blocks"][1]
+        self.assertEqual(b["tokens"], 4 * 15)             # usage says 15, SSE chunks were 12
+        self.assertEqual(b["requests"], 4)
+        self.assertIsNotNone(b["ttft_p50"]); self.assertIsNotNone(b["tpot_p50"])
+
+    def test_w4_w6_open(self):
+        rc, L = self.run_("--workload", "w4", "--conc", "2", "--rounds", "1", "--joiner-at", "0.05", "--joiner-tokens", "400", "--prompt-tokens", "100")
+        self.assertEqual(rc, 0); self.assertIn("joiner_ttft", L["blocks"][0])
+        rc, L = self.run_("--workload", "w6", "--conc", "2", "--rounds", "1", "--share", "0,0.9", "--prompt-tokens", "300")
+        self.assertEqual(len(L["blocks"]), 2)
+        rc, L = self.run_("--workload", "w2", "--open-rates", "20", "--duration", "1", "--drain-s", "5", "--max-tokens", "8", "--prompt-tokens", "100")
+        self.assertEqual(rc, 0); self.assertEqual(L["blocks"][0]["unfinished"], 0)
+
+    def test_analyze(self):
+        rows = []
+        for i, (a, b) in enumerate([(10, 21), (10, 20), (11, 22)]):
+            rows += [("x", {"workload": "w2", "level": 4, "arm": "A", "pair": str(i), "agg_tps": a, "errors": 0}),
+                     ("x", {"workload": "w2", "level": 4, "arm": "B", "pair": str(i), "agg_tps": b, "errors": 0})]
+        r = par_analyze.analyze(rows, "agg_tps")
+        self.assertAlmostEqual(r[0]["geomean"], 2.0, delta=0.06); self.assertEqual(r[0]["pairs_better"], 3)
+
+
+if __name__ == "__main__":
+    unittest.main()