Repository navigation
bench: community report, RTX PRO 4500 x1/x2 + RTX PRO 4000, Threadripper PRO 3975WX, engine 0.1.31 - #417
Closed
mlfather wants to merge 495 commits into
Closed
bench: community report, RTX PRO 4500 x1/x2 + RTX PRO 4000, Threadripper PRO 3975WX, engine 0.1.31#417mlfather wants to merge 495 commits into
mlfather wants to merge 495 commits into
Conversation
…el (Niko1221#179); --draft-vocab en keeps the English/code subset (~110 MiB less VRAM); CUDA 13.0 for RTX 50 builds (Niko1221#220); llama.cpp's web UI is not unpacked (Windows path limit, Niko1221#206); Swift Q2_0 hidden (#171); the earlier data folder and nested Strata-data move into --data-dir (Niko1221#198); whole shards without a finish mark count (Niko1221#173) Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…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
… fewer slots than the resolve block
…eam does for the others
…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)
…rompt/token path)
…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>
…-prefill-borrow) the boundary is searched as the borrowing configuration would; the reserve only shrinks the caches (STRATA_SPLIT_OWN_PLACE=reserve: the search sees it) 9070 XT + R9700: the reserve moved K 28 -> 32 and split decode 58.5 -> 54.8 tok/s; own buffers at K=28 decode 58.4. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…9070 XT 4K prompts 718 -> 1,211 tok/s with it); docs/AMD_HIP.md: the router and WMMA switches, with credit to PR Niko1221#329's equal-speed kernel Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…WN=1, or auto for the 12% rule) - the owner's choice: the default borrows as 0.1.31 did, with the concurrent refill, the 96-blob split ring and own buffers only on a card with free VRAM, so the split's output stays 0.1.31's (2K +27%, 16K +5%, decode unchanged on the R9700 + 9070 XT); docs Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…'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: # setup.py
# 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
…75WX, engine 0.1.31 UD-Q4_K_XL, UD-IQ4_XS and IQ3_S on the same build at 8K/64K/128K, one card with a 100 GiB resident budget, then two and three cards with the layer split. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01V3YDbe7CsrJt7tTWsck1Uq
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01V3YDbe7CsrJt7tTWsck1Uq
Owner
|
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. |
Owner
|
Thanks for the report. It is in 0.1.40.2 as bench/results/2026-10-01-community-rtx-pro-4500, 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. |
Contributor
Author
|
Hi Niko, I spent Fable 5.1 tokens reviewing Strata's engine and server with a concurrency roadmap in mind, and I'd like to share what came out of it. The full write-up (findings by severity with file and line references at 82f46a8, where the time goes, 30 single-stream optimizations with the counter that measures each one, and a four-phase design for serving several requests at once) is here: The one thing worth looking at first: with "parallel": 2, two long prompts and one short one sent close together deadlock every request (serve/server.py:1150-1172, the control lock is held while waiting for a slot and a yielded request holds its slot while waiting for the lock). It reproduces 6 of 6 times against the fake engine in serve/test_parallel.py, no GPU needed. Two related gaps: an ERR during a batch admission holds the lock for 300 s, and the silence watchdog does not exist in batch mode. I can open these as an issue with the reproduction steps if you prefer them tracked separately. The branch also adds tools/load_test.py, a stdlib-only HTTP load generator that produces the first-token and throughput table BATCHING.md measures by hand. If any of it is useful I'm happy to open a PR against main: Thanks for the project |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Community benchmark: RTX PRO 4500 Blackwell (x1, x2) + RTX PRO 4000, Threadripper PRO 3975WX, 512 GB RAM
Adds
bench/results/2026-10-01-community-rtx-pro-4500/(README with method, per-pack tables, split decisions, needle recall, limitations; per-run JSON; server configs; engine logs; per-card nvidia-smi samples; BUILD.json). Extends the report posted in #402 with two- and three-GPU layer-split runs, as requested there.Setup. Engine 0.1.31, source build, commit
9259cad, CUDA 13.3, sm_120. Host: Threadripper PRO 3975WX (AVX2, no AVX-512), 512 GB DDR4-8ch, NVMe. Packs on the same binary: Unsloth UD-Q4_K_XL (imported perdocs/UNSLOTH_Q4.md, GGUF read in place, noexperts.bin), Unsloth UD-IQ4_XS, ISTA-DASLab IQ3_S. Prompts: 2,700-token review, 370-token codegen, and the same padded to 63,997 and 127,999 tokens; greedy, thinking off, 1 warm-up + 3 (short/codegen) or 2 (long) runs; context limits 8,192 / 81,920 / 139,264, one server per limit. Single card:--resident-budget-gib 100(every expert cached, no file reads). Two and three cards:--layer-split autowith the default cache mode (the resident budget cannot be combined with a split).Decode tok/s, median, 1 / 2 / 3 GPUs (1 = one RTX PRO 4500; 2 = two 4500; 3 = two 4500 + RTX PRO 4000):
Prefill tok/s and time to first token, 1 / 2 / 3 GPUs:
Peak VRAM 31.2-31.4 GB per card at every setting; host RAM 101-104 GiB (UD-Q4_K_XL, single card with the resident budget), 74-78 GiB (IQ3_S). GPU power 140-195 W.
Observations. (0) The model's full 262K window is usable on two cards: a 256,011-token prompt reads at 4,884 tok/s (52 s to first token) and decodes at 115 tok/s, 12 % under the 128K figure as the KV cache takes room from the expert cache. (1) With two cards the caches hold ~99.5 % of the routed expert mass (UD-Q4_K_XL K=23 split, 16.6K of 24.6K pairs), which is why the Q4_K pack nearly doubles: its experts otherwise run on ggml's per-token AVX2 path. (2) A third card adds nothing to decode and slows long-prompt prefill (two card boundaries, the slowest card last). (3)
STRATA_KQ256=1on UD-Q4_K_XL: within noise (67.6 vs 61.5 short, 71.7 vs 72.4 codegen). (4) Needle recall at 32K/128K, depths 10/50/90: 6/6 found. (5) UD-IQ4_XS imports on the unpatched 0.1.31 (on 0.1.18 it needed a local patch). (6)docs/UNSLOTH_Q4.md's serve example omits--serve; the binary exits asking for--tokenswithout it.Limitations: single request at a time, greedy only, two runs for long prompts, multi-GPU runs in the default cache mode, no quality suite beyond the needle check. Results-only submission; no engine changes.
🤖 Generated with Claude Code
https://claude.ai/code/session_01V3YDbe7CsrJt7tTWsck1Uq