Repository navigation
bench: community report — RTX 5090, IQ3_S at 1M context (YaRN), agent-style workload - #466
gravitomagnetic wants to merge 534 commits into
Conversation
…es bits Item 8 verification. The parity section now runs all three scaling configs - none, linear 2 (angle shrink only), yarn 2 (ramp + mscale) - and "bit-identical" is asserted on the BIT representations (memcmp of the state words), not float !=, which would pass a NaN-payload or -0.0 flip. Each scaled variant must also differ from the none baseline in pooled-key bits, so the parity cannot pass because both sides ignore the scaling. Green on this tree: 0 of 33409 state words differ in every config; 128 (linear) and 192 (yarn) pooled words move. Co-Authored-By: Claude <noreply@anthropic.com>
v0.1.29's new draft_kv site (the 0.1.26 batched draft prompt pass) kept upstream's float freq_base form through the rebase, which both misses the scaling on the draft K/V and does not compile against the RopeScaling signature the other four sites carry. It passes rope_scaling() now, like the prompt K, Q and indexer-query sites. Co-Authored-By: Claude <noreply@anthropic.com>
…correction (TODO 22) The note fired for every scaled run inside the trained 262144 and its text named YaRN's magnitude correction even for linear, whose resolved mscale is exactly 1. Split in two: the in-window warning now only says that no position needs the extension while the resolved scaling still applies to every angle, and the magnitude explanation became its own note, printed only when YaRN is resolved and its magnitude multiplier differs from 1 (yarn factor 1 inside the window stays quiet; linear never sees it). The main rope scaling line is unchanged, the numbers it prints already carry the resolved mscale. Diagnostics only. Checked on the rebased tree: none prints nothing, linear 2 prints the in-window note without the magnitude line, yarn 2 prints both (mscale 1.069315), yarn factor 1 prints the in-window note only, and yarn 2 past the trained window prints the magnitude note without the in-window one. Co-Authored-By: Claude <noreply@anthropic.com>
Move core diagnostics from the dependent NVMe work (4d73bb3) into the shared-core review branch. Read back authoritative bytes and resident ring pages without rewriting state; remove a redundant GPU fixture refill that could mask restore failures. Bound hash workspace and check synchronization/copy errors. Opt-in diagnostics record batched versus token draft prefill. Validation: current core translation units pass C++ syntax checks; 1188 host transfer checks pass with ASan/UBSan. Real GPU and full-model validation remain pending the benchmark window.
Replay identical requests and require token, known-answer, reuse, resolved-setting and comparable main-state parity. Retain the complete pooled-state diagnostic and add a completed-row fingerprint matching upstream. Add negative controls that remain active under optimized Python.
Grow snapshot buffers in segments, track the first rewritten token, and charge retained active storage to the existing RAM budget. Refresh running state, partial pages and draft boundary state. Add full-capture versus incremental growth/rewind checks and report snapshot sizes with usable pressure-test budget ranges.
Reserve accounted space for later appends without moving prior payloads. Prefer full capture when optional capacity would evict useful conversations, and report the actual additional RAM required. Exercise 4096 small appends and non-mutating admission checks under sanitizers.
…28% on 1-4K-token prompts, the same output; STRATA_PREFILL_STREAM_MIN overrides Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…rd and a before_load hook Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
… sharing the GPU Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…load() (the idle thread cannot unload between load and run); unload() sends QUIT before terminate; the VRAM guard checks the first GPU (docs); Niko1221#163's test engine answers poll() Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…TA_EXECV so the server is PID 1, SIGTERM stops it cleanly, KV/GPUS/LOW_RAM as env vars Merged by hand onto 0.1.29: SIGTERM takes Niko1221#153's Ctrl+C path (the timed sleep loop and the second-Ctrl+C kill stay); SIGINT keeps Python's handler; the image builds RTX 20 (75) too. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
Two gates stand between this engine and a 6.0 card, plus one attribute the card does
not honour. None of them needs the kernels rewritten:
* `CMakeLists.txt` refuses to configure below sm_75;
* `__dp4a` (6.1+) appears 29 times in the i-quant kernels (`native_mmvq.cu`,
`s2_expert_grouped.cu`, `iq_kernels.cu`) and `__nanosleep` (7.0+) in the three
doorbell waits (`elementwise.cu`, `verify_kernels.cu`);
* the Turing port sizes `fused_gr_read_multi`'s slice from
`cudaDevAttrMaxSharedMemoryPerBlockOptin`, which on an sm_60 answers 65536 - a card
whose launch validator enforces the 49152 B per-block limit instead. The capacity
comes out 6 tokens and every 6- or 8-token window fails the launch with
"invalid argument".
## The changes
* **`include/strata/kernels/dp4a.hpp` (new)** - `STRATA_DP4A` and `strata_spin_pause()`.
The dp4a fallback is llama.cpp's own (`ggml/src/ggml-cuda/common.cuh`): both operands
read as SIGNED bytes, accumulated into int32 - which is what `__dp4a` does for exactly
these call sites, the kernels here being transcribed from `vecdotq.cuh`/`mmvq.cu`. The
fallback is bit-exact, not approximate; the nanosleep sites only pace single-thread
doorbell waits, so on sm_6x the loop spins. Both are `__CUDA_ARCH__` selections, so a
HIP build takes neither: its `__dp4a` and `__nanosleep` are `hip_compat`'s own.
* **`CMakeLists.txt`** - `-DSTRATA_EXPERIMENTAL_SM60=ON` relaxes the architecture gate,
and the refusal message names it. The UTF-8 BOM stays.
* **`fused_gr.cu`** - below sm_70 the slice capacity comes from
`cudaDevAttrMaxSharedMemoryPerBlock` (49152 B -> 4 tokens) instead of the opt-in. The
Turing port's slicing, its no-opt-in fallback and the HIP expression are unchanged.
On sm_70 and newer this is a no-op: the opt-in is honoured there, so the capacity is the
value it already was, and `STRATA_DP4A` and `strata_spin_pause()` compile to `__dp4a` and
`__nanosleep` exactly as before. Building for a 6.0 card needs a CUDA release that still
targets Pascal (CUDA 13 removed it; 12.x has compute_60), and `CMAKE_CUDA_STANDARD` has
to stay 17 for it: nvcc 12.9 refuses `-std=c++20` when the host compiler is VS2019's `cl`
and silently falls back to C++14, which rejects the nested
`namespace strata::kernels {` declarations.
Rebased onto 0.1.27 (0.1.21 -> 0.1.24 -> 0.1.27). The 0.1.21 form of this patch clamped
the MTP and verify windows instead, because the fused GR read could not launch a 6-token
window at all. Now that `fused_gr_read_multi` slices the batch itself, the clamp, its
`fused_gr_max_window` helper and the two `generate.cpp` call sites are gone, and the full
window runs:
strata verify: window up to 6 tokens, 69.1 MiB of device buffers
strata verify: captured the 6-token window (upload no error, sync no error)
Measured on a 6.0 card (GP100-class, four dies through the layer split, the model's
IQ3_XXS quant, an 883-token prompt, 256 generated, the shipped `--spec 4`): every one of
the 24576 experts resident, 100% decode cache hit rate, 110 tok/s prefill and 21.3 tok/s
decode (warm median). This port on the 0.1.21 base measured 110.6 and 17.9 in the same
session with the same prompt, so the full window is worth about 19% of decode here. One
RTX 3090 Ti with the release engine measures 305 and 79 tok/s for the same model. The
port is about fitting, not speed: decode is bounded by Pascal's per-layer arithmetic, not
by capacity or memory. Only the 6.0 path was exercised on hardware.
), setup recommends instead of forcing (Niko1221#406 Niko1221#403 Niko1221#364 Niko1221#384), fixes Niko1221#352 Niko1221#365 Niko1221#369 Niko1221#371 Niko1221#375 Niko1221#393 Niko1221#408 Niko1221#414; setup requires it Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…the HIP build) Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
# Conflicts: # serve/test_server.py
# Conflicts: # README.md # docs/AMD_HIP.md # setup.py
…iko1221#325), an MCP server so an AI assistant can set Strata up (Niko1221#133), a shorter README; setup requires it Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
Niko1221#430, Niko1221#431, jkuepker) A non-streamed request wrote nothing until its end, and a streamed one only a keep-alive per prompt chunk (the first write after a hang-up usually still succeeds), so a dropped request kept the engine reading the whole prompt or generating to max_tokens while the next request waited. A watcher per request checks the client's socket for EOF every 0.5 s and cancels; the engine then stops at its next prompt chunk. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
… a resumed download needs room only for the rest (Niko1221#425) Niko1221#429 (jkuepker): the HIP engine has no prompt kernels for its Q4_K/Q5_K experts and it has not been run on AMD - --check says so, and setup asks (--model UD-Q4_K_XL --yes tries it). Niko1221#425 (jctaborda): the finished shards and the .part files already on the disk count against the download's size. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…path instead of aborting (Niko1221#420, qni-live) llama.cpp's tile choice aborts the process ("J_best=0") when no tile size of a type fits the card's shared memory. The plan now asks the same question for every visible GPU (the experts' gate+up and down shapes, the draft layer's Q8_0 projections) and keeps a product with no tile on the existing FP16 path, saying which type, card and limits. Where every product fits (all cards tested here), the plan is the same as before. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
… over A/B/A/B), WMMA +29-34% with TheRock 7.10 (Niko1221#432, jkuepker) 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
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…heck before the download Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…#410, measured) Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
… RAM (Q2_0, IQ2_XS, the Coder) Measured with setup's own install and the README's speed method on the AMD box. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…t (YaRN), agent-style workload Four prompt sizes (4k/32k/128k/512k), three fresh runs + warm repeat + greedy instruct run each. Engine-reported timings throughout. Includes a corrected note: an apparent 65s per-request stall was disproven by a controlled probe and re-attributed to serial scheduling behind a co-tenant 208k-token prefill.
…adline numbers Ran the repo's own tools/needle_bench.py against the same v0.1.31 server (IQ3_S, YaRN 1M). All five haystacks found. Wall times flagged as upper bounds due to co-tenant agent traffic (Note A).
|
Hi @Niko1221 I took a closer look at the ~60-65 s blocking observed when another request is doing a large 200K+ prefill, and I think there may be a middle ground between the current strictly serial serving model and full multi-request / multi-slot concurrency. From what I can see, the current behaviour is expected from the server architecture: "serve/server.py" holds the service FIFO for the whole request, including both prefill and decode. "StrataEngine" also explicitly assumes that requests are serialized, which is why a single stdin/stdout pipe is sufficient. So with a request like: A: 208K cold prefill B cannot start until A finishes its complete request. This creates head-of-line blocking, and a long prompt can therefore delay an unrelated interactive request for tens of seconds. I understand the reasoning in #465 for not implementing full concurrent decoding. Supporting several live sequences at once would require per-slot KV, draft state, checkpoints, scheduling, shared expert-cache handling, etc., and would be a much larger architectural change. However, I wonder whether a much smaller form of cooperative prefill preemption could be useful. The prefill path already has natural chunk boundaries: for (int64_t c0 = 0; c0 < n; c0 += m.T) { and near the end of each chunk there is already a synchronization / callback point: cudaStreamSynchronize(m.cs); if (on_stage_chunk) if (on_chunk) So instead of trying to execute multiple sequences concurrently, the idea would be: request A: request B arrives In other words, the engine would still execute only one sequence at a time. This would be preemption between safe prompt chunks, not cross-request batching. Conceptually: current: proposed: What makes this seem potentially practical is that Strata now already has most of the state save/restore machinery through the conversation cache. A parked conversation snapshot already contains the important persistent state, including:
and there is already a "conversation_snapshot_restore(...)" path. So rather than introducing several permanently active GPU sessions, a long prefill could potentially be suspended at a chunk boundary, its state parked using the existing snapshot infrastructure, a queued request serviced, and then the original request restored. There are some details that would need care before a chunk can be considered a safe preemption point. For example, in "Prefill::run()": ss.ple_prev[0] = prev[0]; is currently committed after the complete prefill. If suspension is allowed between chunks, this state would need to be committed at the boundary before taking the snapshot. Similarly I think a safe yield point would need to guarantee that all asynchronous work associated with the chunk is finished, including at least: compute stream before parking the session. A possible interface could be something like: enum class PrefillAction { std::function<PrefillAction(int64_t position)> on_chunk_boundary; with the normal fast path simply returning "Continue". The scheduler could then behave roughly as: if no other request is queued: if another request is waiting: This should avoid changing the decode execution model at all. One further possibility would be to keep large chunks for throughput while the server is idle, but reduce the scheduling quantum when another request is waiting. For example, with ~3000 prompt tok/s: 8192 tokens ~= 2.7 s So even without true concurrency, the worst-case additional latency for an interactive request could potentially fall from ~60 s to approximately one prefill chunk plus snapshot/restore overhead. I would not suggest releasing the existing FIFO directly in the middle of "Prefill::run()", because the current engine owns one live "SessionState" and a second request would overwrite KV/GDN/QSA/PLE/draft state. The important part would be to make the current sequence explicitly suspendable and restore it before continuing. So the proposal is intentionally narrower than #465: not: but: I think this could preserve Strata's current single-sequence execution model while addressing the specific long-prefill head-of-line blocking seen in the 5090 benchmark. If this direction sounds reasonable, I would be interested in investigating it further and prototyping the smallest version first, probably only allowing preemption during prefill and only at an existing chunk boundary. |
|
Update (2026-10-04): re-probed on 0.1.39 — the regression this report documented is closed, and we measured the preemption path @j-luwierski's analysis targets. Same box, same protocol (unique-salt cold prompts, 3-run medians, engine-reported timings, no other traffic). 0.1.39 source build (commit 6f32ec0, CUDA 13.4, driver 595.91.07,
0.1.39 beats the 0.1.31 golden baseline at 32K (+9%) and 128K (+5%). Decode 110–163 t/s. Parallel requests (
Full config dump in the PR body; happy to add the 0.1.39 row to the results table if you want this report updated in-place rather than as a comment. |
|
Preemption A/B: measured head-of-line blocking vs Follow-up to @j-luwierski's analysis above and the #465 shipping note. Setup: same box as my bench report (5090 32 GB, driver 595.91.07, 0.1.39 commit 6f32ec0, IQ3_S MoE, 46.84 GiB experts RAM-streamed, 1M ctx YaRN×4, kv int8,
Preemption works: 5.3× (thinking) / 6.5× (no-think) reduction in interactive latency. Under FIFO, B waits out A's prefill and most of its decode (rep 1: A wall 62.1 s, B queue 57.1 s) — exactly the prefill+decode hold @j-luwierski described. Under Cost to the big request: prefill 3,471–3,483 → 3,359–3,405 t/s (−2%); decode 115.4–117.3 → 101.7–104.3 t/s (−13%); A's total wall +2 s. Verdict for single-big-GPU agent boxes: preemption converts "the agent ignores you for a minute" into "answers in seconds while a 125K prefill runs" — the felt difference whenever a background/voice/daemon client collides with a big agent turn (this box runs a persistent chat daemon + the agent on one engine; those were the collisions). Raw data: 4 JSON files ( |
|
Heads-up: today we cleaned up this repository's commit history, and GitHub closed this PR automatically when 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 Sorry for the noise, and thank you for the contribution. |
|
Thanks for the report. It is in 0.1.40.2 as bench/results/2026-10-02-community-rtx5090-1m-agent, with your authorship kept. Raw dumps and logs were trimmed (see TRIMMED.md in the folder); the README, summaries and configs are unchanged. The row is in bench/results/COMMUNITY.md. |
Results-only community benchmark per docs/COMMUNITY_BENCHMARKS.md. No engine changes.
Hardware: RTX 5090 32 GB (Gen5 x16), Ryzen 9 7950X, 122 GB RAM, Ubuntu 26.04.1, driver 595.91.07, source commit 9259cad (v0.1.31), CUDA 13.4.2.
Model/config: Qwen3.8-Flash-Next IQ3_S (split GGUF, sha256 recorded), --max-context 1048576 --rope-scaling yarn --rope-scale 4, --kv int8 --kv-resident 32768, --expert-cache 7500 (auto-fills to 9,804 slots), MTP --spec 4 --spec-min-p 0.5, vision on (no images in requests).
What's measured:
Honest caveats (Note A in the README): an apparent ~65 s per-request stall in early drafts was disproven by a controlled probe and re-attributed to serial scheduling behind a co-tenant 208k-token prefill (this machine runs a live 1M-context agent session alongside the benchmark). Client-side TTFT figures are contaminated by that traffic; engine per-request timings are unaffected. The surviving observation — one in-flight 200k+ prefill blocks all other clients for up to ~60 s — may interest whoever works on serve scheduling; happy to open a separate issue with logs if useful.
Files: README.md (report), agent_bench.jsonl (20 runs), needles.json (recall), bench_agent_workload.py (script), strata-iq3_s.json (config, paths genericized).