Conversation
Dependency-free C99 single header (same shape as st.h/quant.h: all-static, tests include it directly), three layers: - scalar stream codec: ryg_rans-style 32-bit-state byte-renormalized encode/decode over the 16-nibble alphabet, ported unchanged from the measured prototype round, plus a checked decode that enforces the invariants every genuine encoder output satisfies (initial state in [L,256L), exact byte consumption, final state == L) so validators can refuse corruption without ground-truth bytes; - chunk record: n_symbols/packed_bytes/stream_offsets[257] header, derived round16 padding, 256 round-robin-interleaved streams per tensor. TRUST-VERIFY-REFUSE reader: every malformation class gets its own named error (truncation, count mismatch, offset malformations, short streams, nonzero padding, trailing bytes), no partial acceptance; - batched decode arms: scalar (branchy, mirrors the reference), scalar_bf (closed-form renorm -- the SIMD algebra kept testable on any host), NEON (16 streams/group, early per-lane loads, vqtbl freq/start lookups, fused vuzp repack; ~1.15x scalar_bf single-thread on an M-series host, and the G=8 variant that measured 0.88x is why the width is 16), AVX-512 F+BW (16 lanes, one vpermd per table lookup, vpmovqb packed store; ported from the prototype kernel measured at 8.7 GB/s batched on Zen 5). Vector arms are byte-identical to scalar by contract -- this is a codec, there is no tolerance -- and each carries the quant.h-style envelope: compile-time ISA gate, first-use round-trip selftest that disables a failing arm loudly, env kill-switch (RANS_NEON=0 / RANS_AVX512=0), and RANS_PATH=<arm> forcing that refuses rather than silently downgrades. The selftest table mixes freq=1 and freq=3 rare symbols deliberately: power-of-two frequencies keep floor(log2 x) on a rigid lattice that never visits some multi-byte renorm bands (measured directly -- a sabotaged renorm threshold passed every comfortable-table test), so the off-lattice f=3 group is what makes the selftest actually bite.
repack_rans.py mints the entropy tier from per-row int4 expert shards: corpus-wide header index (a tensor and its .qs sidecar can live in different shards), ONE global table from the pooled nibble histogram stamped byte-identical into every output shard (decode stays shard-local), every record round-trip-verified byte-exact BEFORE its shard is written, and deterministic output (sorted orders, sort_keys stamps, no timestamps -- two runs diff clean). Writer-side geometry guard mirrors the fp8 tool's discipline: a tensor at the per-group-64 byte signature (weight/scale ratio exactly 32) or without a confirmable per-row .qs is refused by name, never silently coded under this format's identity. rans_verify.py is the TRUST-VERIFY-REFUSE validator: stamp layer (the colibri.fmt stamp is MANDATORY here -- entropy-coded sizes are data-dependent, no byte-arithmetic inference exists, so the stamp is the only dispatch signal, not a cross-check), table layer (JSON shape, freq sum, prefix start, slot banding, crc32), record framing, per-stream decode invariants, and a deterministic re-encode pin. One named refusal class per corruption; exit 0 only on full trust. rans_format.py is the shared reference module: table construction (largest-remainder quantization, unchanged from the measured prototype), record build/parse, minimal deterministic safetensors i/o, and a pure-Python codec fallback. builds tools/librans_c.* (a thin exported-symbol bridge over rans.h, same optional-accelerator shape as tools/libiq3): with it the tools use the C codec, without it they emit identical bytes far more slowly -- the e2e test pins the two implementations byte-equal.
test_rans.c (wired into TEST_BINS): stream round-trip property suite over random and adversarial inputs (all-zero, all-max, single-symbol-freq=M, 15-symbol boundary, empty/odd/edge lengths, low-resolution and rare-symbol tables), checked-decode refusal semantics including an enumerated single-bit-corruption sweep (a corruption that silently PASSes the checked decode is a verifier hole, so it is asserted impossible, not sampled), record framing round trip at the format's real N=256 plus small-N shapes, the full named-refusal battery, and the batched-arm identity sweep: every compiled arm (scalar, scalar_bf, NEON here; AVX-512 on x86 builds) must be BYTE-identical to the packed reference on the whole suite. The rare-symbol case exists because a deliberately sabotaged renorm threshold passed every other case: power-of-two-frequency tables keep the decoder's state off the multi-byte-refill bands entirely (measured), so the f=3 off-lattice group is what makes this suite able to bite. Envelope semantics are asserted too: kill-switches drop the vector arms from selection, and a forced unavailable arm is an error, never a silent downgrade. test_rans_repack.py (unittest, skips without numpy): synthetic per-row int4 fixtures with a cross-shard .qs sidecar; end-to-end byte-exact round trip through real shard files; .qs raw pass-through; MANDATORY-stamp and per-shard-table metadata shape; writer determinism (two runs diff clean, AND a forced pure-Python run emits bytes identical to the C-codec run); the verifier's corruption battery (payload flip, count mismatch, offset break, truncation, dirty padding, dropped stamp, dropped table, wrong table crc, freq-sum break, dangling stamp -- each refused by name); and the writer-side geometry refusals (g64 signature, missing scales, empty selection).
…t claimed) The complete wire contract a third-party consumer needs: safetensors container shape (original tensor names, U8 record blobs, raw .qs sidecars), the chunk record byte layout with derived round16 padding, the 256-stream round-robin interleave and why it makes wide decode output-coalesced, the per-stream codec including the verifiable encoder-output invariants, and the per-shard table-in-metadata blob with its crc field and generation id. The load-bearing section is the MANDATORY stamp: unlike every existing format, entropy-coded byte counts are data-dependent, so no byte-arithmetic inference can identify this format -- the colibri.fmt stamp is the only dispatch signal, and a future engine read path must gate on it BEFORE the existing size-formula inference, which stays untouched. Scope is per-row int4 v1 (the per-group-64 class is excluded with its rationale), and the registry row proposes the NAME only: the numeric fmt ordinal is deliberately left for the maintainer to assign.
…er bound, never-crash validator Review findings (one blocker, four MAJOR): the codec core survived a 4,907-mutant ASan/UBSan fuzz with zero memory errors and zero arm divergence, but the untrusted-input edges around it did not hold. All fixes, each with its regression pinned: - parse hardening: the packed_bytes check is now non-wrapping (n/2 + (n&1) — the additive ceil form accepted n_symbols=UINT64_MAX with packed_bytes=0), and an amplification bound refuses any header claiming more symbols than payload_len*8*M_max could encode under ANY admissible table (E_OVERSIZE) — the decompression-bomb surface E2's loader would have inherited, refused before anything sizes buffers from n_symbols. C and Python refuse identical bytes with identical classes, and that parity is now a TESTED property (same-bytes battery through both implementations), not a documented intention. - encoder bound: the old per-stream scratch rule (n + n/2 + 64) was arithmetically false for freq=1 symbols at scale_bits >= 13 (a symbol costs up to scale_bits bits), so a legitimate rare-symbol-heavy tensor refused as out-of-memory. Correct bound (ceil(n*scale_bits/8)+4+slack), and bound exhaustion is now E_SCRATCH — a logic failure, never conflated with the allocator. Bulk-freq=1 input round-trips at scale_bits 13/14/15. - named encode refusals: input containing a value > 15 or a symbol with table frequency zero refuses as E_SYMBOL_UNCODABLE via a pre-scan, instead of burning the scratch buffer into a misdiagnosed E_NOMEM; the stream primitive masks to the nibble alphabet so no input bytes can index past freq[16] (ASan-verified). - rans_verify.py's "never crashes" is now enforced, not promised: the shard header shape is validated by name (object-ness, entry shape, data_offsets bounds against the file, str->str metadata), record-size allocations refuse as E_NOMEM instead of raising MemoryError, and per-tensor/per-shard catch-alls convert anything unexpected into a named E_INTERNAL refusal. The review round's six traceback classes all print REFUSE lines now. - test vacuum: make test-python builds the ctypes bridge first, the e2e C-vs-Python byte-identity pin skips LOUDLY if the bridge is somehow absent, and a bridge-presence assertion pins that the default run actually used the C codec. - geometry guard exactness: 2-D weights get positive per-row verification from the shape fields; flattened tensors at ratio 64 — exactly the per-group-128 signature, byte-indistinguishable from per-row I=128 — refuse as E_GEOMETRY_AMBIGUOUS instead of being silently stamped with geometry the ratio cannot prove. - n_streams is validated in every public entry (E_NSTREAMS; the ctypes bridge pins the format's 256) — no division by zero, and the generic decode path can no longer return RANS_OK with untouched output. Record buffers must be 4-byte aligned (E_UNALIGNED) and the offsets table is accessed via memcpy on the parse/write side; the little-endian precondition is a compile-time #error and documented. - quantize_freq: stable argsorts (exact fractional ties now break by symbol index on every host — contract 4's cross-host clause) and a provable-termination precondition (m >= present symbols, else ValueError). - multi-shard repack is crash-evident: a non-empty output directory refuses without --force, and repack-manifest.json is written LAST as the completion marker — an interrupted set can no longer masquerade as done. - the selftest's renorm-band coverage is machine-checked (the fixture fails the arm if multi-byte refills stop occurring) and the suite gained the same assertion — the one load-bearing claim that was previously an author assertion. The review fuzzer is adopted as tests/fuzz_rans.c (make fuzz-rans, ASan+UBSan, seeded/bounded; not a test gate). - docs: the encoder renorm emission bound (x_max formula) makes a third-party WRITER doc-buildable; the >=4-bytes-per-stream validity rule, the amplification bound, and header-integer endianness are stated. Sabotage meta-check re-run on the final tree: a t15 off-by-one in the NEON arm still fails both the identity suite and the first-use selftest. Engine object still byte-identical to base.
…heck Per-record validation proves a record is A well-formed artifact of this format; it cannot prove it is THE record a given mint run produced — a well-formed record swapped in from another checkpoint passes every per-record check (the #671 review point). Close that at the build layer: - repack_rans.py records, per output shard, a `records` map — original tensor name -> sha256 of the emitted record bytes exactly as written — hashed inline while the bytes are already in hand (no second read). The record wire format is untouched: the manifest is a sidecar file, still deterministic (sorted keys, no timestamps), still the completion marker. - `--manifest-only <dir>` retro-generates the manifest for already-minted directories by hashing shards in place, no re-encoding. It preserves every field of an existing manifest and only recomputes the digest maps, so retro output is byte-identical to mint output for the same shards (tested byte-for-byte). It lives in the WRITER deliberately: the manifest is the writer's completion/provenance artifact, and rans_verify.py stays side-effect-free — a tool pointed at a stranger's download must never write into it. - rans_verify.py checks each record's bytes against its digest whenever a manifest with digests sits next to the shards, BEFORE record decode (the swap case is exactly what the later checks are blind to). Named refusals: E_DIGEST_MISMATCH (bytes differ from the mint run), E_DIGEST_MISSING (digests exist but not for this shard/tensor — a record the mint never produced), E_MANIFEST_MALFORMED (a manifest that exists but cannot be parsed is evidence of tampering, refused rather than treated as absent). No manifest, or a digest-less manifest from an older mint, is note-and-proceed — never an error. Never-crash holds. - tests: digest round-trip green with "manifest digests checked" pinned; tamper bite (one flipped payload byte in the shard file, manifest untouched -> E_DIGEST_MISMATCH); retro equivalence (byte-identical manifest over a minted dir; from-scratch retro reproduces the digest map); backward compat (stripped digests -> note + green); corrupt manifest -> E_MANIFEST_MALFORMED without traceback; foreign shard -> E_DIGEST_MISSING. The per-record corruption battery now drops the manifest first — that layer must hold even when an attacker removes the whole-artifact evidence. Mint determinism (manifest included) already pinned by the existing two-run byte compare. Build integrity (these bytes came from that mint run) vs container identity (the stamp/registry lineage) is documented in docs/int4-rans256-g0.md; the engine-side load check ships with the consumer PR. Whole-file layer (micro-addendum, same round): each manifest shard entry also carries `sha256` — the complete .safetensors file bytes, hashed while streaming the write before the atomic rename (retro mode hashes files in place; retro-over-minted stays byte-identical to mint). The file hash subsumes what the record map cannot see — .qs sidecars, headers, padding — closing the scale-swap gap; the record map stays for granular diagnosis. Verify checks the file hash FIRST (cheapest whole-artifact gate, E_SHARD_DIGEST_MISMATCH), then the per-record layer, continuing past a file-level mismatch so record digests localize in-record damage. Tests: .qs-region tamper bites as E_SHARD_DIGEST_MISMATCH with NO record-level refusal (the exact gap, demonstrated), record tamper trips both layers, retro identity and mint determinism re-proven, degraded manifests (no digests at all / file-hash-only) note-and-proceed. Validator delta-3 fix (same commit): manifest entry shapes are now type-validated unconditionally BEFORE the digest-evidence scan — an entry whose only digest fields are malformed-typed (sha256 as a number, records as a list, an entry that is not an object) previously zeroed the evidence scan and was silently classed as a digest-less older mint instead of refusing. Wrong-typed digest fields are malformation regardless of whether valid evidence remains elsewhere; both validator corpses are in the shipped battery (E_MANIFEST_MALFORMED, no TRUST line, zero tracebacks). Retro --manifest-only also warns on stderr when `source` provenance is unknowable, so the loss is visible in output, not just by field absence.
fix(win): test_st_mirror stale dirs + two -Wformat-truncation warnings
Convert speedup using parallel workers
serve: wire Host allowlist through coli CLI
…ngine docs: define colibri as an experimental inference engine
entropy: in-tree rANS codec, int4-rans256-g0 container format, offline tools
…ode (#678) With KV_SLOTS=1 (the default) each serve turn runs through ONE spec_decode call, and run_serve_mux's stdin poll only runs between turns. When the gateway's stop filter matched (e.g. a role marker) it sent STOP, but the engine sat inside the turn generating suppressed tokens until max_tokens -- minutes of "active query, no progress" on CPU-only hosts (#675 follow-up). Fix: mux_ctl_poll, a non-blocking stdin poll called from the mux spec emit callback once per emitted token. It consumes pending STOP/CANCEL lines and raises g_mux_stop/g_mux_cancel, which spec_decode checks next to g_intr, so a server stop now takes effect within ~1 token. STOP exits into the normal mux_done (DONE frame); CANCEL mirrors mux_submit's cancel epilogue (ERROR CANCELLED, consistent KV append). The framed protocol stays in sync: a SUBMIT that arrives while the only slot is busy has its payload drained and is refused with SLOT_BUSY, exactly as mux_submit would. On stdin types the platform cannot poll, behaviour falls back to the pre-fix end-of-turn handling. The flags are only ever raised during a mux spec turn, so chat, run and oracle spec_decode callers are unaffected. Tested: CPU build clean; -DCOLI_CUDA syntax clean; tests/test_stops 4/4; tests/test_openai_server.py 100/100; live mux protocol test against the glm_tiny fixture (READY, STOP honored in 0.27s on the batched path, CANCEL -> CANCELLED, unknown id -> NOT_FOUND, follow-up requests complete). The spec path itself needs an MTP-bearing model; verification on real hardware requested from the #678 reporter. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
fix(serve): honor STOP/CANCEL mid-turn in single-slot speculative decode (#678)
rocWMMA's headers static_assert on archs without MFMA/WMMA units (ROCm/TheRock#1944), so any HIP build for MI50/Vega/RDNA1/RDNA2 died at compile. Make rocWMMA optional: the Makefile derives -DCOLI_HIP_NO_WMMA from HIP_ARCH (normalized: comma lists accepted, native resolved via rocm_agent_enumerator so the arch filter sees real gfx names) and the compat header then compiles the WMMA kernels out, taking the same host-dispatch fallbacks as pre-Volta CUDA. The define must reach both hipcc passes: an arch-macro check in the header would leave host dispatch enabled for kernels whose device bodies were compiled out. Also gate COLI_CUDA_TC_INT4 on COLI_GPU_HAS_WMMA like TC_W4A16 already is: grouped_s4_wmma's body is __CUDA_ARCH__>=750-guarded, and launching it empty silently leaves the output buffer unwritten. Missing rocwmma-dev on a WMMA-capable arch stays a hard #error; mixed capable/incapable HIP_ARCH lists warn that WMMA is disabled for all targets (shared host dispatch requires one consistent setting). Verified on gfx906 (MI50) + gfx1030 fat binary: builds clean, engine serves GLM-5.2 with the VRAM expert tier active. make check passes (229 tests). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The engine takes the string "auto" (sized to free VRAM at init), but the launcher did a strict float() on the env var, so coli doctor/chat/serve crashed with "could not convert string to float: 'auto'". Non-numeric values now mean 0 for the planner (planner sizes the budget itself) while the engine keeps reading the original string. Same guard for RAM_GB, where the engine's atof() already tolerated it but the launcher didn't. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
fix(hip): support GPUs without matrix cores (gfx9xx, gfx101x/103x)
ngram_draft() already proposes a continuation by matching the tail of the current sequence against earlier positions in that same sequence. This adds the same lookup over a persistent file of previously generated token ids: at each decode step the engine proposes what followed the longest suffix of the live context found in the corpus (suffix 8..3, most recent occurrence first). Proposals enter the existing batch-union verify forward, so the target model's argmax still decides what is accepted. Unlike MTP (one token per forward) a corpus hit proposes a whole span. On GLM-5.2 int4, 96-token greedy, replaying a prompt whose generation is in the corpus: 1.78 -> 6.00 tok/forward and 3.69 -> 4.52 tok/s on a resident H200 (90% acceptance), 1.76 -> 7.92 tok/forward and 0.82 -> 1.00 tok/s on CPU (100% acceptance). Note the gap: 3-4x fewer forwards buys ~22% of wall clock, because in an MoE every verify row activates its own experts, so speculation amortizes the dense path and attention but never the expert work. docs/corpus-draft.md documents that ceiling rather than leading with the forward multiplier. Rejected drafts cost time, so the source pauses below COLI_CORPUS_MINACC acceptance (default 50%, the measured break-even: 90% gave +22%, 19% gave -25%). A prompt unrelated to the corpus costs nothing -- the suffix match finds nothing and the source stays inert (measured: same forward count as no corpus at all). Off by default: without COLI_DRAFT_CORPUS the corpus pointer is NULL and behaviour is unchanged. Build a corpus from any run with TOKENS=1, which already dumps the generated ids. Validation on this branch: make check 229 tests OK, clean build 0 warnings, oracle 32/32 TF + 20/20 greedy, and 20/20 greedy again with the source armed (a new draft source must be token-exact when ON, not only when OFF). tests/test_corpus_draft.c covers the lookup contract: longest suffix wins, recency wins ties, the -1 span separator terminates a proposal, caps hold, degenerate inputs propose nothing. One caveat, not specific to this source and filed separately: on CUDA a run whose drafts are accepted in long streaks can diverge from the unbatched path by a single near-tie token (1 in 96 measured), and the same corpus accepts at 90% on GPU vs 100% on CPU. The CPU path is byte-exact.
…IN_GB hosts (#686) CUDA_RELEASE_HOST defaulted to ndev>1, chosen when the host copy was the multi-GPU re-upload path. On a single GPU that is also asked to fill RAM (PIN_GB=all, or a PIN_GB large enough that the two tiers compete for the same memory), that default keeps a full host copy of every VRAM-tier expert and the RAM tier starves. Measured on 1x H200 141 GB + 235 GB RAM, GLM-5.2 int4, 96-token greedy, 3 runs, median: default (RELEASE_HOST off) 9,297 resident experts 86.2% hit 1.10 tok/s CUDA_RELEASE_HOST=1 14,951 resident experts 99.4% hit 4.75 tok/s [PROF] for the default config showed 166.9 GB fetched with 87.7 s of read service inside an 87.4 s decode: the whole token budget went to misses the duplicated RAM could have absorbed. The host copy is provably redundant in that configuration — releasing it still reloads from disk if CUDA later fails. So keep ndev>1 exactly as before and add the single-GPU + large-PIN_GB case: - explicit CUDA_RELEASE_HOST always wins (0 disables, as before) - ndev>1 unchanged, so multi-GPU behaviour is byte-identical - single GPU only opts in when an expert tier exists AND PIN_GB is "all" or >= the tier size (under CUDA_EXPERT_GB=auto the tier is sized to the whole card, so any explicit PIN_GB qualifies) - one stderr line when the new default engages, naming the knob to undo it Also extends the existing "[RAM_GB auto] WARNING ... PIN_GB is inflating the resident set" path: on a single GPU still holding host copies, what inflates the resident set is usually those copies rather than PIN_GB itself, so the warning now names CUDA_RELEASE_HOST. That line already fired in exactly this situation without telling the user which knob to reach for. Validation: make check OK, clean build 0 warnings (CUDA=1 and portable), oracle 32/32 TF + 20/20 greedy.
fix(cuda): release host copies of the VRAM tier on single-GPU large-PIN_GB hosts (#686)
fix(cuda): apply grouped scales in ragged attention
feat(spec): frozen-corpus draft source (COLI_DRAFT_CORPUS)
The measured numbers come from replaying a prompt whose generation is already in the corpus — a best case, not a general speed-up. On novel text the corpus has nothing to propose and the gain approaches zero, which is the case a chat assistant is normally in. Says so explicitly so the table isn't read as a blanket +22%. Also cross-links #689: deep drafts from this source are the easiest way to reach the open CUDA verify-batch divergence at S>=8. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…PU (#687) A resident tensor whose upload fails is disabled permanently and silently falls back to CPU. Today that produces one line per tensor: [CUDA] tensor [2048,6144] on device 0 disabled after an error; falling back to CPU Sixty of those scroll past in a second, none of them says what to do, and the end state reads as healthy: the GPU is fully allocated, nvidia-smi looks busy, and the run is at CPU speed. In #687 that cost a debugging session to work out that CUDA_EXPERT_GB=auto had claimed the card to within 4 MiB, leaving nothing for the lazily-uploaded dense tensors. This keeps the per-tensor detail for the first few, then prints one prominent block naming the knob that fixes it, and repeats the count in the end-of-run CUDA stats — by then the original lines are thousands of log lines back, and the count is what explains a CPU-speed "GPU" run. Diagnosability only: no behaviour change, no new env var, and nothing prints on a healthy run. Validation: make check OK, clean build 0 warnings (CUDA=1 and portable), oracle 32/32 TF + 20/20 greedy.
…alar + AVX2 The native layout of compressed-tensors 'mxfp4-pack-quantized' checkpoints: packed [O,I/2] u8 (low nibble = even column, bit3 = sign) + scales [O,I/32] ue8m0, w = v * 2^(s-127). Computes on the vendor bytes directly - QAT weights (Kimi K3 routed experts) are exactly the trained values, so they are streamed and multiplied as-is, never re-encoded. AVX2 decodes a 16-byte group with one pshufb over the doubled-e2m1 int8 LUT; the 0.5 un-doubling rides the group scale. Scalar and SIMD paths use the same bit-trick exponent decode so both agree on the s=0/255 edge cases. Verified against an independent numpy implementation of the spec (rel err 1.2e-7 = f32 noise) and byte-conventions against compressed-tensors' pack_fp4_to_uint8/unpack_fp4_from_uint8.
… vocabs
Two additions for Kimi K3's tokenizer (tiktoken vocab, fancy-regex pattern):
- kimi Split family (sniffed via \p{Han} in the pre_tokenizer regex): the
o200k case-aware rules plus a leading [\p{Han}]+ rule, Han masked out of
the letter classes (fancy-regex &&[^\p{Han}] intersections), and no '/'
tail in the punctuation rule. A Han codepoint can only match the Han rule,
so Han-free text tokenizes exactly like o200k minus the '/' tail.
- rank-BPE: when model.merges is empty, bpe_piece merges the adjacent pair
whose CONCATENATION has the lowest vocab id - tiktoken's own algorithm,
exact by construction. Merge-list recovery from ranks is provably lossy
(the standard max-rank-minimizing reconstruction mis-merges 'newlines'
as new+l+ines), which is why K3's generated tokenizer.json ships no merges.
tests/test_tok_kimi.c + tools/k3_tokenizer.py --ctest cross-check the C
encoder against tiktoken: exact on all 18 corpus cases (Chinese, kanji vs
kana, Korean, CRLF, contractions, emoji, mixed-script boundaries). Existing
families are untouched: merges-bearing tokenizers keep the merges path.
New sibling engine for Kimi K3 (2.8T / 104B active, 93 layers), following the one-engine-per-family pattern (colibri.c untouched). Shares st.h / json.h / tok.h / quant.h. make kimi_k3; ./kimi_k3 <model_dir> "prompt" --ngen N. - 69 KDA layers: SiLU(causal-conv4) projections, L2-normed q/k, per-channel bounded decay gk = -5*sigmoid(exp(A_log[h])*(f_b(f_a x)+dt_bias)), delta rule S=(I-beta k k^T)Diag(e^gk)S + beta k v^T, per-head RMSNorm x sigmoid full-rank output gate. Checkpoint A_log [128] = per-head [96] zero-padded. - 24 gated MLA layers, absorb form, NoPE (nothing is rotated; the 64 extra dims are cached raw), sigmoid output gate before o_proj. - AttnRes: the residual stream is replaced by softmax mixes over block snapshots (layers 0,12,...,84) + the running prefix sum - twice per layer and once before the head, scored against res_norm.w*res_proj.w, raw-entry mixing in fp32. - Stable LatentMoE: sigmoid router + score-correction bias, top-16, raw-score renormalization; latent 7168<->3584 projections around the experts; RMSNorm on the aggregate; 2 fused shared experts; SiTU-GLU (4 tanh(g/4) sigmoid(g) * 25 tanh(u/25)). - Routed experts stream from the safetensors shards as native MXFP4 through a per-layer LRU (K3_EXPERT_GB): one pread per expert (the six tensors are stored back-to-back - measured), loads issued in disk-offset order (experts are NOT id-ordered in the shards - measured), multi-drive shard split via K3_DIRS (st_init_multi). - Non-expert BF16 is quantized at load (K3_BITS int4-g64 / int8 / f32, K3_MLA_BITS, K3_HEAD_BITS) or read pre-quantized from a repacked container (U8 + .qs sidecars, tools/k3_repack.py) - container start measured 30.4s -> 0.6s on two layers. Validation (docs/kimi_k3.md): with injected inputs and f32 weights the engine matches an independent numpy reference (tools/k3_ref.py) on real checkpoint weights across all four layer types to rel-L2 <= 2.2e-6; container-vs-load-time quantization is bit-identical (layer-0 trace equal); tokenizer exact vs tiktoken. Known limits: single-token stepping (sequential prefill), CPU only, no chat template.
- k3_tokenizer.py: tiktoken.model -> tokenizer.json (byte-level vocab, EMPTY merges on purpose -> tok.h rank-BPE, kimi Split regex, 256 added tokens from tokenizer_config). --check replays rank-BPE in python vs tiktoken; --ctest drives tests/test_tok_kimi against tiktoken. 18/18 exact. - k3_ref.py: independent numpy implementation of the K3 stack (KDA recurrence, gated NoPE MLA, AttnRes, LatentMoE with MXFP4 dequant) reading tensors straight from the HF shards; gen/run/cmp subcommands produce and compare the engine's K3_TRACE hidden-state dumps. This is the reference the engine is validated against (rel-L2 <= 2.2e-6 at f32 on real weights). - k3_repack.py: one streaming read of the HF snapshot -> ideal ingestion container, so a disk-to-disk copy IS the conversion. Experts byte-identical in id order (six tensors back-to-back = the engine's slot layout, one pread per expert), big BF16 mats quantized with the ENGINE-EXACT algorithm (int8/row or int4-g64, --bits/--mla-bits/--head-bits) into U8 + .qs f32 sidecars, sensitive smalls passed through, vision dropped, spec-valid safetensors (no holes) + regenerated index, per-shard resume, --verify-full re-reads every expert byte. Kills the engine's 10-15 min load-time quantization; quantized bytes verified bit-identical to the load-time path. ~1.56 TB source -> ~1.50 TB (int8) / ~1.48 TB (int4).
Full-model bring-up on a 62 GiB box exposed three streaming-path issues: 1. Expert loads were sequential buffered preads; with ~55 GiB of resident weights the page cache has no headroom and reclaim throttled reads to ~1.8 GB/s (drive: 7.1 GB/s O_DIRECT, 2.9 buffered — measured). A token's misses now load in parallel into working-set slots and read O_DIRECT through st.h's twin fds (aligned-window + tiny buffered tail; K3_DIRECT=0 reverts to buffered + WILLNEED). Expert reads: 6.3 GB/s. Decode: 26-28 -> 9.2-10.7 s/token on the 93-layer model. 2. The LRU cap floor was topk, silently committing topk*92 slots (~26 GiB) regardless of K3_EXPERT_GB; experts are consumed one at a time, so the floor is now 1 and the budget is honored. 3. Explicit K3_BITS=4 on an int8 container downcasts to int4-g64 at load: ~35 vs ~57 GiB resident (fits next to a desktop session), same decode speed once I/O is fixed; unset K3_BITS keeps container bits (default unchanged). 93-layer sanity (greedy, repacked int8 container): 'The theory of general relativity was developed by' -> int8: 'Einstein when he felt that the theory of special relativity was insufficient...', int4: 'Einstein in 1915, and it'. Init 42 s (int8) / 131 s (downcast pass).
Teacher-forced bit-width comparisons (with --ngen 0): run the same --ids twice under different K3_BITS and diff the logit streams offline. Used for the int8-vs-int4 dense-side quality evaluation on real text.
The float kernel converts every decoded weight nibble to float before the
FMA; at 25.8 GB of packed expert bytes per K3 token that path measured
~4.4 GB/s effective. The doubled e2m1 values are exact int8 (same LUT), so
a 32-group reduces to maddubs(|w|, sign(x,w)) + madd — the dot_i4i8 idiom —
with the group scale folded once per group:
y = sum_g idot_g * 2^(e8-127)*0.5 * xscale_g
Activations quantize per 32-group to int8 (~0.4% group noise, aligned with
the weight groups so scale folding is exact). Unit test vs the float kernel:
rel-L2 3.5e-3 on random data. On the full model: teacher-forced real-text
ppl 2.684 vs 2.694 (float), argmax agreement 98.4%, KL mean 0.003 — free at
text level, 3x faster (expert compute 5.8 -> 1.9 s/token). Scalar fallback
included; I%32!=0 falls back to the float kernel.
Decode 9.4 -> 5.2 s/token (and 3.0 opt-in), on top of the O_DIRECT work: - K3_PIPE (default 1): a small loader-thread pool (K3_LOAD_THREADS, 4) runs the expert preads while the compute loop chews already-loaded experts — expert j's matmuls overlap expert j+1's read. t_eload now reports the UN-hidden wait (2.5 s/token; pure I/O would be 3.8). - K3_IDOT (default 1): expert matmuls via quant.h's new matmul_mxfp4_i8. Real-text gate: ppl 2.684 vs 2.694 float, agreement 98.4% — free. - K3_TOPP (default 0 = off): drop the low-weight tail of the top-16 to cumulative weight p, renormalized (GLM TOPP semantics) — the only lever cutting I/O and compute together. At p=0.7 it keeps ~8.1/16 experts: 3.0 s/token, but ppl 2.811 vs 2.694 (dNLL +0.043) with accuracy unchanged at 73.4% — measurable, so it stays OPT-IN; gate any setting with a K3_LOGITS A/B first. Full-model numbers, greedy sanity text unchanged. make check green.
Prefill now runs layer-major over K3_CHUNK-token chunks (default 32): - every dense matmul batches over the chunk, so each weight matrix streams from RAM once per chunk instead of once per token; - the MoE builds the chunk's expert union per layer and loads each unique expert ONCE (position lists via counting sort), feeding the same pipelined block loader as decode. Measured on real text at C=32: 2.7x dedup — neighbouring tokens share experts far more than QB-flat routing predicts — 9.6 instead of 25.8 GB streamed per token; - the lm_head runs only on the chunk's last token (K3_LOGITS still dumps every position); - sequential state (KDA recurrence, MLA cache append, AttnRes bookkeeping) advances per token inside each layer in the original order, so chunked results are BIT-IDENTICAL to token-at-a-time: the 125-position teacher-forced logit streams at C=32 vs C=1 match exactly (max abs diff 0.0). K3_TRACE forces C=1 (trace rows are token-major by contract). The KDA delta-rule state sweeps get explicit AVX2 (mul+fmadd over the 128-wide state rows, scalar fallback when head_dim%8!=0): attention 0.9 -> ~0.5 s/token; vs the pre-SIMD logits: 100% argmax agreement, ppl 2.6841 unchanged at 4 decimals. Full-model prefill: ~5.3 -> 2.0 s/token (125-token prompt, C=32, int8 dense, K3_LOGITS on — pure generation prefill is slightly better still). Decode unchanged (same C=1 path). make check green.
…hannel
K3 ships its chat renderer as encoding_k3.py: an XTML structure where only
<|open|>/<|close|>/<|sep|>/<|end_of_msg|> are special tokens and tag names /
attributes are ordinary text, encoded as standalone segments. A turn is
<|open|>message role="user"<|sep|>TEXT<|close|>message<|sep|><|end_of_msg|>
and the generation prompt opens the assistant message plus its STRUCTURAL
thinking channel (<|open|>think<|sep|>; K3_THINK=0 opens <response> directly
— in thinking mode the channel exists in every assistant message, which is
the preserved-thinking property).
The C builder mirrors the reference segment-for-segment (segment boundaries
are token boundaries, so per-segment BPE matters) and is verified TOKEN-EXACT
against encoding_k3.py + tiktoken for system/no-system and thinking/
non-thinking prompts. The CLI hides the structure at decode and prints
[think]/[response] banners; <|end_of_msg|> is already the configured eos.
Full-model greedy check ('how many r in strawberry'): the model reasons in
the think channel (spells the word, finds positions 3/8/9), closes it, opens
response, answers '3 times' — native use of the format end to end.
Single system+user turn only for now (no multi-turn CLI, no tool-call
rendering); make check green.
(cherry picked from commit ca3dcc6)
Kimi K3 engine: native-MXFP4 expert streaming, KDA + gated NoPE MLA + AttnRes + LatentMoE (#658)
olmoe: decode goes disk-bound whenever the expert LRU misses
feat(cuda): say it once, loudly, when resident tensors fall back to CPU (#687)
…iner; --chat Three things, all needed to run Inkling on a host that cannot hold its bf16 dense set. 1. CORRECTNESS — the expert cache silently computed with the wrong weights. moe() acquired a slot for every (token, expert) pair up front and kept the pointers until the compute pass. slot_acquire evicts the LRU when the cache is full — including slots already handed out in the same call: the pointer stays valid but the slot now holds a DIFFERENT expert. No crash, no warning, just incoherent output. Correct behaviour therefore required a cache big enough for every distinct expert in the batch: a prefill of 18 tokens at topk=6 needs up to 108 slots per layer. Now the routed experts are processed in ROUNDS of `cap` pairs — acquire, fill, compute, accumulate — so nothing is evicted while in use. The MoE output is a weighted sum, so accumulating in rounds is identical to one pass, and the cache can go down to one slot per layer (more disk reads, memory proportional to `cap` instead of to the batch). Verified on tools/make_tiny_inkling.py's transformers oracle: 36/36 teacher-forced and 24/24 tokens, at cap=4 AND at cap=1 (where the old code used one slot for both experts of a topk=2 model). 2. MEMORY — optional int4-gs64 dense container. The pre-converted container has int4 experts but bf16 dense weights: 49.4 GB measured, resident, and load_w expands bf16 to f32 while loading (~99 GB peak), so hosts below ~64 GB die before generating. c/tools/ convert_inkling_dense_int4.py quantizes just the dense set into a separate container (15.3 GB); the engine auto-detects <snap>/dense-int4g64/ and INK_DENSE_Q4=0 ignores it. Without the container nothing changes: the bf16 path is untouched. Adds matmul_i4g (group-scaled int4) and matmul_i8r (per-row int8) plus a bounds guard in matmul_w, since the container is a file, not an invariant. Measured error vs the real bf16 weights: ~11% rel-L2 on attention/shared/ dense-MLP at int4-gs64 (which is what 4 bits costs on this distribution), ~0.9% on embed/lm_head at int8. Enough for coherent output: the 975B answers "My name is Inkling." on a 25 GB box, 15.1 GB RSS, clean EOS. 3. USABILITY — `--chat` renders Inkling's chat template in the engine, so the direct CLI stops feeding an instruct model raw out-of-distribution text. coli chat/serve/web already render it through the gateway. Speed is unchanged and still disk-bound: ~8 GB of expert cache against a 464 GB expert bank is ~1.7% residency, tens of seconds per token. This makes the model runnable and testable on a small host, not fast — docs/inkling.md says so plainly rather than quoting the best case. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…he-fix fix(inkling): expert cache served the wrong experts; int4 dense container; --chat
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…t it trims
`colibri.c` has had `TOPP` and `kimi_k3.c` has `K3_TOPP`; `inkling.c` was the
only engine without it, and it is the engine where it matters most. Inkling is
disk-bound on a small host: topk=6 experts x 28 MB x 66 layers is ~11 GB read
per token, and 92% of decode time is expert I/O. Every expert TOPP drops is an
expert not read.
Same semantics as the other two engines: sort the routed (weight, id) pairs
descending, keep them until the cumulative weight crosses `p`, drop the tail,
and do NOT renormalise — the dropped weight simply does not contribute.
si[] is ordered by sigmoid(logit)+bias while the weight is sigmoid(logit)
without the bias, so the weights are not guaranteed descending and the pair is
re-sorted before the cut, exactly as colibri.c does.
This is a QUALITY lever, not a free speed-up: it changes the computation. So it
is opt-in, off by default, announced at startup, and — new here — it reports
its own effect at the end of the run:
[topp] 0.30: 36/72 routed used (50.0% trimmed, 1.00 experts/token avg)
Measured on the tiny fixture at TOPP=0.30: 50% of routed experts trimmed and
expert loads down from 47 to 27 (-43%), which is the mechanism the disk-bound
host needs. A number you can A/B instead of a claim you have to believe.
TOPP=0 (default) leaves the computation bit-identical: the transformers oracle
stays at 36/36 teacher-forced and 24/24 tokens, at cap=4 and cap=1. Values
outside (0,1) are refused with a message rather than silently clamped.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
feat(inkling): TOPP — adaptive routed-expert trimming, and report what it trims
The models table said which engines exist but the Run-it section only ever showed GLM, so a reader had no way to know that coli picks the engine from the model's config.json and renders that family's chat template — i.e. that the command line does not change between models. Adds the build targets, the identical chat/web/serve invocations for GLM, Inkling and Kimi K3, the note that non-GLM chat routes through the local gateway (so TUI, API and dashboard share one arch-aware template), and the two per-model caveats that actually bite: Inkling needs --cap on a RAM-tight host, Kimi needs ~1.6 TB for the snapshot. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Owner
Author
|
Two additions since the description above — both already on
|
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
10 participants
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.
Promotes 51 commits from
dev. The headline: colibrì goes from one model family to three MoE families, 744B to 2.8T parameters — and a 975B answers on a 25 GB box.New model families
Kimi K3 (Moonshot, 2.8T total / 104B active, 93 layers) — #676, by @steve-m, with the chat/API/Web integration by @ZacharyZcR. A sibling engine (
c/kimi_k3.c) that streams the QAT-trained MXFP4 routed experts straight from the original Hugging Face shards — never re-encoded — and quantizes the bf16 dense set at load time. Includes the KDA + gated-NoPE-MLA + AttnRes + LatentMoE architecture, an MXFP4 matmul kernel (scalar + AVX2 + int8-activation), chunked prefill (bit-identical, 2.6×), parallelO_DIRECTexpert reads, and--chatfor K3's XTML format.Status: engine and tokenizer validated — the multi-turn wire matches Moonshot's official
encoding_k3.py77/77 token ids — and the C/Python regression suite passes. Full-model generation is still gated on a host with ~1.6 TB of storage for the checkpoint; nobody has run it end to end yet. Said plainly so the people with that hardware know it's theirs to try.Inkling (Thinking Machines, 975B / 41B active) now runs on a small host — #701. Two problems were in the way:
moe()acquired a slot for every(token, expert)pair up front and kept the pointers until compute, but a full cache evicts slots already handed out in the same call. Correctness therefore required a cache big enough for every distinct expert in the batch — up to 108 slots per layer for an 18-token prefill attopk=6. Below that: incoherent output, no error. Now the experts are processed in rounds ofcapand accumulated, so nothing is evicted while in use, and the cache can go down to one slot per layer.Result:
My name is Inkling.— a 975B model, correct answer, clean EOS, 15.1 GB RSS on a 25 GB machine. Verified against the transformers oracle at 36/36 teacher-forced and 24/24 tokens, at bothcap=4andcap=1(where the old code corrupted).Correctness fixes
attention_absorb_ragged_kernelapplied per-row scales to group-scaled (fmt=4) tensors, so the recommended g64 GLM container produced gibberish throughcoli serve/coli webwhilecoli runwas fine. Located to a single line, fixed with the sharedabsorb_scale()helper, plus a two-group regression fixture and the previously-orphaned ragged test wired into CI. Field-confirmed fixed by @terrizoaguimor on his exact reproduction.STOP/CANCELmid-turn in single-slot speculative decode; before, a stop matched by the gateway sat unread untilmax_tokens(minutes of apparent hang on CPU hosts).doctorandresource_planrecognize AMD/ROCm; a working HIP build is no longer reported CPU-only.Performance
CUDA_RELEASE_HOSTnow defaults on for single-GPU hosts with a largePIN_GB, where the duplicated host copy was starving the RAM expert tier. Measured on a single H200: 1.10 → 4.75 tok/s (3.8×), 9,297 → 14,951 resident experts.CUDA_EXPERT_GB=autofilling the card no longer degrades 60× in silence.olmoestopped callingfadvise(DONTNEED)after every expert read, which was dropping pages the LRU would immediately want again (EXPERT_DROP=1restores the old behaviour for RAM-tight boxes).COLI_DRAFT_CORPUS: a frozen-corpus draft source proposing whole spans. +22% wall clock when replaying a prompt already in the corpus; the docs state the scope rather than the headline.Thanks
This release is mostly other people's work, and the hard parts were found by people with hardware we don't have:
Version bumped to 1.3.0 (
c/version.py). Full suite green: 28 C tests + 233 Python tests. Tagv1.3.0after merge to publish the release artifacts.