The DSA lightning indexer selects the top-index_topk (2048) context keys to
attend to by finding the threshold = keep-th largest attention score. It
previously full-qsorted all nk scores per layer per token (O(nk log nk)) just
to read one pivot value, then scanned the original array in position order to
build the kept set.
Replace the qsort with partial_select_desc (Hoare quickselect, median-of-three,
descending): O(nk) average to partition the keep largest into a[0..k), then the
threshold is min of that block. The two position-order scans (>thr then ==thr)
are UNCHANGED, so the kept-position set is bit-identical -- a stronger contract
than #335's sampling heap (which was multiset-only because the heap was unstable
and changed accumulation order). The quickselect pivot IS by definition the
keep-th largest, so the new threshold equals old tmp[keep-1] exactly.
Measured (bench_dsa_select, keep=2048, median of 2000 reps):
nk=2049: 119us -> 5.7us (21x)
nk=8192: 626us -> 43us (15x)
nk=32768: 2.8ms -> 0.28ms (10x)
nk=65536: 6.6ms -> 0.47ms (14x)
The gap widens with context (linear vs n-log-n). DSA only fires past index_topk,
so this is precisely the long-conversation regime where decode latency matters.
Adds test_dsa_select (in TEST_BINS): 129 cases asserting element-wise identical
kept-set vs an independent qsort reference across shapes (random, peaked,
sorted, reverse-sorted, tie-plateau, all-equal) and edges (keep==1, keep==nk).
Also directly checks the partition invariant.
Adds bench_dsa_select (on-demand, NOT a gate): reproduces the table above.
test_topp proves correctness; bench_topp measures the latency claim. It re-implements
the OLD dist_build (full-vocab qsort) inline on a private buffer and times it against
the REAL new dist_build over identical inputs in one process: V=151936, temp=0.7,
3 shapes (realistic / uniform / plateau) x 4 nuc values (0.5/0.9/0.95/0.99), 2000
timed reps each, median reported. Deliberately NOT in TEST_BINS -- it's a microbench,
not a gate. Build on demand: make tests/bench_topp && ./tests/bench_topp
dist_build() sorted the entire 151936-entry vocab by probability (qsort) on every
sampled token whenever 0 < g_nuc < 1 — the serve default — and again per draft
position under rejection sampling. Measured cost: 5.6-8.0 ms/call; the actual work
is finding the few-hundred-token head whose cumulative mass reaches g_nuc.
Replace the full qsort + linear scan with a max-heap partial select:
- Floyd heapify g_pidx over V by descending g_pbuf prob (O(V), cache-friendly)
- pop winners to the array's high end until cum >= g_nuc (k * O(log V))
- the remaining heap prefix IS the tail -> zero it, renormalize the head
Winners land in g_pidx[out..V-1] in descending order, so s2 accumulates in the
same order as before -> head is unchanged on tie-free shapes (ties were already
unspecified under the unstable qsort). All four dist_build/dist_sample contract
properties hold: g_pbuf stays id-indexed, g_pidx stays internal, the tail is
fully zeroed, the head renormalizes to 1.
No API change, no caller change, no new globals.
c/tests/test_topp.c (new): drives the REAL dist_build via the include-glm.c
pattern against an independent double-precision reimplementation of the OLD
algorithm. 123 cases: 6 sizes (1..1519) x 5 shapes (uniform/peaked/geometric/
plateau/sharptail) x 4 nuc values, plus the g_nuc>=1 guard-off paths and V=1.
Tie-free shapes compare head values within 1e-6 rel (float vs double renorm
noise); tie shapes compare value-multisets. No scratch files -> builds clean on
Windows MinGW without the unmerged mkdtemp shim (#352).
Same fix pattern as test_stops: mkdtemp with a CWD-relative template
(MinGW resolves Windows paths; /tmp is not one), and the fork/pipe/
truncate-based truncation subtest compiles only where those exist —
Windows still runs the cross-platform chunk-loop content check.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Two latent bugs in every st.h reader, both hit in the field:
- a single pread caps at ~2^31 bytes on Linux, so any tensor past
2.1 GB (bf16 embed/unembed tensors of large models qualify) came
back silently truncated with perror printing '... : Success'
(errno untouched by a short read) — the same misleading-error
symptom glm.c fixed for its own reads in #236;
- no EINTR retry.
st_pread_full loops in ST_PREAD_CHUNK pieces (1 GB default, override
for tests), retries EINTR, and reports offset + progress on failure.
All five read sites converted; behavior on well-formed files is
byte-identical (GLM oracle re-verified on this branch: 32/32).
tests/test_st_pread builds with -DST_PREAD_CHUNK=7 so a 96-byte tensor
exercises the multi-chunk loop, and forks a child against a shard
truncated after st_init (init's static bounds check means the pread
path only fires when a file shrinks under a live handle) asserting
exit(1) with a 'short read' message and no 'Success'.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
My test_stops.c called mkdtemp("/tmp/coli_stops_XXXXXX"). These binaries are
built by MinGW into NATIVE Windows .exe files, which resolve Windows paths —
"/tmp" is not one, so mkdtemp failed ENOENT and `make check` went red on the
windows job the moment #143 gave us cross-platform CI.
Now a relative template in the CWD, which is what test_compat_direct.c already
does (`#define TMPF "test_direct.tmp"`). test_uring.c uses /tmp but is
Linux-only by construction (Makefile guards it behind $(LINUX)); I copied the
wrong neighbour.
Worth being precise about what this was, because the red job looked scarier
than it was: Windows is fine. `make check` built the engine, ran the whole C
suite, and passed everything else — test_i4_grouped, kv_alloc, compat_direct —
before tripping on my temp path. The CI caught a bug in the test, not in the
port. That is #143 earning its keep on day one.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
The engine armed its stop tokens from config.json's eos_token_id and nothing
else. That trusts metadata written by third-party conversion tooling, which is
a thing we already know goes wrong: the README documents a mirror shipping
int4 MTP heads that silently give 0% draft acceptance. GLM-5.2 declares THREE
eos ids (<|endoftext|>, <|user|>, <|observation|>); a converter that rewrites
config.json with a reduced list leaves the engine stopping on fewer tokens
than the model emits, and the missed ones get detokenized and printed into the
chat as literal text while generation runs past the end of the turn.
Two independent defenses:
- eos_token_id is now unioned with generation_config.json, which is
HuggingFace's authority for generation (config.json often carries a
partial legacy copy). An extra stop is harmless; a missing one is not.
- every added-token the TOKENIZER marks "special":true is armed as a stop,
whatever the configs say. Those are control tokens (<|user|>, <|assistant|>,
<sop>, [gMASK], the image/video/audio markers) and are never legitimate
content in a reply -- GLM itself lists three of them as official eos.
<think>/<tool_call>/<arg_key> are "special":false and are deliberately NOT
swept up: they are real output. tok.h was parsing added_tokens but throwing
the "special" flag away, so the distinction wasn't available to anyone.
On the real per-row checkpoint this takes the armed set from 3 to 18:
[stop] 18 stop tokens: 154820 154827 154829 154821 ... (15 from the
tokenizer's special set)
Honesty about scope: this is hygiene for a class of bug, NOT a fix for the
trailing-junk report on #298 that prompted it. I hypothesised @woolcoxm's g64
checkpoint had lost eos ids in conversion; he checked, and it hadn't -- his
config arms all three correctly. The emit path is also innocent: is_stop() is
checked BEFORE emit() at every one of the four call sites (4215, 4256, 4908,
4987), so a correctly-armed stop cannot be printed. His trailing junk is still
unexplained and is more likely quantization noise. What this commit buys is
that a checkpoint we don't control cannot leak control tokens into a reply,
which was true before and is not now.
tests/test_stops.c covers both defenses: the union, a missing
generation_config.json, BOTH configs mutilated (the tokenizer still stops all
five control tokens while leaving <think> alone), and T=NULL (the validation
path keeps config-only behaviour).
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
@woolcoxm's matmul_i4_grouped_pair reads x once instead of twice for the
gate+up pair. Verified here against his branch (86e91b1 merged onto dev):
correct to ~2e-8 relative vs the double reference, and BIT-EXACT against two
separate matmul_i4_grouped calls on aligned shapes -- which is the shape the
real g64 checkpoints have (I = 2048 / 6144, gs = 64). His kernel is good.
Guarded behind COLI_HAVE_GROUPED_PAIR since the function only exists on that
branch; add -DCOLI_HAVE_GROUPED_PAIR to the test's Makefile rule when #298
lands and the pair cases activate.
The checks are deliberately asymmetric, and the reason is worth recording.
Bit-exactness is asserted ONLY when I % gs == 0: there every group is covered
by the AVX2 body, whose accumulation order matches the unfused kernel, so any
difference is a real bug. With a partial last group the tail falls to scalar
code and the compiler may contract/reassociate the fused body differently,
producing ~1e-7 differences -- rounding, not logic. My first version demanded
bit-exactness everywhere and duly "found" a bug in his kernel that did not
exist; the tell was that only `up` differed and never `gate`, which is FP luck
rather than a code path. Correctness is checked everywhere against the double
reference; identity only where identity is actually implied.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
matmul_i4_grouped is the reference the CUDA fmt=4 port (#298) is expected to
reproduce, and it had no test of its own. @woolcoxm is currently debugging a
CUDA backend against an oracle nobody had verified, which is two moving
targets at once -- and he can't cross-check on CPU, since a 5-prompt run
takes 8 hours on the 744B model.
This checks matmul_i4_grouped against a plain-C reference that dequantizes
nibble -> (v-8)*scale[i/gs] and accumulates in double, over 11 shapes: I a
clean multiple of gs, a partial last group (the glen clamp), odd I (the
scalar nibble tail), gs > I, gs=16/64/128, S>1, and the nibble extremes
0x00/0xFF -- which decode to -8/+7 because the format is offset-encoded, not
two's complement. Reading that backwards turns 15 into -1 and looks like
data-dependent noise rather than a bug.
All 11 shapes match to ~1e-8 relative, so the CPU kernel is exact and can be
trusted as the reference.
One note on the tolerance, because the first draft of this test got it wrong
and "found" a bug that wasn't there: the error is compared against the sum of
|terms|, not against |result|. A dot product of signed terms can land near
zero through cancellation, and then a 1e-6 absolute error -- ordinary f32
accumulator precision -- reads as a 1e-3 relative one. A wrong scale index or
a wrong group boundary shifts the result by a fraction of the terms, so it is
still caught at 1e-6.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
The engine already tracks where each turn's wall time goes (expert-disk
service, I/O wait, expert matmul, attention, lm_head) — it just only spoke
at exit or under PROF=1. Stream it instead:
- glm.c: mux serve emits a per-turn "PROF" protocol line next to TIERS/HITS
(window deltas per request, same convention as the STAT hit%); the phase
window base is now always captured (a few loads per request).
- openai_server.py: parses PROF into a 120-turn rolling window and serves it
at /profile (read-only, same trust level as /health).
- web: new Profiling tab — stat tiles (tok/s, wall, tokens/forward, disk
service), wall-time composition bars for the last turn and the window,
per-turn throughput and stacked phase columns with hover readouts, and a
table of recent turns. Disk service is shown apart from the stack: it
overlaps with compute, so only the I/O wait the compute thread felt counts
inside wall time. Phase colours are a CVD-validated set with gaps + legend
+ table so identity never rides on colour alone.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01WhTmF8yvEBgSkUKSVfZF7P
Select a portable architecture from the compiler target instead of forcing x86-64-v3 on every platform. On macOS, only enable Homebrew OpenMP when its header and library actually exist, preserving the dependency-free fallback.
On Windows the engine self-exec OMP tuning never runs (Linux/FreeBSD-only)
and posix_fadvise readahead is a compat.h no-op, so a stock Windows run
leaves large measured wins on the table. The launcher now setdefaults, on
win32 only, each independently overridable by setting the variable:
- OMP_WAIT_POLICY=active, GOMP_SPINCOUNT=200000, OMP_DYNAMIC=FALSE,
OMP_NUM_THREADS=<physical cores> (parity with the glm.c self-exec block;
COLI_NO_OMP_TUNE disables exactly this block, presence-based like the
engine). OMP_PROC_BIND/OMP_PLACES deliberately omitted and also removed
from environment_for_plan on win32: MinGW libgomp has no affinity support
("Affinity not supported on this configuration").
- DIRECT=1: unbuffered expert reads. Measured on a 9950X3D + Samsung 9100
PRO Gen5 + Win11: iobench 10.68 GB/s O_DIRECT vs 9.03 buffered (warm);
end-to-end REPLAY 0.48 -> 1.02 tok/s. Matches #162 (1.47x same class).
- PIPE=1: load/matmul overlap, byte-identical output; +8% on top of DIRECT
(PIPE_WORKERS untouched at 8 - 4/8/16 swept flat on Gen5).
- PILOT_REAL=1: real cross-layer prefetch, the only working prefetch on
Windows; +11% and expert hit rate +19 points.
Full ladder methodology and numbers: 96-token greedy REPLAY, one lever per
step, medians of 3-4 runs (see the fork tuning doc referenced in the PR).
tests/test_env_defaults.py covers the defaults, explicit-override-wins,
the kill-switch scope, and the non-win32 no-op.
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
vmmlaq_s32 computes a 2x2 int32 tile (2 weight rows x 2 activation
rows) per instruction on 8-deep segments. Tile o and s in pairs,
halving weight traffic and doubling per-instruction work at S>=2.
Four independent accumulators over a 64-deep unroll keep the loop
throughput-bound (a single chained accumulator measures no better
than SDOT: latency-bound). S=1 and all tails (odd o, odd s, I not a
multiple of 16/32) keep the existing SDOT/scalar code, and scales
apply in the same order, so results are bit-identical.
Compile-time gated on __ARM_FEATURE_MATMUL_INT8. The default Darwin
build passes no -mcpu and is byte-identical (still SDOT, IDOT_KERNEL
"neon"). Opt in with ARCH=native (new Darwin Makefile knob, appends
-mcpu=<arch>), which reports IDOT_KERNEL "neon-i8mm". The same gate
lights up on any aarch64 with i8mm (Graviton3+, Grace).
test_idot grows a driver-level exactness check through matmul_qt_ex:
fmt 1 and 2, S in {2,3,4,5,8}, O in {1,2,3,64,65}, I in {16,17,100,
1408}, bitwise float equality against a plain-C reference. Green on
both build flavors.
Measured on an M5 Pro (18 threads, matmul_qt_ex microbenchmark at
GLM-5.2 expert shapes, best of 3 process runs, vs the SDOT baseline):
gateup int4 S=8 499.7 -> 1076.6 GF/s (+115%)
gateup int8 S=8 512.9 -> 1166.3 GF/s (+127%)
down int4 S=8 696.9 -> 1186.0 GF/s (+70%)
S=1 decode rows unchanged (SDOT path untouched)
opencode / the ai-sdk OpenAI-compatible client sends a large default max_tokens
(> the server's --max-tokens cap, 1024 by default), and generation_options
returned 400 "must be an integer between 1 and 1024" — even for a trivial
"hello". OpenAI-compatible servers clamp to their own ceiling rather than
reject. Now max_tokens > limit is clamped to limit; only non-int / non-positive
values are a hard error. Test updated to assert the clamp + keep the <1 reject.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Cut Windows decode disk I/O from 2.06s/tok to 1.70s/tok (budget=4), meeting the
<=2s/tok target. Two changes on the pread expert-load path:
1. compat.h: replace the posix_fadvise no-op with a real WILLNEED cache-warmer
(overlapped ReadFile into a scratch buffer -> populates the standby page cache
so the later synchronous pread faults from RAM). Re-arms the existing
expert_prefetch/PILOT/next-block prefetch chain on Windows. DONTNEED stays a
no-op (matches macOS; Windows standby-list trimming self-regulates).
Measured: hit rate 16.4% -> 27.6%.
2. glm.c: flip PIPE (async expert-load thread pool) from default OFF to default ON
on Windows. Dispatches expert pread onto worker threads so loads overlap the
matmul, instead of blocking serial load-then-compute. PIPE=0 opts out.
Measured: expert-disk 65.9s -> 54.3s (-18%).
Also adds compat_fadvise assertions to tests/test_compat_direct.c (data integrity
after cache-warmer, safe no-op on bad fd / non-WILLNEED).
mmap (CreateFileMapping/MapViewOfFile) was implemented and tested at length but
reverted: it regressed on Windows (RSS bloat from touched mapped pages collapsed
the expert cache via ullAvailPhys — a fundamental Windows-vs-Linux difference).
Full findings + the dead-end analysis recorded in issue_diskio.md.
c/tests/test_schema_gbnf and c/tests/test_compat_direct were committed
as Linux x86-64 ELF executables (slipped in via #111). On any other
platform make test-c considers them up to date and execs them, failing
with OSError: [Errno 8] Exec format error. Remove them and add the two
names to .gitignore alongside the other test binaries already listed,
so each platform rebuilds its own.
kv_alloc guards every KVState free with if(k->Lc) precisely so it can be
called again on the same KVState (context resize, slot re-init). Exercise
that path: allocate, touch the cache, allocate again at a larger size.
Fails at this commit with a double-free (fixed in the next commit):
malloc: *** error for object 0x7: pointer being freed was not allocated
* Fuse CUDA expert MLP execution
* Group CUDA expert transfers by device
* Instrument grouped CUDA expert execution
* Bound grouped CUDA decode scratch
* Execute expert groups across GPUs in parallel
* Release host backing for multi-GPU experts
* Define quality-preserving memory policies
* Overlap cold expert loading with resident compute
* Adapt expert placement with session LFRU
* Fuse q4 expert gate and up dispatch
* Plan CPU work on physical cores
* Batch grouped expert CUDA kernels
* Separate VRAM and RAM expert placement
* Add ragged multi-sequence decode forward
* feat(runtime): add continuous decode scheduler
* Route concurrent API requests through batch scheduler
* Harden multiplex request lifecycle and framing
* Cancel disconnected multiplex requests
* Bind API port before starting the engine
* fix automatic KV slot allocation
* add native int4 Tensor Core grouped GEMM
* add Tensor Core throughput benchmark
* optimize packed int4 low-row kernels
* add asynchronous CUDA staging streams
* document validated six-GPU dense acceleration
* tune six-GPU expert hot set
* raise validated expert hot-set target
* add CUDA MLA absorption core
* fuse grouped expert gate and up projections
* Warn for explicit lossy routing flags
* Add full-resident expert placement mode
* Adapt VRAM expert slots to live routes
* Accelerate int4 matvec on AVX-512
* Reduce AVX-512 and RoPE decode overhead
* Seed every GPU expert layer after prefill
* Limit live GPU swaps during decode
* CUDA batch MLA attention, kv_b head-sharding, fused o_proj, expert-group dispatch, W4A16 kernels
Lab-qualified on the 6x RTX 5090 machine (914-token request benchmark):
- batch MLA absorption kernel (COLI_CUDA_ATTN=1): whole-batch attention on
device, 154.8s -> 102.4s
- attention -> o_proj fusion on the layer device: -> 97.4s
- kv_b head-sharding across cards (COLI_CUDA_ATTN_SHARD=1), no weight
duplication: -> 94.05s
- per-device expert-group dispatch with pinned-buffer async transfers,
W4A16 tensor-core kernels for the shared expert, OMP hot-thread tuning
Negative results (reverted, kept out): GPU-side weighted scatter-add
(atomics + per-layer D2H lose 43.8%), shared-expert fused small-batch
kernel (-38.8%), W4A4 grouped tensor cores (int4 activations corrupt
output). Details in the lab research log.
* GPU resident pipeline: device-resident prefill attention chain, GPU expert groups in prefill, batched router, W4A16 mixed dispatch
COLI_CUDA_PIPE=1 keeps the prefill data plane on the layer home device;
control flow (routing, cache/pin management) stays on CPU. Any CUDA
failure falls back to the unchanged CPU path.
- Device primitives + unit tests (tests/test_pipe_cuda.cu): rmsnorm
(strided), interleaved RoPE, silu-mul, residual add, fixed-order row
merge (no atomics), device-input GEMM, persistent per-device scratch.
All verified against the engine's CPU math on SM120 (worst 1.2e-5).
- attn_pipe_prefill: q_a -> norm -> q_b -> rope -> kv_a -> norm -> rope ->
batch attention -> o_proj in one device chain (q_a/q_b/kv_a colocated
with kv_b); only the final [S,D] and the new KV rows return to host.
Attention 41.2s -> 30.8s on the 1571-token benchmark.
- Prefill batch-union now uses the GPU expert groups (previously gated to
S<=64, leaving all VRAM-resident experts idle during prefill - measured
21ms of GPU expert time in a 148s prefill). Expert phase 78.9s -> 69.0s.
- Router computed as one batched matmul instead of S sequential rows
(bit-identical math).
- W4A16 tensor-core path for expert groups (COLI_CUDA_TC_W4A16=1) with
row-count mixed dispatch: >=16 rows per expert use tensor cores, smaller
batches keep the naive kernel (tensor cores measured negative below
~16 rows). Expert phase 69.0s -> 64.3s, decode unaffected.
Net on the 1571-token prefill benchmark: 148.8s -> 114.3-126.8s
(component timings stable across runs; wall drifts +-3-5s because
.coli_usage placement learning shifts the expert tiers between runs).
PROFILO now also prints the prefill-phase breakdown.
* Skip OMP hot-thread tuning when CUDA is enabled
The active-spin worker team measured 66.9s->20.9s on the CPU-only Zen5
build, but on the six-GPU full-residency workload the spinning workers
contend with the CUDA dispatch threads: ~4x slower prefill with the
process stuck near 1.8 cores. Gate the tuning on COLI_CUDA so each
configuration keeps the behavior it was measured to prefer.
* Inc.2a: sparse layers fully resident on the layer device, residual hops cards at layer boundaries
COLI_CUDA_PIPE=2 keeps the residual stream on the layer home device for
consecutive sparse layers (cudaMemcpyPeer at boundaries): in/post norms,
attention chain, both residual adds and the shared-expert MLP run on
device. Per layer only the post-norm activations (router + CPU-tier
experts + group gather), the new KV rows and, on DSA indexer layers, the
pre-attention norm leave the card. Per-layer transfers drop from ~130MB
to ~70MB. A device-side snapshot at layer entry makes any mid-layer CUDA
failure fall back to the unchanged CPU path idempotently.
1571-token prefill: 127.1s (PIPE=1 control) -> 117.6/118.9s, components
attention 30.8->26.1, other 31.8->22.5-24.5; output verified coherent
against the control.
* Head-sharded attention inside the pipe: negative on PCIe star topology, gated opt-in
Slicing q per card from the home device and collecting ctx back
serializes ~95MB/layer through the home card's PCIe link: attention
26.1s -> 41.4/44.4s on the 1571-token benchmark (two repeats), wall
117.6 -> 135-138s. The standalone host-path sharding won because six
cards uploaded from host RAM in parallel; a home-device star has no
such parallelism without NVLink. Kept behind COLI_CUDA_PIPE_SHARD=1
for interconnects where peer bandwidth does not share one root port.
* Inc.3: device-resident KV shadow for decode attention
Decode re-uploaded the whole latent+rope window per layer per token
(~300MB/token at 1571 context). Each layer now keeps a device shadow of
the compressed KV on its kv_b card, bulk-synced when behind and appended
incrementally; the host cache stays canonical. Invalidation on kv_bind
(slot switch), kv_alloc (resize) and on any overwrite of mirrored rows,
with the legacy full-upload path as fallback.
Measured (COLI_CUDA_PIPE gate): short-context decode 5.48 -> 5.59/5.87
tok/s, 1571-context decode 4.14 -> 4.22 tok/s. Decode remains CPU-expert
bound; the shadow removes the transfer tax, not the compute.
* tools: unified user-experience benchmark (bench_ux.sh)
Two fixed scenarios (short chat, long-document QA), TTFT + decode tok/s
+ first-line drift check, TEMP=0 DRAFT=0 enforced, medians over REPS
runs. Encodes the measurement discipline from the lab record: same
binary per comparison, judge medians because .coli_usage placement
learning drifts wall times between runs.
* tools: bench_ux.sh executable bit
* gitignore compiled test binaries
* tools: expert_atlas.py — measure per-expert topic affinity (#175)
Diffs .coli_usage across 10 themed probe batches (code/math/chinese/
prose/science/law/poetry/structured/translation/casual, 3 prompts each)
driven through a running API server — one engine load total. Every
touched expert gets a topic-affinity vector, entropy, and a specialist/
generalist label; output experts.json feeds the Brain page hover.
* serve: persist .coli_usage after every turn in mux mode, not only at exit
run_serve_mux saved the learning cache once at shutdown; a crash lost
the whole session's routing history, and live consumers of the file
(expert_atlas.py diffs it between probe batches) saw a frozen snapshot.
Now saved per turn like the interactive path (165KB write, negligible).
* web: Brain hover shows measured expert atlas when published
If /experts.json (from tools/expert_atlas.py, #175) is served next to
the app, the tooltip upgrades from the depth heuristic to measured
data: specialist/generalist label, entropy, and the top-3 topic
affinities. Row index maps to real layer (row+3, last row = MTP 78).
Falls back to the heuristic when no atlas is published.
---------
Co-authored-by: JustVugg <JustVugg@users.noreply.github.com>
* 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>
* 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>
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>
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.
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).
* 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>
* 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>
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>
ZacharyZcR