Skip to content

Pipelined windows over the active slots (+ slot allocator), /metrics for Prometheus (vLLM names) + monitoring kit - #793

Closed
blange48 wants to merge 851 commits into
Niko1221:mainfrom
blange48:upstream-metrics-autoconfig
Closed

blange48 wants to merge 851 commits into
Niko1221:mainfrom
blange48:upstream-metrics-autoconfig

Conversation

@blange48

@blange48 blange48 commented Oct 4, 2026 •

Copy link
Copy Markdown
Contributor

Pipelined windows over the active slots (+ slot allocator), /metrics for Prometheus

These are the parts of #559 that didn't go into 0.1.39, rebased on main (6f32ec0) and adapted to "parallel" and its one status per request. Each commit stands on its own.

1. Pipelined group windows over the active slots (src/program/generate.cpp, 8 lines)

0.1.39 runs the one-GPU batch windows over the active slots. With --batch-groups, a group's window still holds every slot of the group, and the idle ones run pad rows through every layer. The server spreads requests over the groups first, so on a 4-group split two requests sit in two windows that are mostly pad rows. Now a group's window holds its slots up to its last active one. This is crazyaimachine's 8119a7f from #559, without the MTP-draft parts.

4 x RTX 5080 (PCIe Gen3), IQ3_S, 262K context, "parallel": 8, --batch-groups 4 --trim-stage-weights, through the HTTP server, T=0.7, 300 tokens per answer:

Concurrent 0.1.39 with this
1 122 117-122 (unchanged path)
2 99 149
4 199 222
8 366 371

Exactness (--pcie-frac 0 --adapt-every 1000000, STRATA_IQ_MT_MIN=1):

  • tools/batch_test.py: identical to solo with 8 slots and with 3 (groups partly filled);
  • tools/parking_test.py: identical, with 141 MB of K/V reused on the second park.

