cortiq-core and cortiq-engine 0.8.68 reached crates.io from a publish stopped halfway; 0.8.69 is the release.

Added

• CUDA: Qwen3.5-family hybrids on the card (qwen3_5: GDN mixers, gated attention with Q/K norms, an MTP block;
  Qwen3.8-27B), with drafts and their verify there too (docs/CUDA.md).

  • Decode and verify run through one tensor-core kernel (mma.sync, bf16 operands, f32 sums) over q4tp weights
    repacked on the card. Its arithmetic does not depend on the window, so the text with drafts is the text
    without them. A verify picks its rows' argmax on the card.
  • A token or a verify window is one CUDA graph. Gate and up are one matrix with SiLU in the epilogue. A GDN
    layer is two launches; its delta rule gives a thread a column of the state (a coalesced 3 MB a layer), and a
    partial accept goes back in two launches for all layers. K/V are f16.
  • A round's MTP drafts are one CUDA graph, the token embeddings gathered and the argmax taken on the card.
    Prompt lookup runs on CUDA in windows of 8 and 16 rows. Rounds that land two tokens speculate without timing
    plain ones first.
  • Kernels are cached as PTX on disk (CMF_CUDA_CACHE), and the repack runs on the card: the 27B plan builds in
    under 4 s.
  • A turbo file's far field runs on the arm too, without drafts past the seal.
  • Qwen3.8-27B q4tp on an RTX 3090, greedy, 512 tokens, against the Vulkan path on the same card:
    • chat with drafts 119.7 tok/s on average over eight prompts (57.7);
    • without drafts 49.1 (40.8);
    • a copy task 201.0 (74.8), an edit task 220.1 (71.6);
    • prompt 983 tok/s at 4 000 tokens (560).

• CUDA hybrids: a chat's next turn goes back to a GDN snapshot the card took at the last prompt's assistant header
  (its last <|im_start|>), instead of reading the conversation again. A Qwen3.5 template drops the think block
  from past answers, so a turn never extended the cached ids, and MTP models had prefix reuse off. The draft head
  keeps its rows before that point too. Three turns over a 10 KB file: turn 2 read 338 new tokens in 0.50 s
  instead of 4 048 in 4.13 s, turn 3 78 in 0.24 s instead of 4.29 s; the same texts. CMF_KV_REUSE=0 reads every
  turn from the start.

• CUDA hybrids: the DFlash 2 drafter (z-lab/Qwen3.8-27B-DFlash2, a .dflash.cmf sidecar, CMF_DFLASH=1) on
  the card. Its block's four-bit matrices run on the arm's tensor-core matvec, and the rest in f16 on cuBLAS
  and in kernels of its own. Its context is made on the card from the trunk's tapped rows, which the graphs
  copy out on their way. Each round goes to the drafter that lands more at the card's prices, read off the
  sidecar's router when it has one (cortiq dflash-route-fit --prices now writes one). Copies and text from
  the high part of the vocabulary stay with the MTP chain. Qwen3.8-27B q4tp on an RTX 3090: chat 128.4 tok/s
  against the chain's 120.2, the same texts; a copy and an edit 8-10 % slower, the first token 0.5 s later.

• CUDA hybrids: DFlash 2 on a trunk without an MTP head (Arete-27B Turbo). Sidecars with a router fitted at an
  RTX 3090's prices are on Hugging Face beside
  Qwen3.8-27B,
  Arete-27B Turbo and
  OrcaCyber-Nano-1.0 Turbo. Arete-27B Turbo on an
  RTX 3090: 49.0 → 70.3 tok/s on a short answer, 49.2 → 102.1 on code; OrcaCyber: 116.8 tok/s on code with
  DFlash, 111.4 with its MTP head; the same texts.

