Skip to content

bench: IQ4_XS on a 64 GB PC - pinned arena vs experts.bin mmap vs the RAM budget - #416

Closed
pipeob0 wants to merge 497 commits into
Niko1221:mainfrom
pipeob0:gh/bench-iq4xs-64gb
Closed

pipeob0 wants to merge 497 commits into
Niko1221:mainfrom
pipeob0:gh/bench-iq4xs-64gb

Conversation

@pipeob0

@pipeob0 pipeob0 commented Oct 1, 2026 •

Copy link
Copy Markdown
Contributor

What

A benchmark entry for the case the docs do not cover: a model whose routed experts are bigger than
the machine's RAM
, on an AVX-2-only CPU. orcarouter/Qwen3.8-Flash-Next-Uncensored-GGUF (IQ4_XS)
has 60.94 GiB of routed experts on a 64 GB PC, and the engine offers three ways to serve them.
I measured all three on the same file, plus the knobs that decide whether the model is usable at all.

Ryzen 7 5700X3D (Zen 3, no AVX-512), 64 GB DDR4-2666, RTX 5060 Ti 16 GB at PCIe 4.0 x8 (the engine's
own probe: 14.1 GB/s host→device), Windows 11, engine 0.1.31 and 0.1.32 built from source.

Headline

path tok/s
--mmap-experts over the pack's experts.bin + --kv-resident 20480 + STRATA_IQ_PREFETCH=16384 37.5 – 39.0
pinned arena (the default, no --mmap-experts) 33.85 warm / 27.90 cold
GGUF read in place + --resident-budget-gib 40 26.0 – 27.5
experts.bin + --resident-budget-gib 40 24.7 – 25.4
GGUF read in place, no budget 9.7
reference: the same model at IQ3_S, arena 51.5

4× between the best and the worst way to run the same file. Four mechanisms behind the numbers,
each quoted from the engine's own log in the README:

  • The arena does not pin, and at this size the ranking flips. The comment in generate.cpp
    measures the arena at 1.79× better than warm mmap, on a 34 GB experts.bin on 63 GB, and it already
    says the arena is not pinned there (cudaHostRegister on 31.64 GiB fails). The same failure here at
    60.94 GiB (SetProcessWorkingSetSizeEx(14814 MiB) failed (error 1450), 37 slices, large pages
    refused too) has a different consequence: the arena pays ~30 s of load on every start and now loses
    to warm mmap. cudaHostRegister succeeds for the IQ3_S arena (PORTABLE ok on 50,295,996,416 B),
    so the arena is right one quantization step down and wrong one step up.
  • The RAM budget costs ~30% against plain mmap, and both costs are visible: adaptive tier 2688 experts swapped into the VRAM tier (every 4 rounds, 7.580 ms/round), and the CPU pool reads the
    page-locked complement at 15.5 GB/s against 27.4 GB/s from the OS page cache.
  • A page-locked budget evicts the page cache the same model reads from. Same arm, same arguments,
    two rounds with two 40 GiB budgeted runs in between: 38.06 tok/s / 27.4 GB/s → 18.59 / 10.6.
  • The prefetch distance is not a lever. Turning STRATA_IQ_PREFETCH off, or setting it to 8192,
    moves nothing: off / 2048 / 8192 land within 0-3% across three models and both expert paths. The
    lever is --kv-resident (20480 instead of the 32768 default, ~15% on this model).

Also reproduced: #403 from the other side (budget 50 refused after a run at the same budget
succeeded — 40 is the ceiling on 64 GB here), and #369 (--expert-cache-per-layer does not start on
a native IQ pack in mmap). STRATA_LOOKAHEAD=0 changes nothing at this size. Context 220000 costs
2.2%.

Method

Offline strata generate --tokens-file <frozen prompt> --max-new 300 --stats, arms interleaved with
the arm order rotated one position per round, warmup passes not counted. Compared as ms per
draft round at equal draft-round counts
, and as total time for the run when the arms differ in
draft acceptance: the GPU hit path rounds differently from the CPU (the engine says so at start), so
the same arm with the same arguments produced 112–115 rounds in five runs and 127–133 in two, which
alone moves tok/s by ~10%.

