[Engine] Turing (sm_75): make the STRATA_SM75_PERF series the default, with expert mode 8 for i-quant packs #2197

Activity

1. rafatxf commented on Oct 13, 2026

   

   rafatxf

   on Oct 13, 2026

   ContributorAuthor

   More actions

   Four more, all decode, from a kernel-node profile of the verify windows on the same card (each stands alone on main a0116b2; the four together build cleanly and pass qsa_prompt_attn_parity):

   PR	what	bitwise vs today	on by default?
   #2220	QSA decode attention for INT8 KV: the query slice in registers, reduce12, cell-major probabilities (kernel 1.5-4.1×)	yes	Turing default, STRATA_ATTN_V3=0 off
   #2221	Q2_0 matvec with one warp per four rows (QSA attn_q and the shared expert; kernel 1.5-2.4×)	yes	Turing default, STRATA_Q2_WARP=0 off
   #2222	two fewer launches per layer: the plan's flag wait folded into its copy, fused q8_1 quantize + interleave	yes	Turing default, STRATA_FUSE_WAIT / STRATA_QIL
   #2223	STRATA_NO_PCIE_GROUP=1: with --pcie-frac 0 the window graph drops the PCIe-share stage that only ever sees zero groups	n/a	opt-in

   Engine, 512 new tokens, two runs each, on top of the PRs above: #2220 + #2221 Spanish chat 86.0 → 87.8 t/s, code 103.1 → 106.1, 32K 75.3 → 78.4; #2223 another +2.2-2.4 % (32K-128K, +0.6 % at 390K); #2222 another ~+1-2 %.

   With everything (the earlier PRs, the settings above, these four): decode after a 32K prompt 65.5 → ~81 t/s, after 128K 64.5 → 76.8, after 389,560 tokens 62.5 → 75.1 (v0.1.43 as released: 52.8).

   🤖 Generated with Claude Code
2. Adamyno commented on Oct 17, 2026

   

   Adamyno

   on Oct 17, 2026

   Last edited by Adamyno

   More actions

   v0.1.45 on real sm_75: 2× RTX 2080 Ti, STRATA_SM75_PERF=0 vs 1

   Report from hardware the team does not have: two RTX 2080 Ti 22 GiB (sm_75) of a
   three-card rig. The third card stayed on other work; GPUs 0–1 were free.

   Hardware: 2× RTX 2080 Ti 22 GiB (sm_75), i9-7980XE (18c/36t, AVX-512),
   125 GiB DDR4-3600, Linux (kernel 7.0.0-38), CUDA 13.0, driver 580.178.04,
   layer split 15 (CUDA0 layers 0–14, CUDA1 15–47).

   Build: v0.1.45 (477ac13c), source,
   -DSTRATA_ENABLE_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=75 -DSTRATA_BUILD_TESTS=ON.
   cuobjdump --list-elf: 62 sm_75 cubins, no sm_86. ctest: 133/135 passed;
   ple_parity needs a Q2_0 table that is not at its fixture path here, and
   expert_cache_segmented_test failed only under -j4 and passed when run alone.

   Model and config: Qwen3.8-Flash-Next UD-Q4_K_XL (pack ud-q4_k_xl, 4 GGUF
   shards), --max-context 262144 --kv int8 --spec 3 --spec-min-p 0.54 --batch 2 --expert-cache auto --trim-stage-weights --vram-reserve-mib 768, and
   STRATA_NO_LARGEPAGES=1 STRATA_BF16_TC=1. The expert cache sized to 5355 slots
   (15.83 GiB) in every arm — identical placement across the A/B.

   A/B, same binary, environment only, alternating order

   tools/ab_engine.py, fresh server per arm, 3 rounds, medians:

   request	PERF=0	PERF=1	Δ
   chat0 decode (tok/s)	40.2	41.4	+3.0%
   chat1 decode	40.8	41.1	+0.7%
   chat2 decode	40.4	40.2	−0.5%
   10,558-token prompt, decode	39.1	39.7	+1.5%
   10,558-token prompt, prefill	598.6	603.1	+0.8%

   32k prompt (two rounds, 24,000-word prompt = 25,298 tokens, medians):

   request	PERF=0	PERF=1	Δ
   chat0 decode (tok/s)	39.6	39.9	+0.8%
   chat1 decode	36.2	37.2	+2.8%
   chat2 decode	40.3	40.8	+1.2%
   25,298-token prompt, decode	38.9	39.5	+1.5%
   25,298-token prompt, prefill	884.4	898.0	+1.5%

   The direction matches the PRs; the size does not. On this config the gain is
   +0.7…+3% decode against the +5–6% the author measured on one card, and it does
   not grow from 10.5k to 25k tokens (both +1.5% decode). The likely reason: this
   config is expert-cache bound — the hit rate was 60–88% with the adaptive tier,
   so 12–40% of the routed expert rows were read from the CPU or the peer card over
   PCIe each token, and the decode-side kernel gains are diluted. The author's
   single-card Q2_0 setup has a different balance.

   Identity

   A text-level gate is not reliable on this rig as configured, and this is worth
   knowing before trusting a 10-prompt gate here:

   • With the default adaptive tier (--adapt-every 4), the same arm twice
     (same binary, same config) differed on 1 of 4 greedy prompts.
   • With --adapt-every 0 (static cache), the same arm twice differed on 1 of 3
     (chat1, first difference at char 217).
   • With MTP on and --adapt-every 0, PERF=1 twice differed on 3 of 6 prompts
     (chat1, math1, long4k). So the earlier PERF=0-vs-1 differences in that
     configuration are not attributable to the series.
   • Without MTP (--spec 2, no --mtp), PERF=0 vs PERF=1 produced identical
     greedy text on 4/4 prompts (the same-arm control was 3/4 — chat1 flips run
     to run on this rig regardless of the switch).

   So: on this machine, a text gate can show a difference that is not the switch,
   and cannot prove a bitwise claim. The clean comparison available to us
   (no MTP, static cache) is consistent with the series being bitwise; a proper
   --window-hashes run with a fixed cache and a deterministic expert path is the
   test that would settle it.

   Bottom line

   On 2× 2080 Ti with Q4_K_XL, STRATA_SM75_PERF=1 is +1.5% decode at 10.5k and
   at 25k tokens, +0.7…+3% on short chats, +0.8…+1.5% prefill — a real but small
   gain, not the author's +5–6%. It is worth keeping as opt-in on this class of
   config until a deep-context number says otherwise; it is not a reason to change
   the default on sm_75 by itself.

   Soak (2 cards, series on)

   Separate run with STRATA_SM75_PERF=1, 2× 2080 Ti, one engine start, 505
   back-to-back requests over 60 minutes: 0 errors, 0 restarts, decode median
   41.3 tok/s, 10-minute medians 39.6 → 41.8 (expert-cache warm-up, no drift), max
   GPU temperatures 43 °C / 49 °C.
