Skip to content

Qwen3.6-35B-A3B and Ornith-1.5-35B-A3B (qwen35moe): a second, smaller model for 8-12 GB cards - #1402

Closed
aflin wants to merge 11 commits into
Niko1221:mainfrom
aflin:qwen36-35b-support
Closed

aflin wants to merge 11 commits into
Niko1221:mainfrom
aflin:qwen36-35b-support

Conversation

@aflin

@aflin aflin commented Oct 7, 2026 •

Copy link
Copy Markdown

Related: #824 (perronemirko's Qwen3.6 engine, closed by the history rewrite). That PR adds a separate
strata-qwen36 engine beside the main one; this one teaches the main engine the model instead, so the server, setup,
the batched prompt path, MTP and the expert cache all apply to it, and the engine picks the model from the GGUF.

Summary

Qwen3.6-35B-A3B is the same family as Flash-Next (DeltaNet layers, full attention every fourth layer, a MoE with a
shared expert) at 35B parameters: its experts take 12-14 GB of RAM instead of 23-50 GB. It is the model for PCs with
16-32 GB of RAM and 6-12 GB cards. Ornith-1.5-35B-A3B (ornith-ai, MIT) is a coding-agent fine-tune of it with the same
architecture.

The engine reads general.architecture and runs qwen35moe beside qwen4exp on one code path. Nothing on the
command line says which model it is. Flash-Next's arithmetic is unchanged (see "How it was checked").

Measured on Linux (Debian 12), RTX 4070 Ti 12 GB, Ryzen 9 5950X, 62 GB RAM, CUDA 13.0; 32K context, 8-bit K/V, MTP on.
"8 GB" is the same card with 4 GiB held by another process (emulated, not a real 8 GB card).

Short chat, writes 8K prompt, reads 8K prompt, writes
Qwen3.6 UD-IQ4_XS, 12 GB 105 tok/s 4,460 tok/s 86 tok/s
Qwen3.6 UD-IQ4_XS, 8 GB 75 tok/s 2,669 tok/s 68 tok/s
llama.cpp b11438, same file, 12 GB (--fit on) 51 tok/s 708 tok/s 49 tok/s
llama.cpp b11438, 8 GB 39 tok/s 490 tok/s 36 tok/s

Greedy tokens match llama.cpp b11438 exactly on UD-IQ4_XS (three short prompts and the 8K prompt, 32 tokens each,
expert cache at 3,000 slots). With a different cache size a near-tie can flip, because experts computed on the GPU
round slightly differently from the CPU's (seen once: llama.cpp's top two at 28% and 26%). More numbers (UD-IQ3_S,
Ornith, small cards, 256K) are in docs/QWEN36.md.

What changed

Engine: model geometry (include/strata/core/layout.hpp, src/core/layout.cpp, gguf_reader.hpp)

  • ModelGeometry gains the architecture, experts per token, and whether the model has PLE, the sparse-attention
    indexer and hyper-connections; geometry_from_gguf fills it and checks that each size fits the kernels' limits
    (all of Qwen3.6's are at or below Flash-Next's, which stay as compile-time capacities).

Engine: decode (src/core/verify.cpp, src/kernels/cuda/*.cu, src/kernels/cpu/*, src/core/expert_source.cpp)

  • One residual stream (an RMSNorm before the mixer and before the MoE, the output added back), the silu(z) DeltaNet
    output gate, and dense attention: every cached cell is attended (qsa_select_all), with no indexer projections,
    snapshots or replay.
  • Kernel template instances for 8 or 12 query heads per K/V head, 32 or 48 DeltaNet heads, and 8 or 10 experts; the
    CPU pool and expert layout take the model's widths.
  • Native Q2_K, Q3_K and Q6_K expert formats (Unsloth's mixes use them); iq_row_bytes covers Q2_K (without it the
    MTP layer's Q2_K experts read as garbage).

Engine: prompt path (src/prefill/prefill.cpp, kernels.cu)

  • The batched path for one stream: DeltaNet kernels for 32 heads and the silu gate, routing and combine for 8 of 256,
    a dense mode for the tensor-core prompt attention. NVIDIA sm_75+; elsewhere (Prefill::supports) the prompt is read
    through the decode windows.
  • Main's deferred shared expert uses the model's sizes.

Engine: MTP (src/core/mtp.cpp)

  • The draft layer is read from the model GGUF's blk.40 (eh_proj split into its embedding and hidden halves,
    native projections and experts). No separate download.

Engine: VRAM on small cards (src/program/generate.cpp, Verifier::dense_attention_bytes)

  • The cache sizing keeps the dense attention's buffers out of the expert cache. They grow with the context and are
    allocated after the cache: without this a 256K context failed on every card.
  • The small-card reserve floor is 500 MiB for models without the indexer (measured: 400 failed in the graph
    instantiation at 8K, 32K and 64K, 450 started) and stays 300 for Flash-Next.
  • A cache too small to work stops the start with how much is short and what makes room, instead of failing later on
    the prompt buffers.
  • Result: UD-IQ3_S with setup's settings starts from 4.17 GB free; with an 8K context and --draft-vocab en, from
    3.9 GB.

Data

  • data/expert-profile-qwen36.bin: 40 x 256 experts ranked by routing counts from 24 varied prompts (server
    --expert-profile-save with a 64-slot cache). Against an unranked profile, on prompts not among the 24 (three
    runs): 12 GB short-chat writes 99 -> 106 tok/s, 8K writes 84 -> 93 tok/s; no difference at 8 GB.

Setup (setup.py)

  • Families qwen36 (unsloth/Qwen3.6-35B-A3B-MTP-GGUF: UD-IQ4_XS, UD-IQ3_S) and ornith
    (bartowski/Ornith-1.5-35B-A3B-GGUF: IQ4_XS, IQ3_XXS), pinned revisions and SHA-256. ModelScope serves the same four
    files with the same SHA-256.
  • Suggested where no Flash-Next size fits the RAM (16-24 GB PCs). On a card under 5.5 GB setup recommends the smallest
    size and an 8K context (a recommendation, never forced).
  • The published engines (0.1.40 to 0.1.40.2) predate qwen35moe and their version cannot tell them apart, so setup
    looks for the support in the engine file itself and compiles the engine when a ready-made one lacks it. Flash-Next
    keeps the ready-made engine.
  • Tests: tools/test_setup_qwen36.py, tools/test_setup_ornith.py; updated test_setup_choices.py,
    test_setup_unsloth.py.

Tools

  • tools/iq_pack.py leaves the MTP layer out of the expert table; tools/make_profile.py --n-layer;
    tools/iq_fixture.py adds Q2_K and Q4_0; tools/strata_mcp.py recommend() follows setup's rules.

Docs

  • New docs/QWEN36.md (setup, what fits, small cards, speed, against llama.cpp, Ornith, differences, limits);
    sections in docs/MODELS.md and README.md.

What it leaves alone

  • Flash-Next's arithmetic and defaults. Its compile-time sizes stay as the kernels' capacities; every new kernel
    instance is a template next to the existing one, and Flash-Next takes the same instance as before. The setup
    families, sizes and defaults for Flash-Next, Swift, the Coder and Unsloth are unchanged. Two things a Flash-Next user
    can see: setup's "less than 12 GB of VRAM" warning adds a pointer to --family qwen36 when no family was named, and
    a start whose automatic expert cache gets no slot at all now stops there with the existing "no VRAM is left" message
    instead of going on to fail at the prompt buffers.
  • The SYCL engine keeps Flash-Next only: its prompt kernels take the shared header's new size parameters and
    stop with a message for any other size (sycl/src/prefill/kernels.dp.cpp).
  • Multi-GPU: a qwen35moe model refuses --peer-device, --layer-split and --expert-cache-device1..3 at the
    start; those paths carry Flash-Next's shapes and were not run with it.
  • Opt-in features from main that are written for Flash-Next's shapes stay Flash-Next only: the prompt path's CPU
    share (STRATA_PREFILL_CPU_SHARE), the chunked GDN recurrence (STRATA_GDN_CHUNKED), the query branch of
    STRATA_DF_BRANCH and the routing debug hook (STRATA_DBG_FEAT).

How it was checked

  • Backends built (AGENTS.md): CUDA (sm_89, with the tests), HIP (gfx1100, gfx1201, gfx1030 and gfx1151, plus the
    STRATA_HIP_GFX906 build for gfx906, with ROCm 7.10 from AMD's TheRock wheels) and SYCL (the strata-sycl-dev image,
    every target). Only the CUDA build was run; there is no AMD or Intel card here.
  • Flash-Next against main (main's own build at fb58e0d, the same PC, Coder IQ1_M, expert cache 900 slots): the
    short prompts give main's tokens (prompts 1 and 2 identical; prompt 0 alternates between the same two outputs in
    both builds). The 8K prompt is not repeatable on main itself on this PC (with --adapt-every 0 --pcie-frac 0, with
    and without the MTP layer, with STRATA_STAGER_SLEEP=0, STRATA_STAGE_PIN=0 or STRATA_PREFILL_STREAM_AHEAD=0:
    three runs gave three different outputs, while 0.1.40.2 repeated exactly), so it cannot be compared byte for byte;
    the merged build's outputs were among main's. I can open a separate issue with the numbers if that is useful.
  • Qwen3.6: greedy tokens = llama.cpp b11438 on UD-IQ4_XS (3 prompts + 8K, cache 3,000 slots); speed, the small-card
    floor and 256K as measured before the merge.
  • Tests: iq_parity (with and without STRATA_PARITY_IQ_MMVQ) and native_grouped_parity 0 failures; setup, MCP
    and profile tests pass except tools/test_setup_engine_hash.py's two, which fail the same way on main.

Extra Notes

  • Not run on AMD (HIP), Intel (SYCL) or Windows (built only, see above). The prompt path for these models is
    NVIDIA-only (elsewhere prompts are read through the decode windows), and setup's compile path for them is untested
    on Windows. Test requests for those could go in docs/TEST_REQUESTS.md.
  • One GPU only for these models; no images yet.
  • Ornith's draft layer is accepted less often than Qwen3.6's (61% vs 80% of all drafts on a chat prompt). A Q8_0
    copy of the layer did no better, so it is the fine-tune's draft layer, not its quantization.
  • Version: QWEN36_ENGINE follows this tree's version (0.1.41). CMakeLists and MIN_ENGINE are untouched.
  • Commits: the work, and merges of main at 0.1.40.2 and 0.1.41 with their conflicts resolved (12 and 10 blocks).
    Squash or rebase as you prefer.
  • More sizes (Qwen3.6 UD-Q4_K_XL, Ornith Q4_K_L, asked for in the comments) will be a separate pull request once
    this one is settled (one change per pull request).

🤖 Generated with Claude Code

aflin and others added 2 commits October 7, 2026 09:13
… model for 8-12 GB cards

The engine reads general.architecture from the GGUF and runs qwen35moe beside
qwen4exp (Flash-Next) on the same code path, switching on the model's geometry:
one residual stream instead of hyper-connections, dense attention (no indexer),
no PLE, a silu(z) DeltaNet output gate, 8 of 256 experts 2048/512 wide. Flash-
Next's compile-time sizes stay as capacities; kernels get template instances
(8 or 12 query heads, 32/48 DeltaNet heads, 8/10 experts), so its code is
unchanged: the same tokens and speed as before (Coder, 3 prompts + 8K, A/B).

- decode (verify windows), batched prompt path (NVIDIA sm_75+; elsewhere the
  prompt is read through the decode windows), MTP draft layer read from the
  model file's blk.40, server, checkpoints, low-RAM mode
- native Q2_K/Q3_K/Q6_K expert formats (Unsloth's mixes); iq_parity covers
  Q2_K and Q4_0
- greedy tokens match llama.cpp b11438 (UD-IQ4_XS: 3 prompts + 8K exactly)
- RTX 4070 Ti 12 GB, 32K, MTP: UD-IQ4_XS 105 tok/s chat, 4,460 tok/s 8K read;
  8 GB emulated 75 / 2,669; llama.cpp on the same PC 51 / 708
- small cards: the expert cache keeps the dense attention (grows with the
  context; 256K failed on every card before) and a 500 MiB floor out of the
  cache; a cache too small to work stops the start with what is short.
  UD-IQ3_S starts from 4.17 GB free at 32K, 3.9 GB at 8K + --draft-vocab en
- data/expert-profile-qwen36.bin: ranked from 24 prompts' routing
  (--expert-profile-save, 64-slot cache); 12 GB writes +7-11% against an
  unranked one

setup: families qwen36 (unsloth/Qwen3.6-35B-A3B-MTP-GGUF UD-IQ4_XS, UD-IQ3_S)
and ornith (bartowski/Ornith-1.5-35B-A3B-GGUF IQ4_XS, IQ3_XXS), pinned
revisions and SHA-256; suggested where no Flash-Next size fits the RAM; on a
card under 5.5 GB the smallest size and an 8K context are recommended. The
ready-made engines predate qwen35moe, so setup looks for it in the engine and
compiles one when it is missing. MCP recommend() follows the same rules.
Tests: tools/test_setup_qwen36.py, tools/test_setup_ornith.py.

Docs: docs/QWEN36.md (setup, what fits, small cards, speed, llama.cpp,
Ornith, limits), MODELS.md, README.

Tested on Linux (Debian 12, kernel 6.1) with an NVIDIA RTX 4070 Ti 12 GB
(driver 580), CUDA 13.0, Ryzen 9 5950X, 62 GB RAM. Not run on AMD or Windows.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Conflicts, resolved for both models:
- setup.py: the Flash-Next draft-layer fetch (now with STRATA_SOURCE) stays in
  the branch the 35B-A3B families skip (their draft layer is in the model file)
- verify.cpp: STRATA_DF_BRANCH's query branch only with the indexer (its side 1
  is the indexer query); qwen35moe takes the plain path
- qsa_decode_attn.cu: the pre-sm_75/HIP kernel takes both template parameters
  (query heads per K/V head G; NC_/DPP_ for gfx103x)
- pool.cpp: the native phases' rows from the format (n_ff/n_embd), tasks from
  phase_tasks (--pool-tasks)
- prefill.cpp: the deferred shared expert and the CPU share's buffers with the
  model's sizes; the CPU share (opt-in) is off for one-stream models (unmeasured)
- generate.cpp: the dense-attention log line and the APU warning both kept

Also: QWEN36_ENGINE = 0.1.40.2 (this tree's version, setup's four-number
versions); docs/QWEN36.md names 0.1.40.2 among the engines without qwen35moe,
and says the llama.cpp match is exact with a fixed cache (a near-tie can move
with the cache's contents). ModelScope serves the four 35B-A3B files with the
pinned SHA-256.

Checked after the merge (RTX 4070 Ti, Linux, CUDA 13.0): Flash-Next Coder
tokens identical to upstream's own build (3 prompts + 8K); Qwen3.6 UD-IQ4_XS
and Ornith IQ4_XS tokens identical to the pre-merge engine at the same cache
sizes; speed unchanged; small-card floor and 256K as before; iq_parity and
native_grouped_parity 0 failures; setup/MCP tests pass except
test_setup_engine_hash.py's two, which fail the same way on upstream/main.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@ohmgloss

ohmgloss commented Oct 8, 2026

Copy link
Copy Markdown

Would it be possible to add UD-Q4_K_XL quant?

Conflicts, resolved for both models:
- prefill.cpp: main's routing debug hook (STRATA_DBG_FEAT, opt-in) and CPU-share
  arms only for Flash-Next (one-stream models keep the CPU share off); the
  peer-device inputs, sums and combine with the model's sizes
- kernels.cu / kernels.hpp: the opt-in chunked GDN recurrence (STRATA_GDN_CHUNKED)
  only for Flash-Next's 48 value heads, VPK kept for it; the peer-device kernels
  beside the templated moe_combine
- mtp.cpp: formatting only

Also:
- generate.cpp: a qwen35moe model refuses --peer-device, --layer-split and
  --expert-cache-device1..3 at the start (their code carries Flash-Next's shapes)
- sycl/src/prefill/kernels.dp.cpp: the prompt kernels take the shared header's
  new size parameters (they were separate overloads: an ambiguous call); any
  size but Flash-Next's stops with a message - the SYCL engine has Flash-Next's
  kernels only
- QWEN36_ENGINE = 0.1.41 (this tree); docs/QWEN36.md names 0.1.41 among the
  engines without qwen35moe

Built: CUDA (sm_89, with tests), HIP (gfx1100, gfx1201, gfx1030, gfx1151; and
STRATA_HIP_GFX906 for gfx906) with ROCm 7.10 (TheRock wheels), SYCL (the
strata-sycl-dev image, every target). Run on an RTX 4070 Ti only.

Checked (RTX 4070 Ti, Linux, CUDA 13.0): Qwen3.6 UD-IQ4_XS = llama.cpp b11438
on 3 prompts + 8K (cache 3,000 slots); Ornith IQ4_XS as before; Qwen3.6 speed,
small-card floor (4.15 GB free starts) and 256K as before; iq_parity and
native_grouped_parity 0 failures; setup/MCP tests pass except
test_setup_engine_hash.py's two (they fail the same way on main). Flash-Next
Coder: prompts 1 and 2 identical to main's own build, prompt 0 alternates
between the same two outputs in both; the 8K prompt varies from run to run in
main's own build too (fixed cache, no adaptive swaps, no PCIe share, with and
without MTP: 3 runs, 3 outputs), and the merged build's outputs are among
main's.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@aflin

aflin commented Oct 8, 2026

Copy link
Copy Markdown
Author

Merged main at fb58e0d (0.1.41) and resolved the conflicts (10 blocks: prefill.cpp, kernels.cu/.hpp, mtp.cpp); the branch merges cleanly again.

Following the new AGENTS.md: built CUDA, HIP (gfx1100, gfx1201, gfx1030, gfx1151, and the gfx906 build) and SYCL (every target). The SYCL prompt kernels needed the shared header's new size parameters (sycl/src/prefill/kernels.dp.cpp); they stop with a message for anything but Flash-Next's sizes. A qwen35moe model now refuses --peer-device, --layer-split and the remote expert caches at the start. The description has a new "What it leaves alone" and "How it was checked".

One thing found on the way, not from this PR: on main's own build (fb58e0d) the Coder's 8K prompt is not repeatable on my RTX 4070 Ti — three runs, three different outputs, with a fixed cache (900), --adapt-every 0 --pcie-frac 0, with or without the MTP layer, and with STRATA_STAGER_SLEEP=0, STRATA_STAGE_PIN=0 or STRATA_PREFILL_STREAM_AHEAD=0. At 0.1.40.2 the same test repeated exactly. Happy to open an issue with the details.

@aflin

aflin commented Oct 8, 2026

Copy link
Copy Markdown
Author

@ohmgloss Yes: Unsloth's Qwen3.6-35B-A3B UD-Q4_K_XL (22.9 GB, with the MTP layer) runs. On an RTX 4070 Ti 12 GB it wrote 82 tok/s (62 with an 8 GB card's VRAM, emulated), against 105 for UD-IQ4_XS, with greedy tokens matching llama.cpp. For Ornith the only UD-Q4_K_XL I found (peculiar-ragdoll) has no MTP layer; bartowski's Q4_K_L has one and runs at 73 tok/s. I'll add both to setup in a separate PR once this one is settled.

@vehystrix

Copy link
Copy Markdown

Could you add KAT-Coder-V2.5-Dev?
I currently use the quality-v2 quant from myric (on llama.cpp), mainly because it has MTP, but I'll be happy with any quality quant (probably around ~20-22gb)
https://huggingface.co/Myric/KAT-Coder-V2.5-Dev-MTP-APEX-GGUF

@ohmgloss

ohmgloss commented Oct 9, 2026

Copy link
Copy Markdown

On a GTX 1080 (Pascal, sm_61) with the experimental CUDA 12 engine (engine-cuda12, 0.1.41, toolkit 12), Qwen3.6-35B-A3B cannot start:

strata serve: prompt chunk auto: 3328 tokens, a 268-slot ring
strata serve: the prompt path borrows 1479 CUDA0 cache slots (2.05 GiB)
strata serve: prefill: geometry differs from the artifact's

The message is misleading. In src/prefill/prefill.cpp, Prefill::supports() returns false for a card below sm_75 when !g.has_indexer:

// Without an indexer the prompt attention is the tensor-core kernel's dense mode: NVIDIA sm_75+ (elsewhere
// the prompt is read through the decode windows)
if (!g.has_indexer) {
... cudaDeviceGetAttribute(&major, ..., dev) ... || major * 10 + minor < 75) { return false; }
}

qwen35moe has no indexer, so on sm_61 this returns false; Prefill::init then reports "geometry differs from the artifact's", which is about the pack and sends you looking in the wrong place. The comment says the prompt is read through the decode windows below sm_75 — the only "elsewhere" in the code is #if defined(STRATA_USE_HIP), which is also return false, and nothing in the caller falls back.

Two asks:

  1. Either implement the sm<75 fallback the comment describes, or drop the comment and give this branch its own error ("this GPU (sm_61) has no prompt path for this model without an indexer; sm_75+ is needed"). Prefill::init currently folds three reasons into one sentence.

  2. In setup.py, the qwen36 family already says "its batched prompt path is NVIDIA-only (sm_75+)" — but nothing checks the chosen card's compute capability, so Pascal is admitted by old_gpus_opt_in and the family is offered anyway, after an 18 GB download. A refusal (or a clear warning) before the download would have saved the whole run.

Config and pack are not involved: --prefill 2048 instead of auto, and a rebuilt pack, both give the same failure. The engine version matches the source (0.1.41).

Workaround for now: Qwen3.8-Flash-Next (kQwen4Exp path, no CC check) runs on this card.

@misthema

misthema commented Oct 9, 2026 •

Copy link
Copy Markdown

The emulated 8 GB rows made me curious how the real thing behaves, so here are numbers from an actual
one: an RTX 2080 SUPER 8 GB on PCIe 3.0 — and it's my display card, sharing VRAM with the desktop the
whole time.

Box: RTX 2080 SUPER 8 GB (display card; ~2 GB desktop + a 350 MB co-resident process beside the
engine), i5-12400F (no AVX-512 — the AVX2 expert pool), 62 GB DDR4-3200, CUDA 13.4, SM75 build with
STRATA_MMQ_KQUANTS=ON. PCIe probe reported 12.8 GB/s H2D.

Model: a community APEX-MTP build of Qwen3.6-35B-A3B with mixed per-layer expert quants
(21×IQ4_XS + 10×Q5_K + 9×Q6_K). The Q6_K expert layers hit "no GPU kernels for" — expected per the
supported() list, but worth knowing APEX-Quality files trip it. I converted the 27 Q6_K expert tensors
to Q5_K with ggml's reference path (repo with the file + conversion tools:
misthemAI/Qwen3.6-35B-A3B-APEX-MTP-Strata-Q5EXPS-GGUF).

Three-turn repo-review workload, same file both engines, llama.cpp b11438-era via LM Studio
(131K ctx, KV on GPU, 41 expert layers in RAM, its MTP speculative on — 0.91–0.96 acceptance, 2-wide):

llama.cpp Strata (this PR)
cold prefill ≈11K tok 234.6 tok/s 308.3 tok/s
second turn 2,738 tok re-prefilled (18.4 s total) 11,024 reused + 1,175 read (9.1 s total)
decode 28.5–32.2 tok/s 36.7–45.7 tok/s

131K ctx / int8 KV / 32K kv-resident fits the 8 GB card with ~420 expert slots (~40–50% decode hit
rate + ~16–18% pulled over PCIe); decode held 37–46 tok/s with the desktop sharing the card.
The conversation cache is the quiet star on agentic loops — second-turn wall time halved.

Two findings for the PR:

  1. Bug: with --max-context 8192 and no --kv flag (which setup's own logic produces, since it only
    adds --kv above 8192), the run exits with qsa: kv_append: geometry is not set up. The plain-KV path
    passes an un-overridden QsaShapes into the qwen35moe QSA layers. Workaround: always pass a --kv value.
  2. iq_pack.py's tokenizer step needs the regex PyPI package; setup installs the pinned one but a
    bare tools invocation on a fresh box dies with ModuleNotFoundError after the 20-minute conversion.
    Cheap fix: check early or add to a tools requirements note.

Also confirmed on this card: SM75 MMVQ/prompt kernels run, the hyper-connection staged read
checks out bit-for-bit, and MTP-from-file works with --mtp <model.gguf> + --mtp-draft-vocab.
Thanks for the PR — it turned my daily driver into a faster one on the same silicon.

- 16-bit K/V (no --kv, as setup writes at 8K) stopped at the first append with
  "qsa: kv_append: geometry is not set up": the size check required the
  indexer's sizes, which a model without one leaves at 0. Dense attention no
  longer needs them; Flash-Next keeps the same check.
- strata serve on a card without the batched prompt path for the model
  (qwen35moe below sm_75, or HIP) failed with "prefill: geometry differs from
  the artifact's". It now says so at start and reads every prompt through the
  decode windows, as one-shot generate already did.
- setup asks before the download on Pascal / Volta cards, as it does on AMD.

Tested on Linux + RTX 4070 Ti + CUDA 13: 16-bit K/V at 8K gives llama.cpp's
tokens on 4 prompts; the fallback with the sm_75 check raised in a scratch
build (server: a 9,802-token prompt in 113 s, sensible answers). Not run on a
real Pascal or AMD card. Flash-Next Coder tokens = main's; HIP and SYCL build.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@aflin

aflin commented Oct 9, 2026

Copy link
Copy Markdown
Author

Thank you both for the reports, and for the detail in them. Both bugs are fixed in 20f4d4b.

@misthema: the 8K crash without --kv came from a size check that expected the indexer's sizes, which this model doesn't have. I reproduced it here; with 16-bit K/V at 8K the output now matches llama.cpp's tokens. Thanks also for the numbers from a real 8 GB card. The regex issue in iq_pack.py is in upstream code that Flash-Next packs use too, so I'll send that as its own small PR.

@ohmgloss: strata serve had no fallback when the card lacks the batched prompt path; only one-shot generate did. It now says so at start and reads prompts through the decode windows instead. Setup now asks before the download on Pascal/Volta, as it does on AMD.

A limitation: I don't have a pre-RTX 20 card or an AMD card. I tested the fallback on an RTX 4070 Ti with a scratch build that treats it as sm_61. The server started and answered correctly, and a 9,800-token prompt took 113 s, against 3.9 s on the batched path. That shows the code path works, but not that the decode kernels run on Pascal for this model, or how fast it is there. If you can retry on the GTX 1080, I'd be glad to hear how it goes.

@misthema

misthema commented Oct 9, 2026

Copy link
Copy Markdown

Follow-up with more numbers, plus a correction to my first comment.

The correction first: my "no GPU kernels for Q6_K experts" note was against an older commit of this
branch (aa7ab97-era). The current head (20f4d4b) already carries upstream's STRATA_Q6K_EXPERTS via the
0.1.41 merge - built with -DSTRATA_Q6K_EXPERTS=ON, it loads mudler's mixed-quant APEX-MTP file
(I-Quality, 21×IQ4_XS + 10×Q5_K + 9×Q6_K experts) unconverted, on Pascal of all things. The converted
file I mentioned still exists, but with this head it's only needed for stock engine builds.

New box: the dual-GPU Pascal machine - Tesla P40 24 GB (headless) + GTX 1080 8 GB, i7-9700K
(AVX2 pool, no AVX-512), 48 GB RAM, CUDA 12.9, SM61 build (STRATA_EXPERIMENTAL_SM60 - works, with the
expected dp4a path; CUDA 13 can't target Pascal so 12.x it is). Same three-turn repo-review workload,
same file in both engines now, llama.cpp b10935 with its draft-MTP (3-wide @ 0.75):

card engine cold prefill 12.6K decode
P40 llama.cpp (-ngl all, 2 CPU-expert layers) 754.5 tok/s 40.4–41.1 tok/s
P40 this PR (8,149 resident experts) 104.1 tok/s 64.9–70.1 tok/s
1080 llama.cpp (--n-cpu-moe 41) 123.6 tok/s 17.7–21.9 tok/s
1080 this PR (1,423 resident experts) 43.8 tok/s 24.7–32.6 tok/s

(With my Q5-converted file instead of the original, the PR's numbers are identical within noise on both
cards - the expert quant choice isn't what moves Pascal throughput.)

Reading: decode wins everywhere we've measured - +50–65% on Pascal, +31–42% on the Turing card from
the first comment - and the P40 at ~80% residency (8,149 of 10,240 experts) holds 65–70 tok/s flat
across short, incremental and 769-token turns. Prefill is the mirror image: on Pascal the engine is
2.8×–7.2× behind llama.cpp (43.8 vs 123.6, and 104.1 vs a frankly absurd 754.5), while on Turing it was
the one ahead (308 vs 235). So for anyone on pre-Turing cards: chat/agentic loads (decode-bound) favor
this PR strongly; one-shot long-document crunching still favors llama.cpp there. The SM61 prefill path
is the one place the emulated rows were optimistic about.

Build notes for Pascal users, in case they're useful:

  • -DSTRATA_EXPERIMENTAL_SM60=ON -DSTRATA_Q6K_EXPERTS=ON -DCMAKE_CUDA_ARCHITECTURES=61, CUDA 12.x.
  • Under CUDA 12.9 + GCC 14 one test binary (dma_batch_parity) fails to link (nvcc host-link misses
    libstdc++ symbols); ninja strata builds the engine clean without it.
  • At 131K ctx the 1080 sized 1,423 expert slots next to a "headless"-desktop ~260 MiB footprint; the P40
    sizes 8,149 with no explicit reserve (its own 700 MiB floor).

Thanks again.

@aflin

aflin commented Oct 9, 2026

Copy link
Copy Markdown
Author

@ohmgloss @vehystrix Thanks for the requests. Both files now run on a separate branch while #1402 is reviewed (it isn't a pull request yet):

https://github.com/aflin/Strata/tree/qwen35moe-more-models

Setup installs them by name from that branch, checking each download against its SHA-256 from Hugging Face:

git clone -b qwen35moe-more-models https://github.com/aflin/Strata.git
cd Strata
./setup.sh --setup --family qwen36 --model UD-Q4_K_XL          # @ohmgloss
./setup.sh --setup --family kat-coder --model i-quality-v2     # @vehystrix

(START-HERE.bat takes the same options on Windows.) Setup compiles the engine on the PC, because the ready-made engines don't have the new kernels.

  • UD-Q4_K_XL: matches llama.cpp on the same file. On an RTX 4070 Ti 12 GB it writes about 84 tokens/s and reads an 8K prompt at about 4,000 tokens/s.
  • KAT-Coder i-quality-v2 (with its MTP layer): matches llama.cpp too, at about 86 tokens/s writing on the same card. The older i-quality file is there as well.

The branch also covers the other Qwen3.6-35B-A3B files from Unsloth and bartowski. Results for every file, how they were tested, and the one that doesn't load are here:
https://github.com/aflin/Strata/blob/qwen35moe-more-models/docs/QWEN35MOE_GGUFS.md

Limits: everything was tested on one PC (Linux, RTX 4070 Ti, CUDA 13). Nothing has been run on Windows, AMD, Intel, or cards older than the RTX 20 series. @ohmgloss, on the GTX 1080 these files would read prompts through the slower decode path, like the fix for your report, and nobody has tried that on a real Pascal card yet. Let me know if there are problems, and good luck.

@vehystrix

vehystrix commented Oct 9, 2026 •

Copy link
Copy Markdown

I just tested out the branch.

Command used:

./setup.bat --family kat-coder --gguf-dir $HOME\.cache\huggingface\hub\models--Myric--KAT-Coder-V2.5-Dev-MTP-APEX-GGUF\snapshots\e4ae9970325413d582ba2c294122fc5567ec1dd7\ --data-dir E:\Strata\ --build

It built the new engine, using my installed Visual Studio 2026 and Cuda 13.3, and then ran some tuning to get the optimal settings for my PC.

I then pointed it at the llama-context.cpp file from llama.cpp and asked "What's in the file?"
The result: prefill of ~2000t/s, generation of ~90t/s
Using an RTX4080SUPER (16gb) and 64Gb DDR5-6000 (but it for sure didn't need that much ram)

Why that file? I needed something to fill up the context a bit to get reliable numbers. With this I got ~43k tokens of prompt.

What I also had to do:
Replace the default chat template with the froggerific one (replacing E:\Strata\packs\kat-coder-i-quality-v2\tokenizer\chat_template.jinja)
With the default template the model got stuck in a nested thinking loop.

For comparison, with the same model on llama.cpp I got ~15t/s generation and ~500t/s prefill...

Running some agentic coding tasks with it now. If the model degrades into loops or something I'll let you know.

@vehystrix

Copy link
Copy Markdown

I know this is not hard evidence, but after using the model for a while I noticed it was significantly worse than when running under llama.cpp

It was missing tool calls left and right (it never actually got the edit tool call to work, and just decided to rewrite the whole file instead...)
Also it ignored several instructions where I had to steer it back on track.

I never had any of this when running the model with llama.cpp, or at least not so blatant and often.

Same sampling settings used, same chat template, same harness (omp).
I used fp16 for the kv cache with --kv-resident 0 - if I used the default int8 and --kv-resident 32768, it was even worse, to the point I just gave up.
On llama.cpp the kv cache was q8_0

I'm happy to provide more info or run a controlled benchmark, if you tell me what to run.

aflin and others added 2 commits October 9, 2026 19:19
…ense attention)

A qwen35moe model attends to every cell, so with KV streaming (--kv-resident, which setup writes from a
64K context) a selection past the resident cells names more blocks than the VRAM slots hold.  The resolve
kept its misses in a list of n_slots entries and wrote past its end; the blocks it could not place stayed
claimed (page -2) for good, and the attention masked every non-resident cell.  From 32K tokens on, the model
lost about 40% of its context: on a 52.7K-token agentic transcript KL against llama.cpp was 0.57-1.16 (int8
and fp16), with wrong tool calls.

- kv_stream.cu: the resolve claims at most n_slots misses; the rest go back to page -1 and count as read
  from RAM.  ctl[3] (overflow) now means "some blocks were read in place", and the end-of-request line says so.
- qsa_decode_attn.cu: a selected cell whose block is not resident is read from the layer's host copy in place
  (QsaAttnPools::host_*, set by qsa_attn_pools for a streamed layer) instead of masked.  Flash-Next's indexer
  selections always fit the slots, so its path never takes it.
- kv_stream_parity: a dense selection of 10,000 blocks over 4,828 slots, twice: bitwise the resident pool's
  output, nothing left claimed, overflow set (int8, fp16, q4_0, k8v4).

Checked: KAT-Coder i-quality-v2 on an RTX 4070 Ti 12 GB, 64K context, against llama.cpp on the same file at
12.8K-52.7K tokens: KL 0.0006-0.009, top-1 98-100%, with --kv-resident 32768 in int8 and fp16, as without
streaming.  Flash-Next IQ3_XXS with --kv-resident 32768 on a 52.7K prompt: the same 143 tokens as the engine
before this change (which itself differs from run to run at token 88).  Built: CUDA, HIP (gfx1100, gfx1201,
gfx1030, gfx1151), gfx906, SYCL.

Speed: past the resident cells every token reads the rest over PCIe.  52.7K-token prompt, 64K context, no
desktop on the card: 27.0 tok/s streamed vs 59.6 with the KV cache in VRAM (12 GB); 23.3 vs 47.0 (8 GB
emulated).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…om 128K, asked otherwise

Streamed, a qwen35moe model reads the KV cache past 32K tokens from RAM for every token.  Measured with
KAT-Coder i-quality-v2 (8-bit KV) on an RTX 4070 Ti 12 GB with no desktop, whole and with 4 GB held:
- 52.7K-token prompt, 64K context: 59.6 / 47.0 tok/s with the KV cache in VRAM, 27.0 / 23.3 streamed;
- 2K-token prompt: as fast or faster streamed (more expert slots);
- 7.9 GB free: 128K and 256K fail at the first prompt unless streamed; 10 GB free: both run.

So for qwen36 / ornith / their fine-tunes from a 64K context (not WSL):
- under 9 GB of VRAM free (nvidia-smi memory.free, so a desktop counts) from 128K: streaming is turned on,
  and setup says why, what it costs and that --context 65536 avoids it; --kv-streaming off is kept with a
  warning that the first prompt fails;
- otherwise setup explains the trade-off and asks, recommending no (--yes: no);
- both say how to change it later (the config's args, or --kv-streaming on|off).
--kv-streaming on|off answers it as before.  Flash-Next is unchanged and never asked.

Tests: tools/test_setup_qwen36.py KvStreaming (needed vs asked, free VRAM not the card's size, the option,
Flash-Next) and FreeVram.  test_setup_engine_hash fails as before this change (2 failures).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@aflin

aflin commented Oct 10, 2026

Copy link
Copy Markdown
Author

@vehystrix - At least one problem with the --kv-resident has been solved. Please try it out and let me know if there are still problems.

@vehystrix

vehystrix commented Oct 10, 2026 •

Copy link
Copy Markdown

I've been using it for a day now, and haven't really noticed an improvement (I still kept --kv-resident 0, but would that not always be better/more reliable?)

I ran the same prompt trough strata and llama.cpp. Same sampling settings, same .gguf, fp16 on strata & q8_0 on llama.cpp for KV cache, same context size. Prompt was to build a html/css/js webapp, you know, the typical one-shot benchmark stuff.
Final result was of comparable quality, but the strata transcript was full with stuff like this:

I need to either:
 1. Add getters for those filter state variables in filters.js, or
 2. Change render.js to not access them directly and instead check if any filter is active differently.

 Option 2 is cleaner - I can just pass the fact that filtering is active as a parameter, or add simple getter functions to Filters module. Let me fix this by adding getters to filters.js since it's simpler.

╭─── ✘ Edit: 🟨 filters.js ─────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────╮
│line 1: payload line has no preceding hunk header. Use `PUT N.=M:`, `CUT N.=M`, or `PUT <N:`/`PUT >N:` above the body. Got "PUT 47*:\\n+  function getSearchText()    { return searchText; }\\n+  function     │
│getLabelFilter()   { return labelFilter; }\\n+  function getAssigneeFilter(){ return assigneeFilter; }\\n+\\n+  return { setSearch, setLabel, setAssignee, matchesCard, getFilteredCards, getAllAssignees,     │
│reset, getSearchText, getLabelFilter, getAssigneeFilter };".                                                                                                                                                   │
╰───────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────╯

 2026-10-10 15:52:54  Δ 4m18s  ⤵ 1.2K  ⤴ 325  💾 45K  ⏱ 3.4s  ⚡ 44.8/s

 The edit tool is having trouble with my syntax. Let me just rewrite filters.js since it's a small file and the fix is straightforward. I'll also need to update render.js to use the new getters instead of
 direct property access.

 Actually, let me try using write for the whole file - that will be cleaner than fighting with edit syntax.

╭─── ✎ Write: 🟨 ~/wksp/strata-kat/kanban/js/filters.js · 59 lines ────────────────────────────────────────

On the llama.cpp transcript I found 2 edit tool faillures, and both times it recovered like so:

I need to fix two things:

 1. In render.js: Wire up the modal close/cancel buttons inside bindModalForm, and add delete button show/hide logic in _openCardModal. Also handle permanent deletion from within the modal when editing a
    card.
 2. In index.html: Remove the broken references to non-existent exported methods (Renderer._closeCardModal and window.Renderer.__openCardModal). The archive buttons are already wired correctly there, but I
    need to remove the bad lines 169-178.

 Let me make both edits now.

╭─── ✘ Edit: 🗒 … ───────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────╮
│input must begin with "[PATH#HASH]" on the first non-blank line for anchored edits; got: "// ── Card modal (create / edit) ────────────────────────────────────". Example: "[src/foo.ts#1A2B]" then edit ops.  │
╰───────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────────╯

 2026-10-10 16:53:31  Δ 19m13s  ⤵ 2.6K  ⤴ 1.5K  💾 43K  ⏱ 4.7s  ⚡ 31.0/s

 I need to use the proper format with [path#hash] on the first line. Let me re-read render.js to get its current hash, then make the edits properly.

 ● Read ~/wksp/llama-kat/js/render.js:315-420
   2026-10-10 16:54:18  Δ 19m21s  ⤵ 1.5K  ⤴ 102  💾 46K  ⏱ 3.2s  ⚡ 12.5/s

 Now I can make the edits properly using the hash #3AF8 from render.js. Let me do two edits:
 1. Replace _openCardModal, _closeCardModal, and update bindModalForm in render.js
 2. Fix index.html bootstrap script

 Let me use PUT ops to replace these sections. I'll need to be careful about the exact line ranges. The section from lines 315-420 needs replacement, but since edit tool requires original content for CUT or
 range-based PUT, let me use a more targeted approach.

 Actually, looking at the edit tool docs:
 - PUT N.=M: replace inclusive N–M with body
 - I can also do line-by-line edits within ranges

 Let me just replace lines 315-420 entirely by specifying that range in PUT format. The hash is #3AF8 from the read output.

╭─── ✎ Edit: 🟨 ~/wksp/llama-kat/js/render.js ⟦+25⟧

With every other edit call landing on the first try.

I don't know what to make of this, but usage is definitely degraded.

Second point I'd like to make:
prefill speed drops off hard as context increases. What starts at >2000t/s drops down to 150t/s after like 100k context:

[strata] reading the prompt: 8,192 of 162,436 tokens, 3 s so far
[strata] reading the prompt: 13,804 of 162,436 tokens, 6 s so far
[strata] reading the prompt: 21,996 of 162,436 tokens, 11 s so far
[strata] reading the prompt: 30,188 of 162,436 tokens, 19 s so far
[strata] reading the prompt: 38,380 of 162,436 tokens, 29 s so far
[strata] reading the prompt: 38,380 of 162,436 tokens, 39 s so far
[strata] reading the prompt: 46,572 of 162,436 tokens, 43 s so far
[strata] reading the prompt: 46,572 of 162,436 tokens, 53 s so far
[strata] reading the prompt: 54,764 of 162,436 tokens, 63 s so far
[strata] reading the prompt: 54,764 of 162,436 tokens, 73 s so far
[strata] reading the prompt: 54,764 of 162,436 tokens, 83 s so far
[strata] reading the prompt: 62,956 of 162,436 tokens, 88 s so far
[strata] reading the prompt: 62,956 of 162,436 tokens, 98 s so far
[strata] reading the prompt: 62,956 of 162,436 tokens, 108 s so far
[strata] reading the prompt: 71,148 of 162,436 tokens, 117 s so far
[strata] reading the prompt: 71,148 of 162,436 tokens, 127 s so far
[strata] reading the prompt: 71,148 of 162,436 tokens, 137 s so far
[strata] reading the prompt: 71,148 of 162,436 tokens, 147 s so far
[strata] reading the prompt: 79,340 of 162,436 tokens, 149 s so far
[strata] reading the prompt: 79,340 of 162,436 tokens, 160 s so far
[strata] reading the prompt: 79,340 of 162,436 tokens, 170 s so far
[strata] reading the prompt: 79,340 of 162,436 tokens, 180 s so far
[strata] reading the prompt: 87,532 of 162,436 tokens, 186 s so far
[strata] reading the prompt: 87,532 of 162,436 tokens, 196 s so far
[strata] reading the prompt: 87,532 of 162,436 tokens, 206 s so far
[strata] reading the prompt: 87,532 of 162,436 tokens, 216 s so far
[strata] reading the prompt: 95,724 of 162,436 tokens, 226 s so far
[strata] reading the prompt: 95,724 of 162,436 tokens, 236 s so far
[strata] reading the prompt: 95,724 of 162,436 tokens, 246 s so far
[strata] reading the prompt: 95,724 of 162,436 tokens, 256 s so far
[strata] reading the prompt: 95,724 of 162,436 tokens, 266 s so far
[strata] reading the prompt: 103,916 of 162,436 tokens, 269 s so far
[strata] reading the prompt: 103,916 of 162,436 tokens, 279 s so far
[strata] reading the prompt: 103,916 of 162,436 tokens, 289 s so far
[strata] reading the prompt: 103,916 of 162,436 tokens, 299 s so far
[strata] reading the prompt: 103,916 of 162,436 tokens, 309 s so far
[strata] reading the prompt: 112,108 of 162,436 tokens, 315 s so far
[strata] reading the prompt: 112,108 of 162,436 tokens, 326 s so far

Strata config:

{
  "exe": "D:\\repos\\aflin-Strata\\engine\\strata.exe",
  "args": [
    "--pack",
    "E:\\Strata\\packs\\kat-coder-i-quality-v2",
    "--native",
    "C:\\Users\\VEH\\.cache\\huggingface\\hub\\models--Myric--KAT-Coder-V2.5-Dev-MTP-APEX-GGUF\\snapshots\\e4ae9970325413d582ba2c294122fc5567ec1dd7\\KAT-Coder-V2.5-Dev-MTP-APEX-i-quality-v2.gguf",
    "--expert-profile",
    "D:\\repos\\aflin-Strata\\data\\expert-profile-qwen36.bin",
    "--expert-cache",
    "auto",
    "--prefill",
    "auto",
    "--spec",
    "4",
    "--mtp",
    "C:\\Users\\VEH\\.cache\\huggingface\\hub\\models--Myric--KAT-Coder-V2.5-Dev-MTP-APEX-GGUF\\snapshots\\e4ae9970325413d582ba2c294122fc5567ec1dd7\\KAT-Coder-V2.5-Dev-MTP-APEX-i-quality-v2.gguf",
    "--mtp-draft-vocab",
    "D:\\repos\\aflin-Strata\\data\\draft_vocab.bin",
    "--max-context",
    "262144",
    "--kv",
    "fp16",
    "--kv-resident",
    "0",
    "--pcie-frac",
    "0.00",
    "--spec-min-p",
    "0.30"
  ],
  "cwd": "D:\\repos\\aflin-Strata",
  "tokenizer": "E:\\Strata\\packs\\kat-coder-i-quality-v2\\tokenizer",
  "model_name": "kat-coder-v2.5-dev-i-quality-v2",
  "log": "D:\\repos\\aflin-Strata\\strata-kat-coder-i-quality-v2.log",
  "lib_dirs": [
    "C:\\Program Files\\NVIDIA GPU Computing Toolkit\\CUDA\\v13.3\\bin",
    "C:\\Program Files\\NVIDIA GPU Computing Toolkit\\CUDA\\v13.3\\bin\\x64"
  ],
  "port": 8080,
  "gpu": 0,
  "gpus_asked": true,
  "open_browser": false,
  "sampling": {
    "temperature": 0.7,
    "top_p": 0.8,
    "top_k": 20,
    "min_p": 0.0,
    "presence_penalty": 1.5,
    "repetition_penalty": 1.1
  }
}

llama.cpp config:

llama.exe serve
      --port 9931
    --webui-mcp-proxy
    -ctk q8_0
    -ctv q8_0
    --flash-attn on
    --parallel 1
    -cram 24576
    -ngl all
    --no-mmproj
    --no-mmproj-offload
      -hf Myric/KAT-Coder-V2.5-Dev-MTP-APEX-GGUF:KAT-Coder-V2.5-Dev-MTP-APEX-i-quality-v2
      --chat-template-file c:\Users\VEH\.llama\froggeric--Qwen-Fixed-Chat-Templates--chat_template.jinja
      -c 262144
      -ncmoe 18
      --spec-type draft-mtp,ngram-map-k4v,ngram-mod
      --spec-draft-n-max 3
      --spec-draft-n-min 0
      --spec-ngram-map-k4v-size-n 12
    --spec-ngram-map-k4v-size-m 48
    --spec-ngram-map-k4v-min-hits 2
    --spec-ngram-mod-n-match 24
    --spec-ngram-mod-n-min 48
    --spec-ngram-mod-n-max 64
      --temp 0.7
      --top-k 20
      --top-p 0.8
      --presence-penalty 1.5
      --repeat-penalty 1.1
      --reasoning on
      --reasoning-preserve

@aflin

aflin commented Oct 10, 2026

Copy link
Copy Markdown
Author

@vehystrix - I ran ornith1.5 with claude code for a task that took about 140k tokens. I noticed the same problem with an edit tool call, and it worked around it. It also got itself into a loop, but broke out of it when interrupted and prompted, and successfully completed the task.
I'll have another go at it and see what I can find.
Thanks for the report.

Conflicts (8 blocks in 5 files), resolved keeping both sides:
- mtp.cpp: upstream's route-resident adjustment (Flash-Next's 512 experts) runs first, then the GGUF-expert /
  blob branch (qwen35moe's draft layer).
- verify.cpp: dense_attention_bytes and upstream's set_logit_bias, side by side.
- qsa_decode_attn.cu: the attention kernel takes both the query-head count G (12 Flash-Next, 8 qwen35moe) and
  upstream's gfx906 QUERY_SWIZZLE / REDUCE12; the gfx906 launches pass G, and the REDUCE12 score path reads a
  non-resident cell from the host copy (load8r) like the other paths.
- prefill.cpp: draft_kv's Flash-Next-only guard uses upstream's decline(); the CPU share keeps qwen35moe out
  (one_stream) and takes upstream's cpu_dead fallback, with the window's real expert count (KK).
- generate.cpp: the verifier's dense-attention reserve beside upstream's split-aware draft binding; the
  "own buffers" line stays off where there is no prompt path, after upstream's lend-sizing report.

Two follow-ups for qwen35moe:
- STRATA_ROUTE_TAIL_SKIP's new CUDA default applies to Flash-Next only (g.has_indexer): it was measured on
  10-of-512 routing; a qwen35moe model (8 of 256) keeps 0.1.41's answers until measured
  (STRATA_ROUTE_TAIL_SKIP=7 turns it on).  docs/DETAILS.md says so.
- setup: QWEN36_ENGINE = (0, 1, 42), this tree's version.

Checked: CUDA build (all 264 targets); kv_stream / kv_hybrid / kv_q4 / kv_q8 / kv_steps / qsa parity pass;
setup tests pass but test_setup_engine_hash (its 2 failures predate this); the 37 serve test modules pass.
Ornith-1.5 IQ4_XS on an RTX 4070 Ti, greedy against the engine before the merge: a 300-token prompt the same
256 tokens; a 12.8K agentic prompt the same first 199 tokens (0.1.42's measured PCIe share moved more experts to
the CPU here, 0.55 -> 0.30, which upstream documents can change answers slightly); no tail-skip line.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@vehystrix

Copy link
Copy Markdown

It's just that I know what to expect from the model from using it a lot in llama.cpp, and it just feels less capable running with strata.
I really want to use it with strata because of the massive speed improvement, but it's strange the quality is so different with the same weights.

@aflin

aflin commented Oct 10, 2026

Copy link
Copy Markdown
Author

@vehystrix Thanks for the side-by-side runs. What I found:

  • Accuracy at long context: with your settings (KAT i-quality-v2, --kv fp16, --kv-resident 0, 262K context) I compared Strata against llama.cpp on the same file at 88K, 118K and 150K tokens of an agentic transcript. KL stays at 0.003 or below, top-1 agreement is 99.6%, and both engines' tool calls are valid at each point. So the engine's numbers hold up at long context. The differences that remain are in the server: llama.cpp's --reasoning-preserve (Strata drops the reasoning of turns before the latest user message), and how tool definitions are rendered. I'm looking at both. One run each at temperature 0.7 is also noisy, so a few repeats would help tell them apart.
  • Prompt speed: confirmed. 93% of the time goes to the prompt attention kernel. It was written for Flash-Next's sparse attention, and for this model every query reads the whole context by itself. A kernel that shares the K/V reads across queries is the fix; I'm on it.
  • The branch is merged with 0.1.42.

The tensor-core prompt attention ran one block per (query, KV head). With
Flash-Next's 2,051-cell selection that is fine, but in dense mode (qwen35moe)
every block read the whole context on its own, so each prompt chunk got slower
as the context grew: at 150K tokens it was 93% of the prompt with fp16 K/V.

A new kernel for dense mode, 8 query heads per KV head, fp16 or int8 K/V,
sm_80+: a block takes 8 neighbouring queries (64 MMA rows, no pad rows), so a
K/V chunk is read once for all 8. The softmax stays in registers and the q.k
fragments become p.v's A fragments. Accuracy is kept as before: q and p as
FP16 hi + lo, int8 codes exact with their scales in FP32, each chunk's p.v
added to the running sum in FP32. cp.async overlaps the next K with p.v and
the next V with q.k.

Flash-Next (12 heads per KV head), Turing, Volta, AMD and Q4_0/K8V4 K/V keep
their kernels. STRATA_PROMPT_ATTN_DENSE=0 goes back to the old one.

RTX 4070 Ti, KAT i-quality-v2, 150K-token transcript, 262K context:
fp16 K/V 771 s -> 133 s, 8-bit 170 s -> 113 s (llama.cpp: 417 s).
qsa_prompt_attn_parity: new dense checks pass (32K: fp16 5.7x, int8 1.6x,
within 4.2e-6 of FP64 like the old kernel). Long-context KL against llama.cpp
at 88K/118K/150K unchanged (fp16 <= 0.0036, top-1 584-585 of 585). Builds:
CUDA, HIP (gfx1100/1201/1030/1151), gfx906, SYCL.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@aflin

aflin commented Oct 11, 2026

Copy link
Copy Markdown
Author

@vehystrix The prompt-speed fix is pushed. The prompt attention now handles 8 tokens per block, so each piece of the K/V cache is read once for all 8 instead of once per token. On a 150K-token agentic transcript (KAT i-quality-v2, 262K context, RTX 4070 Ti):

K/V Before Now llama.cpp (q8_0, -ncmoe 18)
fp16 (your setting) 771 s (99 tok/s at 146K) 133 s (714 tok/s at 146K) 417 s
8-bit 170 s 113 s

Accuracy is the same as before: against llama.cpp at 88K, 118K and 150K tokens, KL stays at 0.0036 or below and top-1 agreement is 584–585 of 585. It needs an RTX 30 or newer card with fp16 or 8-bit K/V; STRATA_PROMPT_ATTN_DENSE=0 goes back to the old kernel if you want to compare.

If you get a chance, I'd like to know how it does at 100K+ on your setup.

@bizhonggeng

bizhonggeng commented Oct 11, 2026 •

Copy link
Copy Markdown

@aflin Thanks for the long-context accuracy numbers — that's solid.

Regarding --reasoning-preserve: I've already implemented this locally and it works end-to-end. The issue is that Strata drops the assistant's
reasoning_content from turns before the latest user message, which breaks the chain-of-thought on stateless sessions.

What I changed (3 files, no engine recompile):

  1. serve/frontend.py — Added **kwargs to the inner render() in _detect_caps(), and passed preserve_thinking=True through the probe so GET /props reports supports_preserve_reasoning: true.

  2. serve/server.py — Injected preserve_thinking into req["chat_template_kwargs"] in with_shared() (only for api == "openai"). The whitelist alone
    only stores/prints; injection is what makes it work.

  3. serve/web/app.js — assistantMessages() now always includes reasoning_content every turn (was conditional on "no body"). Stateless server can only
    preserve what the client sends back.

Verification: 3-message session (user / assistant+reasoning / user). Compare prompt_tokens with/without the kwarg — A==B>C means injection works.

The diff is ready if you want to review it. Happy to submit as a separate PR or include it here.

@vehystrix Thanks for the side-by-side runs. What I found:

  • Accuracy at long context: with your settings (KAT i-quality-v2, --kv fp16, --kv-resident 0, 262K context) I compared Strata against llama.cpp on the same file at 88K, 118K and 150K tokens of an agentic transcript. KL stays at 0.003 or below, top-1 agreement is 99.6%, and both engines' tool calls are valid at each point. So the engine's numbers hold up at long context. The differences that remain are in the server: llama.cpp's --reasoning-preserve (Strata drops the reasoning of turns before the latest user message), and how tool definitions are rendered. I'm looking at both. One run each at temperature 0.7 is also noisy, so a few repeats would help tell them apart.
  • Prompt speed: confirmed. 93% of the time goes to the prompt attention kernel. It was written for Flash-Next's sparse attention, and for this model every query reads the whole context by itself. A kernel that shares the K/V reads across queries is the fix; I'm on it.
  • The branch is merged with 0.1.42.

@bizhonggeng

Copy link
Copy Markdown

Here's the diff:

   diff --git a/serve/frontend.py b/serve/frontend.py
index 1dccb91..ce61280 100644
--- a/serve/frontend.py
+++ b/serve/frontend.py
@@ -66,9 +66,9 @@ class ChatTemplate:
     def _detect_caps(self) -> dict[str, bool]:
         """llama.cpp's capability names, checked at load time against this template and Strata's tool-call format.
         These are rendering hints, not a guarantee that the model will follow a request."""
-        def render(messages, tools=None):
+        def render(messages, tools=None, **kwargs):
             try:
-                return self.render(messages, tools=tools)
+                return self.render(messages, tools=tools, **kwargs)
             except Exception as exc:                # noqa: BLE001 - a hint only: never stop the server for a probe
                 # A template may reject a role or feature (a custom one may fail in any way): the feature is off,
                 # discovery still answers for the others, and the start-up goes on.
@@ -108,7 +108,7 @@ class ChatTemplate:
                 "supports_system_role": "strata_caps_system" in render(
                     [{"role": "system", "content": "strata_caps_system"}, user]),
                 "supports_parallel_tool_calls": calls_supported(2),
-                "supports_preserve_reasoning": "strata_caps_reasoning" in render(history)}
+                "supports_preserve_reasoning": "strata_caps_reasoning" in render(history, preserve_thinking=True)}
 
 
 # ------------------------------------------------------------------------------------------------ requests
@@ -393,6 +393,8 @@ def openai_to_messages(req: dict) -> tuple[list[dict], list[dict] | None, dict]:
             kwargs = {"enable_thinking": False}
         elif k == "reasoning_effort" and "enable_thinking" not in kwargs:
             kwargs.update(effort_kwargs(v))
+        elif k == "preserve_thinking":
+            kwargs["preserve_thinking"] = bool(v)
     return _late_system_to_user(messages), tools, kwargs
 
 
diff --git a/serve/server.py b/serve/server.py
index 553f440..69650b8 100644
--- a/serve/server.py
+++ b/serve/server.py
@@ -2971,6 +2971,12 @@ class Service:
                     req["thinking"] = {"type": "disabled"}
                 else:
                     req["output_config"] = {"effort": effort}
+        if api == "openai" and "preserve_thinking" in s:
+            # No request-level flag: fill it from the shared Chat settings so the launcher's switch reaches the template
+            # (openai_to_messages reads chat_template_kwargs; the request's own value always wins).
+            ctk = req.get("chat_template_kwargs") if isinstance(req.get("chat_template_kwargs"), dict) else {}
+            if "preserve_thinking" not in ctk:
+                req["chat_template_kwargs"] = {**ctk, "preserve_thinking": bool(s["preserve_thinking"])}
         return req
 
     def start_telemetry(self):
@@ -5544,6 +5550,9 @@ def clean_shared_defaults(d) -> dict:
         elif key == "experimental_speed_projection":
             if not isinstance(value, bool):
                 raise ValueError("experimental_speed_projection: true or false")
+        elif key == "preserve_thinking":
+            if not isinstance(value, bool):
+                raise ValueError("preserve_thinking: true or false")
         else:
             raise ValueError(f"unknown setting {key!r}")
         out[key] = float(value) if key in ("temperature", "top_p") else value
diff --git a/serve/web/app.js b/serve/web/app.js
index 69735fc..8c82072 100644
--- a/serve/web/app.js
+++ b/serve/web/app.js
@@ -739,9 +739,10 @@ function assistantMessages(m) {
   // #1392: a turn with no answer text (only reasoning, a stop before the first content token, or an error) still goes
   // back as an assistant turn, so the history keeps alternating and a reasoning model sees that it already answered.
   // The server leaves such a turn out of the prompt itself (serve/frontend.py, #843); reasoning_content rides along.
+  // The reasoning rides along on every turn too, so a server set to preserve_thinking keeps earlier thinking in the prompt.
   if (!ran.length) {
     const msg = {role: "assistant", content: m.text || ""};
-    if (!m.text && m.reasoning) msg.reasoning_content = m.reasoning;
+    if (m.reasoning) msg.reasoning_content = m.reasoning;
     return [msg];
   }
   const out = [];

@vehystrix

Copy link
Copy Markdown

Note that the --reasoning-preserve doesn't actually do anything in llama.cpp, because it's on by default in the chat template. Since I use the same chat template on strata, it's on there as well. The default template does have it off, but I use the froggeric one on both engines.
There's no code in strata to actively remove the reasoning traces, this is all the chat template.

Prefill speed is massively improved. I'm running it trough some more workloads now.
I'm beginning to think it's small rounding errors/differences compounding everywhere. Maybe I need a simpler harness so the model has less chance of tripping up. I'll probably give pi a try next.

@moondrop12138

Copy link
Copy Markdown

Vision works on qwen35moe with zero engine changes (stock tools/vision encoder + --vision)

The PR gates vision off for the qwen36/ornith families ("vision": False, "mmproj": None in setup.py). I tested the multimodal path on this branch and it works end-to-end as-is — the stock encoder accepts the Ornith mmproj and the engine's --vision/GENI path handles the embeddings on the qwen35moe geometry. Maybe worth wiring up here.

Setup (Windows 11, RTX 5070 Ti sm_120, CUDA 13.3):

  • Engine: local build of this branch (0.1.42) with -DSTRATA_Q6K_EXPERTS=ON (needed for the Q6_K experts and Q6_K dense projections of the APEX quant I run)
  • Encoder: cmake -S tools/vision -B build-vision -G Ninja -DLLAMA_DIR=third_party/llama.cpp -DSTRATA_VISION_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=120 -DSTRATA_PORTABLE=OFF → build-vision/bin/strata-vision.exe, copied next to the engine
  • Model: SC117/Ornith-1.5-35B-A3B-Heretic-MTP-APEX-GGUF APEX-I-Quality (23.5 GB, qwen35moe, fused MTP in blk.40) with that repo's mmproj-Ornith-1.5-35B-A3B-Heretic-MTP-BF16.gguf
  • Run config: engine args gain --vision, plus the vision section (same shape setup writes for Flash-Next):
"vision": {
  "exe": ".../engine/strata-vision.exe",
  "mmproj": ".../mmproj-Ornith-1.5-35B-A3B-Heretic-MTP-BF16.gguf",
  "model": ".../Ornith-1.5-35B-A3B-Heretic-MTP-APEX-I-Quality.gguf",
  "gpu": true,
  "max_tokens": 1024
}

Results

  • strata-vision loads the mmproj in 2.8 s and warms up at 1024 image tokens in 0.3 s, no errors — mtmd takes the Qwen3.5-VL tower as-is
  • /v1/models now reports "input_modalities": ["text", "image"]
  • OpenAI-style request with image_url (base64 data URL) works: a synthetic test image (red circle left, green triangle top-middle, blue square right, caption text bottom) was described correctly — shapes, positions, colors (it even said "olive green" for RGB 133,153,0), and the caption was read. 4.7 s end-to-end; the image cost 283 prompt tokens

Caveats

One GPU (sm_120), one model file, one test image, GPU encoder only (CPU mode untested). I have not tried Qwen3.6-35B-A3B's own Unsloth files — the family already points mmproj_hf at that repo, so an mmproj may be there to wire. Happy to run more checks if useful.

aflin and others added 2 commits October 11, 2026 08:27
…ngs or the config

Qwen3.6's, Ornith's and KAT-Coder's chat templates leave out the thinking of turns before the latest user
message unless preserve_thinking is set; Strata never passed it.  Now:

- a request's chat_template_kwargs.preserve_thinking reaches the template (OpenAI, and the Anthropic API as a
  Strata extension, so Claude Code's thinking blocks can stay in the prompt);
- "preserve_thinking" in the shared Chat settings (POST /settings) or in strata-<model>.json fills it for
  requests that leave it out (llama.cpp's --reasoning-preserve); the request wins, then the settings, then the
  config;
- /props reports supports_preserve_reasoning when the template keeps it on request;
- the web app sends reasoning_content back on every turn.

Off by default: nothing changes unless it is set.  Checked with the Ornith template on the mock engine: a
3-message history with earlier reasoning is 33 prompt tokens without it, 114 with it, on both APIs.

Based on bizhonggeng's patch in Niko1221#1402 (OpenAI path), extended to the Anthropic API and the config.

Co-authored-by: bizhonggeng <30041982+bizhonggeng@users.noreply.github.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The stock image encoder (tools/vision) reads Qwen3.6's and Ornith's mmproj as is, and the engine's --vision
path needed no change (moondrop12138, Niko1221#1402).  Setup now asks about images for these families as it does for
Flash-Next (default no; --vision yes|gpu|cpu|no answers it):

- Qwen3.6: bartowski's mmproj-Qwen_Qwen3.6-35B-A3B-bf16.gguf (Unsloth's is named mmproj-BF16.gguf), its repo
  pinned like the others; Ornith: bartowski's mmproj-Ornith-1.5-35B-A3B-bf16.gguf;
- the KV-streaming advice counts the encoder's VRAM (1,254 MiB measured, plus the 700 MiB reserve) out of the
  free VRAM it decides on;
- the "--vision cpu reads prompts ~3x faster" tip (Niko1221#469, Flash-Next) is not given for these models: Ornith read
  a 24K-token prompt at ~4,130 tok/s with or without the GPU encoder.

Measured on an RTX 4070 Ti, 32K context: Ornith IQ4_XS and KAT i-quality-v2 (with Qwen3.6's mmproj) described a
test picture right over both APIs; with the GPU encoder Ornith wrote 135 vs 143 tok/s (mean of 12 replies).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@aflin

aflin commented Oct 11, 2026

Copy link
Copy Markdown
Author

Thanks, both of you. Both additions are in the branch now (qwen36-35b-support 4a0bb25 and 70c5d3e, merged into qwen35moe-more-models).

@bizhonggeng: keeping earlier reasoning. Your patch worked as described on the OpenAI API. Claude Code and other Anthropic-API clients didn't get it, so I extended it (you're credited as co-author):

  • chat_template_kwargs.preserve_thinking is read on both APIs. On the Anthropic API it's a Strata extension, so earlier thinking blocks can stay in the prompt.
  • "preserve_thinking": true in strata-<model>.json turns it on for every request, like llama.cpp's --reasoning-preserve. The request's own value wins, then the shared Chat settings, then the config.
  • It's off by default; nothing changes unless it's set.

Checked with Ornith's template: a 3-message history with earlier reasoning is 33 prompt tokens without it and 114 with it, on both APIs.

@vehystrix you're right that this is the template's choice. With the froggeric template it was already on for you, so it doesn't explain the difference you saw. It matters for the default Qwen3.6, Ornith and KAT templates, which leave the reasoning out unless asked. Glad the prompt speed is better.

@moondrop12138: images. Confirmed here, and setup now offers images for Qwen3.6, Ornith, bartowski's Qwen3.6 and KAT-Coder. Setup asks, the default is no, and --vision yes answers it.

  • Ornith uses bartowski's Ornith mmproj. The other three use bartowski's Qwen3.6 mmproj (Myric publishes none for KAT, and Unsloth's file is just mmproj-BF16.gguf).
  • RTX 4070 Ti, 32K context:
    • Ornith IQ4_XS (GPU and CPU encoder) and KAT i-quality-v2 described a test picture correctly (shapes, colors, positions, caption) over both APIs, at about 288 prompt tokens.
    • A 1920x1200 photo took 1,019 tokens and 6.2 s including the answer.
  • Cost of the encoder on the GPU:
    • It holds 1,254 MiB of VRAM.
    • Generation was about 5% slower: 135 vs 143 tokens/s, the mean of 12 replies.
    • Prompt reading was the same, ~4,130 tokens/s on a 24K-token prompt. So setup doesn't give these models the Flash-Next tip to move the encoder to the CPU.
  • Setup's KV-streaming advice now counts the encoder's VRAM, which matters on 8 GB cards.
  • If you set up without images, ./setup.sh --setup --vision yes adds them: it downloads only the mmproj and compiles only the encoder.

Not tested: Qwen3.6's own files with images, AMD.

Let me know if there are problems.

Conflicts: the route tail skip's default (src/program/generate.cpp, docs/DETAILS.md).  Upstream's new rule is kept
(off from 90% of the experts cached, on for HIP gfx1200/gfx1201) with this branch's exception: qwen35moe keeps it
off unless STRATA_ROUTE_TAIL_SKIP=7.

Also: 0.1.43's adaptive overlap (on by default on Linux CUDA) is not the default for qwen35moe.  Ornith IQ4_XS on an
RTX 4070 Ti, 6 prompts x 500 tokens, 2 starts: 128 tok/s with it, 144 without (the same as before the merge);
STRATA_ADAPT_OVERLAP=1 turns it on.  QWEN36_ENGINE follows the source version (0.1.43).

Built: CUDA sm_89, HIP gfx1100/gfx1201/gfx1030/gfx1151 and gfx906, SYCL.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@aflin

aflin commented Oct 11, 2026

Copy link
Copy Markdown
Author

Merged upstream 0.1.43. The two conflicts were both the route tail skip's default; qwen35moe still keeps it off unless STRATA_ROUTE_TAIL_SKIP=7.

One change beyond the conflicts: 0.1.43 turns adaptive overlap on by default on Linux CUDA, and on qwen35moe that was slower. Ornith IQ4_XS on an RTX 4070 Ti, the same 6 prompts of 500 tokens over 2 starts: 128 tok/s with it, 144 without. So it stays off by default for qwen35moe; STRATA_ADAPT_OVERLAP=1 turns it on. Flash-Next keeps upstream's default.

Built for CUDA, HIP (gfx1100/1201/1030/1151, gfx906) and SYCL.

@vehystrix

Copy link
Copy Markdown

I've been using it for a while now today. Getting it to properly use the edit tool still fails, I'll tell it not to rewrite the entire file and it'll make a python script to rewrite parts of it instead...
Anyone know if Ornith is better?

What I also found:
Using --kv-resident 32768 slows down decode massively when context grows. At ~150k context I was down to ~15t/s,
setting the value back to 0 and resuming the session got the speeds back to ~50t/s

@bizhonggeng

bizhonggeng commented Oct 12, 2026 •

Copy link
Copy Markdown

Update — my earlier i-quant hypothesis does not hold up; it looks like the engine version (and the sampling default), not the expert format.

Some follow-up of my own, on the latest head (cd1daa4, 0.1.43):

What I ran. UD-IQ3_S (i-quant) and UD-Q3_K_M (K-quant), Qwen3.6-35B-A3B, single-turn long context at ~65k, same prompt, same harness (pi code).
Sampling explicitly set to temperature 0.6, top_p 0.95, top_k 20.

Result. Neither model looped. No length truncation, no repetition. The thinking stays short and the turns stop normally.

Why this matters for the hypothesis. The one common factor behind the loops I reported earlier was not i-quant — it was greedy decoding. Strata's default for an API client that sends no sampling is temperature 0 (serve/server.py:596; the engine spells it out: "Absent keys keep today's behavior: greedy, no penalties"), and setup writes no sampling block, so a code agent that sets nothing runs greedy. With greedy, both i-quant and K-quant loop for me; with temperature 0.6, neither does. That is the opposite of "i-quant is the fragile format".

What I cannot separate yet. I changed the engine version and the sampling at the same time, so I cannot tell how much is the version. A clean disentangle would be: latest engine + greedy (should still loop) vs the previous engine + temperature 0.6. If the old engine with sampling is clean, then this is purely sampling; if it still loops, the version fixed something else as well. I'd run that, but I'd rather not over-claim before doing it.

On @vehystrix's report just above (same sampling settings, still degraded): those look like two different things — the edit tool failing on its own anchored-edit input format, and prefill dropping from >2,000 to ~150 t/s past 100k — neither is the repetition loop, and that run was on an older engine. If anything, that pushes me toward "the newer engine addressed the loop", not "it is all sampling".

Either way, one thing seems worth fixing upstream: the default is greedy for any client that sends no sampling, and setup writes no sampling block. Every code agent that connects and sets nothing lands there. A sampling block in the generated config (or a note at first start) would avoid the whole class of reports.


_

Subject: Long-context loops / tool-call failures — common factor across reports, and a fast test

Hi @aflin,

Following the long-context quality discussion here (especially @vehystrix's repetition loops and failing tool calls). This is a hypothesis with a proposed experiment, not a confirmed root cause — but the pattern across independent reports is consistent enough to be worth a look, and the test below is quick to run.

The common factor: every report involves i-quant

Reporter Model / quant i-quant?
@vehystrix (this PR) KAT-Coder-V2.5-Dev MTP-APEX-i-quality-v2 Yes — APEX I-Quality puts IQ4_XS on routed experts for L10–29
@vehystrix (this PR) Ornith-1.5-35B-A3B Yes — setup's ornith family ships IQ4_XS / IQ3_XXS
Several users Qwen3.6-35B-A3B UD-IQ4_XS Yes
lybhb8 (Niko1221/Strata#438) Qwen3.8-Flash-Next GSQ-RCO-IQ1_M Yes
"Tested in Coding", 12h run (r/LocalLLaMA) Qwen3.8-Flash-Next UD-IQ4_XS Yes

No exception so far. The two Flash-Next cases differ from the qwen35moe ones in model, hardware and engine code path — yet the symptom is identical. That is what made me look for a shared ingredient rather than a shared code path.

Also worth noting: both quant sizes the qwen36 family offers in setup are i-quant (UD-IQ4_XS 18.2 GB, UD-IQ3_S 15.3 GB). So the small-card users this PR targets are all on i-quant by default.

Why the format, and why only at long context

Two lines of evidence:

  1. Niko1221/Strata#438 already falsified "which device computes the expert" by ablation. Sweeping --expert-cache 1,500 → 9,019 slots (hit rate 62% → 97%, essentially all experts on GPU) still degenerated at every size. It also ruled out suffix drafting, --spec 2/4, KV type and repetition penalties. So your own note — "experts computed on the GPU round slightly differently from the CPU's" — cannot be the whole story: at a 97% hit rate there is almost no CPU-computed expert left.
  2. The IQ dequantization path is structurally more fragile. k-quant is linear block scaling; i-quant resolves values through a non-linear codebook. ggml's own ggml/src/ggml-cuda/mmvq.cu in ggml-org/llama.cpp says Q3_K / IQ2_ / IQ3_ "regress due to register pressure and lookup table contention." A table-driven path has far more implementation freedom than a multiply-add, which is exactly where an independently written CPU kernel and CUDA kernel can drift apart in the last bits — and Strata maintains both (iq_avx512.cpp / iq_kernels.cu). The APEX technical report (mudler/apex-quant) measures the same from the quality side: "IQ formats underperform K-quants on MoE experts … the near-Gaussian expert weight distribution (kurtosis 3.41) suits the K-quant block structure better."

Per-step the error is tiny — your observation is "llama.cpp's top two at 28% and 26%", a near-tie. It only matters because greedy is the worst case for it: at temperature 0 a near-tie is exactly where a last-bit difference flips the argmax, and one flipped token re-bases everything after it. Over thousands of steps some fraction of those flips land in a repetition attractor.

This predicts something counter-intuitive and quick to check: raising temperature should reduce loop frequency. If it does not, drop this hypothesis.

The test

Unsloth's Qwen3.6-35B-A3B-MTP-GGUF publishes both families for the same model:

  • UD-IQ4_XS (18.2 GB) — IQ, current default
  • UD-Q3_K_M (17.1 GB) — K-quant, and smaller than the IQ4_XS
  • UD-Q4_K_XL (22.9 GB) — K-quant

Same model, same repo, same prompt, same context, only the expert format changes.

One thing worth checking first, at almost no cost: you reported earlier in this PR that UD-Q4_K_XL "matches llama.cpp on the same file" (~84 tok/s on a 12 GB card). Do those runs show any looping or tool-call degradation? If K-quant has been clean all along while i-quant reports keep arriving, that settles the direction without a new experiment.

Caveat, from lybhb8 in Niko1221/Strata#438: "one --kv fp16 run came back clean and looked like the root cause. Re-running that exact config gave 3 clean / 1 degenerate over 4 runs… The clean run was luck." Any comparison here needs ≥5 repetitions per config, ideally with fixed expert cache and --adapt-every 0.

To be clear about what I am not claiming: Strata's i-quant kernels may be perfectly correct, and this may be entirely the format's intrinsic error floor. The test separates those cases — if the K-quant build is clean and both engines agree on the IQ4_XS file, it is the format; if Strata and llama.cpp disagree on the same IQ4_XS file at long context, it is the kernels.
I am going to run this comparison on my side as well.

Compiled with the assistance of an AI model (hy4 preview).

_

@aflin

aflin commented Oct 15, 2026

Copy link
Copy Markdown
Author

Closing this pull request. Thank you all for the testing, reports and suggestions here, and for the patience while it grew: @vehystrix, @ohmgloss, @bizhonggeng, @misthema and @moondrop12138. Thanks also to @Asofwar, whose decode-path and router speedups I've taken. And thanks to @Niko1221 for Strata itself.

I've moved this work to a personal fork called MoMoE ("More MoEs"): https://github.com/aflin/momoe. It isn't meant to compete with Strata. It's my own variant, shaped around how I use it, and almost everything that makes it work is still Strata's. The README credits the project and the people whose changes it uses.

What's different:

  • Platforms: Linux with NVIDIA (CUDA), plus an early Metal engine for Apple Silicon Macs. It runs Qwen3.6 and Ornith-1.5 well, and Flash-Next on Macs with enough memory. Flash-Next is slow there, though: on my MacBook Air (M4, 32 GB) it does about 4–7 tokens/s. It runs, which llama.cpp can't manage on that machine, but it's too slow for agent work.
  • Models: Qwen3.6-35B-A3B and Ornith-1.5 alongside the Flash-Next family, from this PR.
  • No web interface: it's just the API server. momoe setup writes a named config, and momoe start <config> starts the server. Setup compiles the engine from source; I'm not providing binary releases (yet).
  • momoe run claude | codex | opencode | unsloth | lmstudio: starts the server if needed and points the tool at it. It never installs anything; if a tool is missing, it tells you how to install it.

Let me know if there are problems, and good luck with Strata.

@Asofwar

Asofwar commented Oct 20, 2026

Copy link
Copy Markdown

Closing this pull request. Thank you all for the testing, reports and suggestions here, and for the patience while it grew: @vehystrix, @ohmgloss, @bizhonggeng, @misthema and @moondrop12138. Thanks also to @Asofwar, whose decode-path and router speedups I've taken. And thanks to @Niko1221 for Strata itself.

I've moved this work to a personal fork called MoMoE ("More MoEs"): https://github.com/aflin/momoe. It isn't meant to compete with Strata. It's my own variant, shaped around how I use it, and almost everything that makes it work is still Strata's. The README credits the project and the people whose changes it uses.

What's different:

  • Platforms: Linux with NVIDIA (CUDA), plus an early Metal engine for Apple Silicon Macs. It runs Qwen3.6 and Ornith-1.5 well, and Flash-Next on Macs with enough memory. Flash-Next is slow there, though: on my MacBook Air (M4, 32 GB) it does about 4–7 tokens/s. It runs, which llama.cpp can't manage on that machine, but it's too slow for agent work.
  • Models: Qwen3.6-35B-A3B and Ornith-1.5 alongside the Flash-Next family, from this PR.
  • No web interface: it's just the API server. momoe setup writes a named config, and momoe start <config> starts the server. Setup compiles the engine from source; I'm not providing binary releases (yet).
  • momoe run claude | codex | opencode | unsloth | lmstudio: starts the server if needed and points the tool at it. It never installs anything; if a tool is missing, it tells you how to install it.

Let me know if there are problems, and good luck with Strata.

u are welcome)

@aflin
aflin deleted the qwen36-35b-support branch November 6, 2026 05:05
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

7 participants