Three measurement traps worth recording, all hit during this work and all fixed in the final numbers:

  • the first arm of a round pays the page-cache warming on the mapped path, so later arms look better;
  • one warmup pass does not warm a 61 GB file (the first measured round still read at 9.7 GB/s);
  • an arm that differs in two flags at once measures neither. The first version of this entry
    reported the prefetch default as ~9% low; the 2048 B arm in that comparison also ran with the
    default --kv-resident 32768. Holding kv-resident constant, the prefetch effect disappears.

Which version, and what has since moved. Every number here is from 0.1.31 and 0.1.32 built from
source. 0.1.33 changed the --resident-budget-gib path (#403: a clamped budget now leaves a margin and
the safety check re-reads free RAM) and made --expert-cache-per-layer start on native packs (#369).
The budget-mode columns (24.7-27.5 tok/s) therefore describe the code as it was before that fix; the
mapped-mode numbers, which are the point of the entry, do not involve the budget path at all.

Corrections after posting

Two claims in the first version of this entry were wrong, and the history of the branch says so:

  • I wrote that the source comment does not know the arena fails to pin. It does: it reports
    cudaHostRegister on 31.64 GiB failing, and its 1.79x/3.7x numbers are measured at 34 GB of experts
    on 63 GB. What this entry adds is that at 61 GiB the ranking flips.
  • I wrote that the prefetch default is ~9% low on the mapped path. It is not: that arm also differed
    in --kv-resident. Section 4 has the re-run on three models.

Caveats

One machine, one model, offline mode (no server, no KV streaming across requests), 300 tokens per
run, greedy + --spec 4. Reproducible to ±1.5% at fixed draft-round counts; the cold columns depend
on what the previous run left in RAM.

Niko1221 and others added 30 commits September 30, 2026 11:18
…7 left 16 GB cards below the stall line), a cancelled request no longer fails the next one, the API key protects /status; setup requires it

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

df6980d put the software prefetch in iq_avx512.cpp, so it only reaches CPUs with AVX-512; the AVX-2 kernels, which is what a Zen 3 / Haswell / Skylake-client box actually runs the i-quant rows on, had none.  Same env, same default, same hint: STRATA_IQ_PREFETCH is the distance in bytes, 0 = off, default 2048.

Measured on a Ryzen 7 5700X3D (AVX-2 only, 64 GB DDR4-2666, 4 KB pages - large pages refused) with IQ3_S decode, 4 rounds interleaved, outputs identical:
  gate/up phase   13.571 -> 13.024 ms/round   (-4.0%)
  rows            25.1   -> 26.1 GB/s         (+1.0)
  whole round     45.54  -> 44.94 ms/round    (-1.3%)
and with the pool cold (first round after the arena is built) the round is -2.4%.  The same 4 rounds with STRATA_IQ_PREFETCH=0 land within 0.1% of the unpatched binary, so the patch itself costs nothing.  native_expert_parity: 0 failures with the prefetch on and off.

The non-temporal hint was measured worse than T0 at the same distance (-0.7% vs -1.3%, +0.5 vs +1.0 GB/s), so this keeps T0.  4096 is a little worse than 1024/2048, which is why the default stays 2048 as in the AVX-512 file.
…(bit-identical; STRATA_GDN_PIPELINE=0 = old kernel)
…'s queries (bit-identical; STRATA_SCORES_MULTI=0 = old grid)
The sampled path (penalties, top_k, top_p, min_p, temperature, draw) kept one
row on one SM: sampler_kernel excluded the logits already kept by sweeping the
kept list for every logit of every round, O(k^2 x n_vocab) shared-memory
compares per row, and ran the double-precision tail on all 1,024 threads.

Three implementations now pick the same token:

- Split (new default): stage 1 cuts each row into 4,096-logit blocks (61 per
  row at 248,320 logits), each warp keeps the top_k of its 1,024 penalised
  logits in registers, and the block merges its four warp lists. Stage 2, one
  warp per row, merges the block lists and runs the tail on one warp.
