Rebased onto current dev, split into 3 logical parts (all validated):
1. CPU portability (serve-mode _O_BINARY pipe fix — stock main hangs on MinGW without it; RAM detection cap 0->9/layer; POSIX guards for select/mmap/madvise; warmup script).
2. AVX-VNNI 128-bit int8/int4 dot kernel (Alder Lake+/Meteor Lake+), bit-identical to AVX2 (author-verified on Meteor Lake; compiles out to AVX2 elsewhere) + _mm256_extracti128_si256 typo fix that blocked -march=native.
3. CUDA DLL via LoadLibrary, gated behind CUDA_DLL=1 (host never links cudart; silent CPU fallback if absent; author-verified on RTX 5070 Ti).
Validated here: make check 59/59, oracle 32/32 TF, Windows cross-compile clean + glm.exe loads+runs via WSL interop. Fixes the #123 Windows build failure.
On Windows $(OS) is Windows_NT in every shell; check it first so native PowerShell/CMD (no uname on PATH) doesn't fall through to the Linux branch. Non-Windows unchanged (else branch still uses uname). Linux build verified green.
Docs-only. Documents that the OMP active-spin steals SoC power from the Metal GPU on Apple Silicon (default regresses -39%); COLI_NO_OMP_TUNE=1 + PIPE=1 recovers and beats the pre-rebase branch (2.24 vs 2.06 tok/s). Flags a follow-up: Metal builds should default to passive OMP wait.
Extends the ablation grammar to int{2,3,4,8}[-g<N>][-rot][-nohead]. -rot round-trips weights through an orthogonal Hadamard Q=diag(±1)·H/√n on the input dim, measuring the exact weight error of a deployed rotate-activations scheme. Engine-free tool (c/tools/quant_ablation.py only). Verified: syntax clean, scheme parser correct (int3/-rot/-nohead), no unsafe constructs. Findings feed the v2 int3-g64 direction (#81/#108).
#134: olmoe.c stored experts as int8_t but silently accepted any bits argv;
bits=16 (falsely documented as f32) truncated in quantize_rows -> wraparound
garbage experts. Guard bits to 2..8 with a clear error (int8 is token-exact;
f32 experts are not implemented here, unlike glm.c's fmt=0 path). Doc corrected.
#133: shipped ref.json was from a different checkpoint (continued 'The capital
of the United States is Washington'); the intended target is
allenai/OLMoE-1B-7B-0125-Instruct (per convert_olmoe.py), whose greedy oracle
continues 'The official language of France is French'. full_ids/text updated to
the reporter's verified oracle (engine is already token-exact 12/12 vs it).
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
env_for() mapped --gpu/--vram only inside the --auto-tier branch, so the standalone form started a CPU-only engine with no warning (#121 — nearly published as a GPU benchmark). Now: --gpu none disables CUDA; --gpu list/auto sets COLI_CUDA/COLI_GPUS; --vram sets CUDA_EXPERT_GB; and --gpu/--vram on a CPU-only binary exit with a clear 'make glm CUDA=1' message instead of falling back silently. Tested on a CPU-only build: fail-fast fires for --gpu/--vram, --gpu none and default still start the CPU engine (no regression).
step_decode_batch (run_serve_mux) passes per-slot positions[]/kv_start, so nt=pos+1-st0 can exceed the old Tk+1 cap -> heap-buffer-overflow on the first serve request (ASan-confirmed, dnnspaul). sc_cap now = max nt over the batch, counted exactly as the write loop (per-slot positions/kv_start + DSA top-k). Non-batched paths reduce to the correct cap (MTP kv_start=-1 -> Tk+1). Validated: make check 58/58, oracle 32/32 TF, build 0-warning; real-model serve completes with ASan clean.
Two OpenAI-compat tool-calling bugs found against the real GLM-5.2 (dnnspaul): (1) string-typed args coerced to numbers — declared schema type now decides, string kept verbatim, bool rejected as number, schema-less params keep permissive decoding; (2) tool_choice was ignored — none/auto/required/{function} now honored, invalid returns 400. Python-only (openai_server.py + tests), engine untouched. 36/36 tests pass (verified independently in a clean worktree).
Measuring "what does int4 cost?" by comparing colibri's score to a published
model-card number does not work: this harness scores 0-shot log-likelihood while
published numbers are few-shot/CoT, and that protocol gap can swamp the
quantization effect entirely (#108).
This removes the confound by construction: take an fp16 model, push its weights
through colibri's own quantizer (quantize -> dequantize, in place), and score both
with the SAME harness on the SAME questions. The only variable is the quantizer, so
the delta IS the quantization cost. Runs on OLMoE in minutes, so a scheme can be
ranked BEFORE committing to a multi-hour GLM conversion.
Quantizer math is replicated from tools/convert_fp8_to_int4.py (symmetric absmax,
per-row scales) and generalised with an optional group size, so grouped/finer schemes
can be compared directly against what ships today.
Measured on OLMoE-1B-7B, n=200/task (#108):
scheme hellaswag arc_c mmlu mean delta
fp16 77.0% 47.0% 47.0% 57.0% --
int4 (shipped) 74.0% 41.0% 31.5% 48.8% -8.2pp
int4-nohead 73.5% 40.5% 37.5% 50.5% -6.5pp
int4-g128 78.5% 45.5% 38.0% 54.0% -3.0pp
int4-g128-nohead 78.5% 46.5% 38.0% 54.3% -2.7pp
The per-row int4 container costs ~8pp, concentrated on the HARD task: MMLU falls to
31.5% against a 25% random baseline while easy HellaSwag barely moves -- per-row
scales eat the small logit margins that hard questions depend on (the same margin
erosion that flips near-tie tokens in #100). group=128 recovers ~63% of the loss.
Keeping lm_head/embed in fp16 is NOT the fix (+1.7pp alone, +0.3pp atop grouping).
Includes a coverage assert: transformers fuses MoE experts into 3D tensors, so a
ndim==2 filter silently skips every expert and leaves ~85% of the model in fp16 while
appearing to work. The tool fails loudly instead of reporting fiction.
Dev-only tool (torch + transformers); the engine's dependency-free path is untouched.
* docs: Metal expert-matmul backend design (Apple Silicon)
Empirically-validated design for a batched MoE expert-matmul Metal backend.
Microbenchmarks (scratchpad) establish: runtime-compiled Metal needs no Xcode;
V3 (float4 + threadgroup reduction) kernel is correct and fast; synchronous
per-matmul dispatch loses to CPU due to ~150us Metal launch latency, so the win
is batched full-layer dispatch (854us/layer, 707 GFLOP/s) reading expert slabs
zero-copy from unified memory.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: backend infrastructure + kernel-correctness test (M1)
Add backend_metal.{h,mm} — an opt-in Apple-GPU backend built with METAL=1 on
macOS. Runtime-compiled shader (no Xcode needed), zero-copy over unified memory.
Implements coli_metal_matmul (general quantized GEMV, f32/int8/int4/int2) via a
threadgroup-reduction + float4 kernel; batched moe_block is stubbed (returns 0 ->
CPU fallback) for M2. tests/test_backend_metal.mm validates all formats and edge
shapes (odd S, non-mult-4 dims) against a CPU reference (nerr ~2e-6). Makefile
gains a METAL=1 Darwin branch and a metal-test target. Default build unchanged.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: batched moe_block + zero-copy slab registry (M2 backend)
Implement coli_metal_moe_block: gate/up/silu/down for a whole expert block in ONE
command buffer, with GPU memory barriers between stages and BINDLESS gpuAddress
pointers so each expert is read zero-copy from its own RAM slab (exceeds Metal's
~31 buffer-binding limit). coli_metal_register/unregister wrap page-aligned slabs
via newBufferWithBytesNoCopy and resolve interior pointers to GPU addresses.
Per-row ragged expert routing supported; CPU does the final weighted scatter-add.
test_backend_metal validates decode + ragged blocks vs a CPU reference (nerr ~2e-6).
Still gated off in glm.c until the moe() wiring lands.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: wire batched moe_block into glm.c, token-exact (M2 integration)
moe() now dispatches each routed-expert block through the GPU in one command
buffer when COLI_METAL=1, reading expert weights zero-copy from page-aligned
RAM slabs (registered in expert_load). Falls back to CPU per-block on any
unresolved slab or GPU fault. Default build byte-identical (all #ifdef COLI_METAL).
Fixes a heap-corruption crash: expert_load registers slabs from parallel OpenMP
threads, so the slab registry is now mutex-guarded (buffer creation stays outside
the lock). Added command-buffer error checking (fall back to CPU on GPU fault)
and a COLI_METAL_DEBUG one-shot trace.
Validated token-exact vs the CPU path (greedy): identical 12-token output;
expert-matmul time 29.9s -> 21.1s with pinned experts still on CPU.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: instrument moe_block (GPU/CPU split, wall-vs-kernel time)
Add diagnostics printed on the PROFILO line under COLI_METAL: GPU vs CPU-fallback
block counts, experts-on-GPU, and a per-block time split (setup / gpu-wall /
kernel / scatter). Reveals that with a warm cache all experts run on the GPU
(0 fallback) and expert-matmul drops ~1.3x vs CPU, but ~62% of GPU wall-time is
idle/scheduling latency (3.1s kernel of 8.3s wall over 396 sporadic submits) —
the GPU powers down between blocks because attention runs on the CPU per layer.
Points the next optimization at keeping the GPU hot (offload attention).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* docs: Metal backend measured results + next levers
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* docs: Phase 2 fused decode attention plan + absorption-core validated
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: fused decode attention on GPU, token-exact (Phase 2)
coli_metal_attn_decode runs a full S=1 decode attention layer in ONE command
buffer: q_a -> rmsnorm -> q_b -> RoPE ; kv_a -> latent rmsnorm@pos + krot RoPE@pos
(cache write) ; MLA absorption core (qabs/score/softmax/clat/ctx) ; o_proj. The
absorption-core kernels were validated in isolation (nerr ~1e-6) before wiring.
Projection matmuls reuse the mm_gemv kernel; attention weights are uploaded+cached
(serial path, no lock); Lc/Rc caches are page-aligned + registered in kv_alloc for
zero-copy GPU read/write. GLM-5.2 dims compiled in; falls back to CPU for S>1
(prefill/MTP verify), st0!=0, active DSA selection (context>topk), or mismatched
dims. DSA index-key write stays on CPU so future selection still works.
Validated token-exact vs CPU (identical greedy output); attention time 16.5s ->
10.5s (~1.57x), end-to-end 0.20 -> 0.28 tok/s.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* docs: Phase 2 fused attention complete + known limits
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: attention coverage/latency instrumentation + honest results
Add per-layer fused-attention counters (METAL-ATTN line): GPU layer count, gpu-wall
and true kernel time. Measurement (DRAFT=0, all-S=1 decode) shows the fused attention
triggers on all decode layers but is submit-latency-bound: gpu-wall 3.70s vs kernel
0.63s (83% idle latency over 546 sporadic command buffers). Attention time is neutral
vs CPU; the earlier MTP-on "16.5->10.5" was run-to-run variance. Design doc corrected
with the honest result: both offloads are gated by Metal's ~5ms cold-GPU submit
latency; reducing submit count (fuse attention+experts per layer) is the real lever.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: fused attention handles S<=4 (covers MTP verify forwards)
Extend coli_metal_attn_decode from S=1 to S<=4: the core kernels (qabs/score/
smax/clat/ctx) gain a query-row dimension with per-row causal masking (query s
attends keys [0, pos_base+s]); rmsnorm/rope/copy became row-aware; projections
run S rows via mm_gemv. This covers the default MTP config (draft=3 -> S=4 verify
forwards), which previously fell back to CPU attention entirely.
Token-exact vs CPU (identical greedy output, MTP on). Perf is inconclusive at
short context: still submit-latency-bound (attn gpu-wall 5.5s vs kernel 0.9s) and
the measurement is dominated by disk-streaming variance (+/-15s between runs).
Next: measure with a fully-warm cache to isolate compute, then reduce submit count.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* docs: clean warm A/B shows real ~1.4x (experts+S<=4 attention), token-exact
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: interleave attention q/kv paths, 7->4 barriers (iter 2)
The q-path (q_a->rmsnorm->q_b->rope) and kv-path (kv_a->copy->rmsnorm+rope) are
independent until the absorption core, but were serialized by memory barriers.
Interleave them into 4 barrier-separated stages so the GPU overlaps independent
dispatches. Token-exact; attention gpu-wall 3.04s -> 2.73s (~10%).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: zero-copy attention weights + fuse shared expert into GPU block (iter 2)
Dense QT weights/scales now allocate page-aligned + registered (qalloc) under
METAL, so the fused attention reads q_a/q_b/kv_a/kv_b/o zero-copy instead of
uploading ~6 GB of duplicates (RSS -3 GB, upload copies gone). bind_gemv resolves
registered pointers (buffer,offset) with a pre-check guard.
Phase E's shared expert (identical shapes to a routed expert: gate/up [I,D],
down [D,I], same int4 container) is appended to the first Metal moe_block as an
extra expert with rw=1.0 over all S rows — removes 3 CPU matmuls per layer and
fills the same GPU submit. CPU Phase E still runs on any fallback.
Zero-copy validated token-exact: 35.1s -> 29.7s (0.34 tok/s) warm.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* docs: iteration 2 findings
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* docs: iter 2 final ~1.56x + iter 3 plan (disk/GPU overlap)
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: overlap disk loads with GPU compute inside the layer (iter 3)
Split each MoE block into two GPU submits: the RESIDENT experts (pin/LRU hits,
plus the fused shared expert) are encoded and committed BEFORE the missed
experts' OMP pread loop, so the GPU computes while the disk reads; the missed
subset follows in a second (sync) submit once loaded. New two-phase backend API
(coli_metal_moe_block_begin/end) with handle-owned scratch so the async submit
cannot collide with the sync path's static buffers; moe_submit/moe_finish are
shared by both. Per-subset CPU fallback preserved (resident and missed fall back
independently on unresolved slab or GPU fault).
Token-exact. Warm 96GB: expert-matmul 8.96 -> 4.92s (resident compute now hidden
inside the disk window; expert idle latency ~5.7s -> ~0.9s), total 28.97s
(0.35 tok/s) vs CPU 50.2s = ~1.73x.
Note: 'make glm METAL=1' after a default build does NOT rebuild (target looks
up-to-date) — touch glm.c or clean when switching build flavors.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* docs: iter 3 disk/GPU overlap results (~1.73x)
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: keep-alive spinner experiment (env-gated) + latency decomposition
COLI_METAL_SPIN=1 keeps trivial GPU work in flight on a separate queue to probe
whether inter-submit idle is clock ramp-down; thread is detached (a joinable
global thread std::terminate'd the process at exit). First contended A/B was
inconclusive but showed the spinner does NOT collapse attention wall per-call
(~16ms both ways), so ramp-down is not the whole story. METAL-ATTN now decomposes
latency: cpu-sched (commit->kernelStart) vs gpu-sched (kernelStart->GPUStart) vs
kernel execution, to pinpoint where the ~13ms/call goes. Default behavior
unchanged (spinner off).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: standalone regression tests for fused decode attention
run_attn builds full-size fake GLM-5.2 attention weights (int4, page-aligned,
registered), replicates glm.c's absorb-branch math exactly on the CPU (q_a ->
rmsnorm -> q_b -> rope; kv_a -> latent rmsnorm + krot rope -> cache; per-head
qabs/score/softmax/clat/ctx; o_proj), and checks coli_metal_attn_decode against
it at S=1/3/4 and pos_base 0/12/37 — including the Lc/Rc cache write-back, which
end-to-end runs cannot isolate. All pass (nerr ~5e-6, cache ~1.4e-5). The whole
Metal path (gemv, moe_block, fused attention) is now testable without the model.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: route large row-batch matmul_qt GEMMs to the GPU (prefill)
matmul_qt now dispatches to a new coli_metal_gemm when S >= COLI_METAL_GEMM_MIN
(default 16), the weight is int8/int4 and registered (all dense QT allocs are,
via qalloc), and we're not inside an OpenMP region (mirrors the CUDA guard).
Decode-sized matmuls stay on the CPU where NEON wins vs submit latency; prefill's
big GEMMs (kv_b reconstruction at S=Tk, o_proj, dense MLP, step_all's S x vocab
logits) amortize it — microbench showed ~6x over the CPU idot at S=16.
Standalone test: registered int4 GEMM S=64 vs cpu_ref (nerr 2.9e-6).
Machine busy again; end-to-end token-exactness + threshold sweep pending idle.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* README: document the experimental Metal backend (Apple Silicon)
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: 1.5-2.1x faster moe_gemv (simdgroup-per-row + 8-value loads)
Replace one-threadgroup-per-output-row (128 threads reducing via threadgroup
memory) with one SIMDGROUP per output row, 4 rows per threadgroup, and uchar4
loads (8 nibbles / 8 int8 per lane-iteration). Removes the threadgroup barrier
+ shared-memory reduction entirely (simd_sum only) and doubles load width.
Engine-like block-shape microbench (pure GPU time): S=4 block 2548->1739us,
S=1 block 934->437us, big block 4582->3414us — 358-389 GB/s vs 182-264.
Row-bound guard added (NT) since the grid rounds up to 4 rows/TG.
All backend tests pass (moe_block nerr 2.4e-6, attention unchanged).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: mm_gemv simdgroup-per-row + 8-value loads (attention projections, prefill GEMM)
Apply the moe_gemv V2 transformation to the general quantized GEMV: one simdgroup
per output element (4/threadgroup), 8-value loads for i8/i4/f32, no threadgroup
reduction. Same measured 1.5-2.1x class of win; serves the fused-attention
projections (q_a/q_b/kv_a/o), coli_metal_gemm (prefill), and coli_metal_matmul.
All three dispatch sites updated (NT row-bound guard, grid ceil(NT/4) x 128).
Full test suite green, incl. non-mult-of-8 tail paths (2050x6146) and all fmts.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: experimental COLI_MMAP=1 — experts as zero-copy views into mmap'd files
Lazily mmap each safetensors file (PROT_READ, MAP_SHARED, mutex-guarded — expert
loads are OMP-parallel), register the mapping with Metal, and make expert_load a
pointer assignment into the map: no pread, no slab, no copy; the OS page cache is
the cache. Alignment guards fall back to the slab path. Default OFF.
First validation (machine at load 66 + 46GB swap): token-exact, RSS 58 -> 10.5 GB
as designed, but GPU wall exploded (~130 MB/s effective) — the GPU demand-faults
file-backed pages, catastrophic when memory pressure evicts them. Needs an
idle-machine A/B to judge fairly (llama.cpp's identical technique relies on pages
staying resident); possible fixes if slow even idle: CPU pre-touch of missed
experts' pages before the GPU submit, or madvise/mlock windows.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: CPU pre-touch for COLI_MMAP expert pages (fix GPU demand-faulting)
In mmap mode, fault the missed expert's pages in on the CPU inside expert_load
(madvise WILLNEED for async readahead + a page-stride touch): this is pread's I/O
without the copy and without the slab, it runs inside the existing OMP loop that
overlaps with the resident-experts GPU submit (iter 3), and it guarantees the GPU
only ever reads resident pages — GPU demand-faulting of file-backed pages
measured catastrophic (~130 MB/s). Read-only addition: outputs unchanged from the
validated mmap run; perf pending the idle-machine A/B.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* docs: idle-machine suite results (~1.33x same-session; mmap negative result)
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: COLI_METAL_UNTRACKED experiment (negative result, default off)
Env-gated MTLResourceHazardTrackingModeUntracked on registered wraps + scratch to
test whether cross-CB hazard tracking causes the ~10ms/CB gpu-sched delay. Idle
A/B: no effect (gpu-sched 3.9 vs 3.4s, noise), token-exact. Together with the
spinner negative, this pins the attention CB delay as inherent scheduler/wake
overhead on an empty pipeline — removable only by eliminating the CB boundary,
which CPU-side routing at ~58% hit-rate forces. Metal side is at its floor:
kernel 3.5s+0.8s (near BW ceiling), sched ~3.2s, disk ~15s dominant (10 tok).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* docs: loop conclusion — best config DIRECT=1+COLI_METAL=1, 0.42 tok/s (~1.4x)
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: refactor attention into encode_attention()+resolve_attn() (layer-CB prep)
Behavior-preserving: attn_decode is now a thin wrapper; all attention tests
byte-identical. Prepares embedding the chain in a full-layer command buffer.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* metal: full decode layer in ONE command buffer (token-exact)
coli_metal_layer_decode runs the whole layer prelude on the GPU in a single
submit: in_ln rmsnorm -> fused attention -> residual add -> post_ln rmsnorm ->
shared expert (gate/up/silu/down) -> router (f32 simdgroup matvec + sigmoid) ->
exact phase-A top-K selection (greedy argmax over sigmoid+bias with CPU tie
order, --topp truncation, norm_topk, routed_scale) in a serial-per-row kernel.
The CPU's per-layer work shrinks to: read 8 expert IDs, resolve/load, expert CBs
(disk/GPU overlap unchanged), scatter. moe() consumes the precomputed routing
(g_pre_*: skips phase A, keeps eusage/eheat/ereq counters for the learning
cache) and adds the GPU shared-expert output instead of computing phase E.
ld() tensors (norms/router/bias) now allocate registered so the GPU reads them
zero-copy. DSA index keys still computed on CPU from the in_ln-normed x (new
inrm output). Every missing condition falls back to the full CPU layer.
Validated token-exact vs CPU (identical greedy output, MTP on). Profile:
"altro" 3.8s -> 0.53s (12 tok); 0.42 tok/s despite disk-variance headwind.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* docs: Phase 3 full-layer CB results — 0.43 tok/s record, token-exact
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* gitignore: Metal build artifacts, venv, bench datasets
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
* remove internal design docs before PR
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
---------
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Both attention score buffers were fixed stack arrays (float sc[8192]).
The score count nt is only capped at index_topk when DSA selection
covers the layer; without indexer weights in the snapshot (has_dsa=0),
with DSA=0, or on the MTP layer, nt spans the full context. Past
position 8192 every OMP worker wrote beyond sc[] on its own stack:
silent corruption up to the guard page (~9400), then segfault.
Reproduced on a GLM-5.2 int4 snapshot without indexer tensors:
14.7k-token prompt crashed seconds into [prefill] layer 1/78, three
workers faulting simultaneously in attention._omp_fn.2 on the
sc[jj]=a*attn_scale store.
Fix: allocate the scratch once per attention call on the heap, sized
omp_get_max_threads() x (Tk - kv_start) — the true nt upper bound for
both the dense range and the DSA top-k list — and slice per thread.
Non-OpenMP builds get inline fallbacks, preserving the dependency-free
CPU path.
Validated: make check clean; short greedy output token-identical to the
previous binary; 10,232-token prefill segfaults on the old binary and
runs clean on the fixed one (layers 1-9+ verified, remainder is
expert-streaming disk time).
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
* Makefile: support Linux PowerPC (ppc64le) builds
PowerPC GCC uses -mcpu instead of -march, so the Linux branch failed
with unrecognized option -march=native on ppc64le. Detect ppc64le and
ppc64 via uname -m and use -mcpu=$(ARCH) there. The x86-64 path is
unchanged.
Validated on an IBM POWER8 S824 (Ubuntu 20.04, gcc 9.4): make test-c
passes, teacher forcing 32/32 positions and greedy 20/20 tokens against
the transformers oracle, engine reports the scalar idot fallback.
Signed-off-by: Scott <scottbphone12@gmail.com>
* VSX integer-dot kernels for POWER8 (12.8x int8, 7.6x int4 over scalar)
Adds a VSX path to dot_i8i8 and dot_i4i8 using vec_msum, which sums
byte products directly into s32 lanes, so the 16-bit saturation bound
of the AVX2 maddubs trick does not apply. abs(w) is built with a
modulo-subtract select instead of vec_abs so w=-128 wraps to 128
unsigned instead of saturating to 127. Nibble unpack uses
vec_mergeh/vec_mergel, which interleave like x86 unpacklo/unpackhi on
ppc64le (verified on hardware). g_i4s=1 on VSX since the f32 fallback
is plain scalar there: measured 5.5x for int4 IDOT at S=1.
Measured on an IBM POWER8 S824 (gcc 9.4, Ubuntu 20.04 ppc64le),
single thread, 1536x6144:
dot_i8i8 1.48 -> 18.99 Gops/s (12.8x)
dot_i4i8 2.33 -> 17.72 Gops/s (7.6x)
S=1 int4 matmul path: 3.925 -> 0.505 ms/call (7.8x vs scalar build)
Adds tests/test_idot.c: exactness test of the compiled idot kernels
(any arch) against a plain-C reference, covering odd tails and the
w=-128 edge. Passes on avx512-vnni (x86) and vsx (POWER8). The tiny
oracle stays token-exact on the VSX build: TF 32/32, greedy 20/20.
Signed-off-by: Scott <scottbphone12@gmail.com>
---------
Signed-off-by: Scott <scottbphone12@gmail.com>
Co-authored-by: Scott <scottbphone12@gmail.com>
The existing PILOT prefetcher predicts the next layer's routed experts
from the router logits and issues advisory readahead (posix_fadvise
WILLNEED / expert_prefetch) so the page cache is warm when the main
thread reaches that layer. That only warms the OS cache; the main thread
still pays the dequant/slab-build cost of expert_load on the critical
path.
PILOT_REAL=1 upgrades the prefetcher to perform the *actual* expert_load
into the next layer's LRU cache (ecache[L+1]) from the pilot I/O thread,
overlapped with the current layer's compute, so on a hit the main thread
finds the expert already resident and skips the load entirely.
Safety invariant (value-identical output vs OFF):
- MATMUL path: the pilot writes ONLY ecache[layer] for layer >
g_cur_moe_layer; the matmul in moe() reads ONLY ecache[layer]==
g_cur_moe_layer. A barrier at the top of moe() claims the current
layer and waits out any in-flight pilot load already targeting it,
after which the pilot drops every new load <= that layer. So the
matmul and the pilot never touch the same layer concurrently and no
half-loaded slot is ever matmul-ed.
- SCAN path (review fix — Option A): pilot_prefetch also runs on the
main thread and its residency scan reads ecache[lnext]/ecn[lnext] for
the FUTURE layer lnext = current+1 — exactly the layer the pilot
worker is writing. That scan previously ran off-lock (a real, benign
data race) and a stale comment falsely claimed the two threads never
touch the same layer's ecache. Fixed: the scan now takes g_pilot_mx
(the same lock the worker uses), decides under the lock, and enqueues
AFTER unlocking into the lock-free pilot_q ring (no re-entrant double
lock). The comment is rewritten to describe the real two-part
invariant honestly.
- The loading slot is published (eid set, ecn grown) only after the
pread completes; while loading it is hidden (eid=-1) from cache scans.
- The slow pread runs outside the handshake lock; the lock only guards
slot selection/publication. The shared LRU clock (m->eclock) is bumped
with an unconditional relaxed atomic add so the main thread and pilot
can't lose an increment; this is value-identical to the plain ++ with
only a negligible relaxed-atomic cost (no runtime branch for the OFF
case — not worth it).
Reliability (review fix — non-fatal speculative load):
expert_load() aborted the whole process (exit(1)) on any missing tensor,
OOM, short read or pread error. That is correct for the main / on-demand
/ REPIN / pin paths but fatal for a ~28%-mispredicted speculative pilot
load — a wrong guess could kill the server. expert_load now takes a
`fatal` flag: all main callers pass fatal=1 and keep today's exit-on-
error behavior byte-for-byte; the pilot passes fatal=0, so on error it
abandons the load, leaves the slot hidden (never published), bumps
g_pilot_drops, and logs a one-line stderr warning (never a silent
swallow). The main load path is behaviorally unchanged.
Residual gap closed (re-review): the fslab scale-buffer allocation still
went through falloc(), which exit(1)s on OOM regardless of `fatal`, so a
fatal=0 pilot load could still abort the process on a scale-buffer OOM.
expert_load now branches: fatal=1 keeps falloc() (byte-identical exit-on-
OOM for the main path), fatal=0 uses a checked malloc replicating falloc's
anti-wrap guard and, on failure, does the same non-fatal cleanup as the
slab-OOM branch (frees/NULLs s->slab and s->fslab, zeroes their caps,
returns -1) leaving a clean hidden slot (eid stays -1). The qt_from_disk
f32 fallback path is unquantized-only (GLM always has .qs) and is left
as-is, now with a comment noting it's unreachable for the pilot.
Tuning (review fix — PILOT_K default):
Real loads create LRU eviction pressure that hint-only WILLNEED does
not, so a large K thrashes at 28% mispredict. When PILOT_REAL is on and
the user did not set PILOT_K explicitly, K now defaults to 6 (best
measured this session) instead of 8; hint-only PILOT keeps the 8
default.
Default OFF (PILOT_REAL=0); PILOT_REAL=1 opts in and implies PILOT=1. A
per-run stats line reports real cross-layer loads completed vs dropped.
No Makefile change: pure C (stdatomic.h + sched.h).
Claude-Session: https://claude.ai/code/session_01DS7oc65c5RdA9V99otRCwt
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
* openai_server: parse GLM-5.2 tool calls into OpenAI tool_calls
Fills the 'tools not supported yet' stub. Upstream declared nothing and
rejected tools/functions; this makes them real end to end:
render_chat: emit a tools-declaration <|system|> block; replay assistant
tool_calls as <tool_call>NAME<arg_key>K</arg_key><arg_value>V
</arg_value></tool_call>; render tool-result messages as
<|observation|><tool_response>...</tool_response> (one
observation per consecutive tool run).
parse_tool_calls: turn the model's <tool_call> markers back into OpenAI
tool_calls; strip <think> from surfaced content.
generation: non-stream returns message.tool_calls + finish_reason
'tool_calls'; stream suppresses the markers from content
deltas (marker-length hold-back so a split <tool_call> is
caught) and emits tool_calls deltas after generation.
Default-safe: with no tools in the request, render and streaming take the
exact upstream path (byte-identical); tool handling is gated on tools present.
Marker format is authoritative (GLM-5.2 chat_template.jinja).
Optional de-mangler (COLI_TOOL_SALVAGE=1, default OFF) recovers malformed
calls from heavily-quantized models by mapping a lone payload onto the tool's
primary parameter; never rewrites well-formed output, never on by default.
A [api] tool-calls telemetry line reports strict vs de-mangled.
Pure stdlib; no new deps.
* openai_server: COLI_THINK global thinking default (opt-in)
COLI_THINK=1 makes thinking the default when a request sends neither
reasoning_effort nor enable_thinking (a launch-time global switch equivalent
to the old server's --think, useful because reasoningEffort propagation from
openai-compatible front-ends is unreliable). An explicit client value always
wins. Default off => exact OpenAI-standard behavior, so this is inert unless
enabled.
* openai_server: keepalive during long prefill (SSE reasoning pings)
engine.generate() blocks silently through the cold prefill (minutes for a
large prompt), and OpenCode/undici drop the socket after their idle timeout
-> the stream is 'terminated' before the first token. Upstream's SSE path
emits nothing until generation produces text, so it has no heartbeat.
A background pump emits a reasoning_content '.' delta whenever no event has
been written for KA_GAP (10s): the channel that reliably resets the client
timer and lands in the thinking panel, so answer content stays clean. All
wfile writes share a lock so the pump and event() never interleave; a
last-write timestamp gates the pump so it's silent while real tokens flow
(decode). The pump stops as soon as generation returns.
Inert for fast responses (nothing sent if events flow within 10s). No new
deps.
* openai_server: COLI_DEBUG echoes decoded tokens to stderr
Upstream sends decoded tokens only to the client (stdout->SSE), so the
console shows no generation text -- painful when the client has disconnected
or for local debugging. COLI_DEBUG=1 tees each decoded chunk to stderr at the
engine-callback boundary (both the plain and tool-call streaming paths).
Default off; no effect unless enabled.
* openai_server: use GLM-5.2 authoritative tool-declaration block
The tools were declared with a hand-written preamble; GLM-5.2 was trained on
the '# Tools' + <tools></tools> XML structure from chat_template.jinja. With
the wrong preamble the model doesn't recognize the tools as native and
hallucinates other frameworks' syntax (observed: repetitive 'end_action' and
key+value concatenated into one arg). Switch render_chat to the byte-exact
trained format so the model emits well-formed <tool_call> blocks.
PowerPC GCC uses -mcpu instead of -march, so the Linux branch failed
with unrecognized option -march=native on ppc64le. Detect ppc64le and
ppc64 via uname -m and use -mcpu=$(ARCH) there. The x86-64 path is
unchanged.
Validated on an IBM POWER8 S824 (Ubuntu 20.04, gcc 9.4): make test-c
passes, teacher forcing 32/32 positions and greedy 20/20 tokens against
the transformers oracle, engine reports the scalar idot fallback.
Signed-off-by: Scott <scottbphone12@gmail.com>
Co-authored-by: Scott <scottbphone12@gmail.com>
Records the July 12-13 lab findings on one 6x RTX 5090 machine:
- vLLM-Moet comparison and why per-rank expert residency dominates
- colibri full-resident placement: 2.30 -> 6.28-6.84 tok/s, with the
tested-and-rejected directions documented
- MTP speculation: broken int4 head identified (issue #8), int8 head
reaches 69-79% acceptance but MoE verify batches scale expert cost
with S, so speculation loses at every depth; revisit after GPU
grouped GEMM
- AVX-512 int4 kernel qualification: numerically better than the old
order, quality-neutral (ppl delta 0.24%), +4-7% on CPU-heavy routing
README links the record from the honest-numbers section.
For streamed (non-resident) experts the MoE forward is disk-bound: each
64-expert block first blocks on a parallel pread of every miss, then runs
the matmuls serially. Those two phases don't overlap, so the compute
cores idle while the block's weights are still coming off NVMe.
PIPE=1 hands the misses' expert_load() to a small persistent pool of I/O
worker pthreads and lets the main thread start matmul immediately. The
main thread walks the block's experts in routing order and waits on a
per-slot ready flag only for the expert it needs right now, so loads of
later experts in the block hide behind matmul of earlier ones. All
matmul_qt stays on the main thread (it parallelises internally via
OpenMP and gates GPU dispatch on !omp_in_parallel()), so the I/O pool
never competes with the compute team for the matmul itself.
Cross-generation correctness: generation-tagged lock-free cursor
--------------------------------------------------------------------
The batch state (njobs/eids[]/layer/ready[]) is reused in place across
64-blocks, so a straggler or late-woken worker could touch the NEXT
batch's state. Two prior fixes (a drain barrier, then an _Atomic active
counter) each still left a cross-generation window. Both are now removed
and replaced with a single generation-tagged cursor:
_Atomic uint64_t cur = (gen << 8) | index; // gen main-only, index 0..njobs
- dispatch writes njobs/layer/eids[]/ready-reset with RELAXED stores,
then RELEASE-stores cur (bumping gen); that release publishes all
batch state to any worker whose ACQUIRE-load of cur sees the new gen.
- a worker grabs a job by CAS-advancing the index; it reads eids[i]/
layer only AFTER the winning CAS. The CAS comparand carries the
generation, so if a new batch was published the stale CAS fails and
the worker re-reads — it can never grab a wrong-generation job or
read torn state, no matter where it was preempted (wake gap,
post-cursor, anywhere).
- gen is bumped only by the main thread and is monotonic ⇒ no ABA.
- the per-expert pipe_wait(ready[q]) in the matmul loop (kept) makes
every grabbed job complete before the block ends, so no grab
outlives its generation. That is what makes the old `active`
counter and the end-of-block drain barrier unnecessary — both are
removed. ready[] is reset before the publishing release, so no stale
flag survives into the new generation.
- the mutex/condvar now exist ONLY to park and wake idle workers, not
for correctness.
This only reorders I/O, never the computation: greedy decode output is
byte-identical to the blocking path.
Default OFF (PIPE=0) so upstream behaviour is unchanged; PIPE=1 opts in
and PIPE_WORKERS=n sizes the pool (default 8). No Makefile change: pure C
(stdatomic.h + sched.h), builds with the default toolchain.
Claude-Session: https://claude.ai/code/session_01DS7oc65c5RdA9V99otRCwt
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
The MoE forward does many tiny, back-to-back per-expert matmul regions.
Under libgomp's default passive wait policy the worker team is parked
between regions, and the thread re-wake latency dominates the actual
compute. Switching the team to an active wait policy (with a bounded spin
count so long NVMe expert-load stalls still yield instead of burning a
core) keeps the threads hot across those regions.
Measured on the Zen5 build: expert-matmul wall time 66.9s -> 20.9s
(~3.2x on that phase). Output is numerically unchanged — this only
affects thread scheduling, not the computation.
Mechanism: libgomp parses OMP_/GOMP_ vars in a constructor that runs
before main(), so setenv() from main() and continuing is too late
(verified: the already-initialised runtime ignores it). Instead we seed
the winning defaults on first entry with overwrite=0 (so any value the
user already set wins), set a COLI_OMP_TUNED sentinel, and execv() self
once so a fresh libgomp constructor reads them. The sentinel guards the
exec, so we re-exec at most once. The block is the first statement in
main() so argv is passed verbatim to execv().
Discoverability / observability (review follow-up):
- COLI_NO_OMP_TUNE=1 is a documented kill-switch that disables the
whole re-exec + tuning path in one shot (distinct from the internal
COLI_OMP_TUNED re-exec sentinel);
- a one-line stderr breadcrumb is printed before the re-exec
("[OMP] hot-thread tuning: re-exec once (COLI_NO_OMP_TUNE=1 to skip)")
so the self-re-exec is self-documenting;
- execv only returns on failure, so a perror() now follows it
("execv self-reexec failed, running untuned") — a container without
/proc or a deleted inode falls back visibly instead of silently
losing the ~3.2x.
Fully opt-out / overridable and lossless:
- explicit OMP_WAIT_POLICY / GOMP_SPINCOUNT / OMP_PROC_BIND / OMP_DYNAMIC
in the environment win (overwrite=0);
- pre-setting COLI_OMP_TUNED=1 skips the re-exec entirely;
- COLI_NO_OMP_TUNE=1 disables the tuning path entirely;
- guarded by __linux__ (no-op re-exec elsewhere; the setenv defaults
still apply for any later-initialising runtime).
Claude-Session: https://claude.ai/code/session_01DS7oc65c5RdA9V99otRCwt
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
The default ./glm mode validates against ref_glm.json, which is generated
from glm_tiny (a random tiny model). Pointed at the 744B GLM-5.2 it feeds
the model tiny-vocab prompt tokens and compares against tiny references —
0/20 and a degenerate loop, guaranteed on EVERY platform (x86/ARM/Metal),
which reads as an engine bug (it isn't). Detect the mismatch (real vocab +
tiny-range oracle ids) and print the correct commands instead. REF_FORCE=1
overrides. The engine self-test (SNAP=./glm_tiny TF=1) is unaffected.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Adds a Nix flake pinned to nixos-26.05. Builds the glm engine with
gcc+OpenMP (AVX2 portable), wraps the coli Python CLI with pinned deps
(torch, safetensors, huggingface-hub, numpy), and provides a dev shell.
New byte-level GBNF-subset engine (c/grammar.h: parser + set-of-stacks PDA
walker) wired into spec_decode as a third draft source ("metodo F"), tried
before MTP/n-gram. Wherever the grammar admits exactly one legal byte, the
forced span is tokenized and injected as drafts; the existing batch-union
verification confirms them, so a wrong or out-of-sync grammar can never
change the output. Lazy arming skips preambles; adaptive guard (same
pattern as MTP) disables the source below 50% acceptance; grammar-accepted
tokens no longer pollute the MTP acceptance counter.
GRAMMAR=file.gbnf enables it in run and serve modes (also with DRAFT=0 and
with the int4 MTP head from #8); GRAMMAR_DRAFT=n caps the span (default 24).
Measured on M3 Max / int8-MTP container, greedy, MTP=0 DRAFT=0, NDJSON
classification: 0.37 -> 0.50 tok/s (1.60 tok/forward, 81 fw per 130 tok),
100% acceptance (48/48), output byte-identical to baseline.
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
cody
JustVugg