3. agorevski commented on Oct 17, 2026

   

   agorevski

   on Oct 17, 2026

   Contributor

   More actions

   RTX 8000 measurements: real gains, but not every switch is a win

   Measured 2026-10-17 on physical GPUs 2 and 3, two Quadro RTX 8000 48 GiB cards (sm_75). The same-binary STRATA_SM75_PERF=0 versus 1 comparison improved decode by 5.85%, 7.44%, and 11.09% at 4K, 32K, and 128K, respectively, using the median of three paired comparisons. Full returned messages matched across all six runs at each size.

   These are observed gains for this configuration, not a reproduction of the original single-2080-Ti/Q2_0 results or a reason by themselves to change defaults. This machine has an IQ3_S mixed-format model, and the cards reached 87/90 C with unlocked clocks. Thermal/order effects limit the precision, particularly of the large prompt-processing gain.

   What was actually compared

   Source: 477ac13cbe8e582c20484871ce459d20d82df09d / v0.1.45. Twelve of the thirteen implementations listed in the issue and its first comment are already incorporated in this source through integration commits, even though the original PRs remain open. I did not cherry-pick their original heads again.

   The current opt-in umbrella is not exactly the original list: it also includes #2081, #1821, and #1825, and excludes the separate SIMT, v2 prompt-attention, BF16-TC, and fused-expert opt-ins. #2223 was reverted (83426781); STRATA_NO_PCIE_GROUP is inert in this checkout and was not treated as a benchmarkable change.

   Hardware/software:

   • Xeon W-2295, 18 physical cores / 36 threads; 251 GiB RAM; Linux 6.8.0-146-generic.
   • GPUs 2/3 only; PCIe 3.0 x16 maximum links, NVLink present. This engine's layer split does not use NVLink P2P. Default 260 W limits; no clock/power settings changed. GPUs 0/1 were excluded from inference.
   • Driver 580.178.04; CUDA 13.4.92; GCC 13.3; Release source build with CMAKE_CUDA_ARCHITECTURES=75, CUDA and tests enabled.
   • Build deviation: reused existing ggml/llama.cpp commit 434ddbbc0e30522e897670681e503b797c12b7c1, not the repository pin 3cf03257f219afbe7334045ff7c6a06ac68c627d. Identical binary/dependency in every arm; not an official release-binary comparison.

   Model: Qwen3.8-Flash-Next-GSQ-RCO-IQ3_S-00001-of-00002.gguf plus 00002-of-00002, native iq3_s pack. Actual gate/up formats include IQ3_XXS, IQ3_S, and IQ2_S; down formats include IQ4_NL and Q2_0. The separate Q2_0 MTP draft does not make the main model a Q2_0 model.

   Fixed engine settings (local model/profile paths represented by placeholders):

   CUDA_DEVICE_ORDER=PCI_BUS_ID CUDA_VISIBLE_DEVICES=2,3
   strata --serve --pack <iq3_s-pack> --native <first-GGUF-shard>
     --mtp <mtp-runtime-pack> --spec 4 --spec-min-p 0.7
     --max-context 262144 --kv int8 --prefill auto:8192
     --pool-workers 17 --expert-cache auto --expert-cache-per-layer
     --expert-profile <expert-profile.bin> --mmap-experts --ple-io mmap
     --vram-reserve-mib 1024 --vram-reserve-later-mib 1024
     --adapt-every 0 --suffix-draft 0 --prompt-cache 0 --layer-split 24
   

   The server configuration explicitly selected physical gpu=[2,3], gpu_order=as_given; it was bound to 127.0.0.1. Other v0.1.45 settings were unchanged. All inherited STRATA_* selectors were removed before applying each arm's switches.

   End-to-end method and results

   One fresh server/engine per arm; one discarded short warm-up request. Three rounds in AB / BA / AB order, A=STRATA_SM75_PERF=0, B=1. Each server then processed the same frozen 4,096-, 32,768-, and 131,072-token prompts, in that order. Prompts were seeded word archives followed by a long numbered-list instruction; locally tokenized lengths were checked against the API's actual prompt counts.

   Greedy, thinking disabled, seed 7, 256 output-token cap. Every measured response generated 256 tokens; every cache_n was zero. Static expert placement, unchanged profile and capacity; OS page cache left warm, no global cache flush. Model loading/warm-up are excluded from engine PP/TG rates. HTTP wall time was also retained but is not streaming TTFT.

   Rates below are per-arm medians [minimum, maximum] in tok/s. Deltas are medians of paired B/A changes, not ratios of the displayed medians.

   Actual prompt tokens	PP, PERF=0	PP, PERF=1	Paired PP delta	Decode, PERF=0	Decode, PERF=1	Paired decode delta
   4,096	1068.9 [1063.5, 1069.2]	1113.7 [1113.0, 1119.6]	+4.71%	73.4 [72.9, 73.5]	77.7 [77.6, 77.8]	+5.85%
   32,768	1596.5 [1589.2, 1658.9]	1755.4 [1752.1, 1755.7]	+9.75%	65.9 [65.3, 66.5]	71.0 [70.8, 71.3]	+7.44%
   131,072	1520.4 [1518.8, 1552.5]	1889.3 [1878.4, 1894.1]	+23.68%	64.5 [64.5, 64.9]	71.7 [71.5, 72.1]	+11.09%

   Two preceding 32K A/A restart pairs had at most 0.74% paired PP variation and 0.15% decode variation, and identical text. This is a repeatability check, not a confidence interval for the hotter long-context runs. Baseline A's PP rate drifted over time; the cards' temperatures and unlocked clocks prevent interpreting the exact percentages as thermally controlled estimates.

   Startup slot-allocation signatures were identical across arms. Useful startup expert residency was about 83% of all experts; measured request hit rates were 92.6%, 93.3%, and 92.9% at the three lengths, identical between A/B. Within each prompt size, MTP offered/accepted counts also matched across baseline arms. Matching returned text is encouraging but does not prove bitwise logits or general model-quality equivalence.

   All baseline runs, retained rather than just the medians

   Round 2 executed B before A; rows are aligned by round/arm for readability.

   Round	Arm	4K PP / decode	32K PP / decode	128K PP / decode
   1	A	1068.9 / 73.4	1658.9 / 66.5	1552.5 / 64.5
   1	B	1113.0 / 77.6	1755.4 / 71.3	1894.1 / 71.5
   2	A	1063.5 / 73.5	1589.2 / 65.3	1520.4 / 64.9
   2	B	1113.7 / 77.8	1755.7 / 71.0	1889.3 / 72.1
   3	A	1069.2 / 72.9	1596.5 / 65.9	1518.8 / 64.5
   3	B	1119.6 / 77.7	1752.1 / 70.8	1878.4 / 71.7

   Which components helped this model?

   Additional three-pair, alternating 32K leave-one-group-out tests compared A=umbrella with the named group disabled against B=full umbrella. These are conditional restore-group effects, not additive standalone PR gains. All generated 256 tokens with no prompt reuse and identical returned messages.

   Group restored	A override, with PERF=1 in both arms	Paired decode delta, median [min, max]
   #2069 dense IL table	STRATA_SM75_IL=0	+4.41% [+3.98%, +4.44%]
   #2069 grouped experts	STRATA_SM75_EXP=0	-2.62% [-3.02%, -0.71%]
   #2070 hyper-connections	STRATA_GR_UP75=0 STRATA_GR_DOWN_TG=0	+0.28% [+0.14%, +1.29%]
   #2220 int8 decode attention	STRATA_ATTN_V3=0	+2.62% [+1.87%, +3.66%]
   #2222 launch fusions together	STRATA_FUSE_WAIT=0 STRATA_QIL=0	+0.71% [+0.43%, +0.86%]

   The grouped-expert choice is therefore not a universal win on this mixed IQ3_S pack. The dense table and decode attention are the clearest positive component results here. The small HC/fusion deltas should not be oversold. These runs used warm-up then 32K, unlike the main sweep's warm-up/4K/32K/128K history; compare within each pair, not absolute rates between suites. No inference profiler was used, so a small effect is not proof of a particular dispatch.

   Kernel-level checks on the RTX 8000s

   Separate synthetic/real-row harnesses ran concurrently, one on GPU 2 and one on GPU 3, before the sequential two-GPU inference tests. Shared host load and unlocked clocks are additional caveats. Three process-level rounds per timing comparison; CUDA-event measurements, warm-ups and repeated launches. These numbers are kernel results, not whole-engine speedups.

   PR / path	Workload	Baseline -> variant median	Correctness observation
   #2069 dense IL	IQ4_XS, 2560x10240, 4 columns	77.8 -> 32.5 us	Bitwise harness comparison passed; native-vs-selected-IL kernel comparison, not isolated old/new table timing
   #2070 HC combined	BF16, staged down, T=4 / T=6	60.3 -> 56.3 us / 82.9 -> 70.9 us	Equal output hashes; forced token pairs
   #2072 decode halving scorer	131K context, 6 queries	157.8 -> 77.0 us	Equal score hashes
   #2072 split top-k	393K context, 6 queries	0.177 -> 0.049 ms	Same IDs on continuous, quantized, and all-equal scores
   #2093 sampled prompt top-k	393K context, 256 queries	1.386 -> 0.600 ms	Same IDs including ties/fallback cases
   #2071 SIMT1 -> SIMT2	131K context, 256 queries	1.020 -> 0.956 ms	Equal hashes between SIMT1/2, not versus default warp scorer
   #2092 + #2146 int8 prompt attention	32K, 256 queries	3.125 -> 1.374 ms	FP64 tolerance checks passed; hashes differ, as expected

   QSA dispatches were confirmed with launch-only profiling; profiled timings were excluded. For reproduction, the key existing harness commands were:

   # Prefix with CUDA_VISIBLE_DEVICES=3 CUDA_DEVICE_ORDER=PCI_BUS_ID;
   # start each arm with inherited performance selectors cleared.
   QSA_BENCH_DECODE=1 STRATA_SCORES_MULTI=1 STRATA_SCORES_HALVE=0|1 qsa_select_bench 131072 6 500 524288
   QSA_BENCH_HASH=1 STRATA_TOPK_SAMPLE=0 STRATA_SELECT_SIMT=1|2 qsa_select_bench 131072 256 100 524288
   STRATA_TOPK_SAMPLE=0 STRATA_TOPK_SPLIT=0|1 qsa_topk_parity 393216 6 200 524288 0
   STRATA_TOPK_SPLIT=0 STRATA_TOPK_SAMPLE=0|1 qsa_topk_parity 393216 256 100 524288 1
   STRATA_PA_V2T=0|1 qsa_prompt_attn_parity 32768 256 10
   STRATA_PF_FUSED_SM75=1 prefill_fused_iq_test --only="Q2_0 / Q2_0" --no-timing
   

   Here 0|1 and 1|2 mean separate invocations, not literal shell syntax. Alternate AB/BA/AB. HC measurements used fused_gr_bench 300 4 6 8, STRATA_HC_SPLIT=2, and combinations of STRATA_GR_UP75=0/1, STRATA_GR_DOWN_TG=0/2.

   Real mixed-format expert-row tests also found forced mode 4 variable at T=1 and much slower at T=4. That explicit mode bypasses the current automatic down-mode fallback; it must not be presented as the automatic umbrella's engine regression. The leave-one-group-out result above tests the actual automatic policy.

   Separate opt-ins: screening only, not quality approval

   Each received one 32K A/B screening pair, on top of SM75_PERF=1. None exceeded the preselected 3% PP or decode threshold for a further three-pair/deep-context sweep. Single pairs are particularly uncertain; these are not established gains.

   Opt-in comparison	PP delta	Decode delta	Returned message
   PA_V2T=0 -> 1 (#2092 + #2146 together)	+1.13%	-0.57%	Changed
   SELECT_SIMT=0 -> 2 (#2071)	+1.80%	+2.70%	Changed
   XN_F16=0 -> 1, BF16_TC=1 in both (#2091)	+1.79%	+0.43%	Same within this pair

   PA_V2T/SIMT changed MTP acceptance too, so their decode deltas are not pure decode-kernel effects. BF16_TC's outputs differed from the ordinary umbrella baseline; XN_F16 identity applies only with BF16_TC already held on. No long-context KL/perplexity/needle tests were run. This does not support making the answer-changing options defaults.

   Coverage and limits

   • prefill: the fused int8 prompt experts on Turing (sm_75), opt-in (+9 % prefill on RTX 2080 Ti) #2145's synthetic Q2_0/Q2_0 fused-expert hardware parity passed: relative RMS 1.625e-5 against MMQ. Its full-model speed/quality claim remains untested: the available IQ3_S gate/up formats do not qualify, and a canonical full-model Q2_0 pack was unavailable locally.
   • sm_75: Q2_0 matvec with one warp per four rows (bitwise) #2221's Q2_0 matvec did not receive a dedicated isolated timing/dispatch test. Pack labels and Q2_0 down tensors alone are not proof that this kernel was exercised.
   • sm_75: the v2 prompt attention in one stage, three blocks per SM (bitwise, attention 1.66x; stacked on #2092) #2146 cannot be isolated from sm_75: the v2 int8 prompt attention with register staging instead of cp.async (STRATA_PA_V2T=1) #2092 with a current-binary switch; the measured comparison is current v1 versus final integrated v2, not old-v2 versus new-v2.
   • 143 tests were registered; the readiness subset had 20 passes, one architecture skip, no failures. All 34 QSA/parity/profile invocations and 27 dense/expert/HC invocations completed without failures. Q4_0 prompt attention was explicitly skipped on sm_75.
   • Full responses, hashes, per-request timings, startup placement, MTP counts, per-arm telemetry, commands and logs were retained locally. Independently rechecked hashes/counts for 59 completed measured responses. An initial valid pilot was excluded from completed results after a harness TIME_WAIT teardown-check error; the check was corrected and a fresh manifest/output used. No inference crashes or OOMs occurred in the completed suites.
   • No 390K sweep, historical-release-binary comparison, standalone attribution to every PR, all-model validation, thermally controlled repeat, or quality certification. No source changes; GPUs 2/3 were released after measurement.

   Bottom line: the current bitwise umbrella gives a clear, repeatable directional improvement on this two-RTX-8000/IQ3_S setup, with matching text and MTP counts. The dense IL table and int8 decode-attention changes have the strongest isolated engine evidence; grouped expert mode selection deserves format-specific caution. Large synthetic kernel gains do not automatically become large engine gains, and the separate numerically different prompt opt-ins remain unvalidated here.
4. Adamyno commented on Oct 18, 2026

   

   Adamyno

   on Oct 18, 2026

   More actions

   Second IQ3_S pack, a second sm_75 pair: STRATA_SM75_EXP=0 is +2.7% decode; the K-quant control does not move

   Cross-check of @agorevski's leave-one-group-out result (grouped experts −2.62% on his IQ3_S pack), on a different rig and a different IQ3_S pack, plus a K-quant control on the same box.

   Hardware: 2× RTX 2080 Ti 22 GB (sm_75, power limits 210 W / 167 W), i9-7980XE (18c/36t), 125.5 GiB DDR4-3600, Linux 7.0.0-38, CUDA 13.0, driver 580.178.04. Two cards only — the third stayed on other work.

   Build: v0.1.45 (477ac13c), source, -DSTRATA_ENABLE_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=75. Same binary in both arms; only the environment differs.

   Fixed config (both arms): --layer-split 15, 262,144 ctx, --kv int8 --kv-resident 32768, --spec 3 --spec-min-p 0.54 --mtp mtp/rt --mtp-max-t 3 --batch 2 --vram-reserve-mib 768 --trim-stage-weights, STRATA_SM75_PERF=1 in both arms, STRATA_NO_LARGEPAGES=1 STRATA_BF16_TC=1. Sampling pinned greedy (top_k 1, temperature 0) with presence_penalty 0 so the 0.1.45 card set (#2940) does not reach either arm.

   Method: tools/ab_engine.py from the repo, fresh server per arm, 5 rounds, order alternates (AB / BA), warm-up request discarded, three chat prompts (256 new tokens) and one 24,250-token prompt (256 new tokens). Figures are medians of the paired per-round deltas, not ratios of displayed medians.

   IQ3_S (ISTA-DASLab GSQ-RCO, pack gsq_iq3s), A = default (EXP on), B = STRATA_SM75_EXP=0

   request	A decode	B decode	paired Δ median	range
   chat0	67.3	69.1	+2.7%	+1.9 … +2.8%
   chat1	63.8	66.0	+3.5%	+3.4 … +4.2%
   chat2	61.7	63.1	+2.3%	−5.4 … +4.2%
   24,250-token prompt decode	64.3	63.2	−1.7%	−2.6 … +3.6%
   24,250-token prefill	1,378	1,381	+0.2%	—

   All 20 pairs: median +2.7%. Prefill is unchanged.

   The 24k row is not a kernel-only comparison: in 3 of the 5 pairs the two arms stopped at different answer lengths (157 vs 131 tokens), so that row mixes in a drafting difference. The chat rows mostly do not: 13 of the 15 chat pairs generated the same token count, and in the eight pairs where the accepted/offered draft counts also match (rounds 3–4, read from the engine log) the median is +2.45% (+1.93, +2.08, +2.11, +2.80, +3.44, +3.63, +3.91, and one −5.37).

   That −5.37% pair is worth naming rather than smoothing: A and B both generated 205 tokens and both accepted 104 of 153 drafts, so it is not a sampling or drafting difference — it is a timing outlier on my box (A 3,341 ms, B 3,531 ms for identical work). I did not find a cause; every other identical-work pair sits between +1.9% and +3.9%. I am reporting the median with it included.

   Control: Q4_K_XL (K-quant pack ud-q4_k_xl), same box, same binary

   Two rounds. Round 0 gave four pairs with identical token counts and identical accepted/offered drafts:

   request	A decode	B decode	Δ
   chat0	49.1	49.1	0.0%
   chat1	49.2	49.2	0.0%
   chat2	48.3	47.7	−1.2%
   24,250-token prompt	50.0	49.8	−0.4%
   prefill	1,017.1	1,012.6	−0.4%

   Median −0.2%. Round 1 drafting differed between the arms in all four pairs (−0.6, −6.6, +4.9, +4.8%) and is excluded — on this pack the run-to-run drafting noise is larger than any effect. So this control says "no gain on a K-quant pack", not "bitwise": four pairs cannot prove inertness, but they are consistent with the code comment that this switch is the i-quant down layout.

   What I take from it

   • The grouped-expert down layout (STRATA_SM75_EXP, on by default under STRATA_SM75_PERF=1) costs ~2.6–2.7% decode on an IQ3_S pack on this rig, the same direction and size @agorevski measured on his own IQ3_S pack. Two rigs, two packs, same answer — so this looks like the pack class, not his box.
   • On a K-quant (Q4_K_XL) pack the same switch is flat (−1.2 … 0.0% over four identical-work pairs). That points at a format gate (i-quant rows present) rather than turning the switch off on every sm_75 pack.
   • Everything else in the umbrella stays: this was a leave-one-out on one component, not a PERF on/off run.

   Not measured: three-card split, full-power cards, --window-hashes identity, other quant families. Raw results.jsonl per request, the engine logs and the two configs are retained locally — I can attach them or run more rounds if that is useful.
5. coolthor commented on Oct 18, 2026

   

   coolthor

   on Oct 18, 2026

   More actions

   v0.1.45 on a single RTX 2080 Ti 22 GB, IQ3_S: STRATA_SM75_PERF=1 helps, and STRATA_SM75_EXP=0 helps more

   Another sm_75 data point, this time on one card with no layer split. It's also the third rig to see STRATA_SM75_EXP=0 come out ahead on an IQ3_S pack, after @Adamyno and @agorevski.

   Hardware: 1× RTX 2080 Ti 22 GB (sm_75), 200 W power limit, clocks not locked, peak 74 °C. AMD EPYC 7402 (24c/48t), 123 GiB RAM, PCIe Gen3 x16, Ubuntu 26.04, driver 580.173.02, CUDA 12.4. The box's other two GPUs kept serving their usual workloads during the run.

   Build: v0.1.45 (477ac13c), Release, CMAKE_CUDA_ARCHITECTURES=75, built through setup.py. cuobjdump --list-elf: 62 sm_75 cubins. My first build (GCC 15 for C++, GCC 13 as the CUDA host compiler) failed at link time on __cxa_call_terminate. Building everything with GCC 13 (CC=gcc-13 CXX=g++-13 CUDAHOSTCXX=g++-13) worked. I built without STRATA_BUILD_TESTS, so I have no ctest results to report.

   Model and config: Qwen3.8-Flash-Next-GSQ-RCO IQ3_S, single GPU, --max-context 262144 --kv int8 --kv-resident 32768 --spec 4 --spec-min-p 0.7 --lookup-chain 5 --ple-io direct --expert-cache auto --prefill auto --vision --vram-reserve-mib 700, parallel: 1. Expert cache: 7589 slots / 14,737 MiB in every arm. The vision encoder ran on the CPU (gpu: false, 24 threads) the entire time. Common env: STRATA_BF16_TC=1 STRATA_ADAPT_OVERLAP=1.

   Arms (same binary, env only): P0 = STRATA_SM75_PERF=0 · P1 = STRATA_SM75_PERF=1 · P1-EXP0 = STRATA_SM75_PERF=1 STRATA_SM75_EXP=0

   Method: 5 rounds with a fresh server per arm per round, alternating the order each round. I used one warmup request per start and didn't count it. Sampling was greedy (top_k 1, temperature 0, presence_penalty 0, reasoning_effort none), with no prompt reuse (every request got a unique prefix). Decode numbers are the engine's decode_tok_s. Prefill is prompt_read divided by prompt time in seconds. Each cell is the median of 5 runs, with [min, max].

   tok/s	P0	P1	P1-EXP0
   code, 45-token prompt	86.0 [85.7, 86.8]	92.3 [92.2, 92.6]	95.8 [95.3, 95.9]
   Chinese, 48-token prompt	60.0 [59.8, 60.4]	61.8 [61.7, 63.2]	64.5 [64.4, 64.7]
   decode at 4,088 tokens	67.2 [66.9, 67.7]	71.4 [71.3, 77.0]	79.1 [79.0, 79.3]
   decode at 32,789 tokens	63.8 [62.3, 68.9]	72.3 [68.3, 76.9]	78.4 [77.9, 81.2]
   decode at 65,549 tokens	64.7 [64.3, 68.4]	68.8 [68.0, 73.5]	73.1 [69.2, 74.8]
   prefill, 16,406 tokens	1,217	1,318	1,322
   prefill, 65,546 tokens	1,138	1,340	1,340

   vs P0	code	Chinese	4K	32K	64K	prefill 16K	prefill 64K
   P1	+7.3%	+3.0%	+6.3%	+13.3%	+6.3%	+8.3%	+17.7%
   P1-EXP0	+11.4%	+7.5%	+17.7%	+22.9%	+13.0%	+8.7%	+17.7%

   The 65K prompts are beyond --kv-resident 32768, so those rows include KV streaming.

   Output parity (with drafts on, --spec 4):

   • The 256-token greedy continuations at 4K and 32K weren't byte-identical to P0 in either enabled arm. The first differing character was at index 128 (4K) and 103 (32K).
   • P0 isn't repeatable across restarts at long context either. Over 5 fresh starts, P0 produced 5 distinct outputs at 32K and 5 at 64K, and P1-EXP0 produced 4 at each. At 4K, P0 and P1-EXP0 were each stable across all 5 starts (1 distinct output each), but the two outputs differ from each other.
   • I didn't isolate this with --spec 0, so I can't say whether the 4K difference comes from the kernels or from draft acceptance.
   • Short answers that end naturally (a ≤20-word summary at 4K and 32K) matched P0 byte for byte in both arms.

   Smoke checks: a vision test (reading table values) and a structured tool call passed in all three arms. On a 2048-token thinking puzzle, all three arms reached the right final answer, but P1's written explanation contradicted the puzzle's premise in one step. That's a single request, so I'm noting it rather than reading too much into it.

   After this run I switched the card's production service to v0.1.45 with STRATA_SM75_PERF=1 STRATA_SM75_EXP=0. I can rerun a --spec 0 pass, or try any other switch you want checked.
6. rafatxf commented on Oct 19, 2026

   

   rafatxf

   on Oct 19, 2026

   ContributorAuthor

   More actions

   Thanks @Adamyno, @agorevski and @coolthor (and @ioridev in #2069) for the runs. That makes five sm_75 setups, four of them independent of mine. Every PR in this issue was ported into 0.1.45, so I have closed them; #2223 was closed as below the 3 % bar. I'll keep this issue open for the default question.

   STRATA_SM75_PERF=1 across the rigs (decode, against =0)

   who	cards	pack	decode	prompt
   ioridev (#2069)	2080 Ti 22 GB + 11 GB peer	IQ3_S	+7.6 %	+8-9 %
   Adamyno	2x 2080 Ti, layer split	UD-Q4_K_XL	+0.7 to +3 % (expert-cache bound)	+0.8 to +1.5 %
   agorevski	2x Quadro RTX 8000	IQ3_S	+5.9 / +7.4 / +11.1 % at 4K / 32K / 128K	+4.7 / +9.8 / +23.7 %
   coolthor	1x 2080 Ti	IQ3_S	+3 to +13 %	+8 to +18 %
   me (#2069)	2x 2080 Ti, NVLink, layer split	Q2_0 / IQ3_S	+19.1 % / +3.9 %	+6-8 % / +3-7 %

   Every rig saw a gain, and there were no errors in Adamyno's 60-minute soak. The size depends on how much of the decode runs on the GPU. To me that supports making the bitwise umbrella the sm_75 default. The expert layout below is the exception.

   The grouped-expert layout: mode 4 only helps the Q2_0 pack

   Three of you found that turning off mode 4 (STRATA_SM75_EXP=0, which goes back to mode 0) helps on IQ3_S. ioridev and I had compared mode 4 against mode 8 instead (STRATA_EXP_MODE=8). Nobody had put 0 and 8 side by side, so I ran all three modes on the three packs I have.

   Setup: v0.1.45 + STRATA_SM75_PERF=1, 2x 2080 Ti in a layer split, 32K context, --spec 4 --mtp --adapt-every 0.

   Method: engine-direct, 512 greedy tokens per prompt. Each figure is the mean decode tok/s over the es, code, 4K and 16K prompts.

   pack	mode 4 (umbrella default)	mode 0 (SM75_EXP=0)	mode 8 (EXP_MODE=8)
   Q2_0	115.1	111.5 (-3.1 %)	114.6 (-0.5 %)
   IQ3_XXS	69.9	71.8 (+2.7 %)	72.9 (+4.4 %)
   IQ3_S	59.7	60.5 (+1.2 %)	61.4 (+2.7 %)

   The IQ3_S row leaves out the 16K prompt. In the mode-8 arm the drafts went differently there: 238/375 accepted against 262/357 in the other two modes. That run took 13.85 s against 12.37 s, so it mixes in a drafting difference. I don't know the cause: ioridev's native_expert_bench showed 0 differing outputs between modes 0 and 8.

   So mode 4 is the right layout only for the Q2_0 pack. On the i-quant packs, mode 8 is the fastest here, and mode 0 is still better than mode 4. On Adamyno's K-quant control the switch is flat.

   Suggestion, if you make the umbrella a default:

   • Gate STRATA_SM75_EXP (mode 4) on the Q2_0 pack only.
   • Give the i-quant packs mode 8, or mode 0 if you'd rather stay on the CUDA layout.

   If a --spec 0 or --window-hashes run of any of these would help, I can do it.
7. Adamyno commented on Oct 19, 2026

   

   Adamyno

   on Oct 19, 2026

   More actions

   Mode 4 / 0 / 8 on a second IQ3_S pack: mode 8 wins here too (+4.9% decode over the umbrella default), prefill unchanged

   Confirmation of @rafatxf's mode comparison on a different rig and a different IQ3_S pack, in the shape I measured STRATA_SM75_EXP before.

   Hardware: 2× RTX 2080 Ti 22 GB (sm_75, 210 W / 167 W), i9-7980XE, 125.5 GiB DDR4-3600, Linux 7.0.0-38, CUDA 13.0, driver 580.178.04.

   Build/config: v0.1.45 (477ac13c), one binary, environment only. Qwen3.8-Flash-Next GSQ-RCO IQ3_S (gate/up IQ3_XXS 17 layers, IQ3_S 10, IQ2_S 20, IQ4_XS 1; down IQ4_NL 39, Q2_0 9), --layer-split 15, 262,144 ctx, --kv int8 --kv-resident 32768, --spec 3 --spec-min-p 0.54 --mtp mtp/rt --batch 2, greedy (top_k 1, temperature 0, presence_penalty 0), STRATA_SM75_PERF=1 in all three arms. 24,250-token prompt, 256 new tokens, 3 rounds, order alternates (ABC / CBA / ABC).

   request (decode tok/s, median)	mode 4 (umbrella)	mode 0 (SM75_EXP=0)	mode 8 (EXP_MODE=8)
   chat0	58.7	58.8	59.5
   chat1	58.8	61.8	62.5
   chat2	55.3	57.4	60.1
   24,250-token prompt	59.4	61.2	62.5
   prefill, 24,250 tokens	1,324	1,336	1,331

   Paired deltas against mode 4, only the pairs whose token counts match (drafting differences excluded):

   • mode 0: 12/12 pairs, median +2.20% (+0.17 … +5.79).
   • mode 8: 8/12 pairs, median +4.90% (−0.51 … +9.01). Four pairs excluded: chat2 rounds 0 and 2 generated 196 vs 204 tokens, and the 24K prompt rounds 0 and 2 ended at 159 vs 157; their deltas (+8.5, +9.6, +5.2, +3.8%) were in the same direction but mix in a drafting difference.

   Prefill is flat between the modes (±0.9%, within run noise), so this is a decode-layout effect, as expected.

   What it says

   On this pack the order is mode 8 > mode 0 > mode 4, which is the same direction @rafatxf measured on his IQ3_XXS and IQ3_S packs (there: +4.4% and +2.7% for mode 8 over mode 4). Two rigs, three packs (his Q2_0/IQ3_XXS/IQ3_S and this IQ3_S) now agree: mode 4 is the Q2_0 layout only, and the i-quant packs want mode 8.

   Mode 8 changes the arithmetic (int8 rounding), so this is a speed result, not an identity claim — I have no KL or --window-hashes run for it. If a mode 8 identity gate is what would settle the default, I can run one on this pack.
8. Adamyno commented on Oct 19, 2026

   

   Adamyno

   on Oct 19, 2026

   More actions

   Quality check on the same pack: mode 8, mode 0 and mode 4 are value-identical here (4,995 teacher-forced positions)

   Before anyone changes a default: I ran the repo's own yardstick (tools/quant_quality.py) on the same IQ3_S pack, with the expert arena on the GPU so the layout is actually exercised.

   Setup: Qwen3.8-Flash-Next GSQ-RCO IQ3_S, one RTX 2080 Ti 22 GB, data/quant_quality/v1 (code, docs, english, multilingual, agent — 999 positions each, 4,995 total), FP16 K/V, --short-read, --prefill 512, --gpu-experts 1500; expert cache, adaptation, PCIe share and prompt cache fixed by the tool. Reference = mode 4 (the umbrella default), one engine per arm, same binary.

   run	positions	top-1 agreement	mean KL	PPL	PPL ref
   mode 8 (EXP_MODE=8)	4,995	100.00%	0.00000	4.115	4.115
   mode 0 (SM75_EXP=0)	4,995	100.00%	0.00000	4.115	4.115

   Control that the harness does see the GPU expert path: mode 4 with all experts on the CPU (the tool's default --expert-cache 1) against mode 4 with 1,500 GPU slots gives 94.81% top-1, mean KL 0.0265, PPL 4.143 vs 4.115. So moving the experts onto the card does move the logprobs; switching the layout does not. (Mode 4 and mode 0 with experts on the CPU are also identical, as expected.)

   What this means for the default question

   On this pack the mode-8 layout changes where the numbers live, not the numbers: the +4.9% decode I measured is not paid for in output quality at the shapes the yardstick covers. If the choice is made on speed alone, mode 8 is the one to take for i-quant packs.

   Caveats, so this is not read as more than it is: the yardstick reads one token per window (--short-read), while the speed runs use spec-3 verify windows of four rows — and Turing's down projection already falls back to the CUDA layout at 4+ token windows (sm75_down_em), so those windows are only partly covered here. A --spec-follow + --window-hashes run is the test for that shape if you want it. One pack, one card, and no task-level eval (no needle/recall).
9. coolthor commented on Oct 19, 2026

   

   coolthor

   on Oct 19, 2026

   More actions

   Mode 0 vs. mode 8 on a single RTX 2080 Ti (IQ3_S): mode 8 wins here too

   Following up on @rafatxf's mode comparison, I ran mode 0 against mode 8 on the same single-card setup as my earlier comment (one RTX 2080 Ti 22 GB, no layer split, Qwen3.8-Flash-Next-GSQ-RCO IQ3_S, v0.1.45 477ac13c, int8 KV, --spec 4, --kv-resident 32768, CPU vision encoder attached).

   Arms (same binary, env only): both use STRATA_SM75_PERF=1 STRATA_SM75_EXP=0; the mode 8 arm adds STRATA_EXP_MODE=8. In exp_mode() (src/kernels/cuda/iq_kernels.cu), an explicit STRATA_EXP_MODE is read before the GPU default, so STRATA_SM75_EXP=0 doesn't override it.

   Method: same as before. 5 rounds, a fresh server per arm per round, order alternating each round, greedy sampling. Each result is the median of 5 runs. Decode uses the engine's decode_tok_s, with [min, max] shown; prefill is prompt tokens read divided by prefill time.

   tok/s	mode 0	mode 8	change
   code, 45-token prompt	95.6 [94.9, 96.2]	99.2 [98.2, 99.3]	+3.8%
   Chinese, 48-token prompt	64.5 [64.1, 64.8]	66.4 [66.2, 66.4]	+3.0%
   decode at 4,088 tokens	79.0 [78.2, 79.6]	82.0 [81.5, 82.0]	+3.8%
   decode at 32,789 tokens	78.5 [73.6, 80.3]	83.6 [80.7, 84.9]	+6.5%
   decode at 65,549 tokens	73.9 [70.7, 76.7]	78.1 [73.9, 79.4]	+5.7%
   prefill, 16,406 tokens	1,319	1,316	-0.2%
   prefill, 65,546 tokens	1,340	1,347	+0.6%

   Decode is faster at every tested length, with bigger gains at 32K and 64K than at 4K. Prefill is essentially unchanged. That matches what you and @Adamyno saw on IQ3_S.

   The vision and structured tool-call smoke tests passed in both arms. I'm now running mode 8 in production on this card.
10. Niko1221 commented on Oct 19, 2026

    

    Niko1221

    on Oct 19, 2026

    Owner

    More actions

    Thanks for taking the time to report this.

    We're resetting the tracker for 0.2.0-beta. There are more than 1,300 open issues and pull requests, too many for us to read properly, so we're closing everything that isn't being worked on right now. Closed here does not mean rejected.

    To reopen it takes about two minutes:

    1. Re-test on 0.2.0-beta. It may already be fixed. Write down what you see.
    2. Start the title with its area: [Engine], [GUI], [TUI], [Benchmark], [Docs] or [Other].
    3. Fill in the parts of the issue form: what happened, the Strata version, your graphics card, memory and system.
    4. Comment /reopen. A check runs and either reopens the issue or lists what is still missing.

    The rules are in CONTRIBUTING.md.
11. Niko1221

    closed this as not plannedon Oct 19, 2026
12. rafatxf

    changed the title [-]Turing (sm_75, RTX 20) support: tracking issue for the open PRs[/-] [+][Engine] Turing (sm_75): make the bitwise STRATA_SM75_PERF series the default, with expert mode 8 for i-quant packs[/+] on Oct 19, 2026
13. rafatxf commented on Oct 19, 2026

    

    rafatxf

    on Oct 19, 2026

    ContributorAuthor

    More actions

    Re-tested on v0.2.0-beta (3d1fb48, engine 0.1.45.3), source build for sm_75, nothing patched.

    Setup:

    • 2x RTX 2080 Ti 22 GB, layer split.
    • 32K context, --kv int8 --kv-resident 32768 --spec 4 --spec-min-p 0.7 --mtp --pcie-frac 0 --adapt-every 0 --resident-experts.
    • Env STRATA_BF16_TC=1 in every arm.

    Method:

    • One binary. The arms differ only by environment variables.
    • Engine-direct (--serve): a fresh engine per prompt, then 512 greedy tokens.
    • Prompts: es (57 tokens), code (54), a 4K and a 16K document.

    Decode, tok/s

    pack	arm	es	code	4K	16K	mean
    Q2_0	default	95.1	124.4	82.4	78.3	95.1
    Q2_0	STRATA_SM75_PERF=1	112.0	151.3	96.2	92.6	113.0 (+18.9 %)
    Q2_0	+ STRATA_EXP_MODE=8	111.2	150.5	95.6	91.9	112.3 (+18.1 %)
    IQ3_S	default	60.2	72.1	38.5	35.9	51.7
    IQ3_S	STRATA_SM75_PERF=1	63.0	75.8	39.7	37.2	53.9 (+4.3 %)
    IQ3_S	+ STRATA_EXP_MODE=8	64.9	77.7	39.2	37.3	54.8 (+6.0 %)

    • Prompt reading with STRATA_SM75_PERF=1: Q2_0 4K 3.79 → 3.57 s, 16K 13.98 → 12.98 s. IQ3_S 4K 5.13 → 4.95 s, 16K 15.30 → 14.36 s.
    • Drafts: the default and SM75_PERF arms accepted the same drafts (Q2_0 es 262/342, code 330/389; IQ3_S es 262/349, code 341/403), as expected for bitwise kernels.
    • One exception: in the IQ3_S mode-8 arm the 4K drafting differed (242/390 against 249/387), so that one cell mixes in a drafting difference.

    The same picture as on 0.1.45:

    • SM75_PERF gives +19 % on Q2_0 and +4 % on IQ3_S.
    • Mode 8 adds about 1.6 % on IQ3_S and costs 0.7 % on Q2_0, so mode 4 is the right layout only for Q2_0.

    I have updated the title and the body to the feature-request form (proposal, current workaround, environment, evidence).

    /reopen
14. rafatxf commented on Oct 19, 2026

    

    rafatxf

    on Oct 19, 2026

    ContributorAuthor

    More actions

    /reopen
15. github-actions

    reopened this on Oct 19, 2026
16. 14 remaining items

    Load more

17. Adamyno commented on Oct 22, 2026

    

    Adamyno

    on Oct 22, 2026

    More actions

    Small correction on the two @Adamyno lines before they land in the issue body — the runs were wider than the compressed "N/N measurable" counts suggest:

    • IQ3_S, our box: 15 prompts measured. 12 of 15 are byte-identical across all four arms (A/B/C/M8); the other 3 differ only in arm C (the two switches outside the series). The series itself — B vs A and M8 vs B — is byte-identical on every measurable prompt, so series-vs-reference is 15/15. One of the 15 overflows thinking at 512; re-run at 2048 it is byte-identical across all four arms too.

      Paste-ready: @Adamyno's (15 prompts; series-vs-reference byte-identical 15/15; 12 of 15 identical across all four arms, the other 3 arm-C-only)

    • Q2_0, our box: 15 prompts × 5 arms (A/B/C/M8 + the fused D). 11 of 15 identical across all five arms; 3 differ only in arm C (P03_code, P13_reason, P10_translate — C's prefill numerics shorten thinking, so it fits the 512 budget while the other arms overflow); P05 is content-null at 512 and byte-identical across all 10 requests at 2048. A = B = M8 = D on every measurable prompt.

      Paste-ready: @Adamyno's (15 prompts × 5 arms; series + fused prefill identical on every measurable prompt; 11 of 15 identical across all arms, 3 arm-C-only, 1 null at 512 → byte-identical at 2048)

    The verdict doesn't change — the series does not move a single answer; only near-tie wording inside the thinking channel moves, and only under the opted-in numerics switches.
18. yusheng227507 commented on Oct 22, 2026

    

    yusheng227507

    on Oct 22, 2026

    More actions

    The evidence in one place, and the one gap I can close

    @Adamyno's third box and @rafatxf's summary are what I was missing to see the whole picture. Here is the consolidated table from my side, in the same terms, so the three boxes can be compared directly. One item is still open on my side, at the end.

    My box: 2x RTX 2080 Ti 22 GB (sm_75), NVLink NV2, PCIe 3.0 x8 host link (6.5 GB/s), i7-10700K (AVX2, no AVX-512), 94.1 GiB RAM, driver 595.99.02, Ubuntu 24.04.5, Strata 9de04cd4 (v0.2.1.4), engine 0.1.46.3 compiled from source for sm_75 (CUDA 12.8), 262144 context, kv int8, --kv-resident 32768, --spec 4 --spec-min-p 0.5, --layer-split 26, 3 slots. Recipe in every arm: STRATA_IQ_MT_MIN=1 + STRATA_PREFILL_CPU_SHARE=0 + STRATA_ROUTE_TAIL_SKIP=0 + --pcie-frac 0 + --adapt-every 1000000.

    Pack	Arm B vs A (SM75_PERF on/off)	Arm C vs B (+BF16_TC+PA_V2T)	Mode 8 vs default	Inside-arm stability
    IQ3_S	8/8 byte-identical	8/8 byte-identical	8/8 byte-identical	0 unstable
    IQ3_XXS	6/8 byte-identical, 2 near-tie prompts moved	8/8 byte-identical	8/8 byte-identical	0 unstable

    So on this box the series itself moves nothing on IQ3_S, and on IQ3_XXS it moves two near-ties; the numerics opt-ins (arm C) are what change answers. That is the same split @Adamyno found on IQ3_S.

    The two IQ3_XXS prompts that moved, for the record, with the exact texts - both are the other of two equally good answers:

    P3_code   A:  cleaned = ''.join(c.lower() for c in s if c.isalpha())
              B:  filtered = [c.lower() for c in s if c.isalpha()]
    
    P5_list   A:  - 第1至5項為：0, 1, 1, 2, 3
                   - 符合Fibonacci定義的起始序列
              B:  - 0, 1, 1, 2, 3
                   - 1, 1, 2, 3, 5
                   - 0, 1, 1, 2, 3
    

    I also ran the per-kernel screen (16 arms: reference, all on, and the 14 named sub-switches off one at a time) on IQ3_XXS. No single sub-switch brings the reference answer back on either prompt; five of them (GR_DOWN_TG, SCORES_HALVE, TOPK_SPLIT, TOPK_SAMPLE, ATTN_V3) each move P5_list to a third text of their own, and the other nine leave it at the all-on answer. The reason a one-switch screen cannot return the reference is in the code: sm75_default() is what native_mmvq.cu:2584, qsa_decode_attn.cu:449, fused_gr.cu:1207 and :1232, qsa_select.cu:1689, :1704, :2256, :2489 and verify.cpp:310 follow, while only STRATA_SM75_IL, STRATA_SM75_EXP and STRATA_TC_SM75 have a variable of their own. So a near-tie moves if any of several paths rounds differently.

    On the speed side, STRATA_EXP_MODE=8 measured +14 % chat decode / +16 % agent decode on IQ3_XXS and +8.7 % / +13.0 % on IQ3_S against the default layout here, with the prompt path unchanged - and byte-identical text on both packs, so it is not a numerics change on this box.

    The gap

    Every A-vs-B comparison above is on engine 0.1.46.3, because that is what I had built when I ran it. The engine people will actually run is now 0.1.47 (0.2.2). The four standard-protocol runs in PR #3535 are on 0.1.47 and put IQ3_XXS at 21/21 requests and 12/12 sanity (tool calls 4/4), but they do not carry the exact-comparison recipe, so they do not answer the text question.

    So the open item on my side is narrow: does IQ3_XXS still move those two near-ties on 0.1.47, or is that a 0.1.46.3 artifact? I can re-run the same A/B/C/M8 matrix with the same recipe on 0.1.47 and report it here - say the word, or I will run it and post the numbers.

    For the record, what the thread now reads, as far as I can see it:

    Pack	Boxes	Result
    IQ3_S	@Adamyno (12/12 measurable), me (8/8), @rafatxf (KL 0)	byte-identical everywhere
    Q2_0	@Adamyno (11/11 measurable), @rafatxf	byte-identical everywhere
    IQ3_XXS	me (6/8)	equivalent answers; two near-tie prompts out of eight move, and no single sub-switch causes it
    Mode 8	me (both packs), @Adamyno	byte-identical to the default wherever it was compared
19. rafatxf commented on Oct 22, 2026

    

    rafatxf

    on Oct 22, 2026

    ContributorAuthor

    More actions

    @Adamyno thanks for the correction. Your two paste-ready lines replace mine in the issue body: series-vs-reference 15/15 on IQ3_S, and series + fused prefill identical on every measurable Q2_0 prompt.

    @yusheng227507 yes, please run the A/B/C/M8 matrix with the recipe on 0.1.47. That is the engine the default would ship in, and IQ3_XXS is the only pack where the series moved anything. If the two near-ties are still there, the body keeps "equivalent" for IQ3_XXS; if they are gone, it can say so. Your consolidated table matches the body, with Adamyno's counts updated as above.

    Reproduction status: no new measurements in this comment.
20. yusheng227507 commented on Oct 22, 2026

    

    yusheng227507

    on Oct 22, 2026

    More actions

    The gap from my previous comment is closed: on engine 0.1.47, nothing moves.

    Re-ran the A/B/C/M8 matrix on 0.1.47 (source build, sm_75, same box), same recipe (STRATA_IQ_MT_MIN=1 + STRATA_PREFILL_CPU_SHARE=0 + STRATA_ROUTE_TAIL_SKIP=0, --pcie-frac 0 --adapt-every 1000000), fresh server per arm, greedy (temperature 0 / top_k 1 / seed 1), enable_thinking false, max_tokens 2048. 12 prompts × 2 repeats × 4 arms × 2 packs = 192 requests, 0 errors.

    Byte-identical at every position:

    	IQ3_XXS	IQ3_S
    A vs B (+SM75_PERF)	12/12	12/12
    A vs C (+BF16_TC+PA_V2T)	12/12	12/12
    A vs M8 (+EXP_MODE=8)	12/12	12/12
    B vs M8	12/12	12/12

    Every repeat pair is byte-identical too — 96/96 within-arm comparisons, nothing unstable.

    The prompt set covers the two classes that moved on 0.1.46.3 (a palindrome-code and a three-line Fibonacci list), plus translation, word-counting and a second long-doc needle the old set did not have. On the two old flip points, all four arms now resolve to the same one of the two texts they used to flip between (P3 gives the filtered build everywhere; the list gives the per-line numbers everywhere) — so the near-ties did not vanish: with this engine, arms A and B no longer disagree about which candidate to pick, and no other prompt moved.

    Two caveats: the wordings are close to, but not token-identical with, the 0.1.46.3 run (its raw data was overwritten, so no byte-level cross-engine comparison on the originals is possible), and this is one box, one engine, 12 prompts. It does mean that on the engine people will actually run, the series is text-neutral on a widened set, in line with the equivalent answers wording above.

    For completeness, the per-pack speed numbers already in PRs #3535 / #3627 carry over unchanged — none of these settings changed.
21. rafatxf commented on Oct 22, 2026

    

    rafatxf

    on Oct 22, 2026

    ContributorAuthor

    More actions

    @yusheng227507 thanks, that closes the last open line. On 0.1.47, the engine a default would ship in, IQ3_XXS is now 12/12 byte-identical across A/B/C/M8, like IQ3_S.

    I'll update the IQ3_XXS line in the issue body:

    • on 0.1.46.3, two near-tie prompts out of eight moved;
    • on 0.1.47, 12/12 byte-identical in every arm on the same box, with a wider prompt set.

    That leaves "equivalent answers" as the general claim, with no measured difference from the series on any pack on the current engine.

    Reproduction status: no new measurements in this comment.
22. yusheng227507 commented on Oct 23, 2026

    

    yusheng227507

    on Oct 23, 2026

    More actions

    One more sm_75 datapoint, off the main series: DFlash2 vs MTP on this box (engine 0.1.48 self-built sm_75, IQ3_XXS, --spec 4 --mtp rt, 2× RTX 2080 Ti, layer split 26/22).

    Same standard suite on both arms (3 runs each, 33/33 ok), the only difference being the drafter. DFlash2 packed from the pinned revision a22a1744 (760 MB Q8 pack), no fallback lines, the ~1.2 GB drafter fits VRAM fine, and the offer/accept pattern confirms the drafter really was different (offered/req median 98 → 95; offered total 860 → 845):

    metric	MTP	DFlash2	Δ
    short chat decode tok/s	83.2	84.7	+1.8%
    agent session decode tok/s (last-turn)	86.1	85.4	-0.8%
    4K prompt tok/s	910.0	912.5	~flat
    16K prompt tok/s	1600.2	1604.3	~flat
    agent cache hit	69.3%	69.4%	~flat
    draft acceptance (chat+4K)	61.3%	62.4%	~flat

    Net: a wash on Turing — chat slightly faster, agent slightly slower, everything else flat, acceptance about equal. Keeping MTP. If you want a row for the suggestion table in strata_control/drafting.py (currently other: not measured), a paste-ready line:

    sm_75 (Turing, e.g. 2× RTX 2080 Ti): ≈tied (chat +1.8%, agent −0.8%, rest flat) — MTP kept
23. yusheng227507 commented on Oct 23, 2026

    

    yusheng227507

    on Oct 23, 2026

    More actions

    Follow-up on the DFlash2 row: the pack size matters, and it moves the wrong way.

    Ran the same both-arm suite on the (abliterated) IQ3_S pack — same engine 0.1.48 (self-built sm_75), same box, 33/33 ok each arm:

    metric	IQ3_XXS: MTP → DFlash2	IQ3_S: MTP → DFlash2
    short chat decode tok/s	83.2 → 84.7 (+1.8%)	75.2 → 73.1 (-2.8%)
    agent session decode tok/s	86.1 → 85.4 (-0.8%)	80.4 → 75.8 (-5.7%)
    4K / 16K prompt tok/s	~flat	~flat
    agent cache hit	~flat	~flat
    draft acceptance (chat+4K)	61.3% → 62.4%	68.1% → 64.1% (-5.8 pp)
    offered / req (chat median)	100 → 100	97 → 106

    So on IQ3_S DFlash2 is visibly worse, outside the noise band (consistent across all 3 runs), with a clean mechanism: the drafter proposes deeper windows (offered 97 → 106) that verify rejects more (acceptance -5.8 pp). Unconfirmed hypothesis: the drafter was trained on the base model and mismatches the abliterated IQ3_S distribution; the (non-abliterated) IQ3_XXS pack still matches well enough to tie.

    Updated paste-ready line for the suggestion table in strata_control/drafting.py:

    sm_75 (Turing, 2× RTX 2080 Ti): IQ3_XXS about tied (chat +1.8%, agent -0.8%); abliterated IQ3_S DFlash2 loses (chat -2.8%, agent -5.7%, acceptance -5.8 pp) — MTP kept
24. atmJava

    added a commit that references this issue on Oct 23, 2026

    docs+tools: the thinking budget is a small lever, and a bigger one su…

    ...

    0b41dd8
25. rafatxf commented on Oct 23, 2026

    

    rafatxf

    on Oct 23, 2026

    ContributorAuthor

    More actions

    @yusheng227507 before the DFlash2 row goes into the table, a check on what actually ran: in 0.2.3 a served model can't use DFlash2 yet, so I think both of your arms ran MTP.

    Why I think so:

    • The Terminal and app don't pass it. strata_control/drafting.py has ENGINE_SERVES_DFLASH2 = False, and its comment explains why. With "drafting": "dflash2", decide() returns MTP with the note "DFlash2 isn't available in this mode". apply() then only strips --draft/--draft-model, so the engine arguments are the same as the MTP arm. The server prints [strata] drafting: Drafting: MTP (DFlash2 isn't available in this mode) at start, and /metrics carries it as drafting_note. The engine itself prints no fallback line, because it is never asked.
    • The engine would refuse it anyway. Started by hand with --serve, --batch or a layer split, --draft dflash2 logs --draft dflash2 is off: it is for one card and the CLI run (not --serve, --batch or a layer split) in this build; the MTP drafts instead (generate.cpp, around line 5478 at v0.2.3).
    • So the IQ3_S gap needs another explanation. Your earlier comment showed that on an AVX2 CPU without the exact-comparison recipe, output varies run to run, because a token's expert rows round differently depending on how many drafts share its window. That alone changes text lengths and offered drafts. If the arms ran in blocks rather than interleaved, order effects (cache state, thermals) can add to it. Interleaving the arms and checking the drafting line in the server log would settle it.

    Where DFlash2 does run: one card, CLI spec loop.

    • Hardware and build: RTX 2080 Ti 22 GB, i7-6700 (4 cores, AVX2), engine 0.1.48 built for sm_75. The drafter is incoai/Qwen3.8-Flash-Next-DFlash2 packed with tools/dflash2_pack.py (760 MB).
    • Model: Unsloth UD-IQ4_XS, 32K context, --resident-budget-gib 29.
    • Run: 512 greedy tokens per prompt. MTP at --spec 4 --spec-min-p 0.7 against DFlash2 in three settings.
    • DFlash2 was really active: the engine printed DFlash2 drafter on, the drafting cost was about 2.8-3.1 ms per round, and the window-size distribution changed.

    prompt	MTP tokens/round	DFlash2 tokens/round (best setting)	MTP tok/s	DFlash2 tok/s (best setting)
    chat	2.04	2.40 (spec 8, min-p 0.7)	18.7	17.1 (spec 6)
    code	3.05	5.09 (spec 8, draft-conf 0.5)	13.6	13.8 (spec 6)
    4K prompt	1.88	2.05 (spec 8, min-p 0.7)	11.3	12.0 (spec 8)
    16K prompt	2.13	2.20 (spec 8, min-p 0.7)	12.2	12.2 (spec 8)

    DFlash2 accepts more per round, most clearly on code (+45-67 %), but the speed comes out even: −9 % on chat, about flat on code, +7 % on the 4K prompt, flat at 16K. The wider windows it proposes cost the CPU more experts per round. On this 4-core box with a 4-bit pack, the CPU expert pool is what limits decode, so the extra accepted tokens are paid back in verify cost.

    So on Turing today it reads as a wash where it can run (one card), and untested where most of us run (served, layer split) until the engine supports it there. A box with more cores, or a pack whose experts mostly fit in VRAM, might come out differently.

    Reproduction status: measured locally; every number here is from my runs on this box.
26. atmJava

    added a commit that references this issue on Oct 24, 2026

    docs: mode 8 is confirmed and kept, the three upstream comments are p…

    ...

    bb4e0aa
27. rafatxf commented on Oct 24, 2026

    

    rafatxf

    on Oct 24, 2026

    ContributorAuthor

    More actions

    Thanks @atmJava. A 14-core Xeon is a new CPU class for this thread, and both results fit the proposal. I'll fold them into the body:

    • STRATA_SM75_PERF: your box joins the byte-identical list (IQ3_XXS on 0.1.47, gate config, the same answer for =0 and =1).
    • Expert mode 8: +8.5 % / +7.9 % against a 0.7 % A/A2 spread is the largest gain reported so far; the body says +2.7 % to +4.9 % from three setups. I'll widen that line to four setups.
    • --spec on a native pack: good catch. The gate measures a spec window, not plain decode. I'll say so next to the gate recipe.

    On the one first-request divergence: I don't have a confirmed cause. As I read docs/BATCHING.md, the recipe pins what can change within a run: the CPU kernels' grouping, the PCIe share per window, the adaptive tier, and prompt borrowing. It doesn't pin everything that can differ between two cold starts. One candidate I can point to in the code is the prompt chunk. With --prefill auto, the chunk is sized at start from the VRAM room the expert cache leaves (generate.cpp logs prompt chunk auto: N tokens). If your 7.6K-token prompt is longer than the chunk on some start, it is read in a different split, and its KV rounds differently from the first token on.

    Two cheap checks would tell:

    • If you kept the engine log of that start, compare its prompt chunk auto and expert-cache slot lines with a normal start's.
    • Pin --prefill N (N larger than the prompt) and repeat the restarts.

    If neither moves it, a cold-start line in the body ("the first request after a start can differ") is the honest fallback.

    Reproduction status: no new measurement in this comment; the code references are from v0.2.3.
28. atmJava

    added 2 commits that reference this issue on Oct 24, 2026

    ledger+docs: Niko1221#3724 judged by reading (partly reachable, expec…

    ...

    aed495f

    docs: the user closed the upstream route - the three posted comments …

    ...

    d3ed673
29. yusheng227507 commented on Oct 24, 2026

    

    yusheng227507

    on Oct 24, 2026

    More actions

    rafatxf is right — the two DFlash2 comments above are invalid; please disregard their data.

    Verified on my side: engine 0.1.48 src/program/generate.cpp gates DFlash2 to it is for one card and the CLI run (not --serve, --batch or a layer split) in this build, and both my arms were --serve runs on a 2-GPU layer split — so the engine took the MTP fallback in both. The serve logs confirm it byte-exactly: the arms' expert-cache footprints are identical to the plain MTP runs (IQ3_XXS 9701 experts / 14.10 GiB of VRAM; IQ3_S 7534 / 13.66 GiB — a loaded DFlash2 pack would have shrunk both), and no drafter line appears in the log. The offered/accept deltas I leaned on (offered/req 98 → 95, 97 → 106) are run-to-run expert-placement variance on an IQ pack, not drafter evidence.

    So the comparison tables in those comments (84.7 vs 83.2, 73.1 vs 75.2, and the acceptance lines) are MTP-vs-MTP replicates on my box and must not go into the suggestion table. Sorry for the wrong first read — the "really using the drafter" check I did was indirect and I should have looked at the engine's own gate in the source first.

    I will rerun DFlash2 the way the engine actually allows it — CLI, one card, no split (the 22 GB card holds the XXS native pack) — and only post a number if that run gives a clean reading.
30. yusheng227507 commented on Oct 24, 2026

    

    yusheng227507

    on Oct 24, 2026

    More actions

    The clean DFlash2 reading for the suggestion table (CLI, single card, as the engine allowlists):

    Engine 0.1.48 self-built sm_75, RTX 2080 Ti 22 GB, one card, no split, IQ3_XXS native pack, greedy (--temperature 0 --top-k 1), 160 max-new, one run per anchor (the same four upstream prompts), expert cache auto with the shipped env (STRATA_SM75_PERF=1, STRATA_BF16_TC=1, STRATA_PA_V2T=1, STRATA_EXP_MODE=8). The engine confirms the drafter this time: strata generate: DFlash2 drafter on (...pack...), the MTP layer stays as the fallback (992 MiB of draft VRAM, 3.55 ms/round drafting).

    anchor	MTP tok/s	DFlash2 tok/s	Δ
    chat_explain	55.11	49.26	-10.6%
    code_py	61.13	57.96	-5.2%
    math	55.86	55.01	-1.5%
    creative	46.90	42.61	-9.1%

    Tokens per round are about tied (1.86–2.81 across anchors, both arms) — the loss is the drafting cost itself plus the costlier verify on Turing, the same mechanism that loses at wider --spec on this card (what the guard's own "4% to 16% slower" band says). So this joins the ampere row more than the blackwell one: DFlash2 slower on every prompt class here too. Paste-ready for the suggestion table in strata_control/drafting.py:

    sm_75 (RTX 2080 Ti 22 GB, CLI one-card, IQ3_XXS native pack): DFlash2 slower on every prompt (chat -10.6%, code -5.2%, math -1.5%, creative -9.1%) — MTP kept