2. GET /metrics for Prometheus, with vLLM's metric names (serve/prometheus.py, three commits)

  • Format. A Prometheus scrape (Accept: text/plain or openmetrics, or ?format=prometheus) gets the text format. Anything else, the Monitor tab included, keeps the JSON.
  • vLLM names. Dashboards and alerts written for a vLLM server read Strata unchanged. These are the names current vLLM exports:
    • vllm:num_requests_running / _waiting and vllm:kv_cache_usage_perc;
    • the token and request counters;
    • vllm:prefix_cache_queries_total / _hits_total (prompt tokens read / reused);
    • vllm:spec_decode_num_draft_tokens_total / _accepted_tokens_total (the MTP drafts);
    • the histograms vllm:time_to_first_token_seconds, vllm:inter_token_latency_seconds and vllm:e2e_request_latency_seconds.
  • Strata's own facts. What vLLM has no name for comes under strata:, named after its JSON key: strata:live_state, strata:live_slots, strata:live_tok_s, strata:last_hit_rate, strata:gpu_util per card, and so on. The Monitor tab and a Grafana panel read the same value.
  • Source. Everything comes from Service.metrics() plus three latency histograms observed when a request finishes. With "parallel", the histograms read each request's own status, and running / waiting come from live.running / live.waiting.
  • Tests. serve/test_prometheus.py (6 tests): the text format (one # TYPE per family), the histograms, the names against the JSON, and three requests at once against a batch engine double. Documented in docs/DETAILS.md.

We run it in production behind a Prometheus that also scrapes vLLM servers. The same Grafana panels show both.

3. A monitoring kit (docs/monitoring/)

  • prometheus.yml and servicemonitor.yaml scrape /metrics, with the API key sent as a bearer token.
  • grafana-strata.json is a dashboard to import (it asks for the Prometheus data source):
    • its first row uses only vLLM's names, so it reads a vLLM server too;
    • its second row shows the engine's strata: series.
  • Every query was checked against a Prometheus scraping the 4-GPU server above.

Not in this PR any more: the configuration tools. tools/autoconfig.py and tools/tune.py duplicated what setup.py and tools/calibrate.py already do. The useful part, measuring the batch settings with every slot busy, is now a patch to calibrate.py, posted on #907, which changes the same function.

🤖 Generated with Claude Code

Niko1221 and others added 30 commits October 2, 2026 09:38
…he --family/--model to use (Niko1221#444)

With Unsloth's UD-IQ3_XXS or UD-Q2_K_XL in the folder, setup either said only "has no IQ3_XXS model file"
(--family unsloth --model IQ3_XXS) or took the UD-IQ3_XXS file for the GSQ-RCO IQ3_XXS by the size in its name.
Now:
- a chosen shard whose name is a quantization setup has no size for (UD-IQ3_XXS, UD-Q2_K_XL, a K-quant) stops
  with "<file> is UD-IQ3_XXS, a GGUF Strata cannot run" and what Strata runs: ISTA-DASLab's GSQ-RCO files (and
  Swift 1.5's and the Coder's) and Unsloth's UD-Q4_K_XL only;
- a folder without the chosen file but with other GGUFs lists the ones Strata cannot use and the --family/--model
  of the ones it can (an empty folder keeps check_shards' "missing");
- a size the family does not have names the families that do (--family qwen --model IQ3_XXS), and with
  --gguf-dir also what Strata runs.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…e unpruned sizes for that (Niko1221#438)

README and docs/MODELS.md said only "Weaker outside coding". Niko1221#438 measured Chinese answers wrong or looping on the
Coder where English was fine; both now say plainly that this includes Chinese and other CJK text and point to the
sizes that keep every expert (Q2_0, IQ2_XS, IQ3_S).

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…iko1221#461), the low-RAM resident mode on Windows 32 GB PCs (Niko1221#467), fixes Niko1221#460 Niko1221#459 Niko1221#446 Niko1221#447 Niko1221#457 Niko1221#448 Niko1221#444; setup requires it

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
The fill from the GGUF (--stream-experts) read each expert's three slices,
copied the blob and waited for the copy before the next read: ~12,000-18,500
experts with neither the SSD nor PCIe kept busy. Now the slots are admitted
in profile order first (same placement), reads go in file order by up to 8
threads into page-locked batches of 64, and a batch's copies overlap the next
batch's reads. STRATA_FILL_SERIAL=1 keeps the serial fill.

Tested on the B70: cold page cache, Coder fill 23.4 GiB 76.2 s -> 18.8 s;
launch to first token Coder 82 s -> 26 s, IQ2_XS 120 s -> 41 s; outputs
identical to the serial fill.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
… read (Niko1221#471)

The engine's `strata serve: prompt ...` line now says "= 0 reused + 12288 of 98179 read in ... (cancelled)" with the
rate from the tokens read, and DONE gets a 15th field: the prompt tokens actually read. The server's /metrics history
and totals count reused + read for a cancelled read (prompt_total keeps the whole prompt); an older engine's DONE line
(no such field) is read as before.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
Engine: when the MTP draft head's subset does not fit, it prints what the head needs, the VRAM free, and the smaller
shipped subsets (--draft-vocab cyrillic / en, their MiB from the head's real row size) with how to pick one
(START-HERE.bat / ./setup.sh --draft-vocab en, kept in the model's config). Text only, on the failure path.
Server: a start that ended that way carries the engine's hint (or the same advice for an older engine) in its
"the engine exited before it was ready" error. Setup: a card under 14 GB with the default subset gets a note
recommending --draft-vocab en; nothing is changed, an explicit or saved choice gets no note, and the golden
configs are unchanged.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…o1221#475)

In a git clone they run `git pull --ff-only` (a zip copy is told to download the new zip), then `setup.py --update`:
what a plain start does to an install before it starts the model - the Python packages, the ready-made engine when
MIN_ENGINE is newer (a compiled one when its source changed), each model's config upgrades and its draft subset, the
Windows AMD runtime - and nothing else: no question, no model files, no start. Both scripts run from one parsed
block, since the pull can rewrite the running script. Documented in docs/INSTALL.md (Updating) and the README.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…profile-save, Niko1221#477)

`--expert-profile-save P` (--serve) writes the adaptive tier's learning as a profile in tools/make_profile.py's
format, byte for byte (read_expert_profile reads it back): every pair ranked - the experts resident in VRAM now
first, then the routing counted since the start (the decayed usage summed at each adapt step, before its decay),
then the profile the engine started from. On QUIT / stdin EOF and every --expert-profile-save-every minutes
(default 10, 0 = exit only) between requests, through <P>.tmp and a rename. Without the flag nothing is counted or
written: `heat` stays empty and every added loop is over it.

Server: "expert_profile_save" (and "expert_profile_save_every") in the config pass the flags, and the next start
begins from the saved file instead of the config's --expert-profile when its header matches that profile's
layers x experts. tests/core/expert_profile_save_test.cpp (CPU only) checks the order and the bytes.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…, the draft-head hint (Niko1221#474), UPDATE.bat (Niko1221#475), --expert-profile-save (Niko1221#477); setup requires it

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…pt path, grouped on the GPU (opt-in STRATA_PF_FUSED=1)

Strata's own kernels (src/prefill/moe_fused.cu): the layer's input rounded to int8 once per token (a scale per 32
values), the (token, k) pairs grouped by expert on the GPU (counts, offsets, a 64-row tile list - no host sync),
then per batch of up to 128 experts one gate/up launch (mma.sync m16n8k32 s8, the 2-bit codes unpacked in
registers, Q2_0's offset folded into the activation blocks' sums, SwiGLU and H's int8 rounding in the epilogue)
and one down launch into the per-slot rows the existing combine reads.  Experts are read where they are (cache
slot or streamed ring slot, a pointer per expert): no gather.

Used only on the streamed walk (chunks from stream_all_min()) of the canonical Q2_0 layout on sm_80+ CUDA, with
STRATA_PF_FUSED=1; it reuses the MMQ buffers (Xq, H, GU).  Everything else - the default, native/IQ packs, small
chunks, sm_75, HIP - runs the MMQ path unchanged (the old block is only wrapped in the else branch).

tests/cuda/prefill_fused_moe_test (exit 77 below sm_80), RTX 5070: vs a double-precision reference, MMQ rel RMS
1.144e-2 (worst row 4.15e-2), fused 1.144e-2 (4.15e-2); fused vs MMQ 9.0e-5.  One layer at 2048 tokens x top 10
over 512 experts: MMQ path 11.12 ms, fused 3.13 ms.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…e layer's experts) by default

The fused path launches on batches of a layer's streamed experts, so with the 384-slot ring a layer's last batch
waited for slots its own first batch frees.  With STRATA_PF_FUSED=1 (and the canonical Q2_0 layout) the pinned-RAM
ring is 512 slots; without it nothing changes (384 / 96 as before; STRATA_PREFILL_RING still overrides).

RTX 5070, Q2_0 (cfg-q2rel), greedy, 3 alternating pairs, prompt tok/s: 4K 1,275 -> 1,516 (+18.9%), 32K 2,176 ->
2,422 (+11.3%); at 384 slots the fused path gave +4-7% / +3%.  1024 slots: 4K +24-27%, but the 32K chunk falls
from 8192 to 6144 (+6%).  The default path's output is identical to 0.1.35's on both prompts.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
… hold only what the fused kernels need

In a layout whose chunks run the fused experts (STRATA_PF_FUSED=1, the Q2_0 pack, a streamed chunk) GU holds the
grouping tables, H the int8 H rows and Xq the per-token int8 activations, and Hq nothing - instead of MMQ's FP32 GU / H
and per-slot q8_1 rows: ~100 KB a token less in the MoE set, which then no longer sets the shared region's size (the
GDN set does).  A last chunk below stream_all_min() still runs MMQ in the same buffers, so each keeps MMQ's size for
stream_all_min() - 1 tokens.  The ring may grow to 1024 slots with the fused experts (STRATA_PREFILL_RING; the default
stays 512).  Without STRATA_PF_FUSED nothing changes: the same buffers, ring cap (512) and events.

The fused walk now counts the experts it reads from the cache in experts_resident (it stayed 0): every resident expert
of the layer, routed or not - at a streamed chunk's size nearly all are.

RTX 5070, Q2_0 (cfg-q2rel), --prefill auto, 2 alternating pairs, prompt tok/s (borrowed slots):
  4K:   MMQ 1,305 (1621), fused 1,517 (1749), fused + P3 1,512 (1534), + 1024-slot ring 1,605 (2046)
  32K:  MMQ 2,177 (3071), fused 2,453 (3199), fused + P3 2,485 (2641), + 1024-slot ring 2,660 (3153)
  128K (--max-context 131072 --kv-resident 32768): MMQ 2,108 (chunk 8192), fused 2,140 (6144), fused + P3 2,346
        (8192), + 1024-slot ring 2,392 (6144)
A 16K chunk still does not fit the 12 GB card: at 8192 tokens the loan is 2,641 of 3,690 slots, ~2.1K of them per
8K tokens.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…ault (two layers' experts)

With P3's smaller buffers the 1024-slot ring fits beside an 8192-token chunk at 32K (borrowed 3,153 of 3,690 slots,
under the 90% cap).  RTX 5070, Q2_0 (cfg-q2rel), --prefill auto, prompt tok/s at 512 / 1024 slots (2 alternating
pairs): 4K 1,512 / 1,605, 32K 2,485 / 2,660, 128K (KV streaming) 2,346 / 2,392 - there the chunk falls from 8192 to
6144 and is still faster.  Against MMQ (one more pair on this build): 4K 1,299 -> 1,609 (+23.9%), 32K 2,187 -> 2,656
(+21.5%), 128K 2,128 -> 2,454 (+15.3%).  The ring does not change the numbers (the same output as at 512 slots).
Without STRATA_PF_FUSED nothing changes (384 / 96 slots, cap 512).

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…y argmax, the same ids and tokens

- qsa_block_topk: calls of up to 16 queries (decode windows) run on a cluster of 8 CTAs per query
  (block_topk_cluster_kernel): keys in shared memory, digit histograms pushed through distributed shared memory,
  one cluster barrier per radix pass, a parallel digit search, cluster-wide offsets for the ordered emit. The
  threshold and the cut at it are the one-CTA kernels' (same tie rule, NaN keys 0). STRATA_QSA_CLUSTER=0: off.
  RTX 5070 per call: 21.9 -> 15.6 us at 32K, 58 -> 18 at 128K, 200 -> 22 at 262K.
- sample_tokens greedy without penalties: a cluster of 8 CTAs per row (sampler_greedy_cluster_kernel), four loads
  in flight per thread, results merged in CTA 0 (lowest index on ties, NaN never taken, 0 when none).
  STRATA_ARGMAX_MULTI=0: off. 39.6 -> 5.9 us per call at 248,320 logits.
- Both launch through cudaLaunchKernelEx with the cluster attribute (captured in the decode graphs) only where
  they run: CUDA, compute capability 9+ (STRATA_EMULATE_CC respected), the kernel's own PTX/SASS at sm_90+ (an
  older build's JIT copy is a trap), room for one cluster - checked once per device. HIP and sm_75/86/89 keep the
  one-CTA kernels.
- decode_cluster_parity (ctest; skips below sm_90): bitwise vs the one-CTA dispatch and the reference kernel over
  434 top-k cases (contexts 7..262,144, every tail weight, 1-16 queries, ties / NaN / +-inf / -0 / denormal / all
  NaN) and 216 argmax cases, plus a CUDA graph capturing both (first calls inside the capture); --bench timings.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…gate/up IQ2_XXS IQ2_XS IQ2_S IQ3_XXS IQ3_S IQ4_XS, down Q2_0 IQ4_NL; opt-in STRATA_PF_FUSED=1)

src/prefill/moe_fused_iq.cu: the Q2_0 path's design (moe_fused.cu: the grouping on the GPU, a pointer per expert -
cache slot or ring slot, no gather -, SwiGLU and H's int8 rounding in the gate/up epilogue) for the raw GGUF blobs of
the native packs.  The load stage decodes the i-quant blocks to int8 in shared memory (llama.cpp's codebooks and sign
tables, held in shared memory; a 32-value sub-block per thread per 64 values of K, read one stage ahead with 16-bit
loads and two 256-value blocks ahead into L2), with each sub-block's scale (per 16 values for IQ2_XS / IQ2_S, which
multiply at m16n8k16).  The activations are int8 per 32 values in natural order (quantize_act_native); fragments by
ldmatrix; the magic bias 1.5 * 2^23 enters as the mma's C, so a dot becomes a float with one add.  A work item is 256
weight rows x 64 routed rows, or 128 x 128 when the experts have 56-112 rows each (STRATA_PF_FUSED_TILE=64|128).

Wiring (prefill.cpp): a native layer takes these kernels when both its formats are covered, else MMQ (IQ1_M keeps the
FP16 path); the 512-slot ring of the fused path now also for a native pack with such layers.  sm_75, HIP and the
default (no variable) are unchanged: on IQ2_XS, IQ3_XXS and IQ3_S the default output is 0.1.35's (256/256 greedy
tokens, 4K and 32K).

tests/cuda/prefill_fused_iq_test (exit 77 below sm_80), RTX 5070, 256 tokens x 64 experts against a double-precision
reference from ggml's to_float, rel RMS / worst row - MMQ and fused equal to four digits: IQ2_S/Q2_0 1.489e-2 /
2.93e-2 (fused vs MMQ 2.6e-4), IQ2_XXS/Q2_0 1.464e-2 / 2.96e-2 (1.9e-4), IQ2_XS/IQ4_NL 1.465e-2 / 2.41e-2 (3.5e-5),
IQ3_XXS/IQ4_NL 1.458e-2 / 2.73e-2 (1.6e-4), IQ3_S/IQ4_NL 1.470e-2 / 2.65e-2 (6.5e-5), IQ4_XS/Q2_0 1.446e-2 /
2.72e-2 (4.4e-5); both tile shapes.  One layer (512 experts, top 10), ms, MMQ path -> fused at 2048 / 3584 / 8192
tokens: IQ2_S 10.5 -> 5.2, 12.5 -> 9.3, 21.4 -> 17.5; IQ2_XXS 8.8 -> 4.2, 10.7 -> 7.6, 17.9 -> 13.6; IQ3_S 10.6 -> 5.6,
12.3 -> 8.8, 19.8 -> 18.4; IQ3_XXS 10.2 -> 5.1, 11.9 -> 8.3, 19.1 -> 16.5.  MMQ's products alone (no gathers) at
3584 / 8192: IQ2_S 8.9 / 17.9, IQ3_S 8.1 / 15.6: the kernels are about even with MMQ's, the gain is the path around them.
(The test syncs the device after its uploads: a cudaMemcpy from pageable memory on the legacy stream can still be
queued on Windows while a non-blocking stream runs - the grouping then read zeroed ids, every pair on expert 0.
prefill_fused_moe_test's timing part has the same pattern.)

End to end, RTX 5070, speed_135.py's settings, greedy, A (MMQ) / B (STRATA_PF_FUSED=1) alternating, 2 pairs, prompt
tok/s: IQ2_XS 4K 1,228 -> 1,378 (+12%; a cold first A run of 590 left out), 32K 2,107 -> 2,170 (+3.0%); IQ3_XXS 4K
1,001 -> 1,020 (+1.8%), 32K 1,745 -> 1,739 (-0.4%); IQ3_S 4K 927 -> 925 (-0.2%), 32K 1,382 -> 1,362 (-1.4%).  The IQ3
packs' prompts here wait on the host copies of their unpinned experts (IQ3_XXS 4K: 5,062 of 22,573 streamed, ~3.0 s
of 3.4 s), so faster kernels change little.  Each arm repeats itself (256/256; IQ3_XXS 32K B and IQ3_S 32K A split at
162 / 223, the decode's timing); A vs B greedy agreement, identical prefix of 256: IQ2_XS 204 / 88, IQ3_XXS 182 / 89,
IQ3_S 31 / 256.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…(STRATA_PF_FUSED=0 off), the native IQ packs' fused kernels opt-in (=1), cluster decode kernels; the fused test waits for its uploads

The ring rule: 1024 slots for the Q2_0 pack's fused experts (P3), 512 for native packs with fused layers.
Quality (teacher-forced vs the FP16 path, Q2_0): fused as close as MMQ at 8K, closer at 32K (phaseA-tests s18-tf).

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
device, expert_source, native_head, prefill gemm, prefill.cpp, generate.cpp
merged from dpct of aeb35be and d9ab843 with sycl/tools/merge_upstream.py (the
3-way merge restricted to files upstream changed, dpct kernel-name hashes
canonicalized to the port's). Version 0.1.35.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…ls) and their switches

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
The PCIe probe sets the share of the missed experts the GPU reads over
the link (--pcie-frac): the native default 0.55 from 20 GB/s up, less
on a slower link (Niko1221#44). On Windows HIP it timed its copies with events
that do not bracket them there: an RX 6800 read between 3,348 and
25,999 GB/s in 24 engine starts, so every link kept the x16 share. A
standalone copy of the probe on that card: its events measure 0.13 ms
for the 1 GiB, the host clock 43 ms (25.1 GB/s).

On Windows HIP the probe now times the same copies on the host, between
two synchronizes (1 GiB, tens of milliseconds): the same RX 6800 on
PCIe 4.0 x16 reads 27.2 GB/s in every run, which keeps 0.55 there. On
a PCIe 4.0 x8 or 3.0 x16 link the share now shrinks as Niko1221#44 intended.
CUDA and Linux keep the events.

Claude-Session: https://claude.ai/code/session_015Ld5YDzuKW1hgMVwwgRZDP
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…DGET)

On Windows, hipMemGetInfo (ROCclr's PAL backend) reports the card's size
minus this process's own allocations; what the desktop and other
programs hold on the card is not subtracted. On a card that also drives
the desktop the engine counts the ~930 MiB those hold as free, and
--expert-cache auto fills the card past what Windows lets the process
keep in VRAM: WDDM moves part of it out to system RAM. On an RX 6800
(HIP SDK 7.2) with setup's configuration: a 9.83 GiB cache, the card at
16132 of its 16368 MiB, 929 MiB more of the engine in system RAM than
with this change, 30.1-31.0 tok/s decode. On a 0.1.30 build,
--vram-reserve-mib 1500 left the card as full (16037 MiB), since the
reserve is taken from the same figure.

Windows gives every process a video memory budget that does account
for the others (DXGI QueryVideoMemoryInfo). On that card it is 799 MiB
below the card's size while the process is small and about 1.8 GiB
below once it holds 11 GiB; allocations past it still succeed.
cudaMemGetInfo, which the HIP shim maps to hipMemGetInfo, now lowers
hipMemGetInfo's free figure by what the budget withholds (the card's
size minus the budget), on Windows only. What is sized from it then
stays within the budget: --expert-cache auto sizes the cache from it,
and the existing check after the slots are written shrinks it ("only
8 MiB free once the slots are written ... shrinking"): an 8.05 GiB
cache, the card at 15230 MiB, 40.6-41.5 tok/s.
Lowering HIP's own figure, instead of taking the budget minus DXGI's
usage, keeps memory the HIP runtime has freed but still caches counted
as free, as HIP counts it.

dxgi.dll is loaded at run time, so nothing new is linked; the adapter
is found by the device's LUID. Without dxgi.dll, the adapter or a
budget below the card's size, hipMemGetInfo's figure stands. The first
query logs the budget once. STRATA_WDDM_BUDGET=0 turns it off. CUDA
and Linux builds are unchanged.

Claude-Session: https://claude.ai/code/session_015Ld5YDzuKW1hgMVwwgRZDP
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…g forever (Niko1221#481)

A request whose engine prints nothing for engine_silence_s (default 300 s; while a prompt is read, a chunk may take
3x the previous one's time, the first up to its tokens at 50 tok/s more) ends with EngineSilent (an EngineDied): the
engine is killed and the next request starts it again. The STOP drain, under the request FIFO, is no longer an
untimed wait. "engine_silence_s": 0 keeps the old behaviour. Tests: unit (silence, PP chunks, first chunk, STOP
never acknowledged, 0, config) and over HTTP with a fake engine process (non-stream, stream, stop token).

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…0 log lines (Niko1221#496)

Every start failure, not only the draft-head case (Niko1221#474), now carries the engine's own reason: the last 20 lines
this start wrote to the engine log (never an earlier start's). The failed start also closes its pipes and its
handle on the log. Tests: the tail itself, and a fake engine that exits with a reason in its log.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…p-experts (Niko1221#505)

The server window said "loading the experts into RAM (about N GB)" whatever the flags. Now: --mmap-experts maps them
from the model files (the OS file cache reads them, nothing is loaded up front; its "experts via mmap" line triggers
the message when the model has no PLE line), --resident-experts copies the ones the GPU does not hold, and
--resident-budget-gib N copies up to N GiB. setup's start banner says "maps" for the plain mapped mode too. Tests:
the words per flag set, and the narrator over a mapped and an arena log.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…#495)

hf() takes the host from HF_ENDPOINT (default huggingface.co, a trailing slash dropped); the pinned revisions, the
fallback to main and the SHA-256 checks are unchanged. tools/mtp_fetch.py reads the MTP tensors from the same host.
Step 5 prints the folder setup expects the model files in and how to use files downloaded by hand (that folder with
their original names, or --gguf-dir), and the mirror when one is set. INSTALL.md says the same. Tests: the endpoint
in hf()/hf_unpinned()/mtp_fetch, and step 5's words over a mocked install. The golden configs are unchanged.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…ds it; no budget on a --gpus start (Niko1221#498)

The engine refuses --resident-budget-gib with a layer split, so setup always dropped UD-Q4_K_XL to one GPU. With RAM
>= the GGUF size + UNSLOTH_RAM_LEFT_GB (~135 GB) and 2+ usable cards, setup now asks (one GPU stays the default and
--yes keeps it; an explicit --gpus is honoured; an explicit --resident-budget-gib keeps one GPU) and the split writes
"gpu": [..], "layer_split": "auto" without the budget: all the experts loaded into RAM at start, the files through the
OS file cache (the config Niko1221#498 measured: 2x RTX 3090, 31 -> 64-78 tok/s). Latent bug: START-HERE --gpus 0,1 (or the
start-time offer, or a hand-edited config) on a UD-Q4 config kept the budget and the engine exited with code 2; the
split now drops it, or setup stops before saving when the RAM is too small. docs/UNSLOTH_Q4.md says so. Tests in
tools/test_setup_unsloth.py; the golden configs are unchanged.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…hole)

Reported on Niko1221#793 by rhgo1749 (3 x RTX 5070 Ti, batch 6 / groups 3): a pipelined group's window spans its
slots up to its last busy one, but pick_slot's least-recently-used fallback could put a request in a
group's second slot with the first one free ([_ B]), which costs a whole row - 2 requests fell from
~150 to ~78 tok/s on later waves. The engine keeps no pipelined slot as a conversation cache, so with
pipelined groups nothing is worth keeping free: pick_slot now fills a hole below a busy slot of the same
group first, then follows the spreading order. The engine says how many groups it runs (INFO
batch_groups) so a request the engine ran unpipelined keeps the reuse / LRU choice.
serve/test_slot_alloc.py: the spreading order, the reported hole, 2,000 random arrivals and departures
that never open a hole below the pick, and the one-group reuse / LRU order unchanged.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@blange48
blange48 force-pushed the upstream-metrics-autoconfig branch from fbf5054 to 81fcc9e Compare October 5, 2026 18:53
@blange48 blange48 changed the title Pipelined windows over the active slots, /metrics for Prometheus (vLLM names) + monitoring kit, tools/autoconfig.py Pipelined windows over the active slots (+ slot allocator), /metrics for Prometheus (vLLM names) + monitoring kit Oct 5, 2026
@Niko1221 Niko1221 closed this Oct 6, 2026
@Niko1221

Niko1221 commented Oct 6, 2026

Copy link
Copy Markdown
Owner

Heads-up: today we cleaned up this repository's commit history, and GitHub closed this PR automatically when main was force-pushed (that is the "base branch was force-pushed" event above). It is not a rejection.

GitHub cannot reopen it, because its branch is built on the old history. If you want to keep it going: start a new branch from the new main, cherry-pick your own commits onto it (git fetch upstream && git switch -c my-fix upstream/main && git cherry-pick <your commits>), and open a new PR. Or simply leave it: when we take your change for a release we port it onto the new main ourselves and keep you as its author.

Sorry for the noise, and thank you for the contribution.

@blange48

blange48 commented Oct 6, 2026

Copy link
Copy Markdown
Contributor Author

Reopened on the new main as #1179 (same six commits, cherry-picked onto v0.1.40.1 without conflicts; the allocator fix and its measurements are now in the description). Thanks again @rhgo1749 for finding the leading-hole case.

Niko1221 pushed a commit that referenced this pull request Oct 7, 2026
…names (#793)

Asked with Accept: text/plain or ?format=prometheus; every other request keeps the JSON.
Latency histograms, a Grafana dashboard and scrape configs under docs/monitoring.
Only the Prometheus part of #793; the pipelined-window slot allocator needs a 2-GPU check and is left out.
@Niko1221

Niko1221 commented Oct 7, 2026

Copy link
Copy Markdown
Owner

Thank you @blange48. The Prometheus /metrics endpoint with vLLM's metric names, the Grafana dashboard and the scrape configs are in 0.1.40.2 with your authorship kept. The pipelined-window slot allocator is not in this release: it needs a 2-GPU check we cannot run yet.

0xrjman pushed a commit to 0xrjman/Strata that referenced this pull request Oct 7, 2026
…names (Niko1221#793)

Asked with Accept: text/plain or ?format=prometheus; every other request keeps the JSON.
Latency histograms, a Grafana dashboard and scrape configs under docs/monitoring.
Only the Prometheus part of Niko1221#793; the pipelined-window slot allocator needs a 2-GPU check and is left out.
Niko1221 added a commit that referenced this pull request Oct 8, 2026
…t of its group first (#793)

A pipelined group's window spans its slots up to its last busy one, and the engine keeps no pipelined slot as a
conversation cache, so the least-recently-used pick left leading holes ([_ B]) that cost a whole row each.  pick_slot
now fills such a hole first, then follows the spreading order.  The engine reports INFO batch_groups so the server
knows the groups really run pipelined.  serve/test_slot_alloc.py covers the allocator.

Co-authored-by: blange48 <blange48@users.noreply.github.com>
@Niko1221

Niko1221 commented Oct 8, 2026

Copy link
Copy Markdown
Owner

Thanks. The slot allocator part is in 0.1.41, ported by hand: a request takes a free slot below a busy one of its group first. The rest of this PR was already in the tree.

Buddychek added a commit to Buddychek/Strata that referenced this pull request Oct 9, 2026
…e had

DETAILS.md (troubleshooting): the file-tier watchdog tolerance (Niko1221#1407) - what makes the engine print the no-step-finished-but-still-reading note (Linux major faults / read_bytes, low RAM, --mmap-experts on slow storage), the STRATA_WATCHDOG_IO_S allowance (default 10 x STRATA_WATCHDOG_S, 0 = off, no file-tier reading on other systems); the 413 body cap answered before the read (STRATA_MAX_BODY_MIB, 256 MiB); the 400 for wrong-shape JSON bodies and bad Content-Length, traceback in the server log.

BATCHING.md: --batch-groups G > 1 does not combine with --batch-mtp (Niko1221#1413) - one row per slot, the drafters still hold ~0.9 GB VRAM each, the engine warns at start; the slot allocator fills a leading hole in the group first (Niko1221#793); the four unlabelled code fences get json/bash. Every line checked against src/program/generate.cpp and serve/server.py at v0.1.41; markdownlint (fenced-code-language) finds nothing in the four touched pages. Code untouched.
Buddychek added a commit to Buddychek/Strata that referenced this pull request Oct 10, 2026
…e had

DETAILS.md (troubleshooting): the file-tier watchdog tolerance (Niko1221#1407) - what makes the engine print the no-step-finished-but-still-reading note (Linux major faults / read_bytes, low RAM, --mmap-experts on slow storage), the STRATA_WATCHDOG_IO_S allowance (default 10 x STRATA_WATCHDOG_S, 0 = off, no file-tier reading on other systems); the 413 body cap answered before the read (STRATA_MAX_BODY_MIB, 256 MiB); the 400 for wrong-shape JSON bodies and bad Content-Length, traceback in the server log.

BATCHING.md: --batch-groups G > 1 does not combine with --batch-mtp (Niko1221#1413) - one row per slot, the drafters still hold ~0.9 GB VRAM each, the engine warns at start; the slot allocator fills a leading hole in the group first (Niko1221#793); the four unlabelled code fences get json/bash. Every line checked against src/program/generate.cpp and serve/server.py at v0.1.41; markdownlint (fenced-code-language) finds nothing in the four touched pages. Code untouched.
MistVVK added a commit to MistVVK/XeStrata that referenced this pull request Oct 12, 2026
グループの窓は、そのグループの使用中の最後のスロットまでの行を走らせるので、使用中のスロットの
下が空いていると、その行の分だけ無駄になる(upstream Niko1221#793: [_ B] の形)。会話の続きを持つスロットが
なければ、まずそういう空きを埋める。upstream は fc84443 の前の前提(パイプラインのスロットは
キャッシュにならない)のまま空きを最優先にしているが、XeStrata は fc84443 を入れたので、会話の続きを
持つスロットを先にした。upstream のテスト(serve/test_slot_alloc.py)は移していない。

serve の unittest は未実行(マシンを使えないため)。
giridhart pushed a commit to giridhart/Strata that referenced this pull request Oct 15, 2026
…t of its group first (Niko1221#793)

A pipelined group's window spans its slots up to its last busy one, and the engine keeps no pipelined slot as a
conversation cache, so the least-recently-used pick left leading holes ([_ B]) that cost a whole row each.  pick_slot
now fills such a hole first, then follows the spreading order.  The engine reports INFO batch_groups so the server
knows the groups really run pipelined.  serve/test_slot_alloc.py covers the allocator.

Co-authored-by: blange48 <blange48@users.noreply.github.com>
giridhart pushed a commit to giridhart/Strata that referenced this pull request Oct 15, 2026
SuperDistance added a commit to SuperDistance/Strata that referenced this pull request Oct 17, 2026
…onversation cheapest to lose

With no empty slot, pick_slot takes the least recently used one. A short side request (a title, a one-line question)
then takes the slot of a long conversation that is still going, when that conversation's last turn ended before the
other slot's, and its next turn reads the whole history again.

New, off unless set: STRATA_SLOT_EXPIRE_S (seconds, in the server's environment). Set, the order is an empty slot, then an expired one (idle that
long or longer: its client has likely moved on), then the one holding the shortest history, the least recently used
among equals. Unset or 0, every held slot counts as expired and the least recently used goes, as before. A request
still goes to the slot holding its own prefix first, and the pipelined groups' order (Niko1221#793) is unchanged.
serve/test_slot_alloc: one test per rule, the measured case below, random slot tables checked against the rule and
(off) against the old order, and the variable's values. docs/BATCHING.md says the order and the variable.

Measured on one RTX 5090 (engine from main 296b59d, --batch 3, no conversation cache, a third stream decoding all
along, the rule at 600 s): with a 40,248-token and a 1,777-token conversation held, a one-line side request took the
long one's slot on main, and its next turn read all 40,273 tokens again (12.72 and 12.70 s to the first token); with
the rule it takes the short one's slot and the long one's next turn reuses 40,251 tokens (0.41 and 0.34 s).

Replayed on a 21-hour timeline of my own agent traffic (9,771 requests, 4,503 conversations; ids, times and token
counts only) through 4 / 6 / 8 slots, STRATA_SLOT_EXPIRE_S=600 reads 17.8 / 20.6 / 21.1% fewer prompt tokens than the
least-recently-used order; the best possible order, computed knowing the whole timeline, reads 18.5 / 22.6 / 22.4%
fewer. Thinned to 5-50% of its conversations for smaller setups, 2 to 4 slots read 0.3-22.5% less on average; 1 slot
has nothing to choose. Where it loses: a long conversation whose client has gone keeps its slot until it expires
while busy short chats share the other slots.
rifaniponk added a commit to rifaniponk/Strata that referenced this pull request Oct 19, 2026
v0.1.41 brought GET /metrics in the Prometheus format (Niko1221#793) and the
--batch-groups concurrency work (Niko1221#793, Niko1221#857), but the Monitor tab could
not show either: the latency histograms existed only in the text format
the scraper gets, and nothing read live.slots.

 - First token: the median time to the first token, with p90, the
   request count and the gap between tokens underneath, over a
   sparkline of each recent request's own reading.
 - Batch slots: how many of the --batch slots are busy and what each is
   doing (generating, reading, waiting), with the group count when
   --batch-groups splits them. Without --batch it says so and falls
   back to the queue, so the card is never blank.

To feed them the JSON /metrics grew a `latency` block and each finished
request a `ttft_s`. The percentiles come from the same buckets the text
format serves, derived the way promql's histogram_quantile derives
them, so a card and a Grafana panel cannot disagree - a test pins the
JSON's count to vllm:time_to_first_token_seconds_count.

Measured on an RTX 5080 against qwen3.8-flash-next-iq3_s: ttft p50
0.10 s, itl p50 22.6 ms, and the text format's sum (1.403293) is the
two requests' ttft_s (0.083 + 1.32) to the microsecond.

The 3D PC is left alone: its channels each drive a physical part, and
neither a latency nor a slot count has an honest analogue there. The
other eight cards, the telemetry sampler and the Prometheus output are
unchanged.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
SuperDistance added a commit to SuperDistance/Strata that referenced this pull request Oct 20, 2026
…onversation cheapest to lose

With no empty slot, pick_slot takes the least recently used one. A short side request (a title, a one-line question)
then takes the slot of a long conversation that is still going, when that conversation's last turn ended before the
other slot's, and its next turn reads the whole history again.

New, off unless set: STRATA_SLOT_EXPIRE_S (seconds, in the server's environment). Set, the order is an empty slot, then an expired one (idle that
long or longer: its client has likely moved on), then the one holding the shortest history, the least recently used
among equals. Unset or 0, every held slot counts as expired and the least recently used goes, as before. A request
still goes to the slot holding its own prefix first, and the pipelined groups' order (Niko1221#793) is unchanged.
serve/test_slot_alloc: one test per rule, the measured case below, random slot tables checked against the rule and
(off) against the old order, and the variable's values. docs/BATCHING.md says the order and the variable.

Measured on one RTX 5090 (engine from main 296b59d, --batch 3, no conversation cache, a third stream decoding all
along, the rule at 600 s): with a 40,248-token and a 1,777-token conversation held, a one-line side request took the
long one's slot on main, and its next turn read all 40,273 tokens again (12.72 and 12.70 s to the first token); with
the rule it takes the short one's slot and the long one's next turn reuses 40,251 tokens (0.41 and 0.34 s).

Replayed on a 21-hour timeline of my own agent traffic (9,771 requests, 4,503 conversations; ids, times and token
counts only) through 4 / 6 / 8 slots, STRATA_SLOT_EXPIRE_S=600 reads 17.8 / 20.6 / 21.1% fewer prompt tokens than the
least-recently-used order; the best possible order, computed knowing the whole timeline, reads 18.5 / 22.6 / 22.4%
fewer. Thinned to 5-50% of its conversations for smaller setups, 2 to 4 slots read 0.3-22.5% less on average; 1 slot
has nothing to choose. Where it loses: a long conversation whose client has gone keeps its slot until it expires
while busy short chats share the other slots.
rifaniponk added a commit to rifaniponk/Strata that referenced this pull request Oct 20, 2026
v0.1.41 brought GET /metrics in the Prometheus format (Niko1221#793) and the
--batch-groups concurrency work (Niko1221#793, Niko1221#857), but the Monitor tab could
not show either: the latency histograms existed only in the text format
the scraper gets, and nothing read live.slots.

 - First token: the median time to the first token, with p90, the
   request count and the gap between tokens underneath, over a
   sparkline of each recent request's own reading.
 - Batch slots: how many of the --batch slots are busy and what each is
   doing (generating, reading, waiting), with the group count when
   --batch-groups splits them. Without --batch it says so and falls
   back to the queue, so the card is never blank.

To feed them the JSON /metrics grew a `latency` block and each finished
request a `ttft_s`. The percentiles come from the same buckets the text
format serves, derived the way promql's histogram_quantile derives
them, so a card and a Grafana panel cannot disagree - a test pins the
JSON's count to vllm:time_to_first_token_seconds_count.

Measured on an RTX 5080 against qwen3.8-flash-next-iq3_s: ttft p50
0.10 s, itl p50 22.6 ms, and the text format's sum (1.403293) is the
two requests' ttft_s (0.083 + 1.32) to the microsecond.

The 3D PC is left alone: its channels each drive a physical part, and
neither a latency nor a slot count has an honest analogue there. The
other eight cards, the telemetry sampler and the Prometheus output are
unchanged.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
rifaniponk added a commit to rifaniponk/Strata that referenced this pull request Oct 26, 2026
v0.1.41 brought GET /metrics in the Prometheus format (Niko1221#793) and the
--batch-groups concurrency work (Niko1221#793, Niko1221#857), but the Monitor tab could
not show either: the latency histograms existed only in the text format
the scraper gets, and nothing read live.slots.

 - First token: the median time to the first token, with p90, the
   request count and the gap between tokens underneath, over a
   sparkline of each recent request's own reading.
 - Batch slots: how many of the --batch slots are busy and what each is
   doing (generating, reading, waiting), with the group count when
   --batch-groups splits them. Without --batch it says so and falls
   back to the queue, so the card is never blank.

To feed them the JSON /metrics grew a `latency` block and each finished
request a `ttft_s`. The percentiles come from the same buckets the text
format serves, derived the way promql's histogram_quantile derives
them, so a card and a Grafana panel cannot disagree - a test pins the
JSON's count to vllm:time_to_first_token_seconds_count.

Measured on an RTX 5080 against qwen3.8-flash-next-iq3_s: ttft p50
0.10 s, itl p50 22.6 ms, and the text format's sum (1.403293) is the
two requests' ttft_s (0.083 + 1.32) to the microsecond.

The 3D PC is left alone: its channels each drive a physical part, and
neither a latency nor a slot count has an honest analogue there. The
other eight cards, the telemetry sampler and the Prometheus output are
unchanged.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
wannagoxuf pushed a commit to wannagoxuf/Strata that referenced this pull request Oct 26, 2026
上游 pick_slot 的多槽分支明确是"先补组内空洞"(Niko1221#793 的决定, serve/test_slot_alloc.py
有测试编码它), 而"先摊到各组"与之正面冲突 —— 上游的 CI 跑 serve/ 全量测试, 我们那版
改动会让上游测试红(且它还引入了对 self.slot_group 的依赖, 打破了上游测试 fixture 的约定)。

⇒ 本 PR 收回服务端改动, 保持纯引擎改动(PR Niko1221#4121 的价值在引擎侧)。
服务端摊槽保留在我们自己的部署里(实测需要), 并向上游单独说明, 由维护者决定要不要单独收。
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.