* SCHEMA=<file.json>: JSON-Schema -> GBNF compiler for grammar-forced drafts (#48/#70 follow-up)
schema_gbnf.h compiles a practical JSON-Schema subset (strict objects, string/
number/integer/boolean/null, enum/const, arrays with items, nesting) into the
byte-level GBNF subset grammar.h parses, so structured-output workloads get
grammar-forced drafts without hand-writing GBNF. Unsupported keywords fail soft:
the engine runs without a grammar and output is unchanged (drafts are verified,
never constraints - a wrong compile can only cost acceptance, not correctness).
grammar_setup: GRAMMAR= (raw GBNF) keeps precedence; SCHEMA= feeds the compiler
into the same gr_parse path. 8 test groups in tests/test_schema_gbnf.c walk
compiled grammars end-to-end through the PDA (forced spans, enum disambiguation,
nested instances, escapes, leading-zero rejection, fail-closed fallbacks).
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* schema_gbnf: whitespace-tolerant emission (jws at separators)
Measured on GLM-5.2 current main (#146): the greedy continuation writes sloppy
JSON (spaces after colons, fences, long free text) and a compact-only grammar
desyncs at the first stray space, forfeiting every span after it. jws points are
not forced themselves (two legal bytes) but the multi-byte spans around them
keep drafting and the walker survives non-compact output - strictly
acceptance-positive for a verified draft source. Tests re-derived for the new
span boundaries + a sloppy-instance walk.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
---------
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Co-authored-by: JustVugg <JustVugg@users.noreply.github.com>
* cuda-dll: fix Windows build — MSVC host flags, CUDA_PATH default, POSIX setenv shim in the kernel test (#157)
First hardware validation of the #131 CUDA_DLL path (RTX PRO 6000 Blackwell
sm_120, MSVC 14.44 + CUDA 13.2, MSYS2 UCRT64 host build) found three blockers
that made 'make cuda-dll' unbuildable as shipped:
- NVCCFLAGS passed GCC-style -Xcompiler=-Wall,-Wextra to the MSVC host
compiler (hard error D8021). Use -Xcompiler=-W3 on Windows — dash form,
since MSYS make mangles /W3 into a filesystem path.
- NVCC defaulted to $(CUDA_HOME)/bin/nvcc with CUDA_HOME=/usr/local/cuda;
on Windows default CUDA_HOME from the installer's CUDA_PATH and NVCC to
plain 'nvcc' from PATH (CUDA_PATH contains spaces, which the unquoted
recipe checks cannot survive; an MSVC PATH environment is already required).
- tests/test_backend_cuda.cu used POSIX setenv/unsetenv (undefined under
MSVC); add a two-line _putenv_s shim.
Also corrects the stale '11 API symbols' comment (the header exports 15 and
backend_loader.c resolves all 15).
Validated: make cuda-dll (stock flags) + make glm CUDA_DLL=1 ARCH=native →
[CUDA] device init on sm_120, tiny oracle TF 32/32 + greedy 20/20, kernel
suite 'q8/q4/q2/f32 correctness ok', graceful no-dll fallback, plain build
byte-identical CPU behavior.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
* cuda: fix heap corruption in expert_host_release on Windows (CUDA_RELEASE_HOST=1)
expert_host_release() freed the expert slab with plain free(), but the slab
is posix_memalign'd — which compat.h maps to _aligned_malloc on Windows, so
free() corrupts the CRT heap: instant 0xC0000374 crash on the first released
expert. This is the exact pattern the compat.h audit fixed at the original
expert_load site ("l'unico sito che libera memoria aligned e' free(s->slab)");
this call site was added later and reintroduced it. compat_aligned_free is
plain free on POSIX, so non-Windows behavior is unchanged. fslab stays plain
free (malloc/falloc on the CPU path).
Found running the VRAM expert tier on real hardware (80 GB resident on an
RTX PRO 6000, CUDA_RELEASE_HOST=1 to avoid 80 GB of host double-residency —
reproducible crash before, clean generation after).
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
---------
Co-authored-by: lEWFkRAD <186512915+lEWFkRAD@users.noreply.github.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
Co-authored-by: JustVugg <JustVugg@users.noreply.github.com>
* prefill: batch the attention input projections (q_a/q_b/kv_a) over the prompt
During prefill the three attention input projections were called one row at a time
(matmul_qt(..., 1)) inside the loop over the S prompt tokens, while o_proj right
below already runs batched and moe() already does a batch-union. Batching them over
all S rows lets each weight row be read once for the whole prompt instead of once
per token.
Measured on Zen5 + RTX 5090, GLM-5.2 int4, 78 layers, DSA on, OMP tuning on,
prefill run to completion (engine's own timings, upstream/dev @ 6d3ed7e):
prompt projection/RoPE total prefill
256 22.7s -> 17.7s (-22%) 120.9s -> 115.4s (-4.5%)
512 45.5s -> 36.3s (-20%) 231.4s -> 222.7s (-3.8%)
Output is BIT-IDENTICAL to dev: summed log-lik of a fixed 1023-token passage is
-4987.458316 before and after (SCORE mode, CPU, deterministic). Decode (S=1) is
untouched — same calls, same kernel.
matmul_qt_ex(..., allow_idot=0) keeps the projections on the EXACT int4 kernel.
This matters: batching alone would push S past the S>=g_i4s gate (2 on x86) and
silently move them onto IDOT's int8-quantized activations, which costs
+0.169 nats/token on markdown/code and +0.084 on prose (perplexity +18.4% / +8.7%,
measured) for only 1.4% more speed now that #95 made the exact kernel fast. The
gate is a decode-era speed threshold; it should not be crossed by a batching change.
Also: matmul_qt documents the IDOT threshold as "configurable via I4S", but the
getenv was missing and the knob did nothing. Wire it up — it is what isolates the
kernel switch from the batching win.
* prefill: batch the DSA indexer's ix_wk projection too
Same per-token bug as the attention projections, one block below: ix_wk was called
one row at a time inside the prefill loop, so its weights were re-read for every
token on every DSA layer.
matmul_qt_ex(..., 0) keeps it on the exact int4 kernel for the same reason as the
attention projections: batching alone would push S past the S>=g_i4s gate and
silently change the activation quantization.
Output stays BIT-IDENTICAL (summed log-lik over 1023 tokens, in-distribution):
prose -2295.245383 (unchanged)
markdown -3272.146403 (unchanged)
Speed, 256-token prompt, 3 alternating reps each (Zen5 + RTX 5090, dev + the
attention-projection commit as the baseline):
projection/RoPE 17.54s (sd 0.20) -> 16.96s (sd 0.15) -3.3%
total prefill 115.34s (sd 0.82) -> 113.69s (sd 0.35) -1.4%
Both separated beyond 2 sd. Small, but it is free and bit-identical.
ix_wq/ix_wp are left alone on purpose: they sit inside an `omp parallel for` and
are skipped entirely below index_topk context, so batching them is a different
change with a different risk profile.
---------
Co-authored-by: JustVugg <JustVugg@users.noreply.github.com>
Routing at layer L strongly constrains routing at L+1/L+2: measured on GLM-5.2,
co-activation lift over independence is median 1.8x / p99 40x in-domain, and the
structure TRANSFERS across workloads (a coupling table trained on prose/code
keeps +33-46% relative prefetch recall on an unseen NDJSON workload, both
depths) - it is a property of the model, not the session.
Three pieces:
- ROUTE_TRACE=<path>: zero-effect routing dump (one line per position/layer,
top-K ids:gates) from moe() FASE A.
- tools/route_pairs.py: traces -> .coli_pairs table (top-16 co-activated
L+1/L+2 experts per (layer, expert)); tools/route_coupling_report.py:
Frechet-bound screen + marginal-vs-coupled prefetch recall, with a
train-on-A/test-on-B transfer mode.
- COUPLE=<.coli_pairs> (+COUPLE_K, COUPLE_D): scores next-layer candidates by
summed pair counts over the position's routed set and enqueues non-resident
ones into the existing pilot ring (same worker, residency re-check, and
safety invariants; hints only - output byte-identical, verified).
End-to-end on M3 Max (fast NVMe, warm page cache), interleaved
baseline/couple/baseline: K=4 D=1 neutral (0.48 vs 0.48/0.50 tok/s brackets);
K=8 D=2 harmful (0.35 tok/s: ~600 hints/token = ~11 GB/token of readahead
thrashing the page cache). WILLNEED cannot move engine hit% by construction.
So the PREDICTOR is validated, the readahead ACTUATOR does not pay on this
hardware class - default OFF, small K default, and the interesting targets are
matched-latency storage (expert read ~ layer compute) and PILOT_REAL-style
loading on small-RAM boxes, both left to owners of such hardware.
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
* Web dashboard: live expert-tier panel (VRAM/RAM/disk), static hosting, coli web command
- Engine: TIERS protocol line (counts + GB for VRAM/RAM/disk experts)
emitted at READY and refreshed after every served turn, computed from
the live pin/LRU/CUDA state.
- Server: dispatcher tracks the latest snapshot, /health exposes it as
"tiers"; the built web UI (web/dist) is served from the same port
(read-only, traversal-safe, SPA fallback) so the dashboard needs no
second process.
- Web: tier bar + legend in the Runtime panel showing where the 19k
experts live right now, with GB per tier; updates with the existing
health polling.
- CLI: `coli web` = serve + auto-open the browser once /health answers
(the 744B load takes minutes; a poller waits for it), --no-browser to
disable, friendly hint when web/dist isn't built.
* coli web: HERE is a path string, not a Path
* web: same-origin default + auto-connect when served by the engine
The page hosted by coli web talked to http://127.0.0.1:8000/v1 by
default while being loaded from http://localhost:8000 - a different
origin, so the very first probe died on CORS ('Failed to fetch').
When the page is engine-served the default (and a stored stale factory
default) now resolve to window.location.origin + /v1, and the console
auto-probes on load; the Vite dev server keeps the classic default.
The server's own port is also whitelisted in DEFAULT_CORS_ORIGINS for
the localhost/127.0.0.1 cross-name case.
* web: fix auto-connect effect returning Promise (React 19 StrictMode crash)
The async connect() call inside useEffect returned a Promise that React
19 StrictMode stored as the effect cleanup. On re-render, React called
destroy_() on that Promise — 'not a function'. Block body with explicit
return undefined prevents any non-function from leaking into the cleanup
slot.
* web: rich streaming metrics — live token counter, tok/s gauge, TTFT, session totals
During generation: a flashing Zap badge counts tokens in real time.
After: Gauge shows tok/s, Timer shows TTFT (time to first token),
Layers shows prompt→completion token counts, Clock shows queue wait.
Session totals (cumulative prompt + completion) appear in the Runtime
panel. All with lucide icons and tabular-nums for stable layout.
Tier bar gains a smoother cubic-bezier transition and slightly taller
height for visual weight.
* web: hardware environment panel — CPU model, GPU count/VRAM, RAM total/free, cores
Engine emits HWINFO protocol line at startup (CPU name from /proc/cpuinfo,
core count, RAM total/available from /proc/meminfo, GPU count + VRAM from
CUDA). Server parses and caches it, /health exposes as 'hwinfo'. Dashboard
renders it with icons at the top of the Runtime section so the user sees
immediately what the engine is running on.
* web: downgrade to React 18 — React 19 effect cleanup bug crashes the dashboard
React 19 treats the return value of every useEffect callback as a
cleanup function. In practice, any async interaction (streamChat's
rapid onDelta re-renders, the health poller, auto-connect) caused
React 19 to store a Promise as inst.destroy and crash with
'destroy_ is not a function' on the next commit cycle. The bug
reproduced in incognito, without StrictMode, and with every effect
converted to explicit block bodies returning undefined — it is a
React 19 regression in the passive-effect unmount path.
React 18 does not exhibit this behavior. Pin to React 18 until
React 19 is fixed or the codebase migrates to a pattern React 19
handles correctly. All 17 vitest pass, ErrorBoundary retained.
* web: Brain page — the expert cortex, live
A 76x256 canvas grid, one cell per expert (19,456 total): colour =
tier (green VRAM / blue RAM / grey disk), brightness = routing heat
(log-bucketed .coli_usage counts), and a white pulse that flashes on
every expert routed in the current turn and decays — you watch the
model think.
- Engine: EMAP protocol line (1 byte/expert: 2bit tier + 6bit heat,
hex) at READY and after each turn; HITS bitmap of this turn's routed
experts (set where eusage increments, cleared on emit).
- Server: parses both, GET /experts serves {rows, cols, map, hits, seq}.
- Web: Brain tab next to Chat; canvas renderer with rAF pulse decay;
hover tooltip shows layer/expert/tier/heat plus an honest depth-role
heuristic (early = surface features ... late = output shaping, MTP
row labelled as the speculative draft head).
* web: Brain page responsive — cells sized from container via ResizeObserver
Cell size derives from the wrapper's actual client box instead of fixed
1400x900, re-rendering on resize; mobile media query tightens padding,
legend and tooltip. The cortex now fills whatever screen it gets.
* win: direct I/O via FILE_FLAG_NO_BUFFERING + compat_fsize + VirtualLock primitives
compat_open_direct() gives Windows the O_DIRECT twin fd st.h already uses
on Linux/macOS: FILE_FLAG_NO_BUFFERING, same 4K-alignment contract as
O_DIRECT (the engine's DIRECT=1 path already aligns offset/len and slabs
are posix_memalign'd).
Measured on GLM-5.2 744B int4, Ryzen 9 9950X3D / 126 GB / PCIe4 NVMe
(5.8 GB/s at the engine's 19MBx8T pattern), Windows 11, MinGW GCC 16.1,
32-token greedy runs at --topp 0.7, 40 GB pin, current dev HEAD:
buffered: 0.38 tok/s (expert-disk dominates)
DIRECT=1: 0.56 tok/s (1.47x) — byte-identical greedy output vs buffered
compat_fsize() (GetFileSizeEx): CRT lseek(SEEK_END) returns -1 on
NO_BUFFERING fds (measured on UCRT); iobench uses it and gains a
NO_BUFFERING branch so disk numbers are comparable across platforms.
compat_mlock/compat_munlock: VirtualLock with working-set growth (bare
VirtualLock caps at the default working-set minimum, a few hundred KB).
Wired into the engine in the next commit.
tests/test_compat_direct.c covers the alignment contract, data integrity,
fsize on both fd kinds; skips cleanly off Windows.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* win: wire VirtualLock into mem_wire, munlock pairing in expert_host_release
MLOCK=1 was a silent no-op on Windows: pinned experts could be paged out
by working-set trimming under memory pressure. mem_wire now uses
compat_mlock (VirtualLock + working-set growth); expert_host_release
unlocks before freeing, mirroring the POSIX branch.
Validated: 39.6 GB pin wired in 17s on a 126 GB machine, zero failures;
TF oracle 32/32 with MLOCK=1.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
* cuda: thread-local current-device cache in select_ctx
cudaSetDevice on every call is expensive when the serial expert loop
alternates devices. Measured on RTX 5090 + RTX 4090 (Windows, DLL
backend, pre-#68 dispatch): expert-matmul 14.3s -> 25.4s per 32 tokens
going from 1 to 2 devices, entirely per-call context switching. The
current device is per-thread in the CUDA runtime, so a thread_local
cache skips redundant switches; multi-GPU expert serving becomes
positive-scaling instead of negative.
Kernel suite passes on sm_120 + sm_89; TF oracle 32/32 dual-GPU.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
---------
Co-authored-by: olorin <io@zyphyr.co>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
- load_expert_w: detect I8/U8 tensors and read them raw with their .qs scales.
Before this, pointing SNAP at a convert_olmoe.py container crashed with
SIGBUS (st_read_f32 walks I8 data as 2-byte elements). Container misses now
cost half the I/O and skip quantize_rows entirely: 2.08 -> 3.69 tok/s on the
miss-heavy 16-slot ref.json run (M5). Raw bf16 checkpoints keep working.
- matmul_q: Q8_0-style integer path on ARM (per-16 activation blocks, sdot on
dotprod CPUs, vmull fallback). Same IDOT family as glm.c, same semantics:
IDOT=0 stays byte-exact vs the oracle (12/12); default integer path can flip
an argmax tie (documented in glm.c README). End-to-end on M5: 4.5 -> 12 tok/s
warm-cache decode.
- rss_gb: ru_maxrss is bytes on macOS, KB on Linux. RSS lines were reading
'2029 GB' on Macs.
- convert_olmoe.py: snapshot_download(local_files_only=True) raises
LocalEntryNotFoundError when the repo is not cached; the download fallback
was unreachable.
Measured on M5 MacBook (10 cores, 24 GB, macOS 26.5), OLMoE-1B-7B-0125:
ref.json greedy 12/12 with IDOT=0 on both container and raw paths.
Co-authored-by: x <x@Mac.fritz.box>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
run_serve_mux (continuous-batching serve, SERVE_BATCH=1) used select() on
STDIN_FILENO to poll for incoming requests. On native Windows/MinGW,
select() on a non-socket handle routes to winsock and always returns
SOCKET_ERROR (-1), so the batch serve loop compiled and printed READY but
could never accept a request.
Replaced the platform-specific polling with a dual implementation:
- POSIX (Apple/Linux): select() as before
- Windows: WaitForSingleObject on the stdin OS handle (works on both pipe
and console handles), with PeekNamedPipe to confirm bytes are available
The rest of the serve loop is now shared across platforms — no more Windows
stub. The function is no longer guarded behind #if __APPLE__/__linux__.
Co-authored-by: woolcoxm <13604288+woolcoxm@users.noreply.github.com>
Five crash-class bugs found by static analysis, all in glm.c:
C1: forward_all (line ~2709) used a fixed stack buffer float row[8192] for
the final RMS norm. The config checker allows hidden_size up to 1<<20;
any model with hidden > 8192 would smash the stack. Every other hot path
in the file correctly uses falloc(D). Replaced with falloc(D) + free.
C2: read_arr (line ~3334) dereferenced json_get() return without a NULL
check. If ref_glm.json is missing 'prompt_ids' or 'full_ids', this is a
guaranteed NULL-pointer dereference / segfault. Added a NULL check that
returns NULL+0, and the caller in main() now validates the result.
C3: mux_submit (line ~3103) allocated tmp=malloc(maxctx*sizeof(int)) without
a NULL check, then passed it to tok_encode. OOM here writes through NULL.
Added a check matching the sibling allocations in the same function.
C4: serve_ctx_init (line ~3025) allocated s->hist without a NULL check, then
passed it to kv_disk_load which writes token IDs into it. Added a check
matching the falloc() pattern used by kv_alloc.
C5: rope_interleave (line ~896) used a fixed stack buffer float in[256] but
the config checker allows qk_rope up to 1<<16. Added a runtime bounds
check that exits with a clear message instead of silently overflowing.
GLM-5.2 qk_rope=64, well within bounds.
Co-authored-by: woolcoxm <13604288+woolcoxm@users.noreply.github.com>
cfg_load() (line ~911, reads config.json at every startup) and the oracle
reference loader (line ~3846, reads ref_glm.json) both had:
char *b=malloc(n+1); if(fread(b,1,n,f)!=(size_t)n){} b[n]=0; fclose(f);
The empty if-body {} silently ignored a short read. If fread returned fewer
bytes than n (truncated file, disk error, concurrent modification), b[n]=0
wrote a null terminator at position n — but only bytes 0..got-1 were valid.
json_parse then read uninitialized memory between got and n.
Fixed to null-terminate at the actual read position and warn on short read:
size_t got=fread(b,1,n,f); b[got]=0; fclose(f);
if((long)got!=n) fprintf(stderr,"warning: short read on %s ...");
Two instances fixed: cfg_load (config.json) and the ref_glm.json oracle reader.
Co-authored-by: woolcoxm <13604288+woolcoxm@users.noreply.github.com>
Two related crash bugs in the Python planning layer:
P1: Path.read_text() without encoding= defaults to locale.getpreferredencoding()
(cp1252 on most Windows installs). HuggingFace config.json is UTF-8 — if it
contains any non-ASCII byte (Chinese fields, accented chars, emoji in
metadata), read_text() raises UnicodeDecodeError. In doctor.py this is
caught and mis-reported as 'config.json is missing or invalid' (false
negative); in build_plan() called directly it is an uncaught crash.
Fixed: read_text(encoding='utf-8') in both resource_plan.py and doctor.py.
P2: int(cfg.get('kv_lora_rank', 0)) crashes with TypeError if the key is
present but null (JSON null). dict.get() returns the default only when
the key is ABSENT; a null value returns None, and int(None) raises
TypeError. The engine validates against malformed configs in C but the
Python planner did not.
Fixed: int(cfg.get(key) or 0) — coerces both missing and null to 0.
Applied to all 8 int(cfg.get(...)) calls in build_plan().
Co-authored-by: woolcoxm <13604288+woolcoxm@users.noreply.github.com>
P6: discover_gpus() used line.split(',', 3) to parse nvidia-smi CSV output.
If the GPU name contains a comma (e.g. 'Tesla, Inc. V100'), split produces
more than 4 fields, the int() parse fails on the wrong field, and the GPU
is silently dropped — doctor reports 'no NVIDIA device detected' with no
clue why. Fixed by using the csv module (handles quoted fields correctly).
P7: require_auth() used plain string != comparison for the API key
('Authorization' header vs expected 'Bearer <key>'). This enables a
timing side-channel that could leak the key byte-by-byte. Low impact when
bound to localhost (default), but serve() only warns when host is
non-localhost without a key, and users do expose on 0.0.0.0. Fixed by
using hmac.compare_digest (constant-time comparison).
Co-authored-by: woolcoxm <13604288+woolcoxm@users.noreply.github.com>
convert_fp8_to_int4.py could ingest FP8-blockscale, bf16 and f32
checkpoints, but not NVFP4 — the format NVIDIA modelopt emits for
REAP-pruned GLM-5.2 (quant_algo=NVFP4). Those expert weights are U8
(two e2m1 nibbles/byte) with a per-16-block FP8 scale sidecar
(.weight_scale) plus a per-tensor F32 global (.weight_scale_2).
Adds dequant_nvfp4(): e2m1 LUT (verified 1:1 against
ml_dtypes.float4_e2m1fn), low-nibble=even / high-nibble=odd unpacking,
per-block scale repeat-interleaved over group_size=16, times the global.
classify() now consumes the .weight_scale/.weight_scale_2/.input_scale
sidecars, and dequant() routes U8 tensors with a scale sidecar to the
NVFP4 path (keys is now required, so a stray U8 can't fall through to it).
Existing FP8/bf16/f32 paths are untouched.
Guards against silent corruption rather than trusting the input:
* group_size is fixed at 16 and the block-scale column count is
asserted == ceil(I/16); the old code inferred it from the data
(I // ncols), which misaligns silently on a padded/swizzled scale
layout and hard-crashes on a non-multiple-of-16 I — after a multi-GB
shard download. A partial trailing block is handled by slicing to I.
* .weight_scale_2 is asserted < 1: modelopt stores the small global and
MULTIPLIES; llm-compressor/compressed-tensors stores the reciprocal
(>= 1) and DIVIDES. The two are dtype-identical, so without this a
compressed-tensors checkpoint would corrupt every tensor by ~gscale^2.
--selftest-nvfp4 (no network) asserts the 16-code LUT, an encode->dequant
round-trip to <1e-9, and a dequant->colibri-int4 requant bound (<0.30;
the inherent ~0.17 is informational, not a precision claim).
Claude-Session: https://claude.ai/code/session_01DS7oc65c5RdA9V99otRCwt
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
Three crash bugs in convert_fp8_to_int4.py, all on the download path:
P3: 'import fcntl' at line 211 is Unix-only — ModuleNotFoundError on Windows.
The --indir test path returns before reaching it so tests pass, but
'--repo' on Windows hard-crashes. Guarded the import: try fcntl (Unix),
fall back to msvcrt.locking (Windows), skip if neither available.
P4: repo_info retry loop had range(999) — up to ~16 hours of retries on a
bad network, then fell through to line 395 where 'info' was unbound
(NameError). Capped at 10 retries and added an explicit error + return
when exhausted. Also added an early return if no safetensors shards are
found in the repo.
P5: if the shards list was empty (wrong repo, all filtered out), the
'for i, sh in enumerate(shards)' loop never executed and 'i' was unbound
at line 460 (NameError). Now caught by the early return from P4.
Co-authored-by: woolcoxm <13604288+woolcoxm@users.noreply.github.com>
The engine opens/mmaps every safetensors shard of the model (144+ files
for GLM-5.2). On macOS the default soft fd limit is 256, so a stock
terminal session fails while loading with:
<model>/out-00125.safetensors: Too many open files
the engine exited while loading
Raise the soft limit toward the hard limit (capped at 65536) in the
coli launcher before spawning the engine, so users don't need a manual
'ulimit -n' in every shell. No-op on Windows and on shells whose limit
is already sufficient.
Co-authored-by: Harvad Lee <hongyanab@gmail.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
rofl0r noted on #129 that uname describes the host shell, not the build
target, and that a gcc/clang toolchain should be queried with -dumpmachine.
Do that: derive the target triple from $(CC) -dumpmachine (e.g.
x86_64-w64-mingw32 / x86_64-pc-cygwin / x86_64-unknown-linux-gnu /
arm64-apple-darwin / powerpc64le-unknown-linux-gnu) and match
mingw/cygwin/darwin/powerpc64 in it.
The triple follows the toolchain rather than the shell, so detection is
correct under a native-Windows shell (no uname on PATH) and when
cross-compiling (make CC=x86_64-w64-mingw32-gcc on Linux), and it now also
distinguishes cygwin from mingw.
#129's OS=Windows_NT check and uname are kept as ordered fallbacks for the
rare toolchain that does not answer -dumpmachine, so no host regresses. The
CUDA/METAL macOS guards now use the derived DARWIN flag.
Validated on native Windows (WinLibs GCC 16.1.0 x86_64-w64-mingw32,
PowerShell, uname absent): make selects the Windows branch, glm.exe links
-static (no libgcc/libwinpthread/libgomp DLL deps), and all 7 dependency-free
C test binaries build and pass.
Co-authored-by: bopof <285767350+bopof@users.noreply.github.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
The tool tokenized every context with add_special_tokens=False, so pointing it at a
GLM snapshot scored the model out-of-distribution — the same bug @bokiko found in
eval_glm.py (#108).
Measured on this box (GLM-5.2 int4, engine SCORE mode, same 1023 target tokens,
only the conditioning prefix changed):
corpus no prefix +[gMASK]<sop>
natural prose PPL 29.2 PPL 9.4
markdown + code PPL 131.0 PPL 24.5
It does not just depress scores, it distorts sensitivity to numerical changes: an
exact-vs-IDOT kernel A/B measured without the prefix reported a penalty that halved
AND flipped sign on one of two corpora once the prefix was restored (#153). A
quantization-ablation tool that scores OOD is therefore worse than useless — it
produces confident deltas that are artifacts.
The GLM tokenizer does not add the prefix itself (add_special_tokens=True is a no-op
there), so it must be prepended explicitly. Auto-detected from the vocab rather than
opt-in, so it cannot be lost by omission; --prefix overrides. Models with no such
prefix are unaffected: OLMoE (the tool's default) has no BOS at all, add_special_tokens
True/False give identical ids, so the numbers in #108 measured on OLMoE still stand.
Canonical write-up of the GRAMMAR= draft source: mechanism, why it pays in a
disk-streaming MoE specifically, usage/knobs, measured expectations by workload
shape (span-density dependence, from the #146 A/Bs incl. corrections), lossless
guarantees, bench discipline, prior art. Linked from the README feature bullet.
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
* Add make install/uninstall targets
Neither the root Makefile nor c/Makefile had an install target,
so 'build+install' never actually installed anything (fixes#164).
Adds standard PREFIX/DESTDIR/BINDIR-respecting install and uninstall
targets that place glm and coli in $(BINDIR).
* Split install: coli in bin/, engine + support files in libexec/
Addresses feedback on #164 from JustVugg and yurivict: coli goes to
$(BINDIR), glm/olmoe and their Python support modules
(resource_plan.py, doctor.py, openai_server.py, tools/) go to
$(LIBEXECDIR) (default $(PREFIX)/libexec/colibri), matching typical
Unix/FreeBSD-port conventions for a wrapper vs. its internals.
coli now resolves the engine path in this order:
1. $COLI_ENGINE if set (explicit override)
2. glm next to itself (run-in-place from a source checkout, unchanged)
3. $(LIBEXECDIR)/glm (installed layout), also added to sys.path so
the Python support modules still import correctly
Also adds a 'bench' target (builds iobench) since only cuda-bench
existed before.
Tested locally (WSL2/Ubuntu):
- run-in-place: cd c && python3 coli info -> engine ready
- installed: make install PREFIX=$HOME/.local && ~/.local/bin/coli info -> engine ready (found via libexec)
- make uninstall cleans both bin/ and libexec/colibri/ fully
Problem: 'make clean', 'make test-c', and 'make check' use POSIX shell
constructs (for loop, rm -f, rm -rf) that require sh.exe. On native Windows
with WinLibs MinGW (no MSYS2, no Git Bash), there is no sh.exe on PATH.
GNU Make falls back to cmd.exe, which can't parse 'for test in ...; do'
or find 'rm', so these targets fail with 'test was unexpected' or
'CreateProcess error'.
Root cause: the Makefile's recipe lines assumed a POSIX shell is always
available. The IS_WIN detection (from #129) catches the platform but the
shell-dependent targets were never made portable.
Fix: replace the shell-dependent constructs with small Python helper scripts
(Python is already a project dependency for test-python, convert, bench).
This works from cmd.exe, PowerShell, Git Bash, and MSYS2 alike.
Changes:
- tools/run_tests.py (new): runs each C test binary, exits non-zero on the
first failure. Replaces the 'for test in ...; do ./$test || exit 1; done'
shell loop in test-c.
- tools/clean.py (new): removes build artifacts and test binaries. Replaces
'rm -f' and 'rm -rf' in clean. Only removes executables (.exe) and known
artifact names — never .c or .py source files.
- Makefile: PYTHON defaults to 'python' on Windows (not 'python3'); test-c
and clean now call the Python helpers instead of shell constructs.
Verified from native PowerShell (no sh.exe): make clean removes 8-19
files/dirs, make test-c runs all 7 C test suites, source files survive.
Also verified from Git Bash (sh.exe present): behavior unchanged.
Co-authored-by: woolcoxm <13604288+woolcoxm@users.noreply.github.com>
On Windows, os.access(path, os.X_OK) always returns True for any existing
file (NTFS has no execute bit; executability is governed by file extension,
not mode bits). So the test_non_executable_engine test — which chmods the
engine to 0o644 and expects 'fail' — could never pass on Windows.
Fixed doctor.py to use a platform-aware check: on Windows, any existing
file is executable; on POSIX, honor the mode bits via os.access(X_OK).
Fixed the test to assert the correct per-platform expectation: 'pass' on
Windows (chmod is a no-op for executability), 'fail' on POSIX.
Co-authored-by: woolcoxm <13604288+woolcoxm@users.noreply.github.com>
Four jobs, all fast, all free on public repos:
- engine: make glm + make test-c (Linux, CPU, native arch)
- engine-cuda-syntax: nvcc compile-only check (no GPU needed, catches
backend_cuda.cu syntax/type errors before they reach a real machine)
- web: npm ci + build + vitest (Node 22, cached deps)
- python: unittest discover on c/tests/test_*.py (Python 3.12)
Triggers on push/PR to dev and main. No 20-minute pipelines — the whole
matrix finishes in under 3 minutes on GitHub-hosted runners.
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>
Fabio Rovai