71 pull requests since v1.6.2. A sixth model family with its GPU tier, a
rebuilt expert-matmul path, and the CI that would have caught the class of bug
we shipped twice.
A sixth engine: Qwen3.6-35B-A3B β CPU and GPU
- #712 (@kreuzzelg) β hybrid Gated Attention + Gated DeltaNet + streaming
MoE, inc/qwen36.c. Pre-converted containers published (int4-gs64
recommended: cosine to the int8 anchor 0.98777 β 0.99313, KL 0.109 β 0.080
vs per-row). The engine takes any architecture-identical checkpoint
unchanged β KAT-Coder v2.5 runs on it with no code path of its own. - #713 (@kreuzzelg) β CUDA VRAM expert tier with heat-based placement
across GPUs: 1.44 β 10.05 tok/s (7.0Γ) on 2Γ 8 GB cards, output
bit-identical to CPU (cmpover the full 200-token generation), measured
cold with no heat table. Aqt_ready()gate keeps CPU-only builds from
allocating the packed int4 buffers they never read: 7.33 GB saved.
The expert matmul path, rebuilt (all bit-identical)
- #1071 / #1075 / #1076 / #1077 β activation quantization hoisted to layer
level across GLM, Kimi K3 and DeepSeek V4: the same vector was being
re-quantized ~16Γ per layer, serially. Removes ~5.2 ms/token of serial time
and every per-callmallocfrom the hot path. - #1079 / #1086 β K1: plane-nibble int4 layout + unsigned-VNNI dot.
Storing element k and k+32 in one byte deletes the unpack, and since
nibbles are stored unsigned,dot(v,x) = dot(u,x) β 8Β·Ξ£xfeedsvpdpbusd
natively. 8 uops per 64 MACs instead of 32: 1.45β2.65Γ on the IDOT
kernels, zero bytes added. - #1088 β K2: 1Γ4 union tile. The prefill union hands each expert
2β16 rows; the weight block's load+mask is now paid once per four rows
instead of once per row: 2.67β3.10Γ at S=4 (peak 253 GMAC/s). - #1093 β parallel silu and the down-side activation hoist.
- #1094 β K1b: grouped planar IDOT for gs64 containers (
IDOT_GS=1,
opt-in): the recommended container format could not reach the integer
kernels at any batch size before this. - #1082 (@outtodata) β
FUSED3=1opt-in fused AVX2 expert matmul.
Streaming and I/O
- #1097 β DeepSeek V4 expert-loader pool default 3 β 9 lanes:
1.41Γ decode on the real V4-Flash checkpoint (8 interleaved runs on a
quiet 25 GB box).V4_LOADER_LANESstill overrides. - #1056 (@dcutugno) β DeepGEMM sm120 headers fetched at a pinned commit on
first build: 2.5Γ prefill on sm120 with nothing vendored in-tree. - #988 / #1054 / #1055 β DeepSeek V4 CUDA tier and dual-SSD mirror.
Correctness and CI
- #1083 β ARM CI job (
ubuntu-24.04-arm) plus an integer-kernel
bit-exactness gate that runs on both ISAs. Every tiny-oracle job ran on x86
before this, so NEON-divergent paths were invisible β which is how the IDOT
defaults below shipped. Closes #1081. - #1044 / #1080 β IDOT made opt-in in olmoe and inkling: the fast path is
x86-only and quantizes activations, so the same model produced different
tokens on x86 and ARM by default. - #1109 (@SebaWag) β ARM64 dotprod probed by compiling the intrinsic:
GCC 11 defines__ARM_FEATURE_DOTPRODfor a base it cannot emit
vdotq_s32for. Fixes #1104. - #1111 β
__syncwarp()aftergrouped_s4_wmma's store (reported by
@monotophic withcompute-sanitizerevidence). Fixes #1099. - #1073 / #1074 (@bherald) β Kimi K3 cancels prefill between layers
instead of holding the engine for a minutes-long prompt; cancelled requests
no longer count as completed.
Apple Silicon
- #790 β #1113 (@RDouglasSharp) β Metal backend for Kimi K3: KDA state and
window buffers aligned, wrap-once buffer cache, CPU-side MLA KV cache. 1.7Γ/2.4Γ
on the compute-bound phases (KDA attention + projections dispatched to the GPU);
MoE experts stay on the CPU. Kimi K3's first GPU backend.
More correctness fixes
- #1098 (@monotophic) β
__syncthreads()missing from the absorb softmax
reduction, with a determinism test that reproduces the hazard. - #1101 (@monotophic) β allocation and
snprintfresults checked on the
checkpoint-load path (#798). - #1100 / #1108 (@monotophic) β fmt=8/fmt=6 scale-byte accounting in
tensor_bytes/tensor_free, andweights_ownedset before the host-to-device
copy so a failed upload frees its buffer. Each ships with its own regression
test; all four of this contributor's CUDA fixes landed in this release. - #1122 (@ZacharyZcR) β
USAGE_SAVE=0honoured in every engine (#1039): the
history was loaded but written back anyway, which quietly contaminated any A/B
that shared a usage file between arms. - #1121 (@ZacharyZcR) β LRU victim selection now respects a lowered
ecap
(#1034): after an RSS-guard reduction the cache kept evicting against the old
capacity. - #1123 (@ZacharyZcR) β the v1.6.2 warning-cleanup patches landed (#1032).
- #1106 (@monotophic) β duplicate tensor names across indexed shards are now
refused rather than silently resolved to one of them (untrusted containers).
Interfaces
- #829 (@aaristov) β GPU-vs-fallback counters and chat status made visible in
coli serve: the tier's behaviour is now observable instead of inferred. - #1095 (@benmaster82) β OLMoE planner geometry adapter, and #1103
(@SebaWag) β Kimi K3, Inkling and DeepSeek V4 adapters with 23 tests: every
family now has real planner geometry, socoli planstops guessing (#1066). - #1096 (@terrizoaguimor) β DeepSeek V4 serve framing on the shared codec,
completing the codec migration across OLMoE, Kimi K3 and V4. - #1063 / #1068 (@terrizoaguimor) β model families are registry-owned:
coli, the gateway,doctorand the planner read one descriptor table. - #1087 / #1090 / #1096 / #1116 (@terrizoaguimor) β a shared serve framing
codec, now adopted by every engine: OLMoE, Kimi K3, DeepSeek V4 and
Inkling (whose audio payload rides as an opaque extension). Each migration
landed behind a byte-exact wire-transcript freeze, so the gateway contract is
provably unchanged. Byte framing had been duplicated five times, which is how
Windows binary mode silently disappeared from sibling engines (#748). - #1036 (@lineape) β distributed expert workers (LAN, opt-in via
CLUSTER_WORKERS). - Planner: DeepSeek V4 expert naming now recognized, so
coli planand
coli doctorstop counting every routed expert as dense (fixes #1110).