e871fb792d0b6041d437d5c451b0d495de7fd740
78 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
 Vincenzo
|
5d4c3aa11b |
Merge pull request #257 from woolcoxm/windows-optimizations
Windows disk I/O: pread + PIPE + compat_fadvise (1.70s/tok, mmap reverted) |
||
|
Vincenzo
|
4d0438314a |
Merge pull request #246 from woolcoxm/feat/diskio-kv-batching
Disk I/O: KV cache write batching + persistent file handle |
||
|
Vincenzo
|
a6895445a9 |
Merge pull request #254 from woolcoxm/feat/expert-budget
EXPERT_BUDGET=N: miss-aware cap on distinct experts/layer (+75% decode tok/s, 6x prefill on low-RAM hosts) |
||
|
Vincenzo
|
d9266e39e3 |
Merge pull request #248 from ZacharyZcR/serve/mux-kv-diag
serve: emit the [API] KV prefix-reuse diagnostic in mux mode (#153) |
||
|
woolcoxm
|
25219de45b |
windows: PIPE default ON + compat_fadvise WILLNEED cache-warmer (pread path)
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. |
||
|
woolcoxm
|
9c24b1eceb | Merge branch 'feat/expert-budget' into windows-dev | ||
|
woolcoxm
|
dd0d60692d |
experiment: miss-aware EXPERT_BUDGET — keep all cache hits, only drop misses
The original budget dropped experts blindly — even cached ones that cost
zero disk I/O. The miss-aware version pre-scans pin/ecache residency
before applying the budget:
- ALL cache hits are kept (free to compute, no disk I/O)
- Only misses compete for the remaining budget slots
- Miss budget = EXPERT_BUDGET - nhits (min 0)
Results (budget=4, same prompt/config as before):
Original Miss-aware
tok/s 0.33 0.36 (+9%)
hit rate 16.4% 38.2% (2.3x)
prefill 8.9s 5.6s (1.6x faster)
decode 97.4s 88.9s (9% faster)
The hit rate doubling is the key quality signal: the model now gets the
full contribution from all resident experts plus the top-4 new loads,
instead of losing some hits to the budget.
|
||
|
woolcoxm
|
413b370fbf |
experiment: EXPERT_BUDGET=N — cap distinct experts/layer, up to 1.8x faster decode
Add EXPERT_BUDGET env var that caps the number of distinct experts loaded per layer across the batch-union. When the union exceeds the budget, keeps only the highest-aggregate-gate-weight experts and drops the rest from idxs[] so they're never loaded from disk. Complementary to TOPP (per-position) — this trims the cross-position union that multiplies under MTP/prefill. Based on MoE-Spec (arXiv 2602.16052): 'top 32 of 64 experts capture 93% of routing weight.' Measurements (GLM-5.2 744B, 24GB RAM, cap=2, MTP=0, 32 tokens): Baseline (budget=0): 0.18 tok/s, 9.3% hit, 176s decode, 39s prefill EXPERT_BUDGET=12: 0.19 tok/s, 14.0% hit, 171s decode, 12s prefill EXPERT_BUDGET=6: 0.26 tok/s, 21.0% hit, 123s decode, 7s prefill EXPERT_BUDGET=4: 0.33 tok/s, 16.4% hit, 97s decode, 9s prefill Budget=4 nearly doubles decode speed (+83%) and 4x's prefill speed by halving disk reads per layer. Default OFF (EXPERT_BUDGET=0). |
||
|
woolcoxm
|
71c262ce1a |
diskio: KV write batching + persistent handle; generators: unfuse experts
Two independent fixes validated end-to-end on fresh fixtures: 1. KV cache disk I/O (issue_diskio.md opportunities #1 + #4): - kv_disk_append: fopen/fclose every turn -> persistent FILE* kept open for the engine lifetime, lazy open on first append, closed in serve_ctx_free. Eliminates per-turn handle creation overhead. - kv_disk_append: ~157 small fwrites per position -> one contiguous record memcpy'd into a staging buffer then a single fwrite per position. The staging buffer grows on demand via realloc. - kv_disk_truncate: closes the persistent handle before truncating so the file actually shrinks on disc, then reopens lazily. - KVState gains disk_fp, disk_buf, disk_buf_cap fields. - Verified: serve-mode round-trip, write 11 tokens then reload and resume with no re-prefill, then append 8 more and reload to 19. 2. Expert weight unfusing in test-model generators: - The real GLM-5.2-FP8 checkpoint stores routed experts UNFUSED as per-expert 2-D tensors, each with its own _scale_inv. HF fuses gate+up into a single 3-D gate_up_proj for compute efficiency. - The converter and C engine both expect the unfused layout. The fused 3-D tensors were silently skipped by the converter, and the engine crashed with missing-tensor errors. - New unfuse_experts in glm_fp8_emit.py splits gate_up_proj and down_proj into per-expert 2-D tensors. Called after reference generation but before saving, in both generators, both FP8 and bf16. - Also fixed: make_glm_oracle.py FP8 round-trip guard used p.dim()<2 which let 3-D fused experts through and crashed fp8_block_quantize. Changed to p.dim()!=2 to match the converter ndim!=2 guard. Validated full chain on fresh fixtures: generator --fp8 -> 570 e4m3 tensors + 629 scale_inv, was 90 when fused converter --group-size 0 -> per-row int4 fmt=2, engine loads clean converter --group-size 128 -> grouped int4 fmt=4, 8-16x more scales, engine loads clean, fmt=4 auto-detected in both mmap and slab paths dequant error: grouped 1.14-1.22x lower than per-row vs FP8 source |
||
|
ZacharyZcR
|
d04d99e039 | serve: print the KV prefix-reuse diagnostic in mux mode too — [API] KV slot line existed only in the legacy \x02PROMPT path (#153) | ||
|
woolcoxm
|
e2ad6c72e6 |
diskio: KV write batching + persistent handle; generators: unfuse experts
Two independent fixes validated end-to-end on fresh fixtures: 1. KV cache disk I/O (issue_diskio.md opportunities #1 + #4): - kv_disk_append: fopen/fclose every turn -> persistent FILE* kept open for the engine lifetime, lazy open on first append, closed in serve_ctx_free. Eliminates per-turn handle creation overhead. - kv_disk_append: ~157 small fwrites per position -> one contiguous record memcpy'd into a staging buffer then a single fwrite per position. The staging buffer grows on demand via realloc. - kv_disk_truncate: closes the persistent handle before truncating so the file actually shrinks on disc, then reopens lazily. - KVState gains disk_fp, disk_buf, disk_buf_cap fields. - Verified: serve-mode round-trip, write 11 tokens then reload and resume with no re-prefill, then append 8 more and reload to 19. 2. Expert weight unfusing in test-model generators: - The real GLM-5.2-FP8 checkpoint stores routed experts UNFUSED as per-expert 2-D tensors, each with its own _scale_inv. HF fuses gate+up into a single 3-D gate_up_proj for compute efficiency. - The converter and C engine both expect the unfused layout. The fused 3-D tensors were silently skipped by the converter, and the engine crashed with missing-tensor errors. - New unfuse_experts in glm_fp8_emit.py splits gate_up_proj and down_proj into per-expert 2-D tensors. Called after reference generation but before saving, in both generators, both FP8 and bf16. - Also fixed: make_glm_oracle.py FP8 round-trip guard used p.dim()<2 which let 3-D fused experts through and crashed fp8_block_quantize. Changed to p.dim()!=2 to match the converter ndim!=2 guard. Validated full chain on fresh fixtures: generator --fp8 -> 570 e4m3 tensors + 629 scale_inv, was 90 when fused converter --group-size 0 -> per-row int4 fmt=2, engine loads clean converter --group-size 128 -> grouped int4 fmt=4, 8-16x more scales, engine loads clean, fmt=4 auto-detected in both mmap and slab paths dequant error: grouped 1.14-1.22x lower than per-row vs FP8 source |
||
|
Vincenzo
|
3fd47b7bbd |
Merge pull request #242 from woolcoxm/feat/grouped-quant-fmt4
Group-scaled int4 (fmt=4): one scale per 128 elements — fixes incoherent output |
||
|
woolcoxm
|
498ab0c20e |
experiment: group-scaled int4 (fmt=4) — one scale per 128 elements, not per row
Root cause of gibberish output: the int4 quantization uses one F32 scale per output row (2048 scales for a 2048x6144 matrix). The FP8 source has 128x128 block scales — 48x finer. This destroys reasoning while keeping surface fluency. Changes: - Converter: quant_int4_grouped() with --group-size 128 arg. Same nibble packing, but one scale per group of 128 elements along the input dim. - Engine QT struct: added 'gs' field (group size, 0=per-row backward compat) - Engine qt_from_disk: auto-detects fmt=4 when scale array is O*ceil(I/128) elements instead of O. Old per-row models (fmt=2) work unchanged. - Engine matmul_i4_grouped(): AVX2 kernel that applies per-group scales. Accumulator resets at each group boundary: dot(x[grp],w[grp]) * scale[grp]. - Engine matmul_qt_ex: dispatches to grouped kernel for fmt=4 (always exact, no IDOT approximation since the point is quality) - Engine expert_load: both mmap and slab+pread paths detect fmt=4 from scale array size and set gs=128 - qt_bytes: fmt=4 reports correct memory including group scales Backward compatible: existing per-row int4 models work unchanged. The fused gate+up pair path (matmul_i4_pair) falls back to separate matmul_qt calls for fmt=4 — minor perf cost, correctness preserved. |
||
|
Vincenzo
|
a8895f2d84 |
Merge pull request #193 from bokiko/feature/disk-split
stats: opt-in disk-load split by speculation context and layer kind (DISK_SPLIT=1) |
||
|
JustVugg
|
c71e5d1cf8 |
glm: honest short-read reporting in expert pread; Makefile: -pthread for *BSD
Two small correctness fixes surfaced by community reports on non-Linux hosts: #236 — expert_load's buffered pread paths (slab + qs scales) used perror("pread expert") on a short read. Since pread returned a short count (not -1), errno stays 0 and perror prints "Success" — a confusing message right before exit(1) in the score/bench path. New pread_full() helper loops over short reads and EINTR and reports actual/expected bytes and offset, so a truncated shard reads as such instead of "Success". #219 — Linux pulls -pthread in via -fopenmp; the *BSDs do not, so pthread_* fail to link there. Added -pthread to the generic (Linux/*BSD x86-64) and aarch64 CFLAGS/LDFLAGS. No-op on Linux, required on FreeBSD (complements #206). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> |
||
|
JustVugg
|
ddf9212b81 |
glm: PIN_GB=all clamps to the --ram budget instead of pinning every expert (fixes #229)
PIN_GB=all passed gb=-1.0 to pin_load, which took npin=n (all experts) and ignored --ram entirely — the OOM-kill regression from #80. A 92 GB host with --ram 78 was killed mid-generation (anon-rss ~89 GB). Now gb<0 clamps npin to expert_avail() — how many experts fit the RAM budget, same accounting AUTOPIN uses. pin_load already adds the pinned bytes to resident_bytes, so the later cap_for_ram narrows the LRU accordingly with no double count. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> |
||
|
Vincenzo
|
22c91403b7 |
Merge pull request #228 from bokiko/feature/score-prefix-default
engine: SCORE mode defaults the [gMASK]<sop> prefix ON for GLM snapshots (#194 follow-up, #108) |
||
|
Michael Denyer
|
2bdf054622 |
glm: fix kv_alloc double-free on KV re-allocation (stale pre-Metal free block)
kv_alloc had two consecutive if(k->Lc) free blocks. The first freed every k->Lc[i]/k->Rc[i] and both arrays without unregistering from Metal and without nulling k->Lc; the second then re-tested the dangling pointer, called coli_metal_unregister on freed pointers, and freed everything a second time. Safe only when k->Lc is NULL (first call); any re-allocation on the same KVState aborts in the allocator. The first block is the pre-Metal version of the free path: |
||
|
Marquez
|
62419af188 |
routing: opt-in CACHE_ROUTE max-rank prefer (pin∪LRU), default off — 2.4->3.33 tok/s on GB10 (#199)
Paper-style cache-aware MoE selection (arXiv:2412.00099 max-rank): keep true top-J always; fill remaining K slots preferring experts already resident in pin∪LRU within top-M. Default OFF so stock full top-K is unchanged. Env: CACHE_ROUTE, ROUTE_J/M/P/ALPHA, ROUTE_AGREE (auto-on with CACHE_ROUTE). Telemetry: swap%/route_swaps/route_slots, route_agree, route_kl on footer and serve STAT. Complementary to PILOT (prefetch vs selection change). Routing-only PR for clean A/B vs PILOT / #119; no CUDA/fuse stack. See docs/CACHE_ROUTE.md. Closes nothing; for #161 discussion. Co-authored-by: Vincent Marquez <vincentmarquez405@gmail.com> |
||
|
yuri@FreeBSD
|
4a936c8af3 | FreeBSD compatibility: platform guards in glm.c/olmoe.c (#206) | ||
|
woolcoxm
|
6ca9565d2d |
pilot: two-step shared-expert router prediction, +3.1% recall, opt-in PILOT_TWO=1 (#204, #200)
Wires the two-step prediction (kind==2 from experiment/two-step-predict)
into the real pilot_prefetch() path behind PILOT_TWO=1 env var.
Changes:
- la_predict kind==2: computes shared expert (resident, no disk I/O) on
L's post_ln-normalized state, adds output to residual, then runs L+1's
router on the corrected state. Guards n_shared==0 and mloe_inter<=0.
- pilot_prefetch(): when PILOT_TWO=1, computes the same shared-expert
correction before running the router. Workspace allocated once per
call (not per position) to avoid malloc churn.
- LOOKA measurement harness expanded to 4 slots (prev, skip-attn,
PILOT stale, two-step) with updated reporting at both exit points.
- PILOT_TWO env var wired into main().
Measurements (GLM-5.2 744B, 24GB RAM, cap=2):
LOOKA recall: PILOT stale 73.6% -> two-step 76.7% (+3.1%)
End-to-end tok/s: no change (0.16 tok/s) — cache too small (cap=2)
for prediction quality to matter; disk bandwidth saturated regardless.
On higher-RAM hosts (cap>=32) the +3.1% recall would translate to
measurably fewer disk misses.
Prior art: 'Speculating Experts' (arXiv:2603.19289) independently
developed the same idea as a 'quasi-hidden state' using a static
default vector. Our approach uses the actual computed shared expert,
which is input-dependent and more accurate.
Co-authored-by: woolcoxm <13604288+woolcoxm@users.noreply.github.com>
|
||
|
ZacharyZcR
|
ec89136029 |
GPU resident pipeline: batch CUDA attention, head-sharded kv_b, prefill expert groups, W4A16 mixed dispatch (#111)
* 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> |
||
|
JustVugg
|
2ead86a27f |
serve mux: fix Windows request dispatch — BINARY-mode stdin/stdout + PeekNamedPipe-only polling (fixes #195)
Two Windows-only bugs in run_serve_mux left the gateway hanging after READY: 1. No _setmode(_O_BINARY): the CRT collapsed CRLF inside fread() payloads (waits forever for missing bytes) and expanded LF in the READY/STAT sentinels. 2. WaitForSingleObject on an anonymous pipe is undefined (always-signaled or WAIT_FAILED) and PeekNamedPipe fails on file/console handles, so the dispatch gate never opened. New rule: idle -> block in getline (POSIX select(NULL) semantics); active -> PeekNamedPipe poll. Reproduced and verified on real Windows via MinGW cross-compile + WSL interop: old binary writes READY (with CRLF corruption) then hangs forever on a crafted SUBMIT frame; fixed binary answers DONE + STAT and exits cleanly. Linux path untouched (0 warnings, oracle-exact). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> |
||
|
JustVugg
|
e12a4295cc |
glm.c: silence unused-variable 'c' in hwinfo_emit (reported in #148 on Darwin, present on Linux too)
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> |
||
|
Dennis Paul
|
4a0435cbb3 |
glm.c: remove the LMHEAD_EXACT probe accidentally merged with #152 — functional no-op strip (#197)
#152 carried leftover debug scaffolding from the lm_head measurement posted on that PR: a `g_lmhead_exact` global, an undocumented `LMHEAD_EXACT` getenv, and 7 lm_head call sites routed through matmul_qt_ex(..., !g_lmhead_exact). None of it was in the PR description; it rode in because `git checkout -B` carries uncommitted working-tree edits forward. It should not stay: - it does nothing worth having. Pinning lm_head off IDOT measures -0.03% perplexity across 5 corpora (in-distribution, isolated) — nothing. - LMHEAD_EXACT=1 would silently change lm_head numerics, undocumented and untested. - lm_head is fmt=1 (int8), so the IDOT gate applies to it unconditionally on every platform anyway; there is no threshold story here to expose. The flag defaulted to 0 and `!g_lmhead_exact == 1 == matmul_qt`'s own allow_idot, so this removal is a pure no-op. Verified rather than assumed — summed log-lik over 1023 tokens, in-distribution, CPU (deterministic), dev-with-probe vs dev-minus-probe: prose -2295.245383 == -2295.245383 markdown -3272.146403 == -3272.146403 Bit-identical. make check green (6 C suites + 60 python tests), no new warnings. The actual #152 changes are untouched: the three attention input projections and the DSA indexer's ix_wk stay batched and pinned to the exact int4 kernel via matmul_qt_ex(..., 0). |
||
|
bokiko
|
c4d14017df |
engine: default the [gMASK]<sop> prefix ON in SCORE mode for GLM snapshots (#194, #108)
SCORE mode scored raw token streams: GLM sees [gMASK]<sop> at the start of every training sequence, so unprefixed requests run the model out-of- distribution and silently distort logprobs (#108). The eval harness got the text-level fix in #194; contributors driving SCORE directly (e.g. the perplexity work in #153) were still exposed. Same detection rule as tools/eval_glm.py: config.json model_type contains "glm" (case-insensitive). The two ids are looked up in the snapshot's tokenizer.json ([gMASK], <sop>) rather than hardcoded. Requests that already carry the prefix pass through untouched — the patched eval_glm.py sends prefixed streams, so no double-prefixing. SCORE_PREFIX=0 restores stock behavior. One [SCORE] stderr notice when active. Validated on GLM-5.2 744B (M4 Pro, Metal): default run flips the '2 + 2 =' smoke question back to ' 4' (-1.006 vs ' 5' -2.950); SCORE_PREFIX=0 reproduces stock output; pre-prefixed requests give bit-identical logprobs to auto-prefixed ones. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> |
||
|
bokiko
|
b611304166 |
stats: DISK_SPLIT=1 opt-in disk-load split by speculation context and layer kind
Splits the expert DISK LOADS (LRU misses -> expert_load) two ways in the final stats line of 'run': - by context: loads issued while drafting (mtp_draft) vs absorbing verified tokens into the MTP KV (mtp_absorb) vs verify/main forwards - by layer kind: the MTP layer (int8 experts, ~2x bytes) vs the 78 main layers (int4), with exact bytes read (weights + scales) Gated behind DISK_SPLIT=1 (default OFF, parsed once at startup like the other g_* flags): when unset no atomic is ever touched and the stats output is byte-identical to stock. Measurement only - no effect on routing, verification or output. Measured on M4 Pro 48GB (GLM-5.2 int4, METAL, --ram 38): the draft path is ~1.5% of misses and the MTP layer ~4.5% of disk bytes. Requested in #182. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> |
||
|
ZacharyZcR
|
f1fa5bf3c8 |
placement: full-resident experts on large-memory hosts — CUDA_EXPERT_GB=auto, PIN_GB=all, adaptive GPU slots, RoPE cache (#80)
* 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 --------- Co-authored-by: JustVugg <JustVugg@users.noreply.github.com> |
||
|
Fabio Rovai
|
504fcdd930 |
SCHEMA=<file.json>: JSON-Schema -> GBNF compiler for grammar-forced drafts (#148)
* 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> |
||
|
LewfKrad
|
9b08cbc543 |
cuda-dll: make the Windows CUDA build work as shipped — MSVC host flags, CUDA_PATH defaults, POSIX shim in kernel test (#158, #157)
* 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> |
||
|
Dennis Paul
|
12350a494f |
prefill: batch the attention input projections — bit-identical, -4.5% prefill (#152)
* 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 @
|
||
|
Fabio Rovai
|
37ffb61674 |
routing: cross-layer coupling trace dump + pair-table builder + opt-in coupled prefetch source (#176)
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> |
||
|
ZacharyZcR
|
2878c60987 |
web dashboard: live expert-tier panel (VRAM/RAM/disk), one-process hosting, coli web (#126)
* 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. |
||
|
Tom Olorin
|
d47b095875 |
Windows: direct I/O (FILE_FLAG_NO_BUFFERING, 1.47x decode) + VirtualLock pin wiring + select_ctx device cache (#162)
* 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> |
||
|
woolcoxm
|
32629eae7f |
serve: WaitForSingleObject/PeekNamedPipe stdin polling for run_serve_mux on native Windows (#177, fixes #189/#139)
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> |
||
|
woolcoxm
|
d872dd5b4e |
glm.c: harden 5 crash paths — qk_rope buffer guard, OOM-checked mallocs, null derefs (#183)
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>
|
||
|
woolcoxm
|
0f5d5c8541 |
loaders: fix silent short-read (uninitialized memory) in config/oracle readers (#181)
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>
|
||
|
ZacharyZcR
|
9b96228775 |
engine: skip OMP hot-thread tuning on GPU builds (fixes #159 3x CUDA reg + #116 Metal -39%); add COLI_NO_FUSED_PAIR opt-out identifying #163 MTP cause
OMP active-spin (#77) contends with GPU dispatch: skip the re-exec tuning when COLI_CUDA/COLI_METAL set. CPU-only path unchanged (no-op). Also gates the S==1 gate+up fusion behind COLI_NO_FUSED_PAIR — the fusion shifts FP accumulation order and collapses MTP acceptance (#163); default unchanged, opt-out available. Zac-validated on 6x5090 (0.41->1.29 tok/s). |
||
|
woolcoxm
|
2319b942d2 |
Windows native port: serve-mode pipe fix + RAM detection + POSIX guards, AVX-VNNI kernel, gated CUDA DLL (#131, fixes #123)
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. |
||
|
Dennis Paul
|
d439ac8680 |
attention: size per-thread score buffer by the batch's true max nt, not Tk+1 — fixes serve heap-overflow (#117)
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. |
||
|
ZacharyZcR
|
4b1d0e3a57 |
expert dot: AVX-512 int4→float accumulator, runtime-switchable (I4_ACC512), +7% CPU-heavy decode, more accurate than AVX2 order (#95)
AVX-512F/BW build path only; non-AVX-512 builds bit-identical. Qualified on Xeon Silver 4510: max rel-err 2-6× lower than scalar oracle order, ppl 5.99 vs 5.98 (0.24%), runtime escape hatch I4_ACC512=0. Clears the #80/#94 numerical bar. |
||
|
ZacharyZcR
|
cbd599024e |
Unify continuous batching + heterogeneous runtime: decode batching, physical-core planning, disjoint VRAM/RAM placement, topp-policy warning (CPU-validated, CUDA on 6x5090) (#68)
* 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 |
||
|
Code Arranger
|
3716e4006a |
Metal backend (Apple Silicon): batched experts + fused attention on GPU, unified-memory zero-copy, gated behind COLI_METAL — 2.06 tok/s M5 Max (#72, #87, #103)
* 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> |
||
|
John Balis
|
342ceaf368 |
attention: heap per-thread score buffers — fix stack overflow past 8192 ctx on non-DSA snapshots (both absorb + dense paths, oracle-exact) (#110)
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> |
||
|
AutoJanitor
|
bc6bc9c250 |
VSX integer-dot kernels for POWER8 (12.8x int8 / 7.6x int4 over scalar, #ifdef-gated, x86 path untouched) (#98)
* 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> |
||
|
Rodolfo Hansen
|
9bba681ae2 |
perf(pilot): PILOT_REAL real cross-layer expert loads, independent from PIPE pool (data-driven: two pools beat layered/unified on both disk- and matmul-bound hosts) (#78)
The existing PILOT prefetcher predicts the next layer's routed experts
from the router logits and issues advisory readahead (posix_fadvise
WILLNEED / expert_prefetch) so the page cache is warm when the main
thread reaches that layer. That only warms the OS cache; the main thread
still pays the dequant/slab-build cost of expert_load on the critical
path.
PILOT_REAL=1 upgrades the prefetcher to perform the *actual* expert_load
into the next layer's LRU cache (ecache[L+1]) from the pilot I/O thread,
overlapped with the current layer's compute, so on a hit the main thread
finds the expert already resident and skips the load entirely.
Safety invariant (value-identical output vs OFF):
- MATMUL path: the pilot writes ONLY ecache[layer] for layer >
g_cur_moe_layer; the matmul in moe() reads ONLY ecache[layer]==
g_cur_moe_layer. A barrier at the top of moe() claims the current
layer and waits out any in-flight pilot load already targeting it,
after which the pilot drops every new load <= that layer. So the
matmul and the pilot never touch the same layer concurrently and no
half-loaded slot is ever matmul-ed.
- SCAN path (review fix — Option A): pilot_prefetch also runs on the
main thread and its residency scan reads ecache[lnext]/ecn[lnext] for
the FUTURE layer lnext = current+1 — exactly the layer the pilot
worker is writing. That scan previously ran off-lock (a real, benign
data race) and a stale comment falsely claimed the two threads never
touch the same layer's ecache. Fixed: the scan now takes g_pilot_mx
(the same lock the worker uses), decides under the lock, and enqueues
AFTER unlocking into the lock-free pilot_q ring (no re-entrant double
lock). The comment is rewritten to describe the real two-part
invariant honestly.
- The loading slot is published (eid set, ecn grown) only after the
pread completes; while loading it is hidden (eid=-1) from cache scans.
- The slow pread runs outside the handshake lock; the lock only guards
slot selection/publication. The shared LRU clock (m->eclock) is bumped
with an unconditional relaxed atomic add so the main thread and pilot
can't lose an increment; this is value-identical to the plain ++ with
only a negligible relaxed-atomic cost (no runtime branch for the OFF
case — not worth it).
Reliability (review fix — non-fatal speculative load):
expert_load() aborted the whole process (exit(1)) on any missing tensor,
OOM, short read or pread error. That is correct for the main / on-demand
/ REPIN / pin paths but fatal for a ~28%-mispredicted speculative pilot
load — a wrong guess could kill the server. expert_load now takes a
`fatal` flag: all main callers pass fatal=1 and keep today's exit-on-
error behavior byte-for-byte; the pilot passes fatal=0, so on error it
abandons the load, leaves the slot hidden (never published), bumps
g_pilot_drops, and logs a one-line stderr warning (never a silent
swallow). The main load path is behaviorally unchanged.
Residual gap closed (re-review): the fslab scale-buffer allocation still
went through falloc(), which exit(1)s on OOM regardless of `fatal`, so a
fatal=0 pilot load could still abort the process on a scale-buffer OOM.
expert_load now branches: fatal=1 keeps falloc() (byte-identical exit-on-
OOM for the main path), fatal=0 uses a checked malloc replicating falloc's
anti-wrap guard and, on failure, does the same non-fatal cleanup as the
slab-OOM branch (frees/NULLs s->slab and s->fslab, zeroes their caps,
returns -1) leaving a clean hidden slot (eid stays -1). The qt_from_disk
f32 fallback path is unquantized-only (GLM always has .qs) and is left
as-is, now with a comment noting it's unreachable for the pilot.
Tuning (review fix — PILOT_K default):
Real loads create LRU eviction pressure that hint-only WILLNEED does
not, so a large K thrashes at 28% mispredict. When PILOT_REAL is on and
the user did not set PILOT_K explicitly, K now defaults to 6 (best
measured this session) instead of 8; hint-only PILOT keeps the 8
default.
Default OFF (PILOT_REAL=0); PILOT_REAL=1 opts in and implies PILOT=1. A
per-run stats line reports real cross-layer loads completed vs dropped.
No Makefile change: pure C (stdatomic.h + sched.h).
Claude-Session: https://claude.ai/code/session_01DS7oc65c5RdA9V99otRCwt
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
|
||
|
Rodolfo Hansen
|
03d9a23fe4 |
perf(moe): overlap NVMe expert loads with matmul via opt-in PIPE I/O pool (default off, oracle-exact) (#79)
For streamed (non-resident) experts the MoE forward is disk-bound: each
64-expert block first blocks on a parallel pread of every miss, then runs
the matmuls serially. Those two phases don't overlap, so the compute
cores idle while the block's weights are still coming off NVMe.
PIPE=1 hands the misses' expert_load() to a small persistent pool of I/O
worker pthreads and lets the main thread start matmul immediately. The
main thread walks the block's experts in routing order and waits on a
per-slot ready flag only for the expert it needs right now, so loads of
later experts in the block hide behind matmul of earlier ones. All
matmul_qt stays on the main thread (it parallelises internally via
OpenMP and gates GPU dispatch on !omp_in_parallel()), so the I/O pool
never competes with the compute team for the matmul itself.
Cross-generation correctness: generation-tagged lock-free cursor
--------------------------------------------------------------------
The batch state (njobs/eids[]/layer/ready[]) is reused in place across
64-blocks, so a straggler or late-woken worker could touch the NEXT
batch's state. Two prior fixes (a drain barrier, then an _Atomic active
counter) each still left a cross-generation window. Both are now removed
and replaced with a single generation-tagged cursor:
_Atomic uint64_t cur = (gen << 8) | index; // gen main-only, index 0..njobs
- dispatch writes njobs/layer/eids[]/ready-reset with RELAXED stores,
then RELEASE-stores cur (bumping gen); that release publishes all
batch state to any worker whose ACQUIRE-load of cur sees the new gen.
- a worker grabs a job by CAS-advancing the index; it reads eids[i]/
layer only AFTER the winning CAS. The CAS comparand carries the
generation, so if a new batch was published the stale CAS fails and
the worker re-reads — it can never grab a wrong-generation job or
read torn state, no matter where it was preempted (wake gap,
post-cursor, anywhere).
- gen is bumped only by the main thread and is monotonic ⇒ no ABA.
- the per-expert pipe_wait(ready[q]) in the matmul loop (kept) makes
every grabbed job complete before the block ends, so no grab
outlives its generation. That is what makes the old `active`
counter and the end-of-block drain barrier unnecessary — both are
removed. ready[] is reset before the publishing release, so no stale
flag survives into the new generation.
- the mutex/condvar now exist ONLY to park and wake idle workers, not
for correctness.
This only reorders I/O, never the computation: greedy decode output is
byte-identical to the blocking path.
Default OFF (PIPE=0) so upstream behaviour is unchanged; PIPE=1 opts in
and PIPE_WORKERS=n sizes the pool (default 8). No Makefile change: pure C
(stdatomic.h + sched.h), builds with the default toolchain.
Claude-Session: https://claude.ai/code/session_01DS7oc65c5RdA9V99otRCwt
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
|
||
|
Rodolfo Hansen
|
704ed49c16 |
perf(omp): keep OpenMP worker team hot across tiny per-expert matmul regions (seed + re-exec once, fully overridable) (#77)
The MoE forward does many tiny, back-to-back per-expert matmul regions.
Under libgomp's default passive wait policy the worker team is parked
between regions, and the thread re-wake latency dominates the actual
compute. Switching the team to an active wait policy (with a bounded spin
count so long NVMe expert-load stalls still yield instead of burning a
core) keeps the threads hot across those regions.
Measured on the Zen5 build: expert-matmul wall time 66.9s -> 20.9s
(~3.2x on that phase). Output is numerically unchanged — this only
affects thread scheduling, not the computation.
Mechanism: libgomp parses OMP_/GOMP_ vars in a constructor that runs
before main(), so setenv() from main() and continuing is too late
(verified: the already-initialised runtime ignores it). Instead we seed
the winning defaults on first entry with overwrite=0 (so any value the
user already set wins), set a COLI_OMP_TUNED sentinel, and execv() self
once so a fresh libgomp constructor reads them. The sentinel guards the
exec, so we re-exec at most once. The block is the first statement in
main() so argv is passed verbatim to execv().
Discoverability / observability (review follow-up):
- COLI_NO_OMP_TUNE=1 is a documented kill-switch that disables the
whole re-exec + tuning path in one shot (distinct from the internal
COLI_OMP_TUNED re-exec sentinel);
- a one-line stderr breadcrumb is printed before the re-exec
("[OMP] hot-thread tuning: re-exec once (COLI_NO_OMP_TUNE=1 to skip)")
so the self-re-exec is self-documenting;
- execv only returns on failure, so a perror() now follows it
("execv self-reexec failed, running untuned") — a container without
/proc or a deleted inode falls back visibly instead of silently
losing the ~3.2x.
Fully opt-out / overridable and lossless:
- explicit OMP_WAIT_POLICY / GOMP_SPINCOUNT / OMP_PROC_BIND / OMP_DYNAMIC
in the environment win (overwrite=0);
- pre-setting COLI_OMP_TUNED=1 skips the re-exec entirely;
- COLI_NO_OMP_TUNE=1 disables the tuning path entirely;
- guarded by __linux__ (no-op re-exec elsewhere; the setenv defaults
still apply for any later-initialising runtime).
Claude-Session: https://claude.ai/code/session_01DS7oc65c5RdA9V99otRCwt
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
|
||
|
Sidd
|
2416bc9079 |
Translate user-facing runtime output to English, machine prefixes preserved, + CLI output test (#67, #85)
* feat: standardize runtime output in English * test: cover English CLI output |
||
|
JustVugg
|
6e7aa6f92e |
Guard against running the tiny oracle against a real model (#76)
The default ./glm mode validates against ref_glm.json, which is generated from glm_tiny (a random tiny model). Pointed at the 744B GLM-5.2 it feeds the model tiny-vocab prompt tokens and compares against tiny references — 0/20 and a degenerate loop, guaranteed on EVERY platform (x86/ARM/Metal), which reads as an engine bug (it isn't). Detect the mismatch (real vocab + tiny-range oracle ids) and print the correct commands instead. REF_FORCE=1 overrides. The engine self-test (SNAP=./glm_tiny TF=1) is unaffected. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> |