- sampler_one_block_kernel: one block per row, top_k rounds that each take the
  logits strictly after the previous pick in the selection order (value
  descending, id ascending). It is the fallback when the split cannot run
  (stream under graph capture, vocabulary above 262,144, more than 64 rows,
  no scratch) and is selectable with STRATA_SAMPLER_ONE_BLOCK=1.
- sampler_kernel, the previous kernel, kept unchanged as the reference and
  selectable with STRATA_OLD_SAMPLER=1.

The list, its order and the pick are unchanged: the first k of a union lie
within the first k of each part, and a merge of ordered lists is ordered. The
tail keeps the same arithmetic, with the exps and quotients shared across one
warp and the ordered sums and scans run by lane 0.

The block lists live in a per-stream scratch buffer grown on demand. A
grown-out buffer is retired rather than freed, and a failed grow is
remembered so later calls of that size do not retry cudaMalloc.

sampler_parity is the parity test. Fixtures 16 to 18 pin the whole top_k list
position by position under ties (+inf, -inf, NaN, +0/-0, sentinel id 0,
penalised tokens landing on other tokens' values), tied draws through the whole
chain, and the automatic fallbacks (graph capture, row cap). Each pick is an
integer token id compared for equality with a host reference. ctest runs the
binary once per path: sampler_parity (split), sampler_parity_one_block and
sampler_parity_old, so all three paths are held to the same expected ids.

hip_compat/cuda_runtime.h: map cudaStreamIsCapturing, cudaStreamCaptureStatus
and cudaStreamCaptureStatusNone to their hip equivalents, which the new
capture check needs on HIP builds.
… the server

Closing the console window (or a crash, or Task Manager) ends the server without its cleanup, and strata --serve,
strata-vision and the MCP servers kept running with their VRAM and RAM. Each child now goes into a job object that
kills everything in it when its only handle - the server's - closes; what the children start joins it too. The
server itself stays out, so a browser --open starts is not tied to it. A no-op elsewhere.

serve/test_winjob.py: a contained child and its grandchild end when the process that contained them is killed.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…decode captured or not, prefill's pooled16); no box-specific note (review)
…TRATA_GDN_CB) dropped, one kernel on upstream's split (review)
…mlock succeeded else 'loaded (not locked)', neutral comments + --help (review)
Under WSL2 the NVIDIA driver can pin/map only about 1 GiB of host memory in
total (dmesg: "dxgk: create_existing_sysmem: establish_gpadl failed: -122").
With --expert-cache-device1/2 the startup ran out of it twice:

- CUDA2's context is created after CUDA0's weights and MTP, whose mapped host
  buffers have spent the budget: "CUDA2 experts: cudaSetDevice(2) failed: out
  of memory". STRATA_EARLY_REMOTE_CONTEXTS=1 now creates every secondary
  context up front, like CUDA1's (opt-in: the default order is unchanged).
- With three contexts up, the 260 MiB mapped native embedding no longer fits:
  "native embedding: cannot pin 260 MiB". It now falls back to a device copy
  on the current GPU instead of failing the start.

Measured in Docker Desktop (WSL2), RTX 5070 Ti + 2x RTX 5060 Ti, Swift 1.5
IQ3_XXS: starts and serves with both helper caches (5000 experts each).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Display both rates in t/s with independent graph scales and a darker
prefill curve. Expose prefill_tok_s_mean using the engine's existing
rate, excluding cached tokens.

Co-authored-by: Codex (gpt-6-astra-medium) <noreply@openai.com>
…ixed 90 GB

IQ3_S and IQ3_XXS asked for 262K were cut to 128K on any machine under
90 GB of RAM. The limit is now counted: the experts' arena, plus the
context's 8-bit KV (in RAM with KV streaming, ~13.7 KB a token), plus
24 GB of room.