• CUDA: the Fast Code stack (qwen4_exp, Qwen3.8-Flash-Next) on the card: hyper-connection mixes, GDN, QSA, PLE
  and MoE, the experts read from an arena with a layer-order LRU. Speculation rounds roll the GDN state, the conv
  ring and the PLE history back exactly; drafts come from a copy cursor, copy frames and the file's MTP head. RTX
  5090, fcaq256: an edit that copies the prompt 497–548 tok/s (151 over Vulkan), a module written from scratch
  94–111, the same text with and without drafts (#107).

• CUDA, Fast Code stack: a masked route (CMF_QWEN_SELF_PACK=1 CMF_QWEN_MASK=1) refills the arena from the
  prompt's own routing and picks only among the experts it holds, so decoding uploads nothing: a module written
  from scratch 94–111 → 135–160 tok/s on an RTX 5090, HumanEval-164 148/164 against 151–153 unmasked. Decode
  steps and windows of a masked route build their expert units on the card (36.4 → 34.5 ms a verify window).
  q2tp experts and 8-bit tubes (Qwen3.8-Flash-Next-hot-cmf) read on the card, exact against the host path.
  The experts a refill chose are evicted last; a layer's n-gram rows are gathered when the pass reaches it.
  CMF_QWEN_ROUND_LOG=1 times every round, CMF_QWEN_NCU=<n> times every launch of round n (#108).

Fixed

• CUDA: DFlash on a trunk without an MTP head stopped after its first partial accept ("GDN state of layer 0 is at
  row 46, the call at 43"). The round now takes the GDN states back to the accepted row.
• cortiq serve no longer holds the first text of a reply until it has 100 characters (#107).
• cortiq serve with enable_thinking: false decoded without drafts. The <think> token it suppresses counted
  as a penalty in the draft gate. On CUDA a suppress list alone now keeps the drafts, verified on the suppressed
  rows: 44.8 → 82–151 tok/s for Qwen3.8-27B at the time, the same text.

Added

• A native CUDA backend for NVIDIA cards, experimental and off by default: build with --features cuda, select
  with CMF_GPU=cuda (docs/CUDA.md). It runs dense stacks with full or sliding-window attention
  (the Spark-X2.x class, q4tp / q8_2f / f32 weights): decode keeps the weights in their codec, the prompt expands
  them to f16 planes for cuBLAS on the tensor cores. cudarc for the driver API, kernels in CUDA C compiled by
  NVRTC on first use; nothing is linked, and a machine without the libraries falls back to the CPU. Spark-X2.6-4B
  q4tp on an A40 against the Vulkan path: prompt of 1 000 / 4 000 / 16 000 tokens 6132 / 6096 / 4808 tok/s
  (3177 / 3789 / 2949), decode 117 / 112 / 93 tok/s (83 / 80 / 54); logits within 0.02 of the host's with float
  activations, the same greedy ids as Vulkan on six prompts for both files. A turbo file's
  far field runs on it too: a prompt past the seal is read whole on the card and sealed once, and from there a
  token reads the far-field state the card keeps (two more kernels a layer; the near rows go through the exact
  attention's). A 64 000-token prompt: first token in 42.3 s, then 100.7 tok/s in 3.9 GB of the card, against
  60.3 s and 85.9 tok/s over Vulkan; the kernels against the host's from one seal, logits within 0.0004; facts
  planted in 240 000 characters found 8 times of 8. CMF_O1_WHOLE=1 takes the wgpu path the same way. Other
  architectures are not ported: under CMF_GPU=cuda they run on the CPU. cortiq gpu reports the card.

Changed

• The far field's seal takes a few layers at a time, each on its groups' threads (CMF_O1_SEAL_PAR=N; one after
  another on a Mac): 9 layers of 32 000 rows of Spark-X2.6-4B on 7.65 CPUs seal in 2.7 s instead of 4.1.

Fixed

• wgpu: the first token after a prompt of a q8 / q8_2f file came from the hybrid head, whose CPU half works with
  integer activations, so its logits sat up to 0.30 off the host's with float activations (0.015 at the later
  tokens) and a near tie could go the other way. The first token now takes the token graph's head pass, as every
  later token does: 0.014 on Spark-X2.6-4B q8_2f, the same greedy ids as the CUDA path on six prompts, prompt time
  and memory unchanged. CMF_HEAD_GRAPH=0 restores the old head.

Added

• The MNN-Pocket-0.3B family (taobao-mnn): the assistant, its base and SFT checkpoints, the DFlash2 draft and the
  decision model as CMF files (infosave/MNN-Pocket-0.3B-cmf). The q4tp file is rounded from the authors' INT4
  (QAT) weights; the files carry the authors' sampling and a loop stop.
• cortiq decide-pointer MODEL --input request.json: typed decisions (choice / noul / score) with a pointer-head
  decision model (MNN-Pocket-0.3B-Dec): the authors' request serializer, one causal pass over the state and each
  question's branch (the reference's branch mask), the float32 pointer head and calibrated temperature carried as
  decision.pointer.*. Dec-f16 picks the reference's answer on 49/49 questions, probabilities within 0.001.
• DFlash 2 sidecars from SpecForge configs (is_causal absent, block_size at the top level) and for tied bases.

Changed

• Metal: the four-row q8_2f matvec decodes every plain dense plan (it was Spark-X2.5's alone; CMF_Q8_R4=0 closes
  it), and the decode attend runs the split-K kernel in blocks of 16 / 32 / 64 from 2 to 1 536 positions
  (CMF_GQA_SHORT=0 restores the per-head / per-key kernels). M4: MNN-Pocket-0.3B q8_2f 162 → 202 tok/s (the CPU
  154–161), q4tp 214 → 252; token time 0.81x on MiniCPM5-2B, 0.78x on Minico-M1; text unchanged.
• wgpu: the batch graph's q8_2f projections of 1–8 rows and the DFlash block's run the decode kernel batched
  (q8_2f_matvec_bu; each row the decode's to the bit, CMF_Q82_UNR=0 keeps the old kernel): six rows of a 0.3B
  model's FFN 1.6–2x faster. The DFlash block runs on the device for hidden sizes that are multiples of 64 and
  four-bit projections narrower than 2 048 columns. The solo drafter's pay-off is judged on summed round time over
  summed tokens from its third round. MNN-Pocket-0.3B q8_2f + DFlash2 on an RTX 2000 Ada: Python 320 → 415 tok/s,
  text identical to plain decoding.
• The converter no longer writes a tied lm_head.weight that is a byte copy of the embedding (the embedding serves
  the head): MNN-Pocket-0.3B q8_2f 375 → 316 MB.

Fixed

• Streaming output (run, serve SSE, the network split path) no longer prints U+FFFD where a byte-level token ends
  inside a character: the bytes wait for the token that completes it (44 of 192 MNN-Pocket answers, Chinese text
  and subscripts).

Added

• cortiq moe-defrag --keep N cuts a MoE to the N experts a layer that a routing profile (CMF_QWEN_PROFILE_SAVE)
  routes to most, writes the profile in the new numbering, and --merge appends the MTP sidecar, so the cut is one
  file. The cut records the checkpoint expert behind every kept slot (provenance.moe_defrag.kept_ids); hot-bake --experts reads a cut's slots from the bf16 checkpoint by those numbers (--defrag-stats rebuilds the list for
  older cuts). Qwen3.8-Flash-Next Fast Code (256 experts a layer by code and front-end routing): HumanEval 93.3 %,
  prompt 827–844 tok/s and generation 136–166 tok/s on an RTX PRO 6000.
• cortiq import-gguf --template maps a qwen4exp GGUF (Strata Coder) onto the tensors of an existing CMF, each
  tensor's reading chosen by cosine against the template; IQ2_S and IQ3_XXS are dequantized as gguf-py does.
• Expert packs for a Flash-Next-class MoE on a card that cannot hold every expert (CMF_QWEN_PACKS=<dir>): one
  routing profile per task, chosen per request by a patent-1 resonance router over the model's own routing of the
  prompt (cortiq qwen-packs fits it); after the prompt the arena is refilled with that task's experts in 10–40 ms.
  CMF_QWEN_MASK=1 keeps decode inside the arena (opt-in; it costs answer quality).
• Qwen3.5 GDN hybrids on wgpu: f16 K/V mirrors (CMF_GPU_KV16=0 keeps f32), q4tp prompt GEMMs and causal flash
  attention on the matrix units, position-parallel GDN conv and norm, a staged GDN scan, deferred chunk readback,
  flash-decoding with f16 twins. Qwen3.5-4B q4tp on an RTX A5000: prompt 800 → 2900–3260 tok/s at 1K and 607 → 3871
  at 16K, decode 116 → 128 and 87 → 111, level with llama.cpp's Vulkan backend.

Changed

• qwen4 (Flash-Next) on wgpu: 256-token prompt frames with the experts, HC mixes, router and GDN gates on the matrix
  units; a decode token in four chains encoded while the card runs the previous one; radix top-k in the indexer and
  index blocks to 65 536 (exact attention to 262 144 positions). RTX PRO 6000: prompt 551 → 887 tok/s at 4K, generation
  with MTP 122 → 145.
• On a card that cannot hold every qwen4 expert the prompt streams layer by layer (CMF_QWEN_STREAM=0 restores
  per-frame admission): each layer's experts are admitted once per 4096-token chunk (2048 below 11 GiB). The host tier
  of pinned memory the arena copies from by DMA is on by default there. Fast Code (q4tp, 41 GB whole) now runs on
  12–24 GB cards (it could not stream before); cards emulated on an RTX 5090, prompt at 4K / generation: 12 GB 240 /
  30–33 tok/s, 16 GB 389 / 48, 24 GB 768 / 65–72.
• A loaded routing profile (CMF_QWEN_PROFILE) no longer keeps every prompt frame reading its routing back; only
  CMF_QWEN_PROFILE_SAVE does. Resident prompt ×1.7–2.4.
• Prompt frames past 4096 expert winners sort them by counting; CMF_QWEN_Q82COOP=1 runs the q8_2f prompt
  projections on the matrix units.
• cortiq bench warms the prompt with a different text of the same length right before the timed run (the old
  warmup left a reusable prefix and ran before the pair micro-bench).

Fixed

• The converter writes no MTP header when the config declares MTP but the checkpoint ships no MTP tensors.
• f16 K/V mirror capacity rounds to 64: the 64-row tile was refused past 2048 positions.
• The embedded MTP head of qwen4 opens its own model handle (it shared the backbone's arena uid);
  rollback_prefix_len and graph_gdn_layer_count guard qwen4's empty generic layers (a panic in run/serve).
• qwen4 prompt frames run n_experts + 1 expert segments: a frame whose winners took every expert of a cut plus the
  shared one would have skipped its last segment.

Added

• cortiq serve prints each request's speed as llama-server does (prompt eval time, eval time, total time,
  with the prompt tokens reused from the cached prefix) and returns llama-server's timings object: in a
  non-streamed response, and on the usage and finish chunks of a stream.
• A Spark-X2.6-4B Turbo file (far-turbo with scripts/spark26/keep_spark26.bin, exact-row weights fitted for
  its 9 full-attention layers). RTX 2000 Ada, Vulkan, 64 000 tokens: decode 56.2 tok/s against 18.5 and 4.2 GB
  of the card against 7.9.
• cortiq ppl --window-score-from K scores a window's predictions from position K on, the earlier ones
  being context only: with --window-len 512 --window-stride 512 --window-score-from 256 the walk is that of
  llama-perplexity -c 512, so the two numbers compare.
• cortiq import-gguf reads IQ3_S (the 3.4-bit grid codebook inside imatrix IQ4_XS mixes), a port of ggml's
  dequantize_row_iq3_s, checked against gguf-py on a real block.

Changed

• Spark-X2.6 on native Metal reads its prompt through the rows graph (exact GELU, the per-head g_proj gate,
  sliding windows, each layer's own RoPE table), and with CMF_ANE=1 its FFNs go to the Neural Engine: a
  14 000-token prompt on a MacBook Air M4 in 34 s instead of 54.
• Metal KV mirrors hold f16 for Spark-X2.6 (CMF_METAL_KV_F16=0 keeps f32) and grow by a quarter instead of to
  the next power of two; a token's attend takes a kernel with a thread per key from 64 positions (CMF_GQA_KB=0
  keeps the split-K one): decode after a 14 000-token prompt 24.4 tok/s → 30.8, the same text.
• A turbo file's prompt on Metal: the exact prefix on the chunk graph when the rows graph does not hold the
  model, and a sealed layer's near window as GEMMs for a chunk of 16 rows or more.
• A sealed layer's token on Metal runs four kernels instead of eleven (the near window and the far field in
  one pass each, a merge, the commit; CMF_O1C_TOK=0 keeps the stages): Spark-X2.6-4B Turbo after a
  14 000-token prompt on an M4, 29.1 tok/s → 33.0, the same text.
• A dense q4tp layer's q, k and v projections, RoPE, the K/V append and the per-head output-gate dots go in one
  Metal dispatch, the gate applied inside the O projection (CMF_METAL_QKVFUSE=0 keeps them apart).
• The wgpu batch graph takes the per-head output gate and exact GELU and copies a far field's query trace out of
  its per-layer-geometry arms, so a Spark-X2.6 turbo prompt rides it (16 000 tokens: 565 s → 117 s).
• A dense sliding-window stack's prompt goes through the wgpu batch graph with 512-row submits where the graph's
  whole-run attend is tiled (CMF_DENSE_GRAPH_PREFILL=0 keeps the chunked host walk, CMF_BATCH_K the rows).
  The whole-run attend admits heads of 256 (Spark-X2.6) and takes a tile of 16 tokens against each 64-key block
  — a K/V row read once per tile instead of once per token (gqa_attend_part_xt, CMF_ATTEND_XT=0 keeps the
  per-token kernel) — and on a card with cooperative matrices the tile runs on the matrix units with a group of
  blocks per workgroup and an online softmax (gqa_attend_part_xc, CMF_ATTEND_XC=0 keeps the scalar tile).
  Spark-X2.6-4B q4tp on an RTX 2000 Ada over Vulkan, 16 000 tokens: 182 tok/s walked, 390 tok/s with the scalar
  tile, 582 with the matrix-unit tile, the same greedy text after 9 000 and 60 000 characters. A cooperative
  kernel's workgroup scalars are 16-byte vec4<u32>: a u32 beside the tile arrays shifted every cooperative
  load and store by one element on NVIDIA (tests/gpu_coop_probe.rs keeps the probe that shows it).
• Where the cooperative GEMM is on (NVIDIA by default), a chunk of 128 rows or more runs its wide q4tp GEMMs
  through zi_mm, the 128×128×32 tensor-core tile of the image models: the weight dequantized once into an f16 plane, the chunk's
  activations cast to f16 under the device scale, the result scaled back at the store. 12.0 TFLOP/s on an RTX
  2000 Ada against 6.2 for the 64×64 in-kernel-dequant cooperative kernel (CMF_BATCH_MM=coop; v2 its
  pipelined twin; CMF_BATCH_COOP=0 the scalar tile GEMM at 3.1). Spark-X2.6-4B over Vulkan on that card: a
  4096-token prompt at 1020 tok/s (660 on the 64×64 kernel, 563 on the scalar tile), 16 000 tokens at 878
  (575), the same greedy text. Dequantizing into the plane and running the 64×64 f16 kernel stays opt-in
  (CMF_BATCH_F16=1; it lost to the in-kernel decoder). The 64×64 cooperative GEMMs walk their tiles in
  groups of eight batch tiles so a weight tile leaves VRAM once per group instead of once per batch tile.
  tests/gpu_q4tp_mm_bench.rs times every variant on the Spark shapes and holds each to the scalar GEMM's
  answer. The whole-run attend's partial buffers are allocated once per call instead of once per layer (36 ×
  65 MB in flight was an Out of Memory on a shared 16 GB card).
• The whole-run attend of a full-attention layer with heads of 256 reads its matrix-unit operands straight from
  f16 planes (gqa_attend_part_xc2): the chunk's queries and the cache's keys and values are cast once per
  layer, and each tile loads from storage without the staging copies and barriers of gqa_attend_part_xc
  (CMF_ATTEND_XC2=0 keeps it). Spark-X2.6-4B q4tp over Vulkan on an RTX 2000 Ada: 4096 tokens 1012 → 1088
  tok/s, 16 000 tokens 871 → 901, 64 000 tokens 486 → 535, the same greedy text after 9 000 characters.
• A turbo file's far field is rebuilt with one thread per KV group (it was one group after another): Spark-X2.6-4B
  Turbo at 16 000 tokens, 6.5 s → 3.5 s of the prompt, 489 → 539 tok/s, the same greedy text after 60 000
  characters.
• x86-64 CPU prompts: the batched chunk attend runs on x86 too (an AVX2 micro-GEMM, x86_gemm_rm, and an AVX2
  softmax; a sliding window bounds the key range of the whole chunk; CMF_CPU_ATTEND=0 keeps the per-position
  attend); a q4tp batch takes two weight rows against four columns with FMA (CMF_Q4TP_FMA=0 keeps the vertical
  cell); the activation split and the FFN's activation combine (an exact GELU is one erf an element) run on the
  pool instead of the main thread; the activation panel is sized to half the core's L2; a plain dense stack with no
  GPU reads 128 rows a chunk instead of 48. Spark-X2.6-4B q4tp on 8 Zen 2 cores: a 1000-token prompt 16.7 → 31.6
  tok/s, 4000 tokens 11.9 → 30.1.
• x86-64 CPU decode: every vertical q4tp cell (the decode lanes, the fused gate/up lanes, the row-exact matmat
  cells) accumulates with FMA and takes the 8·Σx term once per row through one shared function, so the matvec
  and the row-exact matmat still agree bit for bit; a grouped-query context of 512 rows or more attends split-K
  (KV group × 256 keys a job, each K and V row read once for all the group's heads, an online-softmax merge;
  CMF_ATTEND_SPLITK=0 keeps one job a head).
• On native Metal the MTP drafts are on by default over a q4tp trunk, as on Vulkan (CMF_GRAPH_SPEC=0 turns them
  off, =1 asks for them over other trunks). The batched verify passed the exact-ID gate there: on a MacBook Air M4,
  Qwen3.8-27B Turbo and OrcaCyber-Nano-1.0 gave the same token ids with drafts and without on five prompts each,
  OrcaCyber at 9.0–14.6 tok/s against 6.1–6.3. (The default's eligibility asked for the wgpu token graph, which
  native Metal does not have; it now checks the block shape Metal's own draft step takes.)
• On a Mac the Neural Engine's surfaces (CMF_ANE=1: a layer's FFN and mixer weights in fp16, two sets, about
  1.5 GB on a 27B) go back once the prompt is read; decode never reads them, and beside a 14 GB weight map on a 24 GB
  MacBook Air they pushed weight pages out. OrcaCyber-Nano-1.0 after 16 000 tokens read with the Neural Engine, drafts
  on: decode 6.4 → 8.2 tok/s (page faults 3.6 M → 31 K, the same text); after 3 600 tokens 7.1–7.5 → 7.7.
  CMF_ANE_KEEP=1 keeps them for a process that reads prompt after prompt.
• On Vulkan/DX12 with an NVIDIA card a q4tp prompt reads through the cooperative-matrix (tensor-core) GEMM by
  default (CMF_BATCH_COOP=0 keeps the scalar tile; =1 asks for it on other vendors), and a 128-row chunk takes
  the new q4tp_mm_coop128_s tile: eight subgroups over one weight tile, so each weight is unpacked once a chunk,
  one lane a row and eight k, with the activations read four floats a load (CMF_COOP128=0 keeps the 64-row tile).
  OrcaCyber-Nano-1.0 q4tp (Qwen3.8-27B class) on an RTX 2000 Ada 16 GB, 2000-token prompt: 74.5 → 112 tok/s, the
  same greedy text; the 128-row tile's sums equal the 64-row kernel's bit for bit (test).

Fixed

• On wgpu a dense model the batched graph declines read its prompt one position at a time (0.8.63): Spark-X2.6
  q4tp on an RTX 2000 Ada, 65 tok/s. It reads chunked on the host walk again (195 tok/s), as sliding-window
  stacks always do now.
• The speculation trial left the drafts after two empty rounds in a row and kept plain decode for 128 tokens. A good
  head meets such a pair now and then; at the start of a chat turn it sent the whole turn to plain decode
  (OrcaCyber-Nano-1.0 on an RTX 2000 Ada, the second of five turns: 13.3 tok/s against 23–26 with the drafts). The
  trial now needs three empty rounds, and a plain decision taken on fewer than five timed rounds is tried again after
  32 tokens. One-shot answers unchanged.
• cortiq import-gguf on a dense qwen35 GGUF with an MTP head (Qwen3.8-27B and its fine-tunes, e.g.
  OrcaCyber-Nano-1.0): llama.cpp counts the nextn block in block_count, and the importer made it a 65th trunk
  layer under the MoE arch name. The trunk is now block_count − nextn_predict_layers, the block lands under
  model.mtp.* with its 1+w norms undone like the trunk's, the header declares the head, and dense qwen35 is
  qwen3_5, as cortiq convert writes it.

Added

• CherryPit-1.2-2B (DRUPEs, a MiniCPM5-2B finetune) in CMF: q8_2f, a chat-calibrated GPTQ q4tp with the attention
  at eight bits, and a Turbo file with exact-row weights fitted for its 42 layers (scripts/cherrypit/). Cards
  docs/hf/CherryPit-1.2-2B-cmf.md, docs/hf/CherryPit-1.2-2B-Turbo-cmf.md.
• Minico-M1-1105 (Mincofficial, a fine-tune of LFM2.5-350M) in CMF:
  three files, q8_2f, q4tp and a turbo file with 248 rows
  kept exact (scripts/minico/keep_minico.bin, fitted with scripts/dory/router/q27_feat.py). MacBook Air M4,
  Metal: q8_2f 158 tok/s, q4tp 276; the turbo file after a 64 000-token prompt 106 tok/s against 38.6.
• run --reasoning-effort LEVEL and the server's reasoning_effort (or chat_template_kwargs.reasoning_effort)
  reach any chat template that reads reasoning_effort, as given. Before, only DeepSeek-V4.1 and MiniMax M3 took
  it. Minico-M1's dial is off, low, medium, high, max.
• cortiq serve returns LFM2's pythonic tool calls, <|tool_call_start|>[name(k='v')]<|tool_call_end|>, as
  message.tool_calls, streamed or not. Before, they came back as text.

Changed

• On native Metal a pure-attention model (Llama, MiniCPM5 and the like) reads its prompt through the batched rows
  graph, as the GDN hybrids already did, instead of the host walk: CherryPit-1.2-2B on a MacBook Air M4, an
  8 192-token prompt in 31 s instead of 58 s (q4tp; q8_2f 34 s instead of 55), 2 048 tokens in 4.5 s instead of
  6.0 (q8_2f 5.5 s instead of 9.2); WikiText-2 ppl through the graph equal to the host's to three decimals.
  CMF_METAL_PREFILL=0 keeps the host walk.
• On Vulkan/DX12 a fully resident pure-attention model reads its prompt through the batched graph, as MoE and GDN
  models already did, instead of the chunked host walk that paid a submit, a wait and a readback per projection
  and per attend. CherryPit-1.2-2B on an RTX 5090: 2 048 tokens in 1.2 s instead of 10.9 (q8_2f; q4tp 1.0 s
  instead of 10.5), 8 192 in 7.2 s instead of 49.6, and decode at that depth 68 → 91 tok/s; the same greedy text
  and WikiText-2 ppl. CMF_DENSE_GRAPH_PREFILL=0 keeps the chunked walk.
• Decode attention on Vulkan/DX12 is a flash-decoding kernel (gpu_wgpu/attend_fd.rs) from 16 cached positions on: a
  workgroup of eight 32-lane subgroups takes a chunk of the KV group, a subgroup reads each K and V row whole and
  coalesced, scores every query head of the group with one subgroupAdd and keeps the online softmax in registers;
  the subgroups meet in workgroup memory and gqa_attend_merge joins the chunks as before. The old split kernel gave
  a model with two KV heads 32 workgroups at 4 096 positions and walked 256 V rows in a row. CherryPit-1.2-2B on an
  RTX 5090, decode tok/s old → new: q4tp 110 → 182 at 512 positions, 113 → 164 at 2 048, 106 → 147 at 4 096,
  92 → 117 at 8 192, 72 → 90 at 32 000; q8_2f 103 → 163, 102 → 142, 98 → 143, 91 → 110; chat answers 110 → 136–164
  (q4tp) and 101 → 135–150 (q8_2f) with the same greedy text. Needs 32-lane subgroups and heads of 64, 128 or 256;
  CMF_ATTEND_FD=0 keeps the old kernels.
• A sealed (far-field) layer decodes on Vulkan/DX12 in two dispatches instead of the seven o1_stages
  (gpu_wgpu/o1_fd.rs): the near rows (sinks, then the ring of the window and the exact rows) through the same
  flash-decoding pass, then one workgroup a query head scores its landmarks, builds the far part as o1s_far does and
  writes (Σ near + Σ far) / (Σ near weights + far mass). CherryPit-1.2-2B Turbo on an RTX 5090, decode past the seal:
  56 → 108 tok/s at 16 000 tokens, 67 → 153 at 32 000. CMF_O1_FD=0 keeps the stages.

Fixed

• cortiq serve with "enable_thinking": false replaced a system prompt the chat template writes itself when the
  conversation has none: the direct-answer directive went in as the system message, so CherryPit-1.2-2B lost its
  identity prompt and answered "Who are you?" as "a MiniCPM series model". The directive is now left out for such
  templates (Tokenizer::template_has_default_system); cortiq run --no-think was not affected.
• The f32 GEMM on the CPU (f16/bf16 files, force_f32 tensors) ran on one core whenever the chunk was shorter than 256
  positions, which every prefill chunk on a Mac is: it now splits the output rows across the pool, bit-identical. The
  f16 CherryPit-1.2-2B scores WikiText-2 (12 × 512 tokens) on an M4 at 785 % CPU instead of 100 %.
• cortiq convert on LFM2 configs written by Transformers 5, which state the tied head as tie_embedding: the file
  did not load ("Tensor not found: lm_head.weight").
• A turbo file kept no row exact where the prompt is read on the host (LFM2 on Metal, whose short convolutions the
  rows graph does not take). Each attention layer sealed itself at the end of the chunk that reached the boundary,
  before the pipeline chose the rows. And the rows past the first seal went into the far field unchosen: without
  the device route the host now keeps the prompt exact up to its last 512 rows and chooses among all of them, as the
  archive does on a card. Facts planted in 100 000 characters, Minico turbo: 4 of 8, the standard file 3 of 8, the
  far field without exact rows 0 of 8.

Added

• Qwen3.8-Flash-Next (qwen4_exp) on wgpu: an O(1) far field for the 12 attention layers, after the Dory pattern
  (crates/cortiq-engine/src/qwen4_o1.rs). A sequence runs exact up to the seal; there the layers' keys, values
  and a strided trace of their queries come to the host, nystrom builds a sink + window + landmark state per KV
  head, the states go up as the layers' O(1) states and the exact caches leave the card. Later frames pack their
  rows dense (q4t_o1_pack), run the engine's o1_batch kernels in place of the indexer and the sparse attention
  and pass through the sigmoid output gate (q4t_o1_gate); a verify window attends without committing and the host
  commits the accepted rows. The GDN layers and the experts do not change. Switched on by the far-turbo header
  (cortiq far-turbo, whose --exact-weights are now optional: without them no row is kept exact) or by
  CMF_QWEN_O1=m,w,sink,seal; CMF_QWEN_O1=off keeps a turbo file exact. RTX 5090, the q2tp-v2 file, a
  6 045-token prompt sealed at 4 096 rows: first token 113 → 105 s, decode 26.5 → 29.3 tok/s with the MTP head;
  WikiText-2 over 8 192 tokens: ppl 2.241 exact, 2.323 sealed at 4 096. On this model the far field is not a speed
  device: its attention is 12 layers × 2 KV heads × 256, exact attention over 16 000 rows is cheap, and
  cortiq bench --core exact against sealed at 8 192 reads 59.6 / 54.6 / 37.7 against 62.3 / 44.6 / 25.8 tok/s of
  decode at 1 000 / 4 000 / 16 000 rows, the first token at 16 000 94 against 127 s. The published Flash-Next files
  stay exact; the far field is there for the attention state that stops growing.
• Hot neurons in the routed experts of a MoE model. cortiq hot-probe --experts MODEL TEXT STAT records Σ(w·act)²
  per layer, expert and neuron on the host path; cortiq hot-bake --experts COLD HF_SRC STAT OUT --frac f zeroes
  the busiest neurons' gate and up rows in place in the q2tp/q4tp experts (rung 0, nibble 8) and writes one q8_2f
  tube per expert (mlp.experts.{e}.{gate,up,down}_proj.tube1.weight) from the bf16 checkpoint, read one layer at
  a time from a directory or the Hub; the shared expert, the experts' layout and the expert cache do not change.
  The host MoE paths add the tubes after the triple; the device keeps them resident in three banks per layer and two
  token-wide kernels (q4t_tube_gu, q4t_tube_dn) add every winner's tube to the frame's MoE output.
  --pick top|err|random ranks by the statistic, by the statistic times the row's two-bit error, or at random
  (the control); the tube's down columns are the cold file's own (GPTQ-compensated) unless --bf16-down.
  Measured on Flash-Next q2tp-v2 (GPTQ-calibrated): the busiest 5 % of every expert (32 of 640 neurons, 29–31 %
  of the energy over 8 192 tokens of code and text) move the code ppl only 4.438 → 4.425 for +6.7 GB — on a
  Hessian-calibrated two-bit file the hot rows are already the best-rendered ones.
• The grain the model routes at, whole at eight bits. hot-bake --shared-q8 writes the shared expert as one q8_2f
  tube (mlp.shared_expert.{gate,up,down}_proj.tube1.weight) beside its own triple, whose gate and up rows go to
  zero, so the device arena keeps the triple's shape in its pinned slot and runs the tube with the route's σ
  (steps 722/723); --whole-experts N does the same for the N busiest experts of every layer (ranked by their
  summed statistic; --pick random is the control), the other experts untouched — the device packs a layer's
  tubes and a slot table sends the kernels to the right entry; the host paths add every tube after the triple.
  Flash-Next q2tp-v2: the shared expert alone 4.438 → 4.241 (+0.3 GB); with the 28 busiest of 512 experts a layer
  (43 % of the layers' Σ(w·act)²) 4.167 at 83.81 GB (+9 %), 83 % of the way to the 97 GB q4tp file (4.111); the 28
  experts alone 4.327, 28 random experts 4.388; WikiText-2 12 × 512: 6.062 → 5.782. The tube banks stay resident
  and are refused when the VRAM budget cannot hold them with 2 GB to spare; on the 32 GB card they come out of the
  expert cache (23.4 → 17.1 GB), and the hot file decodes at 27.8 / 29.9 tok/s against the 2-bit file's 59.6 / 54.6
  at 1 000 / 4 000 rows. Qwen3.8-Flash-Next hot is
  the file; an older cortiq runs it without the tubes.

Fixed

• The per-op GPU MoE block (moe_ffn_gpu, the path a model takes when its device graph is off) dropped the experts'
  tubes (hot-bake --experts) silently; it adds them after the block now, as the CPU paths do.
• --no-think closes the think block after a ChatML <|im_start|>assistant line as well (Ripple-65B's template
  takes no enable_thinking): the model answered after 14 000 tokens of thinking in 0.8.60.

Fixed

• The 0.8.60 binaries for Windows and Android did not build: the Vulkan SM-count query added in 0.8.60 (the grid cap
  of the NVIDIA decode kernels) imported ash, which the engine depends on under Linux alone. The query is now Linux
  only; Windows and Android keep the default grid cap. The crates of 0.8.60 are unaffected.

Added

• Ripple-65B Preview (ricursiveAI) runs on the engine: gated attention, a value residual from the first layer, sandwich
  norms, a clamp on the SwiGLU (ffn_clamp in the arch), 16 full and 64 sliding layers over a vocabulary of 100 352.
  cortiq convert reads the checkpoint; the turbo header (cortiq far-turbo) takes the file with 248 rows kept
  exact (scripts/ripple/keep_ripple.bin), and cortiq hotmix keeps the busiest 5 % of each layer's FFN neurons and
  the K and V projections at eight bits in a two-bit file. The three files
  on an NVIDIA A40 (46 GB), q4tp turbo, 4 000 to 64 000 tokens: 12.0 to 11.8 tok/s at 35.3 to 36.6 GB; perplexity
  against the BF16 checkpoint's 6.940: q8_2f 6.974, q4tp 7.089, q2tp 7.432.
• A DFlash 2 sidecar for it (cortiq convert --model DIR --output MODEL.cmf --dflash-sidecar --quant q4tp), and the
  drafter's block as one device submission on wgpu (gpu_wgpu/dflash_dev.rs: norms, kernel projections, grouped
  convolutions, the int8 projections, a block attention over the context rows kept on the card, the head over the
  draft rows). A round of four drafts costs 85 ms against 32 for a plain token on an RTX PRO 6000 (the block 8 ms,
  the drafter's context 5, the verify of five rows 70; the host block took 87). Ripple-65B q4tp, 400 tokens of code:
  30.2 → 46.6 tok/s with CMF_DFLASH=1 (84 % of the drafts kept); prose keeps a third and the pay-off monitor rests,
  30.5 → 28.8. The block runs its first eight rows (CMF_DFLASH_DEV_ROWS=16 for the trained block: the same drafts
  at twice the price); CMF_DFLASH_DEV=0 keeps the host block, CMF_DFLASH_DEV_CHECK=1 runs both and compares,
  CMF_DFLASH_DEV_TS=1 (with CMF_GPU_TS=1) prints the block's stages, CMF_DFLASH_DEV_MM=coop|r the GEMMs that
  measured slower. The pay-off monitor no longer samples a sequence's first round (it reads the prompt into the
  drafter), weighs a round at a tenth, and picks at most five drafts on its own (the verify of eight rows costs 108 ms
  against 70 for five). The text was the plain run's in every run but one: the int8 verify resolved a near-tie the
  other way at token 252 of a 384-token prose answer under a pinned depth of four; CMF_VERIFY_I8=0 keeps the stream
  bit-exact at 29 tok/s on code. Off unless CMF_DFLASH=1.
• wgpu, q8_2f: a prompt chunk of 4 096 rows or more goes through the cooperative-matrix GEMM with an eight-bit tile
  loader (q82_mm_coop_s), f16 operands and f32 sums, the activation scale from the device. Ripple-65B q8_2f on an
  RTX PRO 6000, a 4 000-token prompt: 43 → 156 tok/s, perplexity 6.973 → 6.974. CMF_Q82_COOP=0 keeps the scalar GEMM.
• wgpu, q8_2f decode: q8_2f_matvec4k splits a row block's columns across the workgroup's four sub-blocks and loads
  the activations once for four rows (down 177 → 163 µs a layer on the 65B). CMF_Q82_MV4K=0 keeps the old kernel,
  CMF_Q82_MV16=1 a sixteen-row one that measured worse.
• The card's own timestamps on the token graph: CMF_GRAPH_TS_ALL=1 stamps every stage of every layer
  (CMF_GPU_TS=2), CMF_GRAPH_DEVICE_SPAN=1 the whole token with two stamps, CMF_GRAPH_PROF=1 the host's part.

Changed

• wgpu decode on NVIDIA, all bit-exact against the previous build: the attention's key loop unrolled by eight and the
  value loop by four (49 → 29 µs a layer), the output projection through the warp-per-row kernel (37 → 31), rope and
  the K/V write fused on gated layers, the sandwich norms with shuffle trees and a second phase in waves of eight
  (30 + 25 → 26 + 17), the narrow-matrix grid capped at four workgroups a multiprocessor (the count read from
  VK_NV_shader_sm_builtins), the value-residual blend as one kernel, and a short first piece of the pipelined
  submission (CMF_GRAPH_FIRST). Ripple-65B q4tp on an RTX PRO 6000 at 4 000 tokens: 23.4 → 29.8 tok/s; q8_2f
  16.2 → 18.3; q2tp 27.8. Variants that measured worse stay as knobs with their numbers beside them: CMF_MV_SPLIT,
  CMF_MV_SG_W, CMF_MV_SG_THIN=1, CMF_MV_SG_KS=1.

Added

• A Mac's Neural Engine for a prompt's chunks, off by default (CMF_ANE=1): it takes every layer's FFN, a GDN mixer's
  z projection and an attention layer's q, the GPU takes the rest of the mixer, and the CPU decodes the next layer's
  weights for it. A chunk of 512 rows goes in two halves a layer apart. Qwen3.8-27B Turbo on a MacBook Air M4, first
  token at 3 600 tokens: 73.9 → 28.3 s, the same text; at 10 000 tokens with --o1-spill 217.6 → 113.5 s with a peak
  footprint of 3.4 GB. Against the GPU's numbers z and q differ by 4e-4 and an FFN by 1.4 %. CMF_ANE_MIX=ffn|z|all
  moves the split, ,int8 stages the mixer's weights as int8.
• With it, the model file is read ahead in blocks while the first chunk runs, and the weight arena is cut at tensors,
  so a command prepares what it reads: a cold file's first token 35.3 → 33.6 s. CMF_METAL_PREFETCH, CMF_METAL_WINDOWS
• CMF_ANE_TRACE=1 prints every wait, ANE call and GPU command of a prompt on one clock, CMF_VERIFY_CUT=1 runs each
  stage of the rows graph as its own command, CMF_ANE_CHECK=1 prints how far the ANE's numbers are from the GPU's.
• Native Metal, drafts over q4tp weights (CMF_GRAPH_SPEC=1): a verify of two or three rows runs on a batched matvec
  (q4tp_matvec_b2, _b3) at 95.7 and 82.4 GB/s, where the narrow GEMM reads 60 to 67 at any batch, and its logits
  equal the plain path's to four decimal places. CMF_METAL_MVS=0 for the GEMM.
• There the draft depth follows how often each draft of the chain holds, in place of a fixed seven, and where the
  answer repeats a stretch of the prompt or of itself the tokens that followed it continue the chain up to seven
  drafts. CMF_GRAPH_SPEC_K=n pins the depth, CMF_SPEC_LOOKUP=0 turns the lookup off.
• Qwen3.8-27B Turbo on a MacBook Air M4, greedy, 6.4 tok/s without drafts, the same text with them: an essay 10.1,
  an explanation in Russian 10.7, a Rust function with tests 12.1, quotes from the prompt 15.0, an edit of the
  prompt's code 20.5 tok/s; after a prompt of 3 600 tokens 6.0 → 8.3. Seven drafts gave the essay 6.2.
• CMF_METAL_CMD_TIMES=1 with CMF_METAL_GPUPROF=1 prints a token-graph command's time in the kernel and on the GPU
  and what Metal holds for the process; the batch_kernel_lab test times a batched q4tp matvec from an MSL file
  against the engine's kernels on a model's own weights (scripts/dory/q27_batch_lab.sh).
• cortiq mtp-heal measures and trains a draft as the rounds use it. --data A.tdata@TOKENS,B.tdata takes several
  sets and packs their records into windows; --trunk-cache DIR keeps the trunk's pass on disk for the next run;
  --stride N starts a chain every N rows and lays its steps out as a round reads them; --hard F puts F of a
  target row's mass on the trunk's own token; --codec and --also-output CODEC=PATH write the trained block as
  q4tp, q8_2f or f16 from one training; --epochs 0 measures, and with a codec re-encodes, with no training.
  It prints the share of held-out rows each chain step holds, per data set, the drafts a round keeps
  (h1 + h1·h2 + …), and with a stride longer than the chain the same for rounds taken in order, each starting
  after what the one before kept. The checkpoint written is the one that keeps the most drafts a round.
• mtp-heal --dump DIR writes the chain's output rows and top-4 per step with the tensors a draft reads, and
  --taps DIR=L,L,… the trunk's rows after the named layers, for draft models built outside the engine
  (scripts/dory/q27_ext_draft.py runs DFlash 2 and DSpark on them, q27_res_head.py, q27_res_train.py and
  q27_res_gate.py the resonance drafts).
• CMF_GRAPH_SPEC_PIN=1 keeps speculative rounds on whatever they cost, so a count of how a draft holds includes
  the stretches the trial would decode plain (scripts/dory/q27_spec_hold.sh).
• A DFlash drafter beside a model, off by default (CMF_DFLASH=1): cortiq convert --model DIR --output MODEL.cmf --dflash-sidecar writes a published drafter (Qwen3.8-27B-DFlash2: five layers over five of the trunk's rows) as
  MODEL.dflash.cmf, and a plain greedy round on native Metal can take its seven drafts from one block in place of
  the MTP chain. Which drafter a round goes to is read off the trunk's row at the tip by a router fitted from logged
  rounds (cortiq dflash-route-fit, two rank-16 subspaces of such rows), and text copied from the prompt stays with
  the chain and its lookup. Qwen3.8-27B Turbo on a MacBook Air M4, the same text: a Rust function with tests
  11.7 → 13.3 tok/s, an essay 9.6 → 9.9, the others within the run's spread. cortiq dflash-check runs the block
  against the published drafter's numbers; CMF_DFLASH_PICK=dflash|mtp pins a drafter. On wgpu the chain keeps the
  rounds unless the block is pinned: RTX 2000 Ada, 141 ms a block against the chain's 11 ms of drafts.

Changed

• Metal: attention over a prompt chunk's rows is three GEMMs per KV group in place of a loop over rows, and a GDN
  mixer's a and b projections are one dispatch. Qwen3.8-27B on the M4's GPU alone, 3 600 tokens: 87.0 → 73.9 s.
  CMF_ATTN_GEMM=0 for the loop.
• wgpu: the MTP chain's draft depth follows what a round of each depth costs on the card, timed as the rounds run,
  in place of the share of drafts kept alone, and a plain token's time from a speculation trial is kept for the
  process's next requests. Qwen3.8-27B Turbo on an RTX 2000 Ada 16 GB, five chat turns of 200 tokens in one
  process, turns two to five: 26.5 → 29.8 tok/s; a first answer +2 %. On an RTX 5090: 68.3 → 72.8 and +2.7 %.
  CMF_GRAPH_SPEC_K=n pins the depth as before.
• Dory's turbo file keeps its ceiling while a long prompt is read. Past the first seal at 8192 rows a chunk's rows leave
  the device as the chunk ends, on Vulkan and Metal, and the far field is built anew from them before the prompt's last rows.
  --o1-spill DIR (CMF_O1_SPILL) keeps them in files, on the CPU too. RTX 5090 at 64 000 tokens: peak on the card 10.8 → 3.76 GB,
  host RAM 1.7 GB and 9.6 GB of disk with files, at the same speed. Metal reads a sealed prompt in chunks with o1c, and its seal for a
  file with exact rows moves from 32768 rows to 8192. On the CPU a sealed prompt keeps one row's queries in 8.

Fixed

• Metal: the draft chain's input stage, the embedding and the input projection, is a command of its own. Joined with
  the draft block, which lies at the other end of a 13.6 GB weight buffer, the command waited 150 to 170 ms in the
  driver for 9 ms of work in the first rounds of an answer: a round of two drafts on Qwen3.8-27B 362 → 204 ms.
  CMF_MTP_CHAIN_JOIN=1 joins them.