Conversation
olmoe.c was benchmark-harness-only (ref.json prompt/full ids in, token ids out) — no tokenizer wiring, no way to actually converse with it. Adds: - CHAT=1 env var on main(), bypassing the ref.json harness entirely - tok.h/sample.h wiring (tokenizer load, temperature/top-p sampling, the tokenizer's own special-token set arms the stop condition automatically) - run_chat(): allocates the KV cache once for CTX tokens (default 4096, capped there since attention()'s per-head score buffer is fixed-size) and never reallocates it -- each turn after the first reuses the previous turns' cached keys/values instead of re-processing them, so this is real multi-turn context, not a fresh generate() call per message - OLMoE-Instruct's actual chat template (<|user|>/<|assistant|>, with bos_token/eos_token both being the tokenizer's "|||IP_ADDRESS|||" token -- a genuine artifact of its tokenizer, not a bug here) - /reset to clear context without restarting the process - chat_olmoe.sh: a memory-safe launcher (systemd-run --scope with MemoryMax + MemorySwapMax=0) so a misbehaving run degrades to a clean self-kill instead of taking the whole machine down Validated with a real multi-turn conversation (two follow-up coding questions) -- second turn correctly used context from the first without re-explaining itself, confirming the KV-cache carry-over works.
The Metal GEMV shader (mm_gemv, backend_metal.mm) handled fmt 0-3 only. Anything else -- including dev's grouped-int4 format (fmt=4: packed int4 nibbles, same layout as fmt=2, but one f32 scale per `gs`-element group along I instead of one scale per row; CPU reference matmul_i4_grouped in quant.h, active CUDA tier in #298/#451) -- fell through the shader's unconditional `else` arm and was decoded AS RAW F32. This was a silent-wrong-answer defect, not just a missing- acceleration gap: bind_gemv (used by the fused decode-attention kernels, coli_metal_attn_decode/coli_metal_layer_decode) accepted a per-weight fmt from its caller with no validity gate at all, so any fmt=4 dense/attention tensor reaching that path would "succeed" (ok=1) while silently computing garbage. coli_metal_matmul/coli_metal_gemm had explicit fmt<=3 gates and so safely rejected fmt=4 (0/CPU-fallback); bind_gemv had no such safety net. COMMIT ORDERING (bisect-safety, see PR_BODY.md sec 1 for the full argument): this commit lands FIRST, before the qt_from_disk allocator fix (next commit). At this commit alone, the capability is CORRECT but INERT: qt_from_disk's fmt=4 scale buffer is still allocated with a bare falloc() (unregistered, not page-aligned), so resolve() in backend_metal.mm can never find it, and every real fmt=4 dense/ attention tensor's scale buffer permanently misses bind_gemv/ coli_metal_gemm regardless of what this commit adds to the shader -- safe CPU fallback, not a wrong answer. No commit in this branch is ever "the shader is fixed but a scale buffer resolves to the wrong/stale allocation" -- fixed-but-unreachable, then reachable-because-fixed, in that order, never the reverse. Verified empirically at this exact commit (this branch, this session): both METAL=0/METAL=1 build clean, 0 warnings, and `make metal-test` passes 27/27 -- metal-test's fmt=4 cases construct their own tensors directly via posix_memalign + coli_metal_register (never through qt_from_disk), so none of them depend on the allocator fix landing first; the capability this commit adds is exercised and proven correct here, just not yet reachable from a real loaded model. Shader: unlike fmt 1-3 (single `acc * scale[o]` after the full dot product), fmt=4's group scale must be folded into the accumulation per-group, before summing across groups -- this is matmul_i4_grouped's actual math, not a simplification. New `constant int& gsz [[buffer(9)]]` parameter carries the group size (ignored, harmlessly, for fmt!=4). Each of the 32 SIMD lanes owns one packed byte (2 elements) per 64-wide stride -- memory access stays coalesced and a group never splits a byte (gs is always even, per the loader's candidate list). Deliberately not vectorized like fmt=2's uchar4/float4 path -- see PR_BODY.md sec 7 for why that tradeoff was made on purpose. Host plumbing: gs threaded through every path that reaches mm_gemv -- coli_metal_matmul, coli_metal_gemm, and bind_gemv (-> AttnW -> coli_metal_attn_decode / coli_metal_layer_decode) -- plus the colibri.c call sites (matmul_qt_ex's Metal-GEMM gate now includes fmt==4; attention_rows/layer_forward_rows pass .gs alongside each .fmt already being passed, a no-op for every model not using fmt=4). fmt_bytes and the new fmt_scale_bytes fix coli_metal_matmul's scale-buffer wrap size, which was previously hardcoded to O floats (wrong for fmt=4's O*ceil(I/gs)). kv_b (MLA absorption core) and the batched routed-expert MoE path (moe_gemv/coli_metal_moe_block) are deliberately untouched -- both already gate fmt=4 to a safe CPU fallback today, and extending them is a structurally separate kernel family; see PR_BODY.md sec 7 for the full scope statement. Tests (tests/test_backend_metal.mm): new cpu_ref_grouped (double- precision oracle mirroring matmul_i4_grouped, same construction tests/test_i4_grouped.c already established) and run_grouped harness, using a magnitude-relative tolerance (justified in PR_BODY.md sec 5) instead of the existing ymax-relative convention, since a grouped dot product's cancellation can otherwise misreport as a kernel defect. Covers the spec's shape matrix (I mult/non-mult of 64, I<=64 degenerate, S=1/S>1, outlier-heavy rows) via coli_metal_matmul, one case via coli_metal_gemm (the large-batch path), and two via a new run_attn_grouped harness that exercises the real fused-attention bind_gemv path end-to-end (not just the standalone kernel entry point). (These two attn cases have known blind spots as first committed here -- closed two commits later, see that commit's message and PR_BODY.md sec 10; kept as originally written here for bisect fidelity.) make metal-test: 15/15 (stock e9b3614) -> 27/27 (12 new fmt=4 cases; original 15 unchanged -- the negative control). make glm METAL=0/ METAL=1: clean, 0 warnings, both before and after. make test-c / test-python: unaffected, all pass. Full gates, evidence classes, and DEVIATIONS/UNCERTAINTIES in PR_BODY.md (untracked, worktree root). (cherry picked from commit f0fa5a9) (cherry picked from commit 82984cdeef67ede71d48b32f2ddd4734f4622eba)
Latent defect, found while tracing how per-row/per-group scale buffers reach the Metal backend (prep work for wiring fmt=4 grouped-int4 into the Metal GEMV, previous commit). qt_from_disk's fmt=4 branch allocated the [O,ceil(I/gs)] scale array with falloc() -- a plain malloc -- while every other format (fmt 1/2/3, right above and below this branch) uses qalloc(), which page-aligns the allocation and coli_metal_register's it under COLI_METAL. Failure mode, and why this commit is the ACTIVATION step, not an independent capability: backend_metal.mm's resolve() only finds pointers inside coli_metal_register'd slabs. Before this commit (previous commit's shader/host-plumbing work alone), a fmt=4 tensor's falloc'd scale buffer could never be resolved, so bind_gemv/ coli_metal_gemm always returned false/0 and silently CPU-fell-back -- the correct shader from the previous commit was unreachable dead code for any real fmt=4 dense or attention tensor. This commit is what makes it reachable: after this fix, resolve() succeeds, and the previous commit's already-correct fmt=4 shader branch actually gets exercised on real model weights for the first time. COMMIT ORDERING (bisect-safety, see PR_BODY.md sec 1): this fix lands SECOND, after the shader/host-plumbing commit, specifically so that no commit in this branch is ever "the allocator resolves fine but the shader doesn't know what to do with fmt=4" -- the dangerous ordering, where a scale buffer starts resolving successfully into a codepath that still silently mishandles it as f32. That state never exists in this branch's history; this commit only ever activates an already-correct consumer. Verified: `make glm METAL=1`/`METAL=0` clean, 0 warnings; `make metal-test` 27/27 (unchanged from the previous commit -- metal-test's fmt=4 cases never went through qt_from_disk in the first place, so this allocator fix doesn't change their outcome, only real model loading); `make test-c`/`test-python` pass. (cherry picked from commit 2ba12d8) (cherry picked from commit 880fd7eb1019c01b2bc9692ecfc09640cfbd951e)
Validator finding, both proven, both fixed: (a) The S=1 case used pos_base=0, giving T=pos_base+S=1. Softmax over a single key is identically 1.0 regardless of the score's value, so that case's final output (got vs ref, nerr/cache) was provably independent of q_a's kernel entirely -- not merely insensitive to some subset of defects, blind to ALL of them. Fixed by moving the S=1 case to pos_base=37, mirroring run_attn's own non-grouped S=1 pos=37 case, which exists for the identical reason. (b) Both cases fed q_a's raw GEMV output through RMSNorm before comparing anything downstream. RMSNorm(c*v) == RMSNorm(v) for any positive scalar c, so a whole-tensor uniform-scale-calibration bug in q_a's grouped-int4 kernel would be invisible to the fused-path comparison no matter how pos_base is chosen -- (b) is independent of (a); fixing pos_base alone does not fix this. Fixed by capturing q_a's raw GEMV output before t_rms overwrites it in place and comparing it directly against the CPU oracle (cpu_ref_grouped, same magnitude-relative construction run_grouped() uses) via a standalone coli_metal_matmul call on the exact weight/scale/x data the attention test generated -- reported as qraw=... in the printed line, folded into that case's existing pass/fail (case count stays at 27/27, these two cases are strengthened in place, not duplicated). Verified by re-deriving the validator's three mutation classes (off-by-one group index; swapped nibbles; uniform 1% scale error) as temporary edits to mm_gemv's fmt==4 branch, run against four diagnostic configurations isolating fix (a) from fix (b) -- confirmed the OLD configuration (S=1 pos=0, no qraw) catches none of the three, fix (a) alone catches 2/3 (misses the uniform-scale class, exactly as the RMSNorm-invariance argument predicts), fix (b) alone catches 3/3 regardless of pos_base, and the shipped fix (a+b) catches 3/3. Full detection matrix in PR_BODY.md sec 10 (untracked, worktree root). The diagnostic harness used to produce that matrix was temporary and is not part of this commit; backend_metal.mm is unchanged by this commit (the shader mutations used for verification were applied and reverted before this commit, confirmed byte-identical by diff). COMMIT ORDERING: lands third/last, after both the shader/plumbing commit and the allocator-activation commit, following the same bisect-safety logic applied to this branch as a whole (PR_BODY.md sec 1) -- this is a test-only strengthening with no interaction with either of the other two commits' reachability arguments; it lands last simply because it is a review-round finding on top of an already-complete Stage 1, not because ordering safety requires it here specifically. make metal-test: 27/27 (unchanged case count, strengthened assertions). make glm METAL=0/METAL=1: clean, 0 warnings. make test-c/test-python: unaffected, all pass. (cherry picked from commit fc0f5a5) (cherry picked from commit 6a88a645fa4ba0d25ba6ee4dfddb9f849c06b059)
…oc->qalloc trade-off - test_gemm_largebatch.mm predates the 9-arg coli_metal_gemm prototype; gs=0 selects per-row scales, the exact pre-fmt=4 semantics its fixtures assume. - The fmt=4 scale-buffer qalloc site is outside #ifdef COLI_METAL: on CPU-only builds it trades falloc's checked-exit for qalloc's unchecked malloc, consistent with qsalloc for fmt 1/2/3 -- now said in the comment, with a note that fmt=5's group scales still use falloc (Metal-inert, pre-existing).
Off-by-default research instrument: ABLATE_SCORE=<manifest> runs a teacher-forced prefill per item with a chosen (layer,expert) cell ablated in one of three modes (contribution-zero / route-around / module-swap) and reads out the final logits (ABLATE_OUT, coli-ablate/1). Inert unless enabled (guarded by g_abl.mode); standalone unit gate tests/test_ablate proves inertness + per-mode semantics (18/18). Causal complement to the Expert Atlas (#175).
backend_vulkan.c/.h + shaders/qmatmul.comp: opt-in Vulkan int4/int8 quantized GEMV via RADV, mirroring coli_cuda_matmul. Shader decodes int4 as nibble-8 (offset-binary), numerically consistent with the CPU path. Validated on the RX 9070 (RADV GFX1201, Mesa 26.1): the built-in VK_TEST harness passes all int4/int8 cases vs the CPU reference (maxrel ~1e-4). Bypasses ROCm's dropped Polaris support -> also targets the RX 580. Head-to-head naive int4 GEMV on the RX 9070 (same shapes, synchronous per-call): ROCm 0.115ms vs Vulkan 0.306ms (6144->1536, S=1) -- ROCm ~2.7x faster as-is (naive shader + host-visible memory + heavy submit overhead; not the tuned coopmat path that wins in llama.cpp). Next: optimize the shader/backend to reach or beat ROCm. Not yet wired into glm.c (Makefile/hooks pending).
Optimizations toward matching ROCm on the RX 9070: - backend: cache descriptor set + resubmit the prerecorded command buffer when tensor/shape/scratch are unchanged (single int4 GEMV 6144->1536: 0.306->0.214 ms). - shader: llama.cpp-style mul_mat_vec (MIT techniques, per-row int4/int8): x staged once in shared memory, one subgroup per output row + subgroupAdd (no barrier tree), grid-stride rows. Needs SPIR-V 1.3 (glslc --target-env=vulkan1.2). Correct (maxrel ~1e-4). - VK_TEST: batched throughput probe (N dispatches / one submit). Honest result (noise-controlled, vs colibri's production coli_cuda_expert_group): production ROCm ~0.179 ms/expert (fused dual gate+up + down, batched); our unfused VK ~0.260 ms/expert -- ~45% slower. VK wins the short-reduction down-proj, loses the long-reduction gate/up. Reaching ROCm needs fused dual gate+up + expert batching + a shape-adaptive reduction.
…pert New qmatmul_gate_up.comp + coli_vk_gate_up: computes hidden=silu(gate(x))*up(x) in ONE dispatch, reading x once for both projections (VK equivalent of colibri's grouped_hidden_w4_dual). Second 6-binding compute pipeline; build_pipeline() helper refactors the pipeline setup for both. Correct vs CPU ref (maxrel ~9e-5). The fusion is the win: gate+up drops from ~0.13 ms (two separate matmuls) to ~0.080 ms/expert (fair, 8 DISTINCT experts cycled so weights come from VRAM not L2 — added bench_experts_fair to control for the caching artifact). Full fused expert: VK gate_up ~0.080 + down ~0.050 = ~0.13 ms/expert ROCm production coli_cuda_expert_group = 0.179 ms/expert (stable) -> VK ~25-30% FASTER, holding under cache-controlled measurement (some run variance 0.07-0.10 on gate_up). Reaching this needed the fused dual projection, exactly the optimization the production HIP path already had and our unfused VK lacked.
…one submit) The real engine primitive: K experts, fused gate+up+silu then down, hidden staying on-device, all in one submit. Per-expert descriptor sets (gate_up: 6-binding, down: 4-binding) sliced into packed x/hidden/y via descriptor offsets; one phase barrier between gate_up and down. Mirrors coli_cuda_expert_group; correct vs CPU ref (maxrel ~7e-4 at K<=8). Honest throughput (RX 9070, K=8, distinct experts): GPU-only (resubmit recorded cmd buffer): 0.113 ms/expert -- BEATS ROCm 0.179 (~37%) per-call (as-is API): 1.03 ms/expert -- 5.6x SLOWER than ROCm The GPU compute is genuinely faster (fused kernel delivers), but the per-call HOST setup (80 descriptor updates + recording 16 dispatches every call) dominates. The earlier microbench 0.13 was the GPU-only number; it did not capture per-call cost. Fix (next): cache descriptor sets + command buffer across calls (hot experts are reused across tokens), or a bindless/BDA single-dispatch-multi-expert design like ROCm's grouped kernels. Either drops per-call toward the 0.11 GPU-only floor. (K=32 maxrel 2e-3, slightly over 1e-3 threshold -- fp accumulation, worth a look.)
…s ROCm The per-call cost (1.03 ms/expert) was ALL in reading the output back: eg_y/y/h were allocated in write-combined DEVICE_LOCAL (ReBAR) memory, which the CPU reads at ~40 MB/s. VK_PROF breakdown (K=1): memcpy_x 0.005 | desc 0.027 | record 0.005 | gpu 0.19 | memcpy_y 0.598 ms -- the readback dominated; descriptor updates + recording were negligible (so no descriptor/command caching needed). Fix: pick_memtype_cached() (HOST_VISIBLE|HOST_COHERENT|HOST_CACHED) for buffers the CPU reads back (eg_y, y, h); inputs/hidden stay write-combined (fast CPU writes / GPU-only). Result (RX 9070, K=8, distinct experts): per-call expert_group 1.03 -> 0.117 ms/expert, now == the 0.111 GPU-only floor. vs ROCm 0.179 -> VK ~35% FASTER end-to-end, per-call. The real primitive now beats ROCm; hypothesis confirmed. (K=32 maxrel 2e-3 is fp32 precision on the 6144-elem reduction, fine for greedy argmax.) VK_PROF=1 env-gated.
make glm VK=1 compiles backend_vulkan.o (plain C + vulkan headers), builds the .comp shaders to .spv via glslc (--target-env=vulkan1.2), links -lvulkan. Independent of CUDA/HIP (own -DCOLI_VULKAN). Verified on the RX 9070: glm links libvulkan.so.1, both shaders compiled. glm.c hooks (COLI_VULKAN) come next; default build unchanged.
End-to-end integration of coli_vk_expert_group into the engine: - QT gains a resident ColiVkTensor *vk (+vk_eligible); qt_vk_reset frees it when a slot is reused for another expert (expert_load_impl hook, mirrors qt_cuda_reset), so the LRU never computes with stale weights. g_vk_resident tracks the tier size. - coli_vk_init at startup (COLI_VULKAN=1), shader path COLI_VK_SHADERS. - moe() VK path (decode S<=4): upload routed int4 experts to VK once (capped by COLI_VK_EXPERTS, default 1024), compute the resident ones as one batched coli_vk_expert_group (fused gate+up+silu -> down, on-device), CPU-fallback the rest. - coli_vk_tensor_ensure() backend entry: upload a resident tensor without computing. Verified on the RX 9070 (make glm VK=1): [VK] expert tier active, greedy decode of 'The capital of France is' -> 'Paris.' CORRECT. Default build (no VK=1) unchanged. NOTE: a VK-only build has no GPU dense/attention offload (that is the HIP CUDA_DENSE/ COLI_CUDA_ATTN path), so end-to-end tok/s here is CPU-dense-bound (0.06 cold), NOT a measure of the expert path -- which is the ~35%-faster-than-ROCm primitive. A hybrid HIP-dense + VK-experts build, or a VK dense/attn port, is the next step for throughput.
…d expert on Vulkan vk_matmul_qt(t,y,x,S) routes a resident int4/int8 matmul_qt through coli_vk_matmul (uploads the weight once into t->vk, then reuses it). Wired, with CPU fallback, into the decode attention projections (q_a, q_b, kv_a, o) and the shared expert (sh_gate/up/down). Env COLI_VK_DENSE=1. Together with the expert tier, a Vulkan-only machine now runs experts + dense projections + shared expert on the GPU; only the MLA attention core (absorb/softmax/ RoPE, small in latent space) and I/O stay on CPU. Verified on the RX 9070 (make glm VK=1, COLI_VULKAN=1 COLI_VK_DENSE=1): greedy 'The capital of France is' -> 'Paris.' correct. Default build unchanged. NEXT for a fully-GPU Vulkan path: a dedicated MLA attention-core compute shader (scores/softmax/weighted-values/RoPE) — the last CPU-bound piece.
…te piece on Vulkan-only machines New shaders/attention_absorb.comp runs the whole decode absorb core for one layer in ONE dispatch, one workgroup per (query row, head): absorbed query (q_nope through the int4/int8 kv_b nope rows), scores over the cache window, softmax, weighted latent, and the value-row projection. Subgroup-per-token score dots + subgroup-per-row value projection, q/qabs/clat staged in shared. The KV latent/rope cache is mirrored in persistent per-layer device buffers (coli_vk_kv_ensure/_row/_reset), appended ~2.3 KB/token/layer instead of re-uploading the window each call — same design as the CUDA kv_dev shadow, with the same invalidation points (row rewrite, kv_bind, kv_alloc resize) tracked by a vk_kv_valid watermark in glm.c. Falls back to CPU on DSA top-k selection, ragged KV, the MTP layer, or any backend failure (mirrors COLI_CUDA_ATTN's guards). Opt-in via COLI_VK_ATTN=1. Validated on RX 9070 (RADV): 5/5 CPU-ref harness cases maxrel <= 1.7e-4 (GLM decode shape, kv_start window, S=2 causal, int8, T=2000); engine output byte-identical to the pre-change build under the same env; decode score-softmax-value 1.75s -> 0.61s per 16 tokens (2.9x) with prefill untouched. Remaining VK perf work: resident-on-device layer pipeline and a shape-adaptive long-reduction for the o-projection.
…he parallel CPU loop) Previously budget 0 still entered the vk_active block, sending every routed expert through its serial CPU fallback. With the gate, dense+attention can run on Vulkan while routed experts keep the normal parallel CPU path — measured the best Vulkan-only config on the RX 9070 (1.31-1.37 tok/s vs 1.46-1.50 HIP; the GPU expert tier as integrated is slower, 1.11, because uploads and submits sit serially on the decode critical path). Bench note: Vulkan-only runs on this box need COLI_NO_OMP_TUNE=1 — the OMP self-tune (skipped under COLI_CUDA, so HIP never hit it) sets OMP_WAIT_POLICY=active + GOMP_SPINCOUNT=200000, and 12 pinned spinning threads starve the PIPE I/O pool and pilot worker (CPU expert rows 28 -> 5 GB/s, 0.52 tok/s). A standalone probe confirmed VK init itself is harmless.
…pert (1 submit each)
CORRECTNESS: qmatmul.comp staged x into shared xsh[6144] unconditionally; the
o-projection's input row is H*vh = 16384, so every VK dense o-proj since the
dense port computed with a truncated/undefined activation tail — deterministic,
so greedy output was plausible and stable, which masked it. Rows with
I > 6144 now skip staging and read x from the storage buffer (uniform branch,
coalesced; harness case fmt=2 I=16384 O=6144 added, maxrel 2e-4). With the fix
the VK build's greedy output matches pure CPU exactly ('Paris.') — the earlier
divergence was this bug, not fp tie-breaks. gate_up/expert_group get D<=6144
host guards (their shader shares the pattern; engine dims are within bounds).
PERF (toward HIP parity):
- coli_vk_attention_absorb_project: absorb + resident o-projection in ONE
submit, ctx stays on-device (was: absorb submit + ctx readback + o submit).
Harness: fused 0.51-0.53 ms/call vs 0.73 unfused absorb ALONE, maxrel <= 7e-5.
- Shared expert now runs as coli_vk_expert_group(count=1): fused gate+up+silu
-> down, hidden on-device, one fence instead of three matmul submits.
- expert_group harness threshold 3e-3 for K>=32 (documented fp32 accumulation
on the 6144-length double reduction; was a standing false FAIL at 2e-3).
Both projections read the same x row(s): coli_vk_matmul_pair stages x once,
records both dispatches, and waits one fence — replacing two full submit+wait
roundtrips per layer per token. Harness-validated (maxrel 4.4e-5, 0.123 ms/pair
vs ~0.35 for two singles); engine greedy still matches pure CPU ('Paris.').
Also: expert_group harness threshold 3e-3 at any K (the fp32 accumulation
lives in the 6144-length reduction chain, not in the expert count).
…ch overlapped with CPU rows Replaces the LRU-slot-tied tier (which uploaded on the decode critical path, churned with evictions, and computed ALL routed experts on the GPU while 12 cores idled — measured slower than CPU-only experts). - vk_registry_fill(): at startup, upload the top-COLI_VK_EXPERTS experts by persistent usage history into a (layer,eid) registry, decoupled from cache slots — stable residency like the HIP VRAM tier (RX 9070: 256 experts, 4.83 GB, 1.0s; pinned RAM slots feed uploads directly, the rest stream through one transient slot). - coli_vk_expert_group_issue/_take: the group gets its own command buffer + fence and splits into submit-and-return / join, so moe() ISSUES the GPU batch, computes the CPU share concurrently, drains the pipe + loads the GPU-side slots (cache invariants unchanged), then takes and accumulates. Sync coli_vk_expert_group (shared expert, harness) wraps the same path. - vk_active now keys on a non-empty registry; empty (no history/budget 0) falls back to the normal parallel CPU loop. Harness: issue/take validated bit-identical to the sync path; full PASS. Engine: 'Paris.' correct with the tier active.
…etch entirely The resolve phase now classifies registry-resident experts FIRST (before pin/LRU): they take no cache slot, dispatch no load, and skip the LRU recency bump — so they age out of RAM and free capacity for the CPU-served experts, the same effect CUDA_RELEASE_HOST gives the HIP tier. The pilot worker, cross-layer lookahead, and next-block readahead treat registry residency like RAM residency (decode only; prefill still loads normally since the tier serves S<=4). Counted as a new 'vk' bucket in the hit-rate split. The group block computes GPU entries straight from the registry (no ESlot), attributes its CPU share to t_ecpu/rows and only the take-wait to t_egpu, and on device-loss falls back via a transient ws slot load.
…ath was inverted Without a RAM pin every candidate loads transiently, and the inverted check skipped every SUCCESSFUL load (pin-fed fills masked it: the pins covered the whole top-256). PIN_GB=0 + COLI_VK_EXPERTS=256 now fills from disk in ~1.6s. Also: bail after 64 consecutive load failures, diagnostics on the first few.
256-384 measured flat within noise (1.73-1.78 tok/s); 320 is the best median with ~2.3 GB VRAM headroom left for the long-context KV mirror.
Upstream PR #399 (KV8 fp8 / TQ4 quantized latent KV) leaves Lc/Rc NULL when a quantized tier is active — the mirror sync would deref NULL. Guard on the arrays themselves (env-independent), falling back to the CPU attention path; an fp8-aware VK mirror (upload Lc8+scale, decode e4m3 in the shader, 4x less mirror traffic) is the follow-up once #399 merges.
Upstream #168 (merged as 5e42e70) now carries the engine int3-g64 support this commit used to port, textually near-identical — including the uring_finalize_load qt_resolve_fmt fix we carried separately. What remains ours is the Vulkan side: vk_matmul_qt/pair accept fmt 5, the absorb kv_b/o and shared-expert paths feed fmt-5 tensors, and vk_registry_fill passes the true fmt through xf/d.fmt instead of assuming int4.
qmatmul/qmatmul_gate_up/attention_absorb learn the 24B-per-64-group two-plane layout: lanes stride the 16-value low-plane words, each word's partial dot is multiplied by its group scale before the subgroup reduction, and the per-row scale multiply is skipped. The absorb query pass stops pre-folding the kv_b row scale into q for fmt=5 (scales vary along K) and applies the (row,group) scale per element instead. upload_tensor accepts fmt=5 (rows are already word-aligned: ceil(I/64)*24), sizes the scale buffer O*ceil(I/64) via scale_floats() mirrored in tensor_free's accounting, and the expert-group prep validates up/down fmts (down may differ per-projection; phase 2 pushes downs[0]->fmt). Engine gates opened: vk_matmul_qt/pair, absorb kv_b+o, shared expert, and the pinned tier registry (int4/int3 accepted, fmt passed through). VK_TEST: fmt=5 cases across dense (incl. tail group + unstaged o_proj shape), fused gate_up, expert group K=8/32, absorb (+causal window), int3-vs-int4 fair throughput lines. Shaders validated against glslangValidator vulkan1.2.
- sample.h: COLI_LOGIT_DUMP=1 prints top-5 (id:logit) per pick_tok step — the tool that separated backend error from fp tie-flips (VK matched CPU logits to ~1e-4 at every matched step; run-to-run token divergence turned out to be engine-wide threading jitter, present on int4 CPU-only too). - harness: fmt=5 matmul_pair case (the q_a+kv_a decode path) and count=1 expert group (the shared-expert shape) — the two production paths the first harness round left uncovered. - tier messages made format-neutral (int4/int3-g64).
docs(api): OpenAI server supports tool-calling
Resolved by a maintainer instead of asking for another rebase. The only conflicting hunk was the hand-written TEST_BINS line, which dev no longer has: gates are derived from the build rules, so this branch's tests are picked up by their own rules and the manual list entry is dropped. No other file conflicted and no commit on this branch was rewritten. Verified before pushing: every test this branch adds a rule for is in the gate set, and make test-c passes. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Resolved by a maintainer instead of asking for another rebase. The only conflicting hunk was the hand-written TEST_BINS line, which dev no longer has: gates are derived from the build rules, so this branch's tests are picked up by their own rules and the manual list entry is dropped. No other file conflicted and no commit on this branch was rewritten. Verified before pushing: every test this branch adds a rule for is in the gate set, and make test-c passes. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
formats: self-describing container stamp (TRUST-VERIFY-REFUSE) + FORMATS.md registry (#524)
Add opt-in per-expert causal-ablation harness (ABLATE_SCORE) for moe()
metal: fmt=6 (E8/IQ3) expert decode — GPU path + on-GPU FWHT rotation
…lback-dev serve: trigger Claude Code system-role fallback
The only conflict was the handler dict in c/coli: this branch adds "mirror":cmd_mirror, dev added "tune":cmd_tune in #673. Both are real commands with their own cmd_ function, so both entries are kept -- picking a side would have silently deleted a working subcommand. Verified: coli --help renders, and both 'coli mirror' and 'coli tune' resolve. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
feat: add learned partial mirror planner
dev went red on macOS only, in the test added by #536: AssertionError: PosixPath('/private/var/folders/.../same.safetensors') != PosixPath('/var/folders/.../same.safetensors') On macOS /var is a symlink to /private/var, so tempfile hands back /var/folders/... while mirror_plan.py resolves every path it is given (discover_shards line 115, create_plan lines 197-199). The test compared a resolved path returned by the tool against an unresolved one it built itself. Linux has no such symlink, which is why it passed there -- including in my own pre-merge run, which is how this reached dev. Resolving the temp root in setUp makes every derived path resolved, so both sides match on every platform. Test-only; mirror_plan.py is unchanged. Verified by reproducing the macOS condition on Linux with a symlinked temp dir: the old comparison fails against it and the new one passes. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
fix(test): resolve the temp root so test_mirror_plan passes on macOS (dev is red)
…ngine
inkling and kimi_k3 did not compile under msys2/UCRT64 at all. Their serve loops
polled stdin with fd_set/select directly, which Windows does not provide in that
form:
inkling.c:1396: error: unknown type name 'fd_set'
kimi_k3.c:1431: error: unknown type name 'fd_set'
This is why binary releases have only ever shipped the GLM engine (#720): a user
who downloads v1.3.0 and points it at Kimi K3 finds no engine, and the README
promises four model families.
colibri.c already solved this, and its Windows path carries two bug fixes:
#139 (select() on a pipe handle routes to winsock and always returns
SOCKET_ERROR, so the loop never accepted a request) and #195 (anonymous pipes
are not waitable objects, and PeekNamedPipe fails on file/console handles).
Copying that a third and fourth time would have reintroduced both in engines
where nobody would look for them, so it moves to compat.h as
coli_stdin_readable() -- already reachable from both engines via st.h.
The engines keep their own function names and call sites; only the three-line
body changes. colibri.c is deliberately NOT refactored onto the shared helper:
it works today, 14 open PRs touch that file, and the churn would cost
contributors a rebase for no functional gain.
CI now builds all four engines on Linux, macOS and Windows -- the matrix the
release archives are built for. Before this it built colibri + inkling on Linux
only and kimi_k3 on nothing, which is why the gap stayed invisible for months.
Verified: all four engines build on Linux with the change.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…endency Two mistakes in the previous commit, both caught by CI on Linux. 1. The '#include <sys/select.h>' landed INSIDE the '#ifdef __APPLE__' block at the top of the file -- the string replace matched the first '#include <sys/types.h>', which happens to be Apple-only. So on Linux neither sys/select.h nor unistd.h came in from here, and test_route_trace failed with 'STDIN_FILENO undeclared'. The include now sits immediately above the function it serves, at file scope, guarded only against Windows. 2. STDIN_FILENO comes from <unistd.h>, which this header does not include on every platform. Using the literal fd 0 removes the dependency entirely -- and is what both engines already did before this change. stdin is 0 under POSIX. Also corrected the header's own top comment, which claimed compat.h is 'a total NO-OP on Linux'. That stopped being true the moment a portable helper needed a POSIX branch, and a comment that lies is worse than no comment. The reason it is right anyway is stated there: a portable helper must exist on every platform, or the .c files need their own #ifdef -- which is exactly what this header's rule forbids. All four engines build and make test-c passes, including test_route_trace. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
fix(win): inkling and kimi_k3 build on Windows; CI covers every engine on every release platform
Binary releases contained c/colibri alone. A user who downloaded v1.3.0 and pointed it at Kimi K3 found no engine at all -- while the README's front page promises four model families. That is #720: nothing MichaelFomenko could have typed would have worked. Now builds and packages all four engines, plus tools/k3_tokenizer.py. Kimi K3 ships tiktoken.model rather than tokenizer.json, and that script synthesizes it; without it the engine would be in the archive but coli chat still could not drive it -- the second half of the same issue. No launcher change needed: engine_for() in c/coli already resolves the engine from the model's config.json and looks for kimi_k3/inkling next to itself, which is the layout an archive already has. Cost: 863 KB for all four (colibri 423K, kimi_k3 179K, inkling 154K, olmoe 128K). The verification step is extended rather than trusted. It already unpacked the archive and asserted coli finds the engine, because a packaging mistake is invisible to every other job here -- green build, green tests, unusable artifact. It now asserts the three sibling engines are present and executable, that k3_tokenizer.py is packaged and parses, and that each engine sits where the launcher's own resolver looks for it. This is only safe because CI now proves all four engines build on Linux, macOS and Windows (engines-all-platforms, added with the Windows serve-loop fix). Before that, inkling and kimi_k3 did not compile on Windows at all. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
feat(release): ship every engine in the archives, not just GLM (closes #720)
kimi_k3 and olmoe had no OpenMP tuning of any kind, while colibri.c and inkling.c did. On an SMT machine their team runs at logical-core count, and #718 measured what that costs: 2.3x on a Ryzen 9 5950X (16C/32T) from the thread count alone -- large enough to swamp most tuning deltas people quote. They get the SIZING half only, not the spin-wait half, and that split is the whole point of this change: sizing OMP_NUM_THREADS = physical cores #718: +2.3x on Zen3. spin-wait OMP_WAIT_POLICY=active, GOMP_SPINCOUNT, KMP_BLOCKTIME #707: -2.2x decode on a low-residency host (M1 Max, ~10% resident) #116: -39% on Metal. #159: ~3x on x86+CUDA. #341: 3000% CPU on FreeBSD with an idle team. Kimi is the most disk-bound engine here -- measured at 6.7% expert hit and 891 GB streamed for 32 tokens -- which is exactly the regime where a spinning team starves the I/O pool doing the real work. Giving it the whole block would have been a measurable regression, so it gets the half that helps. No re-exec needed: colibri.c re-executes because OMP_WAIT_POLICY and friends are read by libgomp's constructor before main(), but omp_set_num_threads() is a runtime API. The safe half is also the simple half. Silent and inert unless it helps: does nothing when the user set OMP_NUM_THREADS, when COLI_NO_OMP_TUNE is set, when there is no SMT to avoid, or when the physical-core count cannot be determined -- never a guessed number (#325 is what a silent fallback to 1 costs). macOS reads hw.perflevel0.logicalcpu before hw.physicalcpu: on Apple Silicon the latter counts E-cores too, and with a barrier per matmul the slowest thread paces the team (#707, -4.2%). Intel Macs have no perflevel*, where physicalcpu is right. colibri.c and inkling.c are deliberately NOT touched: they are tuned already, 14 open PRs touch colibri.c, and the churn would cost a rebase for no gain. Verified: all four engines build; detection returns 6 physical of 12 logical here, matching lscpu; the line is suppressed under OMP_NUM_THREADS and COLI_NO_OMP_TUNE; make test-c passes. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
feat(omp): size the OpenMP team by physical cores on kimi_k3 and olmoe (#718)
…e value The OMP tuning guard tested getenv() for PRESENCE on four variables. That is right for two of them and wrong for the other two. COLI_OMP_TUNED and COLI_NO_OMP_TUNE are kill-switches, and presence-based is deliberate there -- c/coli documents it: "impostarla a qualsiasi valore, anche 0, disattiva". Unchanged. COLI_CUDA and COLI_METAL are STATE, where 0 means "no GPU". Testing presence meant the tuning was skipped for exactly the users who have no GPU and most need their CPU team configured. It is reachable rather than theoretical, and `coli` itself triggers it: the launcher writes e["COLI_CUDA"]="0" at c/coli:281 and :360 when the plan does not use CUDA. So on a CPU-only Linux host, a documented `coli run --auto-tier --gpu none` silently lost the tuning that a bare ./colibri receives. Measured before and after on this box: before (dev): COLI_CUDA=0 -> no tuning <- the bug after: COLI_CUDA=0 -> tuning runs after: COLI_CUDA=1 -> skipped (correct, #159) after: COLI_METAL=0 -> tuning runs after: COLI_NO_OMP_TUNE=0 -> still skipped (kill-switch, unchanged) coli_env_on() treats 0/false/off/no as off, so the fix covers a user writing COLI_CUDA=false as well as the launcher's "0". Reported by @fredchu in #707, who read it out of the source and said plainly he could not measure it -- he has no Linux host. The measurement above is that half. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…wrong
The site said colibri runs GLM-5.2. It has run four families since v1.3.0, and
two of the model cards contradicted the project outright:
Inkling 975B MoE - Planned -> Live (docs/inkling.md ships; runs on 25 GB)
Kimi K2 1T MoE - Planned -> Kimi K3, 2.8T MoE, Live
Telling visitors that Inkling and Kimi are on the roadmap, while the README
front page says both run today, is the kind of contradiction someone finds in
thirty seconds.
Hero rewritten rather than merely widened. The old line worked because it put
two incompatible things next to each other -- an enormous model, your machine.
Replacing that with a range ('744B to 2.8T') informs and stops landing; the
contradiction was the message. It now reads:
These models do not fit in your machine. They run in it anyway.
Same rhetorical shape as the copy further down the page ('Weights are not state
to hold. They are data to stage.'), so the page speaks with one voice. The
subtitle now also explains WHY it is possible -- a MoE token touches a small
fraction of the weights -- which was missing entirely and is what turns an
unbelievable claim into an understandable one. The numbers move there, where
they serve the reader who wants detail instead of the one who is skimming.
DeepSeek and Qwen3 deliberately stay 'Planned': #165 and #712 are not merged,
and the site should not promise what the code does not do.
Text only. No CSS, structure or script changes -- the sole markup edits are the
two cards' buttons becoming real links now that both models are runnable.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
fix(omp): COLI_CUDA/COLI_METAL are state, not kill-switches — test the value (#707)
site: four model families, not one — and two cards that were plainly wrong
JustVugg
changed the title
pin_fill loaded the ENTIRE VRAM-ranked prefix into host RAM in one parallel pass, then entered the upload loop that releases each slab. So CUDA_RELEASE_HOST worked exactly as documented and could not help: by the time the first release ran, every slab was already resident. Peak host RSS was the whole CUDA_EXPERT_GB budget, whatever that budget was. @mcollinswisc hit this on a machine with MORE VRAM THAN SYSTEM RAM -- 96 GB on an RTX PRO 6000 against 64 GB of RAM -- and reproduced it for everyone else with a cgroup: CUDA_EXPERT_GB=40 under MemoryMax=24G is OOM-killed during the pin fill, exactly where their log stops. Their hypothesis was right and is what this commit acts on. The code's own comments stated the intended contract and show the gap: 'Load the VRAM-ranked prefix first. Once uploaded its host backing is released before the disjoint RAM-ranked suffix is allocated.' 'npin+=prefix_est; additive: prefix RAM is returned after upload' Both true of the prefix AS A WHOLE. The design treats prefix RAM as transient, and it is -- transient in aggregate rather than per expert, and aggregate is the number that has to fit. It also explains why RAM_GB and RSS_GUARD_GB changed nothing for them: those bound the steady state, and this peak is a startup transient no steady-state budget models. Now staged in rounds: load a round in parallel, upload it, release it, next. Peak host staging becomes stage*expert_bytes instead of the full budget, and the parallel load keeps its bandwidth. Round size is min(4 GB, budget/8), at least one expert and never more than the prefix -- budget/8 keeps rounds large enough to saturate the loader threads, the absolute cap protects the little-RAM, big-budget case that is the report. Same bug family as the Inkling expert-cache defect: acquiring the whole working set up front when the working set is larger than what fits, where processing in rounds is both correct and bounded. Behaviour is UNCHANGED for everyone else. Without CUDA_RELEASE_HOST, or without CUDA, or with no VRAM prefix, stage == prefix and the loop runs exactly once, which is the previous code path. The accumulators (remaining/placed_b/placed_n, gpu_expert_bytes) already lived outside the loop, so they carry across rounds untouched. Verified: builds clean; syntax-checks with -DCOLI_CUDA; make test-c passes; the CPU path is byte-identical in structure with stage == prefix. NOT verified functionally: I have no CUDA hardware. The compile-level and non-CUDA evidence is what I can give; @mcollinswisc has the reproduction. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Third engine onto the shared telemetry from #700, after colibri.c (#716) and kimi_k3.c (#719). The counters, the load and the save move to route_trace.h; the bump sites are untouched, since m->eusage now aliases rt_counts_all(). The transition is the point: rt_load reads both the IKU1 block every previous inkling wrote and the shared text format, so an existing .coli_usage keeps working and is rewritten in the new form. The header checks the old pins_load did inline -- magic, n_layers, n_experts against the live config -- all still fire, now in the reader, which additionally refuses a history written by a different engine. Narrower than before, never wider; and parse geometry does not bend to a trusted path the way identity does. Fixes a latent bug on the way: pins_load documents PIN_N=0 as "seeds the ranking from the history but pins nothing", but the npin guard ran before the memcpy, so the counters started at zero and the next usage_save replaced the accumulated ranking with that run's counts alone. Verified on the tiny fixture: rebuild clean, oracle token-exact (36/36 teacher-forced, 24/24 generated), make check 273 passed / 34 skipped, and the built engine reading a planted legacy history pins the right 9 experts while dropping the dense layer's counts. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
fix(cuda): stage the VRAM tier in rounds so host RSS is bounded (closes #730)
make check was red on Linux, for a reason unrelated to what the test asserts. _cc_flags() adds -fopenmp only on darwin. Everywhere else it compiles colibri.c without it, so every '#pragma omp' becomes a -Wunknown-pragmas warning (-Wall turns that on) -- and the test then asserts stderr is empty: AssertionError: 'In file included from .../colibri.c:419 ...' != '' : e2e harness build produced warnings (production flags require zero) quant.h:99: warning: ignoring '#pragma omp parallel' [-Wunknown-pragmas] The assertion is right: a production-flags build should be warning-free. The flags were not production flags -- they were missing the one the Makefile has. Linux and the BSDs now get -fopenmp like the Makefile does; macOS keeps its existing libomp probe just below, and Windows/MinGW is left alone. The test passes rather than being weakened, so it still guards what it was written to guard. Confirmed pre-existing: the same failure reproduces on unmodified dev, so it is not fallout from anything merged today. make check: 288 tests, OK (skipped=13). Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
inkling: adopt route_trace.h, keeping the IKU1 history readable
fix(test): build the fp8 e2e harness with -fopenmp on Linux — make check was red
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
11 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.
Everything on
devsince v1.3.0 — 156 commits, 29 pull requests, no divergence onmain(a clean fast-forward of history).Two things arrive here that change what a colibrì release is: a third GPU backend, and — for the first time — binary archives that actually contain the models the front page promises.
A third GPU backend
#418 — Vulkan backend (@steve-m). The expert tier, dense projections and the MLA attention core on any GPU with a Vulkan 1.2 driver. This is the only backend for cards the vendor stacks no longer support — an RX 580 works — and it is competitive with ROCm on RDNA4 (#523: measured 19–24% faster on an RX 9070 XT, after the reporter corrected their own confounded first comparison).
It kept growing after the first cut:
fmt=4grouped-int4 andfmt=5int3-g64 decode, 256 MB allocation arenas to remove the tier-size submit tax, priority-classed VRAM allocations with budget-capped fill, a second-device expert tier (COLI_VK_DEV2— e.g. an RX 580 beside the primary card), single-submit fused prologues, spin-then-block fence waits, and an explicit warning when ReBAR leaves VRAM host-invisible instead of failing obscurely.#705 — Vulkan
fmt=7MXFP4 + Kimi K3 GPU tier (@steve-m). e2m1 nibbles with per-32 group scales decoded on the GPU; shared experts resident, fill-once routed-expert tier. Kimi K3 gets GPU acceleration without leaving its native QAT MXFP4 format — no requantisation, no conversion step.Binary releases finally ship all four engines
Until now every archive contained
c/colibrialone. A user who downloaded v1.3.0 and pointed it at Kimi K3 found no engine at all, while the README front page promised four model families (#720). Three changes close that, in order:#736 —
inklingandkimi_k3did not compile on Windows. Not "were not shipped" — did not build. Their serve loops polled stdin withfd_set/select, which msys2/UCRT64 does not provide in that form. Measured across the release matrix beforehand: Linux ✅, macOS ✅, Windows ❌. The portable version now lives once incompat.h, carrying the two bug fixescolibri.chad already absorbed — #139 (select()on a pipe handle routes to winsock and always returnsSOCKET_ERROR) and #195 (anonymous pipes are not waitable objects;PeekNamedPipefails on file/console handles).The same PR adds
engines-all-platformsto CI: every engine, on Linux + macOS + Windows. Before it, CI builtcolibri+inklingon Linux only andkimi_k3on nothing — which is exactly why a whole engine could stop compiling on two platforms and stay invisible for months.#737 — the archives ship every engine, plus
tools/k3_tokenizer.py. Kimi shipstiktoken.modelrather thantokenizer.json, and without that script the engine would be present and still undrivable. 863 KB for all four. The release job's verification step — which exists because a packaging mistake is invisible to every other job here — now asserts each engine is present and executable, that the tokenizer script parses, and that each sits where the launcher's own resolver looks for it.#739 —
COLI_CUDA=0no longer disables the OpenMP tuning. The guard testedgetenv()for presence on four variables. That is right for the two kill-switches and wrong forCOLI_CUDA/COLI_METAL, which are state where0means "no GPU" — so the tuning was skipped for exactly the users who have no GPU and most need their CPU team configured. Reachable rather than theoretical:coliitself writesCOLI_CUDA="0"when the plan does not use CUDA. Reported by @fredchu in #707, who read it out of the source and said plainly he could not measure it; the before/after measurement is in the PR.Performance and correctness
fmt=4decode on Apple Silicon.fmt=6(@michael-denyer): E8/IQ3 expert decode with on-GPU FWHT rotation.kimi_k3andolmoe, which had no tuning at all. Zen3 data: SMT collapse (2.3× at 16 threads), ±15% run variance, PIN_GB starves LRU — HIP/gfx1100 + NVMe #718 measured 2.3× on Zen3 from the thread count alone. They get the sizing half deliberately, not the spin-wait half: [Bug]: OpenMP spin-wait tuning is inert on macOS — and applying it costs 2.2x decode on a 32 GB host #707 measured −2.2× from spin-wait on a low-residency host, and Kimi is the most disk-bound engine here.PILOTprefetch.Shared facilities
.coli_usagehad two mutually unreadable writers under the same filename and two engines that produced nothing;route_trace.his now the single format, andkimi_k3has routing telemetry for the first time. Phase C (inkling,olmoe) is still open in The learning cache is engine-specific: .coli_usage has two incompatible writers and two engines that cannot produce it at all #700 — this is not closed yet, and the issue says so.Fixes and docs
#728 —
/home/vincenzo/glm52_i4removed from the shipped code (@terrizoaguimor). Every error message named a maintainer's home directory;tools/download_glm52.pydefaulted a 400 GB download there.grep -rn vincenzo c/is empty. The author found and fixed a regression in his own patch before it landed, by sweeping all eleven subcommands rather than the three he had changed.#511 the tiny-oracle expectation now says
~30-32/32, which is what it does (@CooperSheroy). #620 documents OpenAI tool-calling (@jeswr). #723 converter selftests without a model source. #709 / #740 README and website now describe four model families instead of one.#731 → #733 —
TEST_BINSderives from the build rules.c/Makefilehad one hand-written line listing every gate, so any two PRs that each added a test conflicted by construction — the file appeared in 26 of 40 open PRs. The first attempt (#731) globbedtests/test_*.cand was wrong in the other direction: it promoted files that deliberately have no build rule, which passed ondevand broke on a contributor branch. #733 derives from the rules instead, which is the honest definition of a gate. There is no shared list left to conflict on.About the contributor queue
Twenty-five contributor PRs merged. Eight of them were unblocked by a maintainer resolving the conflict and pushing to the author's branch rather than requesting another rebase; three of those needed a real merge rather than picking a side, because both sides carried a genuine change. No commits were rewritten.
That was a debt, not a courtesy. #386 was rebased eleven times in thirteen days — the eleventh caused by merging the Makefile fix before it instead of after. New rule, stated on the PRs: a green, unconflicted PR merges ahead of anything opened after it. The reason people rebased repeatedly is that newer, smaller PRs kept jumping the queue and resetting them.
Suggested tag
v1.4.0. A new GPU backend and releases that ship four engines instead of one is not a patch.
Credit goes mostly outward: @steve-m (Vulkan backend, Kimi GPU tier), @monotophic (Metal, formats, container stamp), @ZacharyZcR (CUDA tier, autotuning, E8), @terrizoaguimor (shared telemetry, hardcoded paths, CUDA fixes), @michael-denyer, @bherald, @jeswr, @opxyc, @t83714, @CooperSheroy — and to everyone who filed a measured bug report this cycle, several of which are the only reason the Windows and packaging gaps were found at all.