- A 64 GB PC still gets 128K for both models, as before.
- A box with ~84 GB keeps the 262K it asked for. A Colab A100-40G
  (83.5 GiB) runs IQ3_S at 262K, with images, in 56 GiB.
The canonical path zeroes session state from put_input at position 0, but a
native pack breaks out of that loop before the first put_input and runs verify
windows instead, leaving cudaMalloc'd state uninitialized. On a reused
allocation that produced NaN/overflow and degenerate identical-token output.
Zero the state explicitly when native_pack is set.

Also removes an unused argmax() helper that emitted a compiler warning, and
ignores the local .pi/ harness directory.
…gative control (Niko1221#148; CMake on the backend variables, CUDA-only)

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…mode (Niko1221#167 follow-up); no .pi/ ignore line

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…g until the watchdog (Niko1221#224); CMake warns when sm_120 is built with CUDA < 13.0 (Niko1221#220, Niko1221#224); iq_pack writes native_experts.txt under a temporary name, so a stop part-way leaves no partial pack (Niko1221#172)

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…iko1221#197), faster prompt kernels (Niko1221#187, Niko1221#188, Niko1221#207), correctness fixes (Niko1221#154); setup requires it

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

The artifact pins theta 1e7 with no rope.scaling keys, so the engine had exactly
one rope: inv[i] = theta**(-2i/n_rot), no freq_scale, no correction.  Contexts
past the trained 262,144 had no coherent path.

rope_scaling.hpp adds the one process config - none | linear | YaRN, llama.cpp's
types - and build_rope_table grows an overload that folds the whole scaled angle
(ggml's rope_yarn: interpolation, the corr_dims ramp along the pairs, and the
mscale magnitude correction) into the float64 host table.  The rotation kernels
cannot tell a scaled table from an unscaled one, so nothing downstream changes:
for the default decode path and the indexer's pooled key, scaling IS the table.

The math is transcribed from the vendored ggml (ggml.c corr_dims, ggml-cuda/
rope.cu rope_yarn/ramp, MIT headers kept), not re-derived; the one deliberate
deviation is bookkeeping - the ramp takes the PAIR index ggml's i0/2 would
produce, because every Strata rope kernel speaks in pairs.

For type None the new builder calls the original loop verbatim, and the parity
test holds it bit-identical; linear(4) is checked bit-exact against a float64
reference and equal row-for-row to none at p/4; YaRN against a float64
transcription of the ggml spec, plus the structural checks a tolerance cannot
fake: cos_tab[0] == mscale (position 0 is angle 0 for every pair), first pair
extrapolates / last interpolates at a far position (cos for the first, sin for
the last - near zero sin IS the angle, which is where interpolation happens),
and an observability assertion that the fixtures see scaling at all.

Co-Authored-By: Claude <noreply@anthropic.com>
… sites

The table path got scaling for free (it lives in the table contents); the three
ANALYTIC paths computed the angle in-kernel and so need the config explicitly.
Each wrapper now takes the RopeScaling and passes the resolved constants -
freq_scale, the corr_dims pair, ext_factor, attn_factor - as kernel arguments.
They are process constants set before session_init, so baking them at CUDA-graph
capture is correct by the same argument as every other kernel argument: only
per-token values must ride device memory, and positions still do (st.pos_dev).

The shared float32 arithmetic is rope_scaled_angle (rope_scaling.hpp), a
line-for-line rope_yarn: one definition for native_rope, the native indexer's
pooled-key rotation and prefill's rope, instead of three transcriptions that
could drift.  The spare indexer key keeps its zero angle; under YaRN the mscale
rides in through cos(0) on both pooled keys and indexer queries, so top-k is
unmoved.  rope_parity adds the check this whole wiring answers to: native vs
table on the same input, yarn and none alike - which caught the first draft
passing scaling.mscale() into the helper that applies the log term itself, a
double magnitude correction visible as ~7% on every rotated dim.

Co-Authored-By: Claude <noreply@anthropic.com>
… defaults

Nine flags, llama.cpp's names: --rope-scaling none|linear|yarn, --rope-scale,
--rope-freq-base, --rope-freq-scale, --yarn-orig-ctx, --yarn-ext-factor,
--yarn-attn-factor, --yarn-beta-fast, --yarn-beta-slow.  Ranges die in the
validation block at second zero; the config itself resolves after the model file
has had its say and before session_init builds the table or captures a graph
(rope_scaling.hpp's one-writer rule).  Precedence: CLI over the model file's
qwen4exp.rope.* keys over the struct defaults - today's artifact ships no rope
keys, so that channel is a no-op defaults path for future fine-tunes.

A scaled run prints its resolved summary (type, factor, freq_scale, base,
mscale, context vs trained) to stderr, which is what the bench logs quote, and
warns when scaling is on with the context still inside the trained range -
YaRN's magnitude correction applies everywhere, so that run is not the stock
model.

Co-Authored-By: Claude <noreply@anthropic.com>
CONTEXTS grows 383K/512K, labelled in the menu, and the setup refuses a context
past the trained 262,144 without --rope-scaling (the engine would refuse it
too).  The factor defaults to the smallest that covers the chosen context
(1.5 at 393K, 2 at 524K) and lands in the engine args like every other flag.

Co-Authored-By: Claude <noreply@anthropic.com>
DETAILS.md gains the 'Context extension' subsection - what linear and yarn do,
the setup path (383K/512K in the context menu pick the factor themselves), the
engine flags with llama.cpp's names and their defaults, why the scaling is fixed
for the whole run (K sits in the cache post-RoPE), and the one honest caveat:
inside the trained range a scaled run is a slightly different model, and the
recall numbers are text-only at one depth.  Troubleshooting grows the row for
the refusal a context past 262,144 now raises without the flag; README points
the context question at it.

Co-Authored-By: Claude <noreply@anthropic.com>
Niko1221 and others added 17 commits October 1, 2026 14:34
…'s code exactly (template <bool TAB>)

With STRATA_ROPE_TABLE unset the table was skipped at run time, yet the compiled default kernels changed their
results: the model-level parity run (Swift IQ2_XS, tools/conversation_cache_parity.py, STRATA_STATE_HASH) gave the
same tokens but other GDN / indexer / K/V state hashes than 0.1.31's (bisected to Niko1221#280). native_rope's apply,
the indexer's append / append_first / append_blocks and the prompt path's rope_kernel are templates on TAB now;
the host launches <false> whenever no table applies, and that instantiation has no table read at all. Checked:
the same tokens and every state hash as round 1 with the whole branch.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…nking renders without it only with "anthropic_thinking": "on_request" in the config - the default thinks as 0.1.31 did (a client that never asks would otherwise lose the thinking on every turn); docs

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
# Conflicts:
#	CMakeLists.txt
#	serve/server.py
#	src/program/generate.cpp
… from the start of __init__ until READY - Niko1221#333's lazy start set ended=False before READY, so a request could again plan with context 0 (RestartWindow test)

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…, AMD router +12% decode (bit-identical), RDNA4 matrix-core prompt attention (opt-in), Unsloth UD-Q4_K_XL in setup (experimental) with 2.8x faster prompts, fixes Niko1221#326 Niko1221#327 Niko1221#342 Niko1221#344 Niko1221#331, the PR batch; setup requires it

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

tools/hip/gfx1201-hipblaslt-100500.txt: 16 dense GEMM geometries (bf16 and f16) at T=4096 and T=8192, 32 rows,
calibrated on an R9700 with tune_hipblaslt (16 candidates per row, relative L2 <= 1e-4, max abs <= 1e-2, no padding
writes) against rocm-nightly-gfx120x-all 10.2.0a20260914, hipBLASLt 1.5.0 (version number 100500).

The engine refuses a table of another hipBLASLt version, so the existing gfx1201 table in Niko1221#262 (100401, hipBLASLt
1.4.1) is not used on this ROCm; setup picks tools/hip/<arch>-hipblaslt-<version>.txt by the installed version, so both
can be shipped side by side. Compared with Niko1221#262's 24 rows this table has the same 24 shapes (with solution ids chosen
for hipBLASLt 1.5.0: 13 of the 24 ids are the same) and 8 more, all bf16 (N=1/K=2560 at T=8192, N=4/K=10240, N=512/K=2560
at T=4096, N=2560/K=2560 at T=4096, and N=10240 with K=2560 and K=320 at T=4096/8192).

docs/AMD_HIP.md lists the shipped tables and how to check one.
…with bf16_from_f32()

Both helpers built their BF16 inputs with hip_bfloat16(x).data. Compiled as host C++ (no __HIPCC__) the SDK's
hip_bfloat16 is a struct holding a uint16_t, and the conversion from a float is not the engine's round-to-nearest-even
BF16 encoding (kernels/bf16_bits.hpp has that one).

Observed effect on gfx1201 with ROCm 10.2.0a, parity test with the 100500 table: the T=37 beta=1 BF16 case printed
relative_l2=0 max_abs=0 with the old encoding and relative_l2=3.6e-07 max_abs=2.6e-06 with this one; the T=4096 BF16
case prints 0 either way, and the f16 cases are unchanged. I did not inspect the old encoder's bit patterns directly,
so the exact fault is inferred from that difference, not shown.

Use strata::kernels::bf16_from_f32() in both, so the BF16 rows of gfx1201-hipblaslt-100500.txt were calibrated on
intended inputs and anyone regenerating a table gets the same. No change to the engine or to the f16 path.
Clarify usage of hip_prefill_hipblaslt_gemm and tuning process for different 1.5.0 builds.

Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
…rding

docs/AMD_HIP.md: the RDNA4 validation note still said there was no gfx1201 table and none was shipped; it now says
that holds for hipBLASLt 1.4.1 and points to the 1.5.0 table and its measured gain. The "Tuning table" text no longer
says the parity test "exits non-zero when the table does not fit". It describes what the test checks (architecture and
version through the loader, two of the 32 rows against hipBLASEx) and what it does not: a solution id the library
rejects falls back to hipBLASEx and the test still passes, so it says to read STRATA_HIPBLASLT_VERBOSE's
"fallbacks=" summary and to recalibrate for another 1.5.0 build.

tests/hip/prefill_hipblaslt_gemm.cpp: the header comment, the SKIP and missing-row messages and the case comment now
say it is a smoke test over two rows, and it prints "smoke test: 4 cases on 2 of the table's N rows" at the end.
Test behaviour and exit codes are unchanged.
dp4a through the plain signed v_dot4_i32_i8 (RDNA2 has no sudot4).
setup knows the card and the gfx103X-all TheRock index.
docs carry a community report (RX 6900 XT 16 GB: complete build, 28/28
ctest, 38-42 tok/s decode at 131,072 context, 330-339 tok/s prefill)
…ch that follows the prompt, not the capacity

Two changes to the prompt path's QSA selection (qsa_select.cu), measured on a Radeon AI PRO R9700 (gfx1201) with
--max-context 262144:

1. qsa_block_topk picked its register kernel by max_blocks, a CAPACITY (--max-context / 4 + 2). With a 262144-cell
   capacity every prompt, even a 4K one, took the slow reference top-k (up to ~10x slower at 128K+). The prompt path
   now passes the batch's active block count (the argument qsa_block_scores already takes); decode passes none and
   keeps the capacity rule and the original register width, so it is unchanged. On HIP the register kernel is
   instantiated for 66 blocks per thread (contexts up to 270,336 cells, the 262144 --max-context) and, below
   ~28K cells, where it measured 0.2-0.9x of the reference on gfx1201, the reference is kept. Same ids as the reference
   (checked 256/256 at every size).

2. qsa_block_scores_tc on gfx1200/gfx1201: block_scores_wmma_kernel on v_wmma_f32_16x16x16_bf16 (wave32). Each operand
   is split into three bf16 slices (hi / mid / lo are exact 8-bit mantissa slices of the fp32, so no scaling is
   needed) and the six products of order <= 2 are summed; the hi*hi chain keeps its own accumulator and the five
   corrections another, added once at the end (a single shared accumulator rounded the corrections at the big
   accumulator's ulp: 8e-7 of the score scale, failing the harness's gate at short context; split: 1.6e-7, the warp
   kernel's own error is of the same order). The tail block n_bid is the warp kernel's arithmetic. Like the TF32
   kernel it is not bitwise with the warp kernel (a near-tie can select differently). Gated at compile time
   (__gfx1200__/__gfx1201__) and run time (gcnArchName); other HIP targets keep the warp kernel; CUDA is unchanged.

qsa_select_bench now builds on HIP, takes the engine's capacity as a 4th argument, checks the scorer against an FP64
reference (gate: no worse than 4x the warp kernel's error, floored at 1e-6 of the score scale), exits non-zero on a
failed gate, and is registered with ctest (32K context under a 262144-cell capacity).

R9700, synthetic bench, 256 queries, capacity 262144: scores 0.81 -> 0.18 ms at 32K, 3.5 -> 0.65 ms at 131K, 6.8 ->
1.5 ms at 262K; top-k 3.9 -> 0.51 ms at 131K, 12.9 -> 1.3 ms at 262K. End to end (Swift 1.5 IQ3_XXS, --kv int8,
prefill timing, same binary with STRATA_SELECT_OLD=1 STRATA_TOPK_OLD=1 as the old arm): 107.7K-token prompt 79.5 s ->
59.5 s (1.34x; qsa select 23.3 s -> 3.4 s); 30.9K 1.02x; below 10K about 1.01x. ctest 44/44.
…pt-in (STRATA_SELECT_WMMA=1) - it selects slightly differently from the warp kernel; the top-k dispatch by active blocks (the same ids) stays on; docs

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…cks on HIP only - CUDA keeps the capacity rule (the 64K prompts read 1-3% slower with it on an RTX 5070; Niko1221#337 was measured on RDNA4)

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
…ings (0.1.32, Niko1221#304 Niko1221#301)

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

60.94 GiB of routed experts on a 64 GB machine (Ryzen 7 5700X3D, AVX2 only, RTX 5060 Ti 16 GB,
64 GB DDR4-2666). All three ways the engine can serve them, plus the knobs that decide whether the
model is usable: 4x between the best and the worst way to run the same file.

The arena never pins at this size (SetProcessWorkingSetSizeEx error 1450, 37 slices, large pages
refused too) and loses to a warm mapped run; the RAM budget costs ~30% against plain mmap (7.58
ms/round of adaptive swaps, and the CPU pool reads the page-locked complement at 15.5 GB/s against
27.4 from the page cache); a 40 GiB page-locked budget evicts the page cache the same model reads
from. STRATA_IQ_PREFETCH wants 8-24 KB on the mapped path, not the 2048 B default. The AVX-2
IQ4_XS gate/up kernel is worth ~3% end to end here (4 of 4 pairs).

Method: offline generate, arms interleaved with the order rotated per round, compared as ms per
draft round at equal draft-round counts and as total time when the arms differ in draft acceptance.
@pipeob0
pipeob0 force-pushed the gh/bench-iq4xs-64gb branch from c22bb21 to 9f8c762 Compare October 1, 2026 19:35
…nning fails (31.64 GiB on 63 GB); what flips at 61 GiB is the ranking and the per-start load cost
…v-resident; with kv-resident fixed, off/2048/8192 are within 0-3% on three models and both paths
… per round from NVMe is a different state, not tested
@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.

Niko1221 pushed a commit that referenced this pull request Oct 7, 2026
… RAM budget (#416)

Bench folder ported from PR #416 by pipeob0, authorship kept.
Trimmed for 0.1.40.2: 0 raw dumps/logs/binaries left out (listed in TRIMMED.md of each folder).
@Niko1221

Niko1221 commented Oct 7, 2026

Copy link
Copy Markdown
Owner

Thanks for the report. It is in 0.1.40.2 as bench/results/2026-10-01-iq4xs-64gb, 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.

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.