Repository navigation
perf: Q8 resident adaptation (+24.5–28.2% rotation gain on RTX PRO 6000) - #898
Closed
CC-David-CC wants to merge 859 commits into
Closed
CC-David-CC wants to merge 859 commits into
CC-David-CC wants to merge 859 commits into
Conversation
…o1221#475) In a git clone they run `git pull --ff-only` (a zip copy is told to download the new zip), then `setup.py --update`: what a plain start does to an install before it starts the model - the Python packages, the ready-made engine when MIN_ENGINE is newer (a compiled one when its source changed), each model's config upgrades and its draft subset, the Windows AMD runtime - and nothing else: no question, no model files, no start. Both scripts run from one parsed block, since the pull can rewrite the running script. Documented in docs/INSTALL.md (Updating) and the README. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…profile-save, Niko1221#477) `--expert-profile-save P` (--serve) writes the adaptive tier's learning as a profile in tools/make_profile.py's format, byte for byte (read_expert_profile reads it back): every pair ranked - the experts resident in VRAM now first, then the routing counted since the start (the decayed usage summed at each adapt step, before its decay), then the profile the engine started from. On QUIT / stdin EOF and every --expert-profile-save-every minutes (default 10, 0 = exit only) between requests, through <P>.tmp and a rename. Without the flag nothing is counted or written: `heat` stays empty and every added loop is over it. Server: "expert_profile_save" (and "expert_profile_save_every") in the config pass the flags, and the next start begins from the saved file instead of the config's --expert-profile when its header matches that profile's layers x experts. tests/core/expert_profile_save_test.cpp (CPU only) checks the order and the bytes. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…, the draft-head hint (Niko1221#474), UPDATE.bat (Niko1221#475), --expert-profile-save (Niko1221#477); setup requires it Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…pt path, grouped on the GPU (opt-in STRATA_PF_FUSED=1) Strata's own kernels (src/prefill/moe_fused.cu): the layer's input rounded to int8 once per token (a scale per 32 values), the (token, k) pairs grouped by expert on the GPU (counts, offsets, a 64-row tile list - no host sync), then per batch of up to 128 experts one gate/up launch (mma.sync m16n8k32 s8, the 2-bit codes unpacked in registers, Q2_0's offset folded into the activation blocks' sums, SwiGLU and H's int8 rounding in the epilogue) and one down launch into the per-slot rows the existing combine reads. Experts are read where they are (cache slot or streamed ring slot, a pointer per expert): no gather. Used only on the streamed walk (chunks from stream_all_min()) of the canonical Q2_0 layout on sm_80+ CUDA, with STRATA_PF_FUSED=1; it reuses the MMQ buffers (Xq, H, GU). Everything else - the default, native/IQ packs, small chunks, sm_75, HIP - runs the MMQ path unchanged (the old block is only wrapped in the else branch). tests/cuda/prefill_fused_moe_test (exit 77 below sm_80), RTX 5070: vs a double-precision reference, MMQ rel RMS 1.144e-2 (worst row 4.15e-2), fused 1.144e-2 (4.15e-2); fused vs MMQ 9.0e-5. One layer at 2048 tokens x top 10 over 512 experts: MMQ path 11.12 ms, fused 3.13 ms. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…e layer's experts) by default The fused path launches on batches of a layer's streamed experts, so with the 384-slot ring a layer's last batch waited for slots its own first batch frees. With STRATA_PF_FUSED=1 (and the canonical Q2_0 layout) the pinned-RAM ring is 512 slots; without it nothing changes (384 / 96 as before; STRATA_PREFILL_RING still overrides). RTX 5070, Q2_0 (cfg-q2rel), greedy, 3 alternating pairs, prompt tok/s: 4K 1,275 -> 1,516 (+18.9%), 32K 2,176 -> 2,422 (+11.3%); at 384 slots the fused path gave +4-7% / +3%. 1024 slots: 4K +24-27%, but the 32K chunk falls from 8192 to 6144 (+6%). The default path's output is identical to 0.1.35's on both prompts. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
… hold only what the fused kernels need
In a layout whose chunks run the fused experts (STRATA_PF_FUSED=1, the Q2_0 pack, a streamed chunk) GU holds the
grouping tables, H the int8 H rows and Xq the per-token int8 activations, and Hq nothing - instead of MMQ's FP32 GU / H
and per-slot q8_1 rows: ~100 KB a token less in the MoE set, which then no longer sets the shared region's size (the
GDN set does). A last chunk below stream_all_min() still runs MMQ in the same buffers, so each keeps MMQ's size for
stream_all_min() - 1 tokens. The ring may grow to 1024 slots with the fused experts (STRATA_PREFILL_RING; the default
stays 512). Without STRATA_PF_FUSED nothing changes: the same buffers, ring cap (512) and events.
The fused walk now counts the experts it reads from the cache in experts_resident (it stayed 0): every resident expert
of the layer, routed or not - at a streamed chunk's size nearly all are.
RTX 5070, Q2_0 (cfg-q2rel), --prefill auto, 2 alternating pairs, prompt tok/s (borrowed slots):
4K: MMQ 1,305 (1621), fused 1,517 (1749), fused + P3 1,512 (1534), + 1024-slot ring 1,605 (2046)
32K: MMQ 2,177 (3071), fused 2,453 (3199), fused + P3 2,485 (2641), + 1024-slot ring 2,660 (3153)
128K (--max-context 131072 --kv-resident 32768): MMQ 2,108 (chunk 8192), fused 2,140 (6144), fused + P3 2,346
(8192), + 1024-slot ring 2,392 (6144)
A 16K chunk still does not fit the 12 GB card: at 8192 tokens the loan is 2,641 of 3,690 slots, ~2.1K of them per
8K tokens.
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…ault (two layers' experts) With P3's smaller buffers the 1024-slot ring fits beside an 8192-token chunk at 32K (borrowed 3,153 of 3,690 slots, under the 90% cap). RTX 5070, Q2_0 (cfg-q2rel), --prefill auto, prompt tok/s at 512 / 1024 slots (2 alternating pairs): 4K 1,512 / 1,605, 32K 2,485 / 2,660, 128K (KV streaming) 2,346 / 2,392 - there the chunk falls from 8192 to 6144 and is still faster. Against MMQ (one more pair on this build): 4K 1,299 -> 1,609 (+23.9%), 32K 2,187 -> 2,656 (+21.5%), 128K 2,128 -> 2,454 (+15.3%). The ring does not change the numbers (the same output as at 512 slots). Without STRATA_PF_FUSED nothing changes (384 / 96 slots, cap 512). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…y argmax, the same ids and tokens - qsa_block_topk: calls of up to 16 queries (decode windows) run on a cluster of 8 CTAs per query (block_topk_cluster_kernel): keys in shared memory, digit histograms pushed through distributed shared memory, one cluster barrier per radix pass, a parallel digit search, cluster-wide offsets for the ordered emit. The threshold and the cut at it are the one-CTA kernels' (same tie rule, NaN keys 0). STRATA_QSA_CLUSTER=0: off. RTX 5070 per call: 21.9 -> 15.6 us at 32K, 58 -> 18 at 128K, 200 -> 22 at 262K. - sample_tokens greedy without penalties: a cluster of 8 CTAs per row (sampler_greedy_cluster_kernel), four loads in flight per thread, results merged in CTA 0 (lowest index on ties, NaN never taken, 0 when none). STRATA_ARGMAX_MULTI=0: off. 39.6 -> 5.9 us per call at 248,320 logits. - Both launch through cudaLaunchKernelEx with the cluster attribute (captured in the decode graphs) only where they run: CUDA, compute capability 9+ (STRATA_EMULATE_CC respected), the kernel's own PTX/SASS at sm_90+ (an older build's JIT copy is a trap), room for one cluster - checked once per device. HIP and sm_75/86/89 keep the one-CTA kernels. - decode_cluster_parity (ctest; skips below sm_90): bitwise vs the one-CTA dispatch and the reference kernel over 434 top-k cases (contexts 7..262,144, every tail weight, 1-16 queries, ties / NaN / +-inf / -0 / denormal / all NaN) and 216 argmax cases, plus a CUDA graph capturing both (first calls inside the capture); --bench timings. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…gate/up IQ2_XXS IQ2_XS IQ2_S IQ3_XXS IQ3_S IQ4_XS, down Q2_0 IQ4_NL; opt-in STRATA_PF_FUSED=1) src/prefill/moe_fused_iq.cu: the Q2_0 path's design (moe_fused.cu: the grouping on the GPU, a pointer per expert - cache slot or ring slot, no gather -, SwiGLU and H's int8 rounding in the gate/up epilogue) for the raw GGUF blobs of the native packs. The load stage decodes the i-quant blocks to int8 in shared memory (llama.cpp's codebooks and sign tables, held in shared memory; a 32-value sub-block per thread per 64 values of K, read one stage ahead with 16-bit loads and two 256-value blocks ahead into L2), with each sub-block's scale (per 16 values for IQ2_XS / IQ2_S, which multiply at m16n8k16). The activations are int8 per 32 values in natural order (quantize_act_native); fragments by ldmatrix; the magic bias 1.5 * 2^23 enters as the mma's C, so a dot becomes a float with one add. A work item is 256 weight rows x 64 routed rows, or 128 x 128 when the experts have 56-112 rows each (STRATA_PF_FUSED_TILE=64|128). Wiring (prefill.cpp): a native layer takes these kernels when both its formats are covered, else MMQ (IQ1_M keeps the FP16 path); the 512-slot ring of the fused path now also for a native pack with such layers. sm_75, HIP and the default (no variable) are unchanged: on IQ2_XS, IQ3_XXS and IQ3_S the default output is 0.1.35's (256/256 greedy tokens, 4K and 32K). tests/cuda/prefill_fused_iq_test (exit 77 below sm_80), RTX 5070, 256 tokens x 64 experts against a double-precision reference from ggml's to_float, rel RMS / worst row - MMQ and fused equal to four digits: IQ2_S/Q2_0 1.489e-2 / 2.93e-2 (fused vs MMQ 2.6e-4), IQ2_XXS/Q2_0 1.464e-2 / 2.96e-2 (1.9e-4), IQ2_XS/IQ4_NL 1.465e-2 / 2.41e-2 (3.5e-5), IQ3_XXS/IQ4_NL 1.458e-2 / 2.73e-2 (1.6e-4), IQ3_S/IQ4_NL 1.470e-2 / 2.65e-2 (6.5e-5), IQ4_XS/Q2_0 1.446e-2 / 2.72e-2 (4.4e-5); both tile shapes. One layer (512 experts, top 10), ms, MMQ path -> fused at 2048 / 3584 / 8192 tokens: IQ2_S 10.5 -> 5.2, 12.5 -> 9.3, 21.4 -> 17.5; IQ2_XXS 8.8 -> 4.2, 10.7 -> 7.6, 17.9 -> 13.6; IQ3_S 10.6 -> 5.6, 12.3 -> 8.8, 19.8 -> 18.4; IQ3_XXS 10.2 -> 5.1, 11.9 -> 8.3, 19.1 -> 16.5. MMQ's products alone (no gathers) at 3584 / 8192: IQ2_S 8.9 / 17.9, IQ3_S 8.1 / 15.6: the kernels are about even with MMQ's, the gain is the path around them. (The test syncs the device after its uploads: a cudaMemcpy from pageable memory on the legacy stream can still be queued on Windows while a non-blocking stream runs - the grouping then read zeroed ids, every pair on expert 0. prefill_fused_moe_test's timing part has the same pattern.) End to end, RTX 5070, speed_135.py's settings, greedy, A (MMQ) / B (STRATA_PF_FUSED=1) alternating, 2 pairs, prompt tok/s: IQ2_XS 4K 1,228 -> 1,378 (+12%; a cold first A run of 590 left out), 32K 2,107 -> 2,170 (+3.0%); IQ3_XXS 4K 1,001 -> 1,020 (+1.8%), 32K 1,745 -> 1,739 (-0.4%); IQ3_S 4K 927 -> 925 (-0.2%), 32K 1,382 -> 1,362 (-1.4%). The IQ3 packs' prompts here wait on the host copies of their unpinned experts (IQ3_XXS 4K: 5,062 of 22,573 streamed, ~3.0 s of 3.4 s), so faster kernels change little. Each arm repeats itself (256/256; IQ3_XXS 32K B and IQ3_S 32K A split at 162 / 223, the decode's timing); A vs B greedy agreement, identical prefix of 256: IQ2_XS 204 / 88, IQ3_XXS 182 / 89, IQ3_S 31 / 256. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…(STRATA_PF_FUSED=0 off), the native IQ packs' fused kernels opt-in (=1), cluster decode kernels; the fused test waits for its uploads The ring rule: 1024 slots for the Q2_0 pack's fused experts (P3), 512 for native packs with fused layers. Quality (teacher-forced vs the FP16 path, Q2_0): fused as close as MMQ at 8K, closer at 32K (phaseA-tests s18-tf). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
device, expert_source, native_head, prefill gemm, prefill.cpp, generate.cpp merged from dpct of aeb35be and d9ab843 with sycl/tools/merge_upstream.py (the 3-way merge restricted to files upstream changed, dpct kernel-name hashes canonicalized to the port's). Version 0.1.35. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…ls) and their switches Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
The PCIe probe sets the share of the missed experts the GPU reads over the link (--pcie-frac): the native default 0.55 from 20 GB/s up, less on a slower link (Niko1221#44). On Windows HIP it timed its copies with events that do not bracket them there: an RX 6800 read between 3,348 and 25,999 GB/s in 24 engine starts, so every link kept the x16 share. A standalone copy of the probe on that card: its events measure 0.13 ms for the 1 GiB, the host clock 43 ms (25.1 GB/s). On Windows HIP the probe now times the same copies on the host, between two synchronizes (1 GiB, tens of milliseconds): the same RX 6800 on PCIe 4.0 x16 reads 27.2 GB/s in every run, which keeps 0.55 there. On a PCIe 4.0 x8 or 3.0 x16 link the share now shrinks as Niko1221#44 intended. CUDA and Linux keep the events. Claude-Session: https://claude.ai/code/session_015Ld5YDzuKW1hgMVwwgRZDP Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…DGET)
On Windows, hipMemGetInfo (ROCclr's PAL backend) reports the card's size
minus this process's own allocations; what the desktop and other
programs hold on the card is not subtracted. On a card that also drives
the desktop the engine counts the ~930 MiB those hold as free, and
--expert-cache auto fills the card past what Windows lets the process
keep in VRAM: WDDM moves part of it out to system RAM. On an RX 6800
(HIP SDK 7.2) with setup's configuration: a 9.83 GiB cache, the card at
16132 of its 16368 MiB, 929 MiB more of the engine in system RAM than
with this change, 30.1-31.0 tok/s decode. On a 0.1.30 build,
--vram-reserve-mib 1500 left the card as full (16037 MiB), since the
reserve is taken from the same figure.
Windows gives every process a video memory budget that does account
for the others (DXGI QueryVideoMemoryInfo). On that card it is 799 MiB
below the card's size while the process is small and about 1.8 GiB
below once it holds 11 GiB; allocations past it still succeed.
cudaMemGetInfo, which the HIP shim maps to hipMemGetInfo, now lowers
hipMemGetInfo's free figure by what the budget withholds (the card's
size minus the budget), on Windows only. What is sized from it then
stays within the budget: --expert-cache auto sizes the cache from it,
and the existing check after the slots are written shrinks it ("only
8 MiB free once the slots are written ... shrinking"): an 8.05 GiB
cache, the card at 15230 MiB, 40.6-41.5 tok/s.
Lowering HIP's own figure, instead of taking the budget minus DXGI's
usage, keeps memory the HIP runtime has freed but still caches counted
as free, as HIP counts it.
dxgi.dll is loaded at run time, so nothing new is linked; the adapter
is found by the device's LUID. Without dxgi.dll, the adapter or a
budget below the card's size, hipMemGetInfo's figure stands. The first
query logs the budget once. STRATA_WDDM_BUDGET=0 turns it off. CUDA
and Linux builds are unchanged.
Claude-Session: https://claude.ai/code/session_015Ld5YDzuKW1hgMVwwgRZDP
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…g forever (Niko1221#481) A request whose engine prints nothing for engine_silence_s (default 300 s; while a prompt is read, a chunk may take 3x the previous one's time, the first up to its tokens at 50 tok/s more) ends with EngineSilent (an EngineDied): the engine is killed and the next request starts it again. The STOP drain, under the request FIFO, is no longer an untimed wait. "engine_silence_s": 0 keeps the old behaviour. Tests: unit (silence, PP chunks, first chunk, STOP never acknowledged, 0, config) and over HTTP with a fake engine process (non-stream, stream, stop token). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…0 log lines (Niko1221#496) Every start failure, not only the draft-head case (Niko1221#474), now carries the engine's own reason: the last 20 lines this start wrote to the engine log (never an earlier start's). The failed start also closes its pipes and its handle on the log. Tests: the tail itself, and a fake engine that exits with a reason in its log. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…p-experts (Niko1221#505) The server window said "loading the experts into RAM (about N GB)" whatever the flags. Now: --mmap-experts maps them from the model files (the OS file cache reads them, nothing is loaded up front; its "experts via mmap" line triggers the message when the model has no PLE line), --resident-experts copies the ones the GPU does not hold, and --resident-budget-gib N copies up to N GiB. setup's start banner says "maps" for the plain mapped mode too. Tests: the words per flag set, and the narrator over a mapped and an arena log. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…#495) hf() takes the host from HF_ENDPOINT (default huggingface.co, a trailing slash dropped); the pinned revisions, the fallback to main and the SHA-256 checks are unchanged. tools/mtp_fetch.py reads the MTP tensors from the same host. Step 5 prints the folder setup expects the model files in and how to use files downloaded by hand (that folder with their original names, or --gguf-dir), and the mirror when one is set. INSTALL.md says the same. Tests: the endpoint in hf()/hf_unpinned()/mtp_fetch, and step 5's words over a mocked install. The golden configs are unchanged. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…ds it; no budget on a --gpus start (Niko1221#498) The engine refuses --resident-budget-gib with a layer split, so setup always dropped UD-Q4_K_XL to one GPU. With RAM >= the GGUF size + UNSLOTH_RAM_LEFT_GB (~135 GB) and 2+ usable cards, setup now asks (one GPU stays the default and --yes keeps it; an explicit --gpus is honoured; an explicit --resident-budget-gib keeps one GPU) and the split writes "gpu": [..], "layer_split": "auto" without the budget: all the experts loaded into RAM at start, the files through the OS file cache (the config Niko1221#498 measured: 2x RTX 3090, 31 -> 64-78 tok/s). Latent bug: START-HERE --gpus 0,1 (or the start-time offer, or a hand-edited config) on a UD-Q4 config kept the budget and the engine exited with code 2; the split now drops it, or setup stops before saving when the RAM is too small. docs/UNSLOTH_Q4.md says so. Tests in tools/test_setup_unsloth.py; the golden configs are unchanged. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
The engine already leaves --vram-reserve-mib MiB of VRAM free (default 700). INSTALL.md now says what it does, its default, and how to set it (the flag, or "--vram-reserve-mib", "2048" in the config's args). Setup's new flag writes it into the config only when given (one value: it replaces the images' own 700, with a note when it is less than the GPU image encoder needs); on a start of an installed model it is saved into that model's args, and a new copy set up from an earlier install keeps it. The default config is unchanged (golden test). Tests in test_setup_choices.py. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…are, the samples logged The probe timed one 1 GiB burst; an x16 PCIe 4 laptop (RTX A3000) read 18.5, 6.9 and 5.8 GB/s in three starts, and the low reading set pcie_frac (0.12-0.15 where calibrate.py found 0.35). The same 1 GiB is now four bursts of 256 MiB, each timed on its own (events between them on the stream; on Windows HIP the host clock, Niko1221#377), and the fastest is the link's figure: a disturbance only ever slows a copy, while the median would still follow one lasting through half the bursts. The probe copies no more than before. The share was base above 20 GB/s and base * bw / 26 below (19.9 GB/s -> 0.42, 20.0 -> 0.55; 4.0 -> 0.08, 3.9 -> 0). It is now base * min(1, bw / 20): continuous, unchanged from 20 GB/s up (an x16 PCIe 4/5 link reads ~26-28 GB/s, so the default stays 0.55 there), in proportion to the link below, so the link's time on its share stays about what the x16 share costs. The rule is one function for CUDA0 and the layer split's stages. The log names every burst: "PCIe probe: 28.9 GB/s host->device (best of 28.9 28.9 28.9 28.9)" (the RTX 5070 here; pcie_frac 0.55 as before). Noted, not changed: a given --pcie-frac is every stage's share and skips the stages' probes. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…#380) and the host-timed PCIe probe (Niko1221#377) Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…MD counts the desktop's VRAM (Niko1221#380 Niko1221#377 Niko1221#497), a steadier PCIe probe (Niko1221#485), fixes Niko1221#496 Niko1221#495 Niko1221#498 Niko1221#505 Niko1221#493; setup requires it Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…iko1221#374 Niko1221#363 Niko1221#407 Niko1221#413 onto 0.1.35 One dpct run of upstream main + the seven PRs, 3-way merged into the files they touch (sycl/tools/merge_upstream.py); their header changes are forked into sycl/include until upstream merges them. - Niko1221#385 stager waits for a buffer's previous DMA on every job (port: polled seq) - Niko1221#463 decode waits for an adaptive swap before reading residency - Niko1221#453 drafter's batched K/V on the KV-streaming ring - Niko1221#374 first chunk's PLE rows read beside layer 0; port fixes: the next read starts at c0 + T (the port's first chunk is shorter), and the upload is marked by a polled sequence number instead of a host event wait, which deadlocked prompts past ~4K tokens under the Level Zero v2 adapter - Niko1221#363 group stride + fused SwiGLU/q8_1 for the PCIe expert call, re-applied to the port's lane kernels (fp contract off, as upstream's HIP form) - Niko1221#407 --adapt-tuned (opt-in), Niko1221#413 key-head DeltaNet recurrence (opt-in via STRATA_GDN_KEYHEAD=1: 8% slower on the B70; its NVIDIA gate and SM-holding parity bench not ported) Tested on the B70: native_grouped_parity, iq_multi, gr, qsa, kv_hybrid, quantize_act, qsa_prompt_attn parity pass; Coder 19/2,184-token, IQ2_XS and 40K outputs identical to 0.1.33; Coder 2,184-token prompt 784 -> 825 tok/s; 128K int8 KV streaming 895 / 67.0 tok/s; IQ2_XS decode +1.7% (Niko1221#363). Co-authored-by: sergqwer, constantindjonkam, architectds, BlueKingMuch (upstream PR authors) Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…er_intel) - engine (serve mode): GEN key logprobs=K (0..20) writes "LP logprob id:lp ..." after each "T id", from the verify window's head logits (Verifier::copy_logits; before temperature, penalties and sampling); absent key, no LP lines - server_intel.py: pairs LP lines with their tokens (StrataEngine._pump), passes logprobs/top_logprobs to the engine, and fills OpenAI's choices[].logprobs (content / reasoning_content) in streamed chunks and the collected reply; the stop token is not listed. serve/ itself is unchanged. Tested: a fake engine on hardac (streamed, not streamed, without logprobs, stop token excluded); on the B70 with the IQ2_XS: same answer with and without, "Paris" -0.0019 with top-1 equal to the chosen token, no measurable time cost. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
… wait, released by one event In the streamed walk every expert put a cross-stream wait (its copy) and an event record (its ring slot back) on the compute stream around a ~3 us gather. Under WDDM each pair left ~10 us of GPU idle even when the copy had long landed: an nsys trace of a 32K prompt showed 226 ms of idle between back-to-back gathers while the host was 3-11 ms ahead of the GPU. Now a group of 16 is gathered by one launch (copy16_group_kernel) after one wait on its last streamed copy - the copy stream is in order - and one event releases all of its ring slots: the issuer waits on used[used_of[slot]]. A ring entry the routing skipped inside an open group first gathers what the group holds so far, so no more than a group's entries are ever held back from the issuer. The same bytes in the same group slots. STRATA_PREFILL_GROUP_GATHER=0: one at a time, as before. IQ2_XS, RTX 5090, 262K context: 32K prompt -8.6%, 8K -10%, 2K -12.6%; first-token logits bit-identical at 2K, 8K, 16K and 32K and with a 96-slot ring. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01VZy1yKaDDiA8a7svdwaHio
…S.md with the measured quality Changes default output on long prompts of native packs against 0.1.39 (expected; gate notes in the S40 report). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…8, Linux 525); the wheel pins match the 12.9.1 toolkit Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
… keep 0.1.39's ring - a prompt that fits 0.1.39's auto chunk keeps 0.1.39's ring (Prefill::set_ring_budget's small_max): 4K prompts were 1-17% slower with the smaller ring (Coder -15%, IQ3_S -5%, IQ3_XXS -2%; 5 interleaved pairs each); - ring slots are given up for a bigger chunk only where 0.1.39's rule held the chunk under 6144 (ring_cap_for); from 6144 on the ring keeps its 0.1.39 size (the Coder 6144/384 -> 7936/199 was -12% at 32K); - the unpinned arm's cap stays 96 (the byte budget made it ~49 on IQ3_S); - the scan sets its own ring budget instead of set_ring_override, so a layer split's 96-slot override (Niko1221#340) is no longer clobbered by the scan. RTX 5070 vs STRATA_RING_BYTES=0: 4K -0.7..+1.1% (bit-identical); 32K IQ3_XXS +18.5% (cache 1500) / +9.4%, Coder +3.1%, IQ3_S +0.2%. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01QJJwUv9g3cD2ztX6KWS8rk
# Conflicts: # setup.py
# Conflicts: # src/core/expert_source.cpp # src/kernels/cuda/iq_kernels.cu
# Conflicts: # src/core/verify.cpp # src/program/generate.cpp
…l); UD-Q4_K_XL stays experimental - "experimental" is per model now: MODELS["UD-Q4_K_XL"] keeps it, the unsloth family and UD-IQ4_XS drop it. The first menu lists the family without [experimental], with a one-line about (UD-IQ4_XS: a 94 GB download; under ~80 GB of RAM part of its experts are read from the SSD). Its size menu lists UD-IQ4_XS first (the default) and UD-Q4_K_XL second; an experimental size always sorts last. --model names are unchanged. - UD-IQ4_XS: no EXPERIMENTAL warning, --check verdict without it. Images are asked (off by default), per model: the engine's image path has no restriction for this pack, which uses the original model's image encoder; not yet run with images. UD-Q4_K_XL keeps images off. - MCP: the fallback tables mirror setup's (per-model experimental/vision), sizes in setup's order, per-size "experimental" and "images", the vision check per model. - Docs: README and its six translations, MODELS, UNSLOTH_Q4, AI_SETUP, MCP_SERVER. - Tests: the menus, the default, --model by name, images, --check, the pack/tokenizer paths; the golden configs are unchanged. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01QJJwUv9g3cD2ztX6KWS8rk
…e-mib 1000 if a request stalls With the image encoder on the GPU, the 700 MiB reserve can leave ~200 MiB free on a 12 GB card (UD-IQ4_XS end-to-end check on an RTX 5070). Setup prints a one-line recommendation; the written config is unchanged. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01QJJwUv9g3cD2ztX6KWS8rk
In the resident RAM mode (--resident-experts) an adaptive round runs on a
thread the decode loop joins after the draft: choose, copy the evicted slots
back with a stream sync, copy the new experts in from the RAM copy (through
the driver's staging when it is pageable). With ~90 swaps that was ~24 ms
every 4th window (measured on our fork). --adapt-async 1 splits a round into
four steps that advance between windows on a helper thread, so no window
waits for one:
1. choose (adapt()'s rule, over snapshots of the counts and the residency)
and copy the evicted slots back (D2H), each from the cache of the card
that owns its layer
2. stage the exchanges, mark the evicted experts non-resident (the CPU reads
them from the exchange buffers), copy the new experts in: straight from
a page-locked RAM copy, else through a pinned bounce buffer
3. mark the new experts resident, move the evicted blobs into the RAM places
the new ones left
4. flip the copy's offsets
A slot is overwritten only after its old expert stopped being planned for the
GPU, and a RAM place is reused only once its expert is resident. The residency
tables on the devices are uploaded at steps 2 and 3, as apply_pending does
(the verify graph's doorbell reads them). A request starts with the round
drained, before its prompt lends any slot. --expert-profile-save counts the
routing at the start of each round, as the blocking tier does. Swaps that
need no exchange (the evicted expert is in the RAM copy already, or the new
one is not: a RAM budget) take the same steps without a copy back; a new
expert outside the RAM copy is read from the files on the helper thread.
Default off: which window first computes a swapped-in expert on the GPU
depends on when its copy lands, so the tier is not bit-exact from run to run.
Without the flag nothing changes. With it, the blocking tier still runs
outside --serve, without the resident RAM copy, and with --batch slots,
--peer-device or --remote-expert-opt (said once at startup). The exchange
buffers are twice as many with it (the bounce half): 96 more blob-sized
pinned buffers, ~130 MiB with 1.38 MB blobs.
Layer split (with the resident mode on a split, Niko1221#848): each swap's copies run
on the owning card's refill stream under that card's device; the homes are
looked up the way Niko1221#848's resident_stage_swaps does.
Measured: Swift 1.5 IQ3_XXS at 160K (q4_0 KV), i9-14900KF, 32 GB of RAM,
Windows 11, greedy, decode tok/s, blocking -> --adapt-async 1:
- RTX 4060 Ti + RTX 5080 layer split with the resident RAM mode (Niko1221#848),
--trim-stage-weights, the stock draft layer, five prompts x two rounds,
two interleaved pairs: 79.7 / 78.8 -> 87.6 / 86.2 (+10%).
- the same split with a fine-tuned draft layer and --spec-min-p 0.65, four
prompts x three rounds:
81.0 / 81.8 -> 108.1 / 107.0 (+32%).
- RTX 5080 alone with --resident-experts: 26.7 / 26.6 -> 29.5 / 28.5 (noisy).
Without the flag the greedy output is the same as main's, byte for byte (4/4
prompts, 500 tokens, STRATA_IQ_MT_MIN=1 --pcie-frac 0 --adapt-every 0). With
it, 22 varied requests at temperature 0.6 on the split (43,825 experts
swapped in over 622 rounds) gave no degenerate answer.
(cherry picked from commit 18d3da4)
(cherry picked from commit 513395d)
A decode window is a long chain of small kernels, and on a fast card many of them are over before the next one has fetched its weights. With programmatic dependent launch the window's quantizations and dense projections (multi-column MMVQ, multi-row BF16 GEMV) start while the kernel before them is still running: they load their first weights, then wait (griddepcontrol.wait) for its output. Nothing is written before the wait, so every value is the one an ordinary launch computes. - pdl.hpp: pdl_wait / pdl_trigger and launch_pdl. The attribute is used only inside the verify window's capture (pdl_scope), with a CUDA 12.3+ runtime, on sm_90+ (STRATA_EMULATE_CC counts), for a kernel whose loaded code is sm_90+ (ptxVersion and binaryVersion >= 90: older PTX, JIT-compiled, has no wait and would race), and only when every node the capturing stream's next node depends on is a kernel. So there is never a programmatic edge after a memcpy, memset, host or event node. After a join of branches each joined kernel gets a programmatic edge (the wait covers them all); STRATA_DF_PDL=2 allows a single predecessor only. - On CUDA the PDL kernels' activations and outputs are no longer __restrict__ (the compiler could otherwise load them before the wait); the weights keep it. For every architecture: the host pass must see the device pass's signature, or MSVC rejects the template stubs. - HIP: ordinary launches, no-op device functions, the kernels' plain loops. - copy_rows_strided replaces the q/gate split's cudaMemcpy2DAsync on the batched QSA path: a memcpy node would end the PDL chain there. A copy: the same bytes, PDL on or off. - The profiler's stamps pass the early launch on inside the window. - STRATA_VERIFY_NODES also counts the programmatic edges and those that are not kernel to kernel (0 expected). - STRATA_DF_PDL=0: ordinary launches. Byte-identical: PDL changes when a kernel starts, not what it computes. Off on pre-sm_90 cards (the 4060 Ti) and on HIP. Checked: pdl_parity (new test) captures a verify-shaped chain (quantize, small-K and large-K MMVQ, BF16 GEMV, the shared Q8_1 buffer rewritten, a memcpy node, a fork and a join) with and without PDL: bitwise equal outputs over 200 replays with every buffer NaN-poisoned before each; programmatic edges only kernel to kernel, and never into a node that also follows a non-kernel node; present on sm_90+, absent elsewhere. In our fork the same kernels gave identical windows over full requests (per-layer hashes of the MoE inputs and the final residual). Measured together with the next commit (PDL is on the 5080 only: the 4060 Ti, sm_89, has none); the numbers are there. (cherry picked from commit 8e1324a)
Extract only the graph-branch portion of Hardin22 PR Niko1221#904. The GDN and QSA sides run the same kernels on the same inputs with explicit fork/join edges; input metadata transfers can overlap the embedding path. Q4 KV batching and fused PLE post-ops are deliberately outside this focused integration. Original author retained; extraction by David. Derived-from: 1d4f833 Source-PR: Niko1221#904
perf: integrate opt-in PDL and graph branches into Q8 adaptation
Owner
|
Heads-up: today we cleaned up this repository's commit history, and GitHub closed this PR automatically when GitHub cannot reopen it, because its branch is built on the old history. If you want to keep it going: start a new branch from the new Sorry for the noise, and thank you for the contribution. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Draft integration: Q8 on RTX PRO 6000 Blackwell 96 GB, 32K input + 1K output, MTP T4, FP16 KV.
Rotation OFF → ON with async + duplex + per-layer admission fixed: 113.224 → 140.979 tok/s (+24.513%), reverse pair 111.568 → 142.975 (+28.151%). Both pairs have identical output IDs and recorded work. Full-request improvement is 2.88–3.19%; PCIe bytes unchanged, 18.0 GB host-copy payload removed.
Fresh upstream comparison: main + Q8 reader 106.1–106.8 tok/s → full stack 141.7–141.8 tok/s. The reader is required to load this Q8 PLE; the full configurations have different output/work trajectories. See the report for both paired results and full request times.
Depends on #864, #865 and Hardin22's #876 (author preserved) plus duplex/integration work. All flags opt-in; presets unchanged. Seven logical implementation commits. Three CUDA memcheck fixtures: zero errors; default-off parity and Q2 lifecycle checks passed. Broader Q8 state/cancellation/long-input coverage remains, so please keep this as a draft.
Full results, upstream TPS, limitations and history
Immutable evidence and token/work verifier