Repository navigation
Conversation
… 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>
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>
- 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>
…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>
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>
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>
…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>
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>
…main 0.1.43.1 Conflicts in the prompt path and the prompt-path init: kept main's pinned-ring copies (Niko1221#1996) with the branch's names (NE_, KK).
…one residual stream) The verify path called cvec_apply with the hyper-connection write (2*sigmoid(inj/hc)) for every model; a one-stream model never writes inj2_, so a steered qwen35moe layer would be added with a garbage gate. Add bo_ first, then apply the vector, as the prompt path already does, and stop the steered range one layer before the head on one-stream models.
…e attention and the 256 x 8 router - expert_source: a STRATA_ROUTE_TAIL_SKIP row is cleared with the model's width (HN), not Flash-Next's 2560: on qwen35moe the memset ran into the next token's top expert row and, on the last row of a full window, past the end of h_ymiss_. - qsa_decode_attn: the merge walks only the chunks that hold selected cells (step[kStepWidth]) instead of max_context / 64 partials per token; skipped chunks were -FLT_MAX before, so the result is the same bits. - verify / mtp: routers other than the native 512 x 10 (qwen35moe's 256 x 8) route a window in one multi-column projection and one top-k launch (moe_route_window, already bitwise the per-token calls), on CUDA too. - verify: dense attention writes its identity ids once per window, at the stage's first attention layer, not on every attention layer. - kv_stream: with dense ids the resolve walks blocks (width / page_size) instead of every cell. - layout: qwen35moe requires exactly 16 DeltaNet key heads, the count the GDN kernels are built for.
router_top10_fast_kernel (the parallel sum with the serial fallback, the selection in registers) was HIP-only, but its argument is IEEE arithmetic and its intrinsics are CUDA's. It is now router_top10's default on CUDA for 64 < n_expert <= 512, k <= 32 (STRATA_ROUTER_OLD=1: the portable kernel). tests/cuda/router_fast.cpp runs the HIP test on the CUDA runtime: 0 rows differ on 65,536 rows of 512 x 10 and 8,192 of 256 x 8, the serial path included; 15.5 -> 8.9 us per 256 x 8 call on an RTX 5080. Ornith-1.5 decode, strata_bench short chat: 235.7 -> 246.7 tok/s.
…ATTN_SPLITS) The decode attention keeps one partial (acc, max, sum) per 64-cell chunk, so with dense attention (qwen35moe) the partials grow with the context: at 262144 tokens and 16 query heads one query needs 67.6 MB, and the verifier keeps one per window row (393 MiB reserved before the expert cache on an RTX 5080 at --spec 4). STRATA_ATTN_SPLITS=P (16..65535, default 0 = off, the old path unchanged) caps the partials at P per KV head once a capacity exceeds P chunks: partition b reads chunks b, b + P, ... with the chunk kernel's own per-chunk math and folds them into a running max / sum / acc; the merge then walks P partials. qsa_decode_attn_scratch_floats() returns the smaller size, so every caller (verify, mtp, prefill, layer, the expert-cache reserve) shrinks with it. While a selection fits P x 64 cells every partition holds at most one chunk and the output is the default's bit for bit; past that the fold rounds differently. New test attn_splits_parity (FP16 and INT8 pools, 8 and 12 query heads per KV head, step and batched forms, P = 64): 16 of 16 pass, bitwise at 4,091 of 36,901 cells, max relative difference 1.3e-6 at 36,901 cells; kv_hybrid_parity and kv_stream_parity still pass. Built for CUDA sm_120 only (no HIP/SYCL build); decode speed and the VRAM it frees in the engine are not measured yet.
…n an RTX 5080 (Ornith 1.5 IQ4_XS, 262K context)
Owner
Author
|
Всё содержимое этого PR вошло в main (порт на upstream 0.2.1.1, коммит a8e2954). Закрываю. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
qwen35moe: Qwen3.6-35B-A3B and Ornith-1.5-35B-A3B) into upstreammain0.1.43.1. Conflicts insrc/prefill/prefill.cpp(pinned-ring uploads from 0.1.43 fixes the #1712 reproduction, but Memory access fault after 31 min under multi-slot load with image requests (2x RX 7900 XTX) Niko1221/Strata#1996 kept, with the PR's runtimeNE_/KKwidths) andsrc/program/generate.cpp(mem_markkept before the optional prompt path).qwen35moe: a control vector in the decode path is applied as a plain add on the one residual stream (the hyper-connection write readinj2_, which this model never writes); the steered range stops one layer before the head on one-stream models.Tested