cpu: IQ3_S / IQ2_S AVX2 decode without the per-index shift/mask/or (1.07-1.6x per row) by sanastasiou · Pull Request #930 · Niko1221/Strata

sanastasiou · 2026-10-05T11:09:06Z

The AVX2 multi-token kernels in src/kernels/cpu/iq_avx2.cpp are compute-bound on the codebook decode, not on the per-token arithmetic: a hot (in-cache) and a cold (256 MiB) working set measure the same ns/row. This PR makes the IQ3_S and IQ2_S decode cheaper. The arithmetic is untouched (vpsignb, vpmaddubsw, vpmaddwd, exact int32 per QK_K block), so the results are bit-identical.

What changes (one file, IQ3_S and IQ2_S only)

Grid indices. The qh byte indexes a compile-time table that spreads its high bits one per index byte. One punpcklbw puts them next to the 8 low index bytes, so each lookup is a 16-bit load plus the grid load instead of a shift, a mask and an or per index.
Sign vectors. The half's four sign bytes are broadcast straight from the block (vpbroadcastd from memory, a load uop) and expanded with ONE pshufb, instead of sgn_vec's two pshufb + vinserti128. On Intel cores those are three uops on the single shuffle port.
Scales. The (2s+1) vectors come from compile-time tables: one load instead of shift/or/broadcast.
No static constructor. The tables are constexpr (like EvenSigns), so the AVX2 TU still has no static constructor (iq_avx2: build the IQ2 sign table without a variable shift (BMI2 \shlx\ traps on pre-Haswell CPUs) #391). static_asserts pin the table layout.
Gather opt-in kept, branch hoisted. IQ3_S's STRATA_IQ256_GATHER path becomes its own instantiation, chosen once per call in iq256_rows / iq256_gu_rows. With the check left inside decode(), the branch sat in the innermost loop and made the new IQ3_S decode 25-35 % slower at nt=2-3 (default path, Zen 3).
IQ3_XXS unchanged. The same hoist measured -5..+5 % there, which is not worth it.

Measured (ns per row of all nt tokens, hot, one pinned thread, best of 3 rounds with base and PR alternated; base = main @ 6f32ec0; the kernels compiled with -march=haswell):

CPU	type	nt=1	nt=2	nt=3	nt=4
Ryzen 9 5950X (Zen 3)	IQ3_S	256.1 → 213.7 (1.20x)	304.4 → 267.5 (1.14x)	360.7 → 313.7 (1.15x)	399.8 → 374.1 (1.07x)
	IQ2_S	208.0 → 144.1 (1.44x)	225.5 → 192.9 (1.17x)	255.6 → 222.7 (1.15x)	279.6 → 271.2 (1.03x)
Ryzen 5 3600 (Zen 2)	IQ3_S	337.7 → 270.4 (1.25x)	386.4 → 320.5 (1.21x)	451.4 → 378.5 (1.19x)	524.8 → 437.3 (1.20x)
	IQ2_S	244.5 → 172.7 (1.42x)	278.3 → 222.8 (1.25x)	320.0 → 265.8 (1.20x)	360.7 → 327.0 (1.10x)
Core i7-7700 (Kaby Lake, AVX2, no AVX-512)	IQ3_S	405.6 → 354.6 (1.14x)	616.6 → 382.8 (1.61x)	507.2 → 426.5 (1.19x)	573.3 → 492.4 (1.16x)
	IQ2_S	310.5 → 208.4 (1.49x)	359.8 → 254.6 (1.41x)	404.8 → 313.4 (1.29x)	440.1 → 365.9 (1.20x)

On the i7-7700, main's IQ3_S at nt=2 is slower than at nt=3. This reproduced on two cores and two runs; with this PR the curve is monotonic again.

I have not measured end to end in the engine, and I have not measured on a Haswell Xeon. The Kaby Lake core has the same single-shuffle-port layout, so I expect Haswell to land in the same range, but that is not measured. Numbers from an E5 v3 owner (#913) would close that gap.

Tests

Parity harness. A standalone harness (our private test bench; happy to contribute it as a CPU-only tests/ target if useful) compiles this exact iq_avx2.cpp and checks iq256_rows against ggml's _generic references and a float64 oracle.
- Every type, nt 1..9, row ranges with odd counts and matrix borders: 844 tests shuffled, pass on Zen 3.
- The strata_iq256 subset: 180 tests, pass on Zen 2 and Kaby Lake.
- The same 844 tests also pass with STRATA_IQ256_GATHER=1.
Mutation-checked. Breaking the IQ3_S high bits or the IQ2_S sign offset fails the parity tests. Swapping the IQ2_S scale nibbles fails at compile time on the static_assert.
Upstream build. CPU-only cmake -DSTRATA_WERROR=ON -DSTRATA_BUILD_TESTS=ON builds clean.
- ctest: 15/16 pass. expert_multi_test fails on this Zen 3 box because it requires AVX-512 ("this CPU cannot run the expert kernel"), not because of this change.
- native_expert_parity needs CUDA plus the llama.cpp oracle and does not cover the IQ formats; I did not run it.
No AVX-512. Nothing in the shipped object uses EVEX or zmm; this is checked by disassembly.

🤖 Generated with Claude Code

… configs are written atomically (Niko1221#459) A fresh copy of Strata set itself up like the newest strata-*.json of another Strata folder and crashed with JSONDecodeError when that file was empty or cut off. previous_config now takes the newest one that parses as a JSON object; each one that does not gets a one-line warning naming it, and with none the copy is set up as a fresh install. Every place setup writes a run config now goes through write_config (a .tmp file, then os.replace), so a setup stopped half-way no longer leaves an empty config behind. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

…s falls back to the wheels (Niko1221#446) rocm_root took /opt/rocm when bin/hipcc and lib/libhipblas.so* were there and the version was 7.0 or newer. A runtime-only ROCm 7 (no -dev packages) has both, and the engine's cmake configure then failed on the missing hip-lang package. The system ROCm now also needs lib/cmake/hip-lang/hip-lang-config.cmake (lib, lib64 or the multiarch lib folder, where CMake looks) and include/hip/hip_runtime.h; without them setup warns, naming what is missing, and takes the TheRock wheels path, as it already does for a ROCm older than 7.0. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

… a failed tuning says why (Niko1221#447) tools/calibrate.py read any "gpu" list as a layer split and added --layer-split auto, so a one-card config with a helper card for the expert tier ("gpu": [0] + --expert-cache-device1) was refused by the engine before it read a weight, a "0,2" string never got its split, and split_skip_if_fits was ignored. The tuner now builds the arguments with serve.server.engine_args, the same as a normal start. When the tuning fails, setup prints the engine's own last error line from the log (only lines written during the tuning), and says the failure again where it is not lost: before --calibrate starts the model, and in the "All set." summary of an install that tuned. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

…etup recommends the big card alone (Niko1221#448) - engine (layer splits only): when a stage's cache cannot fund the chunk CUDA0 alone would choose, a WARNING names it (CUDA index, name, cache slots, the slots the chunk borrows) and points at the helper cache (--expert-cache-device1, docs/SECOND_GPU.md). The chunk choice itself is unchanged: fits() is split into fits_one() with the same tests; one GPU never reaches the new log. - setup: a pair whose second card is under 12 GB beside a first card of 12 GB or more is still offered, but the first card alone is recommended (the default with --yes and Enter; offer_together asks "n" by default), with a note. A named --gpus pair is kept. Equal pairs (2x 8/16/24 GB) are unchanged. - The automatic drop of the small card from an auto split is not done: stages load their weights before the boundary search and the exact prompt-buffer bytes need the sessions made after it. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

…the totals (Niko1221#457) The server already read drafts_offered / drafts_accepted from the engine's DONE line but only returned them in a response's timings (draft_n / draft_n_accepted). Each /metrics request entry now carries both (null when the engine did not report them) and the totals keep their running sums since the server started, so a dashboard such as llm-visuals can show the acceptance. Additive: no existing field changes. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

…he --family/--model to use (Niko1221#444) With Unsloth's UD-IQ3_XXS or UD-Q2_K_XL in the folder, setup either said only "has no IQ3_XXS model file" (--family unsloth --model IQ3_XXS) or took the UD-IQ3_XXS file for the GSQ-RCO IQ3_XXS by the size in its name. Now: - a chosen shard whose name is a quantization setup has no size for (UD-IQ3_XXS, UD-Q2_K_XL, a K-quant) stops with "<file> is UD-IQ3_XXS, a GGUF Strata cannot run" and what Strata runs: ISTA-DASLab's GSQ-RCO files (and Swift 1.5's and the Coder's) and Unsloth's UD-Q4_K_XL only; - a folder without the chosen file but with other GGUFs lists the ones Strata cannot use and the --family/--model of the ones it can (an empty folder keeps check_shards' "missing"); - a size the family does not have names the families that do (--family qwen --model IQ3_XXS), and with --gguf-dir also what Strata runs. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

…e unpruned sizes for that (Niko1221#438) README and docs/MODELS.md said only "Weaker outside coding". Niko1221#438 measured Chinese answers wrong or looping on the Coder where English was fine; both now say plainly that this includes Chinese and other CJK text and point to the sizes that keep every expert (Q2_0, IQ2_XS, IQ3_S). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

…iko1221#461), the low-RAM resident mode on Windows 32 GB PCs (Niko1221#467), fixes Niko1221#460 Niko1221#459 Niko1221#446 Niko1221#447 Niko1221#457 Niko1221#448 Niko1221#444; setup requires it Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

The fill from the GGUF (--stream-experts) read each expert's three slices, copied the blob and waited for the copy before the next read: ~12,000-18,500 experts with neither the SSD nor PCIe kept busy. Now the slots are admitted in profile order first (same placement), reads go in file order by up to 8 threads into page-locked batches of 64, and a batch's copies overlap the next batch's reads. STRATA_FILL_SERIAL=1 keeps the serial fill. Tested on the B70: cold page cache, Coder fill 23.4 GiB 76.2 s -> 18.8 s; launch to first token Coder 82 s -> 26 s, IQ2_XS 120 s -> 41 s; outputs identical to the serial fill. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

… read (Niko1221#471) The engine's `strata serve: prompt ...` line now says "= 0 reused + 12288 of 98179 read in ... (cancelled)" with the rate from the tokens read, and DONE gets a 15th field: the prompt tokens actually read. The server's /metrics history and totals count reused + read for a cancelled read (prompt_total keeps the whole prompt); an older engine's DONE line (no such field) is read as before. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

Engine: when the MTP draft head's subset does not fit, it prints what the head needs, the VRAM free, and the smaller shipped subsets (--draft-vocab cyrillic / en, their MiB from the head's real row size) with how to pick one (START-HERE.bat / ./setup.sh --draft-vocab en, kept in the model's config). Text only, on the failure path. Server: a start that ended that way carries the engine's hint (or the same advice for an older engine) in its "the engine exited before it was ready" error. Setup: a card under 14 GB with the default subset gets a note recommending --draft-vocab en; nothing is changed, an explicit or saved choice gets no note, and 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

…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>

…ted count opt-in (STRATA_RING_BYTES=1) Default: 0.1.39's 384/1024-slot ring and bytes_needed count, so the default loan and output stay as they were. The auto chunk scan (the PR's second commit) is not taken. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

…the better buy `--prefill auto` walked a fixed list of chunk sizes - 32768, 16384, 8192, 6144, ... - and took the first that fitted, with the ring a fixed 384 slots taken off the top before the chunk was ever considered. Both halves of that were wrong on a rig like this one. The list is coarse exactly where it matters. `bytes_needed` is a sum of (T x positive constant) terms plus a max of such sums, so it rises monotonically with T, and so does every test the scan applies - which makes the largest chunk that fits a bisection on the 256-token grid the prompt path already works on. Seven probes against the list's ten, at one `bytes_needed` each. It matters because this rig affords ~8,700 tokens and was handed 8192; 8704 is not on the list. The order was wrong because chunk and ring spend the same borrowed VRAM, and a ring slot is worth far more than a chunk token. Measured here (2x RTX 3060 + 2x RTX 5060, CUDA3 lending at its 90% cap, 120K prompt): 8960 tokens with the 17-slot ring that leaves reads at 963 tok/s, 8192/130 at 1,008, 7168 with the ring full at its 199-slot byte budget at 1,000, 5632/199 at 915. A ring slot is worth ~0.53 tok/s and a chunk token ~0.05, so the 69 slots between a full ring and 8192's 130 are worth more than the 512 chunk tokens they cost - and once the ring IS full, a smaller chunk buys nothing. So: the largest chunk that still leaves the ring full, and only a rig where no chunk can afford one falls back to the old rule with a kRingMin floor. The room is counted in BYTES, not slots. A ring slot is max_blob, but a cache slot holds its own layer's blob - 2.15 MiB against a 2.54 MiB max_blob here. In slots the ring looked 18% cheaper than it is, a chunk whose ring fitted only after that discount was refused outright, and the scan stopped a step short. `bytes_from_slots` is the inverse of the existing `slots_from_bytes`, and `cache_slots_for` factors the latter out for the serve path's per-participant version. Both inverses clamp at k <= 0: the lend budget is `min(slots - 128, pct * slots / 100)` and goes negative on a cache under 128 slots, where `slots - k` is off_[slots] - the end of the prefix sum, so that cache lends 0 bytes rather than reading past its offsets. Upstream's per-participant `fits`/`fits_one` and the Niko1221#448 diagnostic are kept: the ring room is taken over the smallest participant, `alone` is the chunk CUDA0 alone would read in, and the small-card warning is unchanged. Verified on the rig: the scan picks 7680/199 against the old rule's 8960/17, 1018.8 and 1021.1 tok/s on two cold 120K prompts against 963.1 - +5.9%. Q8_0 picks 5632/101, its 506 MiB budget over a 5,222,400 B blob. Co-Authored-By: Claude Code <noreply@anthropic.com> 0.1.39b (Niko1221): behind STRATA_RING_BYTES=1 with the byte-budget ring; the default keeps 0.1.39's chunk list (plan_lend and the serve scan), its log line gains the ring. The INFO fields (the Monitor commit) are not taken. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

…mbers and quality check Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

's fused_gr carveout on HIP) Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

Added where a config goes to several GPUs (setup --gpus / a multi-GPU install, start --gpus, "use both" at start). Single-GPU configs are not touched. Opt-out: setup --no-remote-expert-opt, or "remote_expert_opt": false in the config (a user key, kept on re-runs). The engine uses it only with a helper cache (--expert-cache-device1..3), so a layer split runs as before. tools/test_setup_remote_opt.py; the 2x16 GB golden configs gain the flag. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

…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

….07-1.6x per row) The AVX2 multi-token kernels spend most of their time decoding, not in the per-token arithmetic. For IQ3_S and IQ2_S three parts of that decode get cheaper; the arithmetic (vpsignb, vpmaddubsw, vpmaddwd, exact int32 per QK_K block) is untouched, so the results are bit-identical: - grid indices: the qh byte indexes a compile-time table that spreads its high bits one per index byte; one punpcklbw puts them next to the 8 low index bytes, so each lookup is a 16-bit load plus the grid load instead of a shift, a mask and an or. - sign vectors: the half's four sign bytes are broadcast straight from the block (a load uop) and expanded with ONE pshufb, instead of sgn_vec's two pshufb + vinserti128 (the single shuffle port on Intel cores). - scales: (2s+1) vectors come from compile-time tables, one load instead of a shift/or/broadcast. The tables are constexpr, so the AVX2 TU still has no static constructor (Niko1221#391). IQ3_S's opt-in gather path (STRATA_IQ256_GATHER) becomes its own instantiation chosen once per call: testing it inside decode() put a branch in the innermost loop that made the new IQ3_S decode 25-35 % slower at nt=2-3. IQ3_XXS is unchanged (the same treatment measured -5..+5 %, not worth it).

… shift per indice PR upstream Niko1221#930 (sanastasiou 9233f99, aperta): indici della griglia da tabella costante (un punpcklbw invece di shift/mask/or per indice), vettore dei segni con un solo pshufb, scale (2s+1) da tabella, percorso gather di IQ3_S come istanza separata scelta una volta per chiamata. Aritmetica invariata. Identita': iq256_dump (972 casi, sei formati, 1-9 token) uguale byte per byte prima e dopo, anche con STRATA_IQ256_GATHER=1; mutazione di un bit IQ3_S rilevata. Teacher forcing a posizionamento fisso TFX-pr930 identico bit per bit a TFX-enc (code ed explain). Velocita': gate/up a un thread IQ3_S 1,25-1,40x, IQ2_S 1,00-1,31x; pool reale (7 worker) strato 20 IQ3_S -25..-40%, strato 21 IQ2_S -20..-33% per finestra. Nel modello 23 strati su 48 hanno gate/up IQ3_S o IQ2_S. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Sorgente congelata 14433b0 (r2: Niko1221#904, Encapsulate, Niko1221#910, Q2_0 esatto, Niko1221#930, Niko1221#949, Niko1221#965), eseguibile shipping b0edf8455031; manifest 0.1.39 conservato come -0139, manifest del 5/10 storici, check_source legge prima source-manifest.json. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Merge di upstream v0.1.40 (1735d64) nel main del fork: 26 file in conflitto, 92 blocchi risolti tenendo entrambi i lati. Dove upstream ha rifatto una funzione che il fork aveva gia' (checkpoint di coda, accodamento K/V del draft, gather IQ3_S, ckpt per richiesta Niko1221#861) vince la versione upstream; restano del fork il porto SM75 e la sua quota di riserva, la cache elastica col nucleo (il flag interno si chiama ora elastic_vmm_ per non scontrarsi con vmm_ di upstream), il kernel Q2 esatto come predefinito (STRATA_Q2_LEGACY sceglie quello di upstream), --pool-tasks, la sonda gate-first e le sonde di laboratorio. Il probe del fork sugli scambi residenti usa la nuova firma con la scheda proprietaria del layer; il test conversation_split_failure_test di upstream ha un cast esplicito che MSVC richiede. Verifiche: CTest 99/101 passati (ple_parity cerca il modello in un percorso relativo, expert_multi_test chiede AVX512), q2_exact_parity 0 differenze su 570634 righe, iq256_dump identico alla PR Niko1221#930 (normale e gather); teacher forcing a posizionamento fisso identico byte per byte a upstream 0.1.40 puro su explain; su code il fork resta diverso da upstream come gia' in 0.1.39 (KL 6,5e-4 in entrambe le versioni) e rispetto alla 0.1.39 del fork cambia in 2 posizioni su 169, quante ne cambia upstream tra 0.1.39 e 0.1.40 (argmax e top-5 invariati). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Niko1221 · 2026-10-06T10:07:57Z

Heads-up: today we cleaned up this repository's commit history, and GitHub closed this PR automatically when main was force-pushed (that is the "base branch was force-pushed" event above). It is not a rejection.

GitHub cannot reopen it, because its branch is built on the old history. If you want to keep it going: start a new branch from the new main, cherry-pick your own commits onto it (git fetch upstream && git switch -c my-fix upstream/main && git cherry-pick <your commits>), and open a new PR. Or simply leave it: when we take your change for a release we port it onto the new main ourselves and keep you as its author.

Sorry for the noise, and thank you for the contribution.

…iant knobs (opt-in) Port of Niko1221#930 (sanastasiou) onto the gathered-variant layout of Niko1221#863: the scalar IQ3_S and IQ2_S decodes take their grid indices from a constexpr high-bit table (one punpcklbw instead of a shift, mask and or per index), their sign vectors from one broadcast load and one pshufb, and their (2s+1) scales from constexpr tables. Bit-identical: iq_avx2_parity passes (every variant, 1..8 tokens, against ggml-cpu). On a Ryzen 7 5700X (one thread, IQ3_S gate/up rows, ms an expert) the scalar decode goes 0.403 -> 0.278 at one token, 0.393 -> 0.370 at three; the gathered decode (Fmt32<121>, unchanged) stays ahead at 2, 4 and 5 tokens. Two opt-in knobs to measure the mix (unset: the behaviour so far): - STRATA_IQ256_GATHER_NT=2,4,5: the gathered decode only for those token counts; - STRATA_IQ_DOWN_MT_MIN: the IQ4_NL down rows' multi-token rule apart from STRATA_IQ_MT_MIN (ggml's one-token dot is faster for down on this CPU). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

The high index bits of a 32-value half are spread one per byte by a table and punpcklbw puts them next to the low index bytes, so each grid lookup is a 16-bit load plus the table load. The sign vector is one broadcast load plus one pshufb, and the (2s+1) scale vectors are table loads. Arithmetic is untouched (same words in the same lanes). The decode is ported onto the 0.1.40 layout of iq_avx2.cpp: the gathered IQ3_S/IQ2_S forms (Fmt32<121>, <122>) and the VNNI rows are left as they are. Co-authored-by: Stefanos Anastasiou <stef.anastasiou@protonmail.ch> (PR #930)

Niko1221 · 2026-10-07T12:27:18Z

Thanks, sanastasiou. Your IQ3_S / IQ2_S AVX2 decode is in 0.1.40.2 (ported onto the 0.1.40 iq_avx2.cpp; the gather and VNNI forms are untouched), credited as co-author. On a Xeon E5-2690 v4 the gate/up rows ran 1.07 to 1.47x faster, 10 interleaved pairs, parity tests pass. On by default.

sanastasiou · 2026-10-07T12:29:16Z

Thanks man! I thought I broke something lmao... I'll keep pushing for more performance on sm_86 and CPU!

The high index bits of a 32-value half are spread one per byte by a table and punpcklbw puts them next to the low index bytes, so each grid lookup is a 16-bit load plus the table load. The sign vector is one broadcast load plus one pshufb, and the (2s+1) scale vectors are table loads. Arithmetic is untouched (same words in the same lanes). The decode is ported onto the 0.1.40 layout of iq_avx2.cpp: the gathered IQ3_S/IQ2_S forms (Fmt32<121>, <122>) and the VNNI rows are left as they are. Co-authored-by: Stefanos Anastasiou <stef.anastasiou@protonmail.ch> (PR Niko1221#930)

Niko1221 and others added 30 commits October 2, 2026 09:31

Merge branch 'f-135-eng' into rel-0.1.35

d8d4ad3

Merge branch 'f-135-small' into rel-0.1.35

7b600d5

Merge upstream main (0.1.35) into b70

e5c2836

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA

Merge branch 'f-136-fused' into rel-0.1.36

bb47c5c

docs: 0.1.36's speed (fused Q2_0 prompt experts, cluster decode kerne…

78f071c

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

docs: the cluster decode kernels run on sm_90+ (RTX 50), not RTX 40

36fa455

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

Niko1221 and others added 17 commits October 4, 2026 03:53

docs: STRATA_RING_BYTES=1 (Niko1221#583, opt-in) with its RTX 5070 nu…

06586e3

…mbers and quality check Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

hip_compat: cudaFuncAttributePreferredSharedMemoryCarveout (Niko1221#646

f166564

's fused_gr carveout on HIP) Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

setup: the CUDA 12 engine's driver floor per OS (Windows 527.41 -> 52…

bc0fc49

…8, Linux 525); the wheel pins match the 12.9.1 toolkit Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

Merge branch 'f-139b-cpu' into f-139-rel

c6275a7

Merge branch 'f-139b-old' into f-139-rel

baed889

# Conflicts: # setup.py

Merge branch 'f-139b-arc' into f-139-rel

8e08550

Merge branch 'f-139b-spd' into f-139-rel

cb97e6c

# Conflicts: # src/core/expert_source.cpp # src/kernels/cuda/iq_kernels.cu

Merge branch 'f-139b-par' into f-139-rel

1189218

# Conflicts: # src/core/verify.cpp # src/program/generate.cpp

docs: README refined + six translations

8b36fa0

Niko1221 closed this Oct 6, 2026

Niko1221 force-pushed the main branch from 1735d64 to 82f46a8 Compare October 6, 2026 09:48

sanastasiou deleted the perf/avx2-iq3s-iq2s-decode branch October 8, 2026 06:54

txiki739 mentioned this pull request Oct 12, 2026

CPU i-quant kernels from upstream: IQ4_XS on AVX-2, the gathered IQ3 decode, AVX-VNNI, an AVX2 Q8_K quantizer eddoursul/Strata#17

Closed

Navigation Menu

Repository navigation

cpu: IQ3_S / IQ2_S AVX2 decode without the per-index shift/mask/or (1.07-1.6x per row) - #930

cpu: IQ3_S / IQ2_S AVX2 decode without the per-index shift/mask/or (1.07-1.6x per row)#930
sanastasiou wants to merge 846 commits into
Niko1221:mainfrom
sanastasiou:perf/avx2-iq3s-iq2s-decode

sanastasiou commented Oct 5, 2026 •

edited

Loading

Uh oh!

Niko1221 commented Oct 6, 2026

Uh oh!

Niko1221 commented Oct 7, 2026

Uh oh!

sanastasiou commented Oct 7, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

20 participants

Conversation

sanastasiou commented Oct 5, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Niko1221 commented Oct 6, 2026

Uh oh!

Niko1221 commented Oct 7, 2026

Uh oh!

sanastasiou commented Oct 7, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

20 participants

sanastasiou commented Oct 5, 2026 •

edited

Loading