Repository navigation
Conversation
A request continues from the longest prefix of its prompt the engine still holds: the live sequence, checkpoints of the fixed-size sequence state (GDN recurrence and conv history, QSA indexer tail, PLE history) in host memory, or a conversation another request replaced, stashed in RAM with its K/V cells. Checkpoints are taken before a prompt's last <|im_start|>, at the prompt's end and every --cache-every cells of a long prompt. Short prompt remainders go through verify windows instead of the batched prompt path. The decode loop no longer verifies past the output limit or a drafted end-of-turn token, so the live state ends where the emitted text does. The server reports reused tokens (OpenAI prompt_tokens_details.cached_tokens, Anthropic cache_read_input_tokens) and drains the engine to DONE inside the request lock. New engine options: --prompt-cache N, --cache-ram MIB, --cache-every N, --feed-max N.
The engine now runs unsloth's Qwen3.8-Flash-Next quants (UD-IQ1_S up to UD-Q6_K_XL and Q8_0) as native packs, besides the ISTA-DASLab files. - Split models: --native takes the first shard and finds the others by name (gguf_split_paths, GgufModel). The expert arena is filled from the shards by tensor name and checked against native_experts.txt, whose lines no longer carry shard-1 offsets. The native head and embedding are found in any shard, and --ple-gguf defaults to the shard that holds per_layer_token_embd. - GPU kernels: expert dot products and dequantizers for Q4_K, Q5_K, Q6_K, Q5_1, Q8_0, IQ1_S and IQ4_XS (llama.cpp's), one dispatch list for all formats, and a startup check that every layer's formats have kernels. The Q5_1 min and IQ1_S delta terms use the quantized activation sums, as ggml-cpu does: 1.1% instead of 1.9% relative error per expert. - PLE: the table may be IQ4_NL or Q8_0 and may share its shard with other tensors; the native PLE key may be any native MMVQ type (Q8_0 here). - tools/iq_pack.py: all shards of a split model, and every float tensor in the form the engine takes. The files' F32 router, SSM gate and injection tensors hold BF16 values (exact); their Q8_0 hyper-connection and PLE value weights are re-rounded to BF16. - native_expert_parity: split models, a --synthetic mode for formats with no model file at hand, and a check of the GPU dequantizers. - README: the unsloth files, their RAM needs and how to pack them. UD-Q4_K_XL on an RTX 3090 (160 GB RAM): prompts 529-594 tok/s, generation 49-57 tok/s. Greedy tokens equal llama.cpp's argmax on the same GGUF at 98.2% (code) and 95.6% (thinking) of positions.
A batched prompt call costs a fixed 4.7-6.6 s on an RTX 3090 behind a PCIe 4.0 x4 link: it lends and refills expert-cache slots and streams every expert its tokens use, however few tokens it holds. Verify windows cost ~15 ms (IQ3_XXS) and ~21 ms (UD-Q4_K_XL) per token, so the two break even at ~480 and ~530 tokens. - --feed-max defaults to 512 (was 256). - A last prompt chunk of up to --feed-max tokens goes through verify windows as well. A prompt just past a multiple of --prefill used to run its last tokens as a chunk of their own (3.6 s on UD-Q4_K_XL). - Chunks and checkpoint segments count from the request's first new token, so a follow-up that starts inside a segment no longer gets a short first chunk. --cache-every is rounded down to whole --prefill chunks. - Image requests always take the batched path, the only one that reads image rows. A request of up to --feed-max tokens with an image used to lose the image. Prompt speed with the prompt cache off, 16K-131K tokens: IQ3_XXS 703-766 tok/s (was 674-747), UD-Q4_K_XL 570-603 tok/s (was 529-594).
- MSVC: _CRT_SECURE_NO_WARNINGS (fopen, sscanf and getenv are standard C), /wd4211 for the host stubs nvcc generates for kernels with internal linkage, and C4324 off around ExpertPool, whose alignas padding is the point. - Unused locals, device functions, a kernel and a constant removed; shadowed locals renamed; the router clamp bound is an if constexpr. - qsa_layer's prologue, a single 917-character line, split into lines.
The draft head copies rows of the model's output head, and the copy took rows in 16-byte units only. Q6_K rows of 2,560 values are 2,100 bytes, so MTP exited at startup with the Q6_K heads of unsloth's UD-IQ3_XXS, UD-IQ4_XS and UD-Q3_K_XL. Rows that are a multiple of 4 bytes are now copied in 4-byte units.
With --spec-follow, a generate run emits a given continuation instead of the window's argmax and accepts the drafts that match it. Runs with different --spec, --spec-min-p or draft-layer settings then process the same tokens and route the same experts, so their speeds compare within ~1% instead of the ~4% two identical free runs differ by once their text diverges at a near tie. The usage text also lists the MTP flags.
…n --stats - The serve loop and the speculative loop each carried the same adaptive expert tier (ranking, refill copies, admission once landed); it is now strata::core::AdaptiveTier. - --stats prints the share of the routed entries the N most-routed experts of the run take, which sizes what a larger expert cache could serve.
The GPU that holds the expert cache is one per-layer bottleneck and the CPU pool, which computes the missed experts, the other. With --second-gpu N, device N holds more experts (the profile's pairs the first cache lacks, then empty slots its own adaptive tier fills with the conversation's) and computes its share of each verify-window layer while the CPU pool computes the rest. It never spins on host memory: per layer the host launches one captured graph (the plan copied in by a kernel, the activations read from mapped memory, the grouped expert kernels of the first GPU's hit path, the rows written to mapped memory) before the pool and collects the rows after it, so a GPU that also drives the display keeps its latency. A layer whose misses are below --second-gpu-min-mb (4) stays on the CPU, which is then quicker than the ~100 us round trip. UD-Q4_K_XL, RTX 3090 + RTX 5070 Ti (12.6 GiB of experts): generation +9% on a 3000-token thinking answer, +14-19% on code; IQ3_XXS, whose rounds wait on the first GPU, gains nothing. The expert arena is now one cudaHostAlloc block per layer instead of a registered reservation. With a second CUDA context in the process, Windows pins at most ~42 GiB through cudaHostRegister, and past that every CUDA allocation fails; allocated blocks reach ~75 GiB with one GPU or two. The layers past the budget share one locked, pageable range as before.
…e; no PCIe stage without a PCIe share
…ert gate (one weight read, one launch)
…refetched and double-buffered down tiles, two up rows per warp (bitwise the single-token read)
…shes RoPE (native_rope_apply_tokens), the int8 K/V append (kv_append_q8_steps), the indexer append (native_qsa_indexer_append_multi), the q/gate split, the q/k norms and the gate now run once per window instead of once per token. Each token's arithmetic is unchanged: verify_parity checks the new kernels bitwise against the per-token ones, and --window-hashes writes a hash of each window's final residual rows so two builds can be compared end to end.
…nk; --prefill-profile The PLE block ran once per prompt token (a 52 MB key GEMV each); its key and value projections are now GEMMs over the chunk, followed by the decode path's per-token arithmetic in one pass (native_ple_postops_tokens, bitwise n_tok native_ple_postops + ple_history_advance calls; verify_parity checks it). The indexer appends of a chunk are one native_qsa_indexer_append_multi launch. 16K-token prompt, IQ3_XXS: 22.8 -> 18.3 s. --prefill-profile prints the GPU time of each section of the prompt path; --window-hashes lines carry the window's first position.
…tion in one pass per query qsa_block_scores launched a warp for every block of the full context (max_cells / 4, 50K blocks at 200K) for every query, nearly all of them idle. Its warps now go round over the blocks: the verify window and the token path use a fixed grid of 1024 per query, the prompt path the chunk's largest n_bid + 1. Each score is computed as before (bitwise). qsa_prefill_attn: the prompt path's attention, a block per (KV head, query) walking the query's selected cells in tiles of 64 with an online softmax, instead of the decode kernel's split-K partial sums through memory. FP32; int8 codes are converted without I2F. verify_parity checks it against the decode kernel (relative difference 3.9e-6). 16K-token prompt: IQ3_XXS 18.3 -> 15.9 s, UD-Q4_K_XL 24.8 -> 22.4 s.
…rrence The conv + L2 norm walked the chunk one token at a time per channel (320 warps); the prompt path now uses the verify window's gdn_conv_l2_multi over all T tokens plus a small history update. The recurrence keeps the decode kernels' arithmetic (four 32-row partial sums per column, added in order) but puts a column's row groups in adjacent lanes, so the sums are shuffles; each thread holds two columns, a head is split over four blocks, and each token's inputs arrive through a cp.async ring 7 tokens ahead: one barrier per token. The output norm is a separate pass. verify_parity checks the prompt path's GDN bitwise against the verify window's kernels (1-37 tokens). 16K-token prompt: recurrence 1.36 -> 0.46 s; IQ3_XXS 15.9 -> 15.2 s, UD-Q4_K_XL 22.4 -> 21.7 s.
… GPU A prompt chunk uses nearly every expert, and the 3090 streamed all it did not hold over its PCIe 4.0 x4 link. With --second-gpu the prompt path now hands them to the second GPU (x16): per layer the 3090 copies the chunk's MoE input (FP16) to pinned host memory while its router and shared expert run; the second GPU computes its experts from its own cache or streamed from the pinned arena, and returns each token's weighted sum (FP32), which the 3090's combine adds to its cached experts' rows (prefill/offload.cpp). Experts go through in batches of consecutive ids (up to 16 experts and 4096 rows): one staging copy per run of adjacent arena blobs, one dequant launch (iq_dequant_experts_f16), one gather, per-expert GEMMs (cuBLAS does not group FP16 -> FP32), one SwiGLU and one ordered per-token add (moe_gather_add, the same sums as one moe_scatter_add per expert). --no-prompt-offload streams them to the main GPU as before. 16K-token prompt, UD-Q4_K_XL: 21.6 -> 11.9 s.
…ots too; one SlotLoan for both GPUs and both modes
… a layer's weights dequantized once; the routed experts of both GPUs through one batched runner; buffers sized to each prompt (16K: 1.4 GiB of cache slots lent instead of 10.9)
… steps run; per-expert GEMMs over side streams; the second GPU's sums return as FP16; pinned uploads
…verify windows at ~185 (IQ3_XXS) and ~45 tokens (UD-Q4_K_XL)
… RAM is idle (--second-gpu-prefetch 4), its rows out in full PCIe transactions, the second tier's moves paced behind the prefetch Each layer's router also runs on the previous layer's FFN input on a side branch of the verify window's graph; once the CPU's rows of a layer are in, the host copies the next layer's likeliest experts that no GPU holds into prefetch slots on the second GPU (~0.3 ms of idle RAM per layer), and that layer's share takes them from there. The second GPU's down kernel wrote its rows one float at a time into mapped host memory: 0.25 ms a layer, and the CPU pool read RAM at 38 instead of 48 GB/s. UD-Q4_K_XL code, fixed text: 71.6 -> 77.4 t/s (rows) -> 86.1 (prefetch 4); thinking 70.2 -> 73.6.
…arted on the CPU's (its launch took the host ~15 us before them); UD-Q4_K_XL code +0.5%
…nched and waited for alone, 0.67 ms a round); UD-Q4_K_XL code +0.9%
… (28.8 GB in ~7 s), direct reads until then; a 16K-token prompt of new text 0.5 s faster
… idle RAM window holds 4 UD-Q4_K_XL or ~7 IQ3_XXS experts (IQ3_XXS code 98.9 -> 108.3 t/s with the second GPU)
…lone (it failed to start); its budget always fills, the rest as empty slots (without the adaptive tier IQ3_XXS had nothing there and lost the prompt offload)
…rupt core, copies in batches Windows sends each GPU's interrupts to one logical processor (here the 5070 Ti's to the first, the 3090's to the second, one MSI each), and the host thread ran on the first: generating IQ3_XXS thinking text the 5070 Ti raised ~12,000 interrupts a second, each ~13 us in dxgkrnl's ISR and ~3 us in DPCs, a fifth of that processor. A copy raises one when it lands, so the prefetch's copies, one per expert, raised most of them. The host saw a layer's ring up to 30-110 us late (p99), its call submitting the 5070 Ti's share took 18-24 us (14-17 on another core), and in the pool's gate/up phase it finished last in 42% of the layers with a share. - The host thread on the last physical core, the pool's workers on the others; with a second GPU the first core stays with the OS and its interrupts (6 workers here, from 7). core_plan() decides; the startup log says where the threads run. - The 5070 Ti's prefetch copies and each adaptive-tier call's copies in one cudaMemcpyBatchAsync (copy_blobs): the first processor's interrupt time 10% -> 4% on IQ3_XXS thinking text. Window hashes identical to 7105fcc's (--adapt-every 0, both models); pool_test --stress passes. Fixed text, 4 runs each in both orders, t/s: both GPUs IQ3_XXS code +1.2%, thinking +0.5%, after 16K +1.3%, edit -0.7%; UD-Q4_K_XL +5.9 / +0.7 / +1.7 / -0.1%; the 3090 alone IQ3_XXS code +4.9%, thinking +0.9%, UD-Q4_K_XL +3.2 / +2.1%. 16K prompt times unchanged (IQ3_XXS 8.11 -> 7.93 s, UD-Q4_K_XL 8.53 -> 8.48 s). Claude-Session: https://claude.ai/code/session_019XxRxQQNyk7Ycdr4ComSs1
Windows moves device interrupts between processors as the load shifts (the power plan's interrupt steering, "Default" here): the GPUs' went to logical processors 0 and 1 one morning and to 1 and 2 that night, beside a network card's ~26,000 a second on 0, so the 3090's landed on a pool worker (up to ~45% of its processor in ISRs and DPCs). Moving that worker off its processor mid-run left the interrupts where they were (two traced runs: to its core's other processor, and to another core). So CorePlacement starts the threads where 5137b33 put them and, once a second between windows, reads each processor's interrupt and DPC time: a thread whose processor spent over 5% there two seconds in a row moves to a quiet processor, a free core's, else its own core's other one. Each move is logged; --stats counts them. Window hashes identical to 5137b33's (--adapt-every 0, both models); pool_test --stress passes. Fixed text, 2 runs each in both orders against 5137b33: no thread had to move (the interrupts stayed on the first core), speeds within the runs' spread (IQ3_XXS +0.2 / +0.2 / +0.2% on code / thinking / edit, UD-Q4_K_XL +7.8 / -0.4 / -0.4%). With the threshold lowered for a test, a worker on the first core moved to its other processor on the 3090 alone. Claude-Session: https://claude.ai/code/session_019XxRxQQNyk7Ycdr4ComSs1
…server fixes Upstream's 236 commits since 6da1f66, with this branch's engine kept wherever both sides do the same thing (the prompt cache, prompt lookup, the verify window's batched kernels, the second GPU tier, the prompt path, the CPU pool's phase protocol). Taken: - serve/: upstream's server (web app with chat, monitor and MCP tools, Claude Code's /v1/messages?beta=true, llama.cpp's discovery endpoints and timings, --fit-max-tokens, LAN access with an API key, an engine that dies is restarted by the next request); the wait loop sleeps so Ctrl+C stops it on Windows. - Sampling: temperature, top_p, top_k, min_p, seed and the penalties per request, on the GEN line. The verify window samples its head after the graph (Philox(seed, position)), with one penalty-history row per window row; windows that read a prompt skip it. spec_min_p sets a request's draft floor. - The serve protocol: INFO before READY, PP lines per prompt chunk, DONE with the drafts accepted and offered, the reused tokens and the expert-cache hits of either GPU. - CPU experts: the AVX-2 multi-token i-quant kernels for CPUs without AVX-512, IQ4_NL down rows of 2+ tokens decoded once for all of them (STRATA_NO_IQ4NL), the i-quant rows prefetched 2 KB ahead (STRATA_IQ_PREFETCH). - --prefill auto (setup writes it): 16384-token chunks, halved until their buffers fit the lendable slots. - setup.py and its tools, tools/needle_bench.py, draft_vocab.py, make_profile.py, calibrate.py; docs. Kept from this branch: data/expert-profile.bin (the 8,000-slot cap; upstream's full ranking would give IQ3_XXS's 3090 cache ~1.3 GiB more, measured no faster). Not taken: upstream's conversation cache and prompt lookup, the helper-GPU caches and layer split, KV streaming and the q4/k8v4 KV, control vectors, the low-RAM mode, pruned expert counts, the HIP and Turing builds (their files are in the tree, not built). Window hashes on fixed text (--adapt-every 0) equal the previous build's on both models (IQ3_XXS with STRATA_NO_IQ4NL=1: the IQ4_NL kernel sums in another order). Claude-Session: https://claude.ai/code/session_019XxRxQQNyk7Ycdr4ComSs1
data/expert-profile.bin is upstream's ranking of all 24,576 (layer, expert) pairs (it ranked 8,000). The 3090's cache takes the pairs it fits, the 5070 Ti's tier the next ones: with 8,000 pairs that tier started with 523 of IQ3_S's 6,619 slots filled (UD-Q4_K_XL: 3,463 of ~4,300) and its adaptive tier filled the rest over thousands of tokens. Fixed text, both GPUs, 4 runs each in both orders, t/s against the 8,000-pair profile: IQ3_S code 173.5 (+15.8%), thinking 130.2 (+4.6%), after 16K 138.6 (+6.3%), edit 236.4 (+9.6%) UD-Q4_K_XL code 115.4 (+4.8%), thinking 102.3 (-1.1%), after 16K 119.8 (-0.2%), edit 161.8 (-1.1%) Claude-Session: https://claude.ai/code/session_019XxRxQQNyk7Ycdr4ComSs1
The verify window's head ended with sampler_greedy_kernel, one 1024-thread block a token over the 248,320 logits, and each draft pass with it and row_top_prob over the 40,525-token draft head, one block each. argmax_rows scans a row in up to 128 blocks (4,096 logits each), the row's last block merging their picks (the larger value, on equality the lower index: the one-block kernel's order); row_top_prob_split computes row_top_prob's 32 warps' sums in 8 blocks, the row's last block adding them in order. Both keep their counters and partials in per-row slots of one scratch, whatever a launch's row count. Bitwise: verify_parity (ties, NaN, -inf, 1-8 rows on one scratch), and the window hashes on fixed text equal the previous build's (IQ3_S, UD-Q4_K_XL; thinking and code text). CUPTI, IQ3_S thinking text: the head's argmax 84-86 -> 9.7 us a window, a draft pass's argmax and top probability 15.1 + 10.1 -> 5.6 + 4.3 us; ~105 us a round (~0.7%). Claude-Session: https://claude.ai/code/session_019XxRxQQNyk7Ycdr4ComSs1
MtpDrafter::capture_round ran the draft layer's front (the embedding, the fc projections, the attention hyper-connection read, the K/V projections and appends) over the window's T cells, then the whole layer again for row a, repeating row a's front. record_forward is now record_front (T rows) and record_rest (one row: the attention from the q projection, the MLP, the final mixer, the head); the round graph runs the front once, copies row a's residual, injection and mixed input to row 0 (copy_row_to_first, the row read on the device) and runs the rest there. The multi-row front is per row the one-row front's arithmetic: the window hashes on fixed text equal the previous build's (IQ3_S, UD-Q4_K_XL; thinking and code text). CUPTI, IQ3_S thinking text (--adapt-every 0): the round graph 721 / 747 / 789 / 784 / 898 -> 655 / 662 / 689 / 692 / 812 us after windows of 1 / 2 / 3 / 4 / 8 tokens. Claude-Session: https://claude.ai/code/session_019XxRxQQNyk7Ycdr4ComSs1
A one-token window always keeps its token, so its graph now advances the sequence state itself: the GDN step kernel writes the state (n_keep = 1, the arithmetic the commit's replay used), the conv kernel writes the history it read (gdn_conv_l2_multi's commit flag), and its indexer append and PLE history were final already. Verifier::commit launches no commit graph after it: the draft round graph then runs alone instead of beside the commit's 98 kernels, which read and write each GDN layer's 3 MB state. Window hashes on fixed text equal the previous build's (IQ3_S, UD-Q4_K_XL; thinking and code text); cache_test.py's regenerated turn and stash round trip reproduce their output. CUPTI, IQ3_S thinking text (--adapt-every 0), the 35% of rounds after one-token windows: the window 11,270 -> 11,380 us, the time from its end to the next window's start 3,706 -> 3,325 us (the round graph 655 -> 450 us, the draft starting 150 us sooner): -270 us a round there, ~-95 us over all rounds. Claude-Session: https://claude.ai/code/session_019XxRxQQNyk7Ycdr4ComSs1
- quantize_q8_0_scaled runs a warp a 32-value block (it ran a thread a block: 80 threads for the draft layer's input): the block's |max| over the same values, then each value's bytes as before. 6.2 -> 2.4 us and 9.2 -> 1.6 us for the Q2_0 draft experts' two quantizations (CUPTI). - The draft graphs' inputs from mapped memory in one copy kernel (copy_from_mapped_multi): 7 launches in the round graph, 2 in a step, became one each. - The router in one kernel (verify_router: the logits, the top 10, and the native draft experts' q8_1 rows), bitwise the gemv, top 10 and quantizer it replaces; as fast as the gemv and top 10 were. - The shared expert on a branch of the draft graphs beside the router and the routed experts, joined before the combine. Window hashes on fixed text equal the previous build's (IQ3_S, UD-Q4_K_XL; thinking and code text; IQ3_S with the IQ4_NL draft experts too). CUPTI, IQ3_S thinking text (--adapt-every 0): the round graph 450 / 661 / 687 / 693 / 824 -> 411 / 592 / 620 / 627 / 760 us after windows of 1 / 2 / 3 / 4 / 8 tokens, each step 435 -> 405-410 us. Claude-Session: https://claude.ai/code/session_019XxRxQQNyk7Ycdr4ComSs1
After the last layer the head streamed the whole output matrix on the 3090 (IQ3_S's Q6_K 521 MB, UD-Q4_K_XL's Q8_0 675 MB). With --second-gpu the 5070 Ti now holds its last rows (--head-split, default 0.45 of them: 234 / 290 MiB, which its expert tier gives up and the 3090's cache gets back): the window's graph writes the head's input rows to mapped memory and raises a count, the host launches the 5070 Ti's part once it sees it (the same quantization, projection and argmax kernels, so every logit is the one-GPU head's), and keeps each token's larger logit, on equality the 3090's. argmax_rows can return the pick's logit. A sampled request or --window-logits gets whole rows: the 5070 Ti's part then copies its logits to mapped memory and the 3090 joins them to its own (join_rows). The drafter gathers its draft head's rows before the 3090 keeps only its part (MtpDrafter::bind_head); without a draft head the head stays whole. split_head.hpp/.cpp. With both GPUs' caches held at one size (--expert-cache 3000 --second-gpu-gib 10) the window hashes and the --window-logits output equal the previous build's (IQ3_S); with the configs the caches' sizes change with the head's. Fixed text, 2 runs each in both orders, t/s against d3d48b1 (at 0.5 / 0.45 of the rows on the 5070 Ti): IQ3_S code +5.3 / +5.8%, thinking +2.9 / +2.5% UD-Q4_K_XL code -6.3 / -0.4% (the 0.5 build ran second of three in both rounds, where UD's code runs are slowest: 106 against 110-117 t/s), thinking +1.4 / +2.7% The 5070 Ti's part starts ~30 us after the 3090's (the host sees the count, launches under WDDM): balanced near 0.46. Claude-Session: https://claude.ai/code/session_019XxRxQQNyk7Ycdr4ComSs1
The hyper-connection read before the head ran gr_read's kernels once per token of the window (lm_head_mix), and once per draft pass in the draft layer. Both now go through fused_gr_read_multi, the layers' read (no injection weight: the final mixer), the window's tokens in one launch. Its sums are gr_read's: window hashes on fixed text and the distributions along llama.cpp's texts (scratch/parity_new.sh: IQ3_S and UD-Q4_K_XL, code, thinking, after 16K) equal the previous build's. CUPTI, IQ3_S thinking text (--adapt-every 0): the time between windows 3,679 -> 3,578 us a round, each draft step 6-9 us shorter. Claude-Session: https://claude.ai/code/session_019XxRxQQNyk7Ycdr4ComSs1
…mixer's write with the FFN norm Upstream's F-1 and F-2 (882bb6d, b046845) in this prompt path. A read normalized its rows into an FP32 copy for the mix and a BF16 one for the projections: gr_norm_rs now writes the row scales and the BF16 image only, and gr_mix_r recomputes each normalized value from the residual rows, in gr_norm's order. The mixer half's write runs with the FFN half's norm (gr_write_norm_rs: one pass over the rows, the reduction as gr_norm's), so the FFN read starts at its projections. The FFN half's write covers the whole chunk after the experts while the next layer's reads run per 2048-token sub-chunk, so it stays apart; the draft layer's prompt pass keeps the FP32 rows (its up projection overwrites its source rows). Window hashes after the 16K prompt equal the previous build's (both models). 16K-token prompt (bench strata_gen.sh, alternating builds): IQ3_S with both GPUs 7,816 -> 7,638 ms (-2.3%, 3 runs each), on the 3090 alone 11,691 -> 11,633 ms (-0.5%: there the dense steps overlap the expert copies); UD-Q4_K_XL with both GPUs 8,399 -> 8,163 ms (-2.8%, 2 runs each). Claude-Session: https://claude.ai/code/session_019XxRxQQNyk7Ycdr4ComSs1
…ayer's FFN write deferred to the next layer's read A layer's expert phase ended with the second GPU's FP16 sums for the whole chunk crossing the 3090's x4 link (80 MiB for 16K tokens, ~13 ms a layer) and one write over the chunk, with the 3090 idle meanwhile. The sums now come back in 2048-token pieces on a stream of their own, and each sub-chunk's FFN write runs at the start of that sub-chunk in the next layer, after its piece has landed and with the next read's norm in the same pass (gr_write_norm_rs takes the partial sums): the next layer starts once the first piece is in, and the rest land beside its sub-chunks. Before the PLE block at layer 1 and after the last layer the writes run without the norm. The sums buffer leaves the buffers' union for the chunk's (the next layer's steps would overwrite it): the prompt borrows 80 MiB more of the caches for a 16K chunk. Window hashes after the 16K prompt equal the previous build's on the 3090 alone, and with both GPUs against it with the sums buffer moved the same way (IQ3_S, UD-Q4_K_XL); against the previous build as it was, the borrowed slots differ, so which GPU computes some experts, and with it the rounding (1 of 100 argmaxes differs on IQ3_S, a near-tie). 16K-token prompt (bench strata_gen.sh, alternating builds): both GPUs, IQ3_S 7,623 -> 7,104 ms (-6.8%, 3 runs each), UD-Q4_K_XL 8,237 -> 7,656 ms (-7.1%, 2 runs each); the 3090 alone 11,755 -> 11,638 ms (-1.0%: the FFN write with the next norm). Claude-Session: https://claude.ai/code/session_019XxRxQQNyk7Ycdr4ComSs1
…step 2b) on both GPUs The prompt path dequantized each batch of experts to FP16 (iq_dequant_experts_f16) and multiplied them with cuBLAS. Layers whose expert types llama.cpp's MMQ covers now keep the weights quantized: a batch's experts' GGUF blocks are gathered side by side in the dequantization buffers' place (mmq::gather_native, a zeroed tail after them), its rows are gathered in FP32 (gather_rows16_f32) and rounded to q8_1, and gate/up and down run as one MMQ launch each over the batch (int8 tensor cores), with SwiGLU between; the weighted scatter-add is as before. Upstream's moe_mmq.cu, and its host shim for ggml-cuda, are built into strata_mmq with the instances of every expert type of ISTA-DASLab's and unsloth's files (the K-quants, Q5_1 and IQ1_S added to its switch; IQ1_M keeps FP16). A single expert with more rows than a batch keeps the FP16 path in pieces; STRATA_PREFILL_MMQ=0 takes it for every layer. The identity row map MMQ writes through is filled with each layer's uploads: the 3090's runner shares the prompt buffers' union with the dense steps, which overwrite it. The activations are rounded to 8 bits, as the decode path's expert kernels already take them. Along llama.cpp's text after the 16K prompt (--adapt-every 0), against llama.cpp: IQ3_S 99.33% / KL 1.11e-3 -> 99.67% / 1.05e-3, UD-Q4_K_XL 98.00% / 3.80e-3 -> 98.67% / 3.95e-3; MMQ against FP16 99.00% / 9.2e-4 and 99.33% / 1.3e-3, the size of llama.cpp's own CPU-against-CUDA difference. 16K-token prompt (bench strata_gen.sh, 2 runs each, alternating): IQ3_S with both GPUs 7,100 -> 6,759 ms (-4.8%), on the 3090 alone 11,626 -> 11,047 (-5.0%); UD-Q4_K_XL with both GPUs 7,772-8,500 -> 7,178 ms (-7.6% or more), alone 17,662 -> 16,828 (-4.7%). Claude-Session: https://claude.ai/code/session_019XxRxQQNyk7Ycdr4ComSs1
…ry, its sums back in pieces The second GPU copied a chunk's MoE input in (80 MiB for 16K tokens) once the first GPU's last sub-chunk had reached host memory and the host had grouped the layer's experts, and started computing ~5.6 ms after the dense steps' end (CUPTI). Each sub-chunk's rows now go on to it as they arrive (ExpertRunner::stage_input, queued behind the previous layer's sums on its stream, which pass through the same buffer), so the input is there when its plan is. Its sums leave in the same pieces (converted to FP16 and copied to host memory piece by piece, piece_done(k)), so the first piece starts over the 3090's link before the last one leaves the 5070 Ti. Window hashes after the 16K prompt equal the previous build's (IQ3_S, UD-Q4_K_XL). 16K-token prompt, both GPUs (bench strata_gen.sh, alternating builds): IQ3_S 6,849 -> 6,686 ms (-2.4%, 3 runs each), UD-Q4_K_XL 7,233 -> 7,044 ms (-2.6%, 2 runs each). Claude-Session: https://claude.ai/code/session_019XxRxQQNyk7Ycdr4ComSs1
…ile the one before runs Each prompt chunk began with its PLE rows read from the SSD (one batched request, ~0.4 s for 16K tokens) and dequantized on the host, both GPUs idle meanwhile (CUPTI: two stalls of 0.36-0.52 s in a 16K prompt). The read now runs on a thread into one of two pinned buffers: the first chunk's from the prompt's start while its embeddings and layer 0 run, each next chunk's from the layer 1 of the chunk before; the PLE block at layer 1 waits for it and the rows go up asynchronously. The rows' token history for the reads runs ahead of the one the prompt cache keeps, which still advances chunk by chunk. Window hashes after the 16K prompt equal the previous build's (IQ3_S, UD-Q4_K_XL). IQ3_S with both GPUs (bench strata_gen.sh, alternating builds, 2 runs each): 16K prompt 6,785 -> 6,357 ms (-6.3%), 64K 25,391 -> 24,372 ms (-4.0%). UD-Q4_K_XL 16K 6,723 ms (7,044 before). The two buffers pin 336 MB more (UD's 71.7 GiB arena stays pinned whole). Claude-Session: https://claude.ai/code/session_019XxRxQQNyk7Ycdr4ComSs1
…+ RTX 5070 Ti (docs/COMPARISON.md) docs/COMPARISON.md holds the current numbers for IQ3_S and UD-Q4_K_XL: speed through the API (prompts 16K-128K, generation on code, thinking, an edit and after each prompt, a follow-up turn, the time to the API, VRAM and RAM), generation on fixed text, agreement with llama.cpp's next-token distributions with their rounding baselines, llama.cpp's perplexity of the files, and the three engines' features. llama.cpp's and upstream's columns were measured once; this fork's are kept current. The README's fork note and DETAILS' second-GPU and unsloth sections point to it instead of carrying numbers of their own. data/draft_vocab.bin is this fork's 40,525-token draft subset again (the merge took upstream's 106,299-token one with the Chinese, Japanese and Korean tokens: a 2.6x larger draft head, English text 1-2% slower by upstream's measure). Claude-Session: https://claude.ai/code/session_019XxRxQQNyk7Ycdr4ComSs1
… setup's layout The engine arguments of the configs COMPARISON.md measures, with paths relative to the Strata folder: the engine in CMake's build/ (setup's engine/ holds upstream's ready-made engine), the model files, packs and MTP draft layer in setup's data folder ../Strata-data. The README says how to build and start them. Both start verbatim from such a layout and answer (checked through junctions to this PC's files). Claude-Session: https://claude.ai/code/session_019XxRxQQNyk7Ycdr4ComSs1
… the MTP layer come from run-iq3_s.bat (port 8081) and run-ud-q4_k_xl.bat (8082) start the server from the Strata folder with their config, as this fork's own launchers do. Both start and answer from a setup-style layout (checked as copies on another port beside them). The README: a pack is tools/iq_pack.py's output from the model's GGUF files (setup makes IQ3_S's, step 3 UD-Q4_K_XL's), the MTP draft layer setup's conversion of the original checkpoint's MTP tensors. Claude-Session: https://claude.ai/code/session_019XxRxQQNyk7Ycdr4ComSs1
Owner
|
Thank you for opening this - your fork is the reference we're porting UD-Q4_K_XL from. As one PR it's too big to review or gate (94 commits, 121 files, many default-path changes). We're taking the Unsloth pieces (8029fa9's shard discovery, per-role expert sources, Q4_K/Q5_K/Q5_1/Q8_0 kernels, the packer forms; later 7706c1b's MMQ dispatch) into upstream under your name, gated so existing models stay byte-identical. The other speedups (second-GPU tier, verify-window rewrites) are welcome as focused PRs, one change each with an A/B. Would you be up for comparing our port against your fork on your 160 GB machine once it runs? Closing this one. |
This was referenced Sep 30, 2026
Closed
Niko1221
added a commit
that referenced
this pull request
Oct 1, 2026
… the mapped mode without experts.bin, --resident-budget-gib, the tier counters - docs/UNSLOTH_Q4.md: what UD-Q4_K_XL is (Hugging Face revision 38bb39e, the four shards with sizes and SHA-256), disk/RAM/GPU needs (111.3 GB download, 1.4 GB pack, no experts.bin), the iq_pack.py command and why --compat-bf16, the numerical policy (native experts/embedding/head/projections/PLE; the converted small tensors listed in conversions.json), the MTP draft layer, a server config (--native shard 1, --ple-gguf found by name, --resident-budget-gib N with N = RAM minus 20-24 GB), measured speed on an RTX 5070 / 64 GB PC (7-8.5 tok/s at 40 GiB; ~60 s to the first answer), the switches (STRATA_LOOKAHEAD=0, STRATA_KQ256=1, STRATA_PARTIAL_PIN=1, STRATA_FETCH_THREADS), scope and validation, credits (eddoursul #245, gopinath87607 #255, jagsan-cyber #247). Experimental, not in setup. - docs/DETAILS.md: the mapped low-RAM mode reads a native pack's experts from the GGUF when it has no experts.bin (checked, identical answers on the Coder); --resident-budget-gib and the routing-aware prefetch; where the tier counters are (--stats, the server log, /metrics' ram_blobs / file_blobs / file_mb); a pointer from "Which model?". Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
Niko1221
added a commit
that referenced
this pull request
Oct 1, 2026
…) - the four shards pinned at 38bb39e with their sizes and SHA-256, a RAM budget of the PC's RAM less 24 GB, packed with --compat-bf16, engine 0.1.32 or newer
- FAMILIES["unsloth"] (marked [experimental], not the default) and MODELS["UD-Q4_K_XL"] ("budget": True,
48 GB of RAM or more); HF_REVISIONS pins unsloth/Qwen3.8-Flash-Next-GGUF; families may have more than two
shards ("shards": 4).
- The download resumes as before; then verify_sha256 checks each shard's size and SHA-256 once (remembered in its
finish mark; a wrong file is deleted so the next run downloads it again).
- resident_budget_gib: round(RAM) - 24 GiB (40 on 64 GB, the measured setting), at most the 71 GiB of experts, at
least 8; a KV cache streamed to RAM (64K+) comes out of it. --resident-budget-gib replaces the low-RAM mode for
this model (no experts.bin; --low-ram is ignored with a warning); one GPU, no images, no speed projection; 8K
context recommended on GPUs under 14 GB; no --ple-gguf for a 4-shard file (the engine finds the PLE shard).
- The engine version is checked right after the engine step, before the 111 GB download.
- choices_from_config reads a size with a dash (strata-unsloth-ud-q4_k_xl.json).
- tools/test_setup_unsloth.py: pins, budget, SHA-256 handling, config round trip, and setup.main() end to end with
every outside effect mocked (URLs, pack flags, engine args, old engine, too little RAM, two GPUs + images).
- README and docs/UNSLOTH_Q4.md / DETAILS.md describe it.
The model support itself follows eddoursul/Strata (PR #245).
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
Niko1221
added a commit
that referenced
this pull request
Oct 6, 2026
… the mapped mode without experts.bin, --resident-budget-gib, the tier counters - docs/UNSLOTH_Q4.md: what UD-Q4_K_XL is (Hugging Face revision 38bb39e, the four shards with sizes and SHA-256), disk/RAM/GPU needs (111.3 GB download, 1.4 GB pack, no experts.bin), the iq_pack.py command and why --compat-bf16, the numerical policy (native experts/embedding/head/projections/PLE; the converted small tensors listed in conversions.json), the MTP draft layer, a server config (--native shard 1, --ple-gguf found by name, --resident-budget-gib N with N = RAM minus 20-24 GB), measured speed on an RTX 5070 / 64 GB PC (7-8.5 tok/s at 40 GiB; ~60 s to the first answer), the switches (STRATA_LOOKAHEAD=0, STRATA_KQ256=1, STRATA_PARTIAL_PIN=1, STRATA_FETCH_THREADS), scope and validation, credits (eddoursul #245, gopinath87607 #255, jagsan-cyber #247). Experimental, not in setup. - docs/DETAILS.md: the mapped low-RAM mode reads a native pack's experts from the GGUF when it has no experts.bin (checked, identical answers on the Coder); --resident-budget-gib and the routing-aware prefetch; where the tier counters are (--stats, the server log, /metrics' ram_blobs / file_blobs / file_mb); a pointer from "Which model?".
Niko1221
added a commit
that referenced
this pull request
Oct 6, 2026
…) - the four shards pinned at 38bb39e with their sizes and SHA-256, a RAM budget of the PC's RAM less 24 GB, packed with --compat-bf16, engine 0.1.32 or newer
- FAMILIES["unsloth"] (marked [experimental], not the default) and MODELS["UD-Q4_K_XL"] ("budget": True,
48 GB of RAM or more); HF_REVISIONS pins unsloth/Qwen3.8-Flash-Next-GGUF; families may have more than two
shards ("shards": 4).
- The download resumes as before; then verify_sha256 checks each shard's size and SHA-256 once (remembered in its
finish mark; a wrong file is deleted so the next run downloads it again).
- resident_budget_gib: round(RAM) - 24 GiB (40 on 64 GB, the measured setting), at most the 71 GiB of experts, at
least 8; a KV cache streamed to RAM (64K+) comes out of it. --resident-budget-gib replaces the low-RAM mode for
this model (no experts.bin; --low-ram is ignored with a warning); one GPU, no images, no speed projection; 8K
context recommended on GPUs under 14 GB; no --ple-gguf for a 4-shard file (the engine finds the PLE shard).
- The engine version is checked right after the engine step, before the 111 GB download.
- choices_from_config reads a size with a dash (strata-unsloth-ud-q4_k_xl.json).
- tools/test_setup_unsloth.py: pins, budget, SHA-256 handling, config round trip, and setup.main() end to end with
every outside effect mocked (URLs, pack flags, engine args, old engine, too little RAM, two GPUs + images).
- README and docs/UNSLOTH_Q4.md / DETAILS.md describe it.
The model support itself follows eddoursul/Strata (PR #245).
anon761
pushed a commit
to anon761/Strata-upstream
that referenced
this pull request
Oct 6, 2026
… the mapped mode without experts.bin, --resident-budget-gib, the tier counters - docs/UNSLOTH_Q4.md: what UD-Q4_K_XL is (Hugging Face revision 38bb39e, the four shards with sizes and SHA-256), disk/RAM/GPU needs (111.3 GB download, 1.4 GB pack, no experts.bin), the iq_pack.py command and why --compat-bf16, the numerical policy (native experts/embedding/head/projections/PLE; the converted small tensors listed in conversions.json), the MTP draft layer, a server config (--native shard 1, --ple-gguf found by name, --resident-budget-gib N with N = RAM minus 20-24 GB), measured speed on an RTX 5070 / 64 GB PC (7-8.5 tok/s at 40 GiB; ~60 s to the first answer), the switches (STRATA_LOOKAHEAD=0, STRATA_KQ256=1, STRATA_PARTIAL_PIN=1, STRATA_FETCH_THREADS), scope and validation, credits (eddoursul Niko1221#245, gopinath87607 Niko1221#255, jagsan-cyber Niko1221#247). Experimental, not in setup. - docs/DETAILS.md: the mapped low-RAM mode reads a native pack's experts from the GGUF when it has no experts.bin (checked, identical answers on the Coder); --resident-budget-gib and the routing-aware prefetch; where the tier counters are (--stats, the server log, /metrics' ram_blobs / file_blobs / file_mb); a pointer from "Which model?".
anon761
pushed a commit
to anon761/Strata-upstream
that referenced
this pull request
Oct 6, 2026
…) - the four shards pinned at 38bb39e with their sizes and SHA-256, a RAM budget of the PC's RAM less 24 GB, packed with --compat-bf16, engine 0.1.32 or newer
- FAMILIES["unsloth"] (marked [experimental], not the default) and MODELS["UD-Q4_K_XL"] ("budget": True,
48 GB of RAM or more); HF_REVISIONS pins unsloth/Qwen3.8-Flash-Next-GGUF; families may have more than two
shards ("shards": 4).
- The download resumes as before; then verify_sha256 checks each shard's size and SHA-256 once (remembered in its
finish mark; a wrong file is deleted so the next run downloads it again).
- resident_budget_gib: round(RAM) - 24 GiB (40 on 64 GB, the measured setting), at most the 71 GiB of experts, at
least 8; a KV cache streamed to RAM (64K+) comes out of it. --resident-budget-gib replaces the low-RAM mode for
this model (no experts.bin; --low-ram is ignored with a warning); one GPU, no images, no speed projection; 8K
context recommended on GPUs under 14 GB; no --ple-gguf for a 4-shard file (the engine finds the PLE shard).
- The engine version is checked right after the engine step, before the 111 GB download.
- choices_from_config reads a size with a dash (strata-unsloth-ud-q4_k_xl.json).
- tools/test_setup_unsloth.py: pins, budget, SHA-256 handling, config round trip, and setup.main() end to end with
every outside effect mocked (URLs, pack flags, engine args, old engine, too little RAM, two GPUs + images).
- README and docs/UNSLOTH_Q4.md / DETAILS.md describe it.
The model support itself follows eddoursul/Strata (PR Niko1221#245).
anon761
pushed a commit
to anon761/Strata-upstream
that referenced
this pull request
Oct 6, 2026
… the mapped mode without experts.bin, --resident-budget-gib, the tier counters - docs/UNSLOTH_Q4.md: what UD-Q4_K_XL is (Hugging Face revision 38bb39e, the four shards with sizes and SHA-256), disk/RAM/GPU needs (111.3 GB download, 1.4 GB pack, no experts.bin), the iq_pack.py command and why --compat-bf16, the numerical policy (native experts/embedding/head/projections/PLE; the converted small tensors listed in conversions.json), the MTP draft layer, a server config (--native shard 1, --ple-gguf found by name, --resident-budget-gib N with N = RAM minus 20-24 GB), measured speed on an RTX 5070 / 64 GB PC (7-8.5 tok/s at 40 GiB; ~60 s to the first answer), the switches (STRATA_LOOKAHEAD=0, STRATA_KQ256=1, STRATA_PARTIAL_PIN=1, STRATA_FETCH_THREADS), scope and validation, credits (eddoursul Niko1221#245, gopinath87607 Niko1221#255, jagsan-cyber Niko1221#247). Experimental, not in setup. - docs/DETAILS.md: the mapped low-RAM mode reads a native pack's experts from the GGUF when it has no experts.bin (checked, identical answers on the Coder); --resident-budget-gib and the routing-aware prefetch; where the tier counters are (--stats, the server log, /metrics' ram_blobs / file_blobs / file_mb); a pointer from "Which model?". Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
anon761
pushed a commit
to anon761/Strata-upstream
that referenced
this pull request
Oct 6, 2026
…) - the four shards pinned at 38bb39e with their sizes and SHA-256, a RAM budget of the PC's RAM less 24 GB, packed with --compat-bf16, engine 0.1.32 or newer
- FAMILIES["unsloth"] (marked [experimental], not the default) and MODELS["UD-Q4_K_XL"] ("budget": True,
48 GB of RAM or more); HF_REVISIONS pins unsloth/Qwen3.8-Flash-Next-GGUF; families may have more than two
shards ("shards": 4).
- The download resumes as before; then verify_sha256 checks each shard's size and SHA-256 once (remembered in its
finish mark; a wrong file is deleted so the next run downloads it again).
- resident_budget_gib: round(RAM) - 24 GiB (40 on 64 GB, the measured setting), at most the 71 GiB of experts, at
least 8; a KV cache streamed to RAM (64K+) comes out of it. --resident-budget-gib replaces the low-RAM mode for
this model (no experts.bin; --low-ram is ignored with a warning); one GPU, no images, no speed projection; 8K
context recommended on GPUs under 14 GB; no --ple-gguf for a 4-shard file (the engine finds the PLE shard).
- The engine version is checked right after the engine step, before the 111 GB download.
- choices_from_config reads a size with a dash (strata-unsloth-ud-q4_k_xl.json).
- tools/test_setup_unsloth.py: pins, budget, SHA-256 handling, config round trip, and setup.main() end to end with
every outside effect mocked (URLs, pack flags, engine args, old engine, too little RAM, two GPUs + images).
- README and docs/UNSLOTH_Q4.md / DETAILS.md describe it.
The model support itself follows eddoursul/Strata (PR Niko1221#245).
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
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
2 participants
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.
After a week of profiling and benchmarking I doubled throughput on 3090 + 5070 Ti.
See the comparison tables.
Use JSON configurations to run https://github.com/eddoursul/Strata/blob/custom/examples/README.md