bad64d1b06cbda80dbdfb1f8f502370d0407bcbb
56 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
 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). |
||
|
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> |
||
|
Fabio Rovai
|
cec7d6b648 |
Grammar-forced speculative drafts: GBNF grammar as a third draft source, guaranteed-accepted forced spans, lossless + opt-in (#48, #70)
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>
|
||
|
nalepy
|
89d95fc73b |
Windows 11 native port, phase 1: MinGW-w64 static build, compat shims, setup + docs (#40)
* fix: Windows port audit fixes — _FILE_OFFSET_BITS guard, O_BINARY st.h, getrusage peak, oracle diagnostic, setup.sh wmic
Audit remediation (all MEDIUM issues fixed):
- compat.h: compile-time #error if _FILE_OFFSET_BITS < 64 on _WIN32
- compat.h: COMPAT_O_RDONLY macro (O_RDONLY|O_BINARY on Windows, belt-and-braces)
- st.h: use COMPAT_O_RDONLY in both open() call sites (plan §1 row 1)
- compat.h: getrusage shim now uses PeakWorkingSetSize (ru_maxrss = peak, not current)
- glm.c: oracle mismatch diagnostic — prints position/expected/got on TF failures
- setup.sh: replace deprecated wmic with /proc/meminfo (MSYS2 provides it)
- .gitignore: *.exe, glm_tiny/, olmoe_hf/, olmoe_i4/
LOW issues addressed:
- _FILE_OFFSET_BITS guard prevents silent 32-bit off_t wrap at >4GB offsets
- coli Windows venv path (Scripts/python.exe) fixed earlier
- posix_fadvise do{}while(0) kept intentionally — no caller checks return code
Verification: oracle 32/32, 27/27 Python tests, rename EEXIST, >4GB pread offset.
* docs: add Windows 11 native port section to README
- Toolchain: MinGW-w64 (winlibs or MSYS2), GCC 16.1 tested
- Build instructions: make glm.exe, tiny oracle verification
- Runtime: SNAP=..., coli chat, coli serve all work
- Status: Phase 1 complete (compiles, correct, static-linked)
- Update platform requirements to include Windows 11 natively
|
||
|
0xhis
|
7d8d5f3109 |
Reduce allocation overhead in quantized matmul and MoE routing (thread-local scratch, trimmed temporaries) (#43)
* Reuse quantization scratch buffers * Trim MoE routing temporaries --------- Co-authored-by: 0xhis <0xhis@users.noreply.github.com> |
||
|
ZacharyZcR
|
9e6e1ac327 |
Isolated sequence KV contexts: KVState extraction, KV_SLOTS serve contexts with per-slot persistence, budget-aware pool accounting (#29)
* Add isolated sequence KV contexts * Log projected multi-slot KV footprint |
||
|
ZacharyZcR
|
1453dab7ae | REPIN follows live session heat (decaying heat map; .coli_usage stays the persistent signal) + disk→RAM→VRAM promotion (#26) | ||
|
ZacharyZcR
|
3e4d08b6bf |
OpenAI-compatible HTTP API: stdlib-only gateway over SERVE with KV prefix reuse across stateless requests (#21)
* Add OpenAI-compatible HTTP API * Support browser API clients * Handle missing KV cache during rewind |
||
|
JustVugg
|
ea5f6fb914 |
Harden against hostile model files: config dim validation at parse + overflow-safe falloc (PR #25 report, done right)
Model files come from untrusted mirrors. One choke point in load_cfg bounds every dimension (hostile config.json now dies with a clear message instead of reaching allocation sites); falloc guards the n*sizeof(float) wrap. No hot-path calloc churn — the report's patch was declined in review, the valid kernel of the idea is taken here. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> |
||
|
ZacharyZcR
|
13e8f09ffc | Organize project tools and local workflows: c/tools, c/scripts, c/tests, root Makefile (flat C core untouched) (#22) | ||
|
JustVugg
|
8f99e12b5e |
KV-to-disk: conversations reopen WARM across engine restarts (.coli_kv, KVSAVE=0 opts out)
serve mode persists the compressed MLA KV-cache incrementally after every turn (~182 KB/token appended, header count written last = crash-safe) and resumes it at startup: the model remembers the whole conversation and zero re-prefill happens. :reset and context-full restarts truncate the file. The MTP layer's KV row is not saved; kv_start=-1 re-arms its decode window. Validated: split-session answer byte-identical to an uninterrupted session (tiny oracle, TEMP=0), and on the real 744B model a restarted chat resumed 58 tokens in 0.0s and recalled a fact from the previous session while prefilling only the new question. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> |
||
|
JustVugg
|
077894210e |
LOOKA: routing predictability counters; PILOT: router-lookahead disk prefetch (async I/O thread); cap auto-raise to RAM budget (#12)
Measured on GLM-5.2 (48 tok, greedy): next-layer routing is 71.6% predictable one full layer ahead (79.4% skipping attention only; 41.3% previous-token). PILOT=1 issues next-layer expert WILLNEED from a dedicated I/O thread while the current layer computes — inline fadvise BLOCKS ~0.5ms/call on a saturated disk queue (+92s/48 tok, measured), hence the lock-free ring + worker. Neutral-in-noise on this dev box (disk already ~80% duty); expected to pay on balanced machines (#12: 43% disk / 46% matmul) — opt-in, default off. cap_for_ram now RAISES the LRU cap up to the RAM budget (ceiling n_experts, CAP_RAISE=0 opt-out): big-RAM machines were silently running with cap=8 (#12: 128GB box using 22GB of a 110GB budget; #13: 92GB box, same). DRAFT=3 on cold cache measured locally: 1399s vs 880s baseline for the same 48 tokens (acceptance 16%, experts/token 1809 vs 800) — confirms #8; DRAFT re-evaluation belongs to warm-cache serve sessions. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> |
||
|
ZacharyZcR
|
57706a0200 |
Tiered CUDA acceleration for routed experts (opt-in, CPU default untouched) + REPLAY fixture harness (#16)
* feat: add experimental CUDA backend for resident tensors * feat: promote pinned experts to a bounded VRAM tier * feat: preload the GPU expert tier at startup * fix: harden CUDA backend failure handling * feat: add deterministic multi-GPU tensor placement * test: add deterministic CUDA benchmark fixture * perf: make routed experts the default CUDA path |
||
|
RDouglas
|
ab6950876b | RFC: live re-pin of the hot-store between turns (REPIN=n, opt-in, default off) (#11) | ||
|
RDouglas
|
3daceaf5ec | IDOT: NEON single-token int4 gate (g_i4s/I4S) — SDOT makes S=1 decode pay on Apple Silicon (+14%, expert-matmul -16%) (#10) | ||
|
Dog279
|
8040dac645 |
IDOT: AVX-512 VNNI kernel paths + batch-size gate for single-row int4 (+20% decode on Xeon) (#9)
- dot_i8i8/dot_i4i8: guarded __AVX512VNNI__ branches using vpdpbusd (64 bytes/iter, s32 accumulation, no 16-bit intermediate). AVX-512 has no vpsignb, so the sign trick becomes abs(w) + mask-negate on x. - dot_i4i8: nibble unpack stays 256-bit; activation blocks are permuted to match the per-lane value order instead of re-sorting weights. - The int4 IDOT gate (was hardcoded S>=2, from AVX2 measurements) becomes g_i4s with a VNNI-aware default of 1: single-row decode goes through the integer path where it now pays. I4S=n overrides for A/B; I4S=2 restores the old behavior exactly. AVX2/NEON/scalar builds are unchanged. - Startup banner reports the active idot kernel (avx512-vnni/avx2/neon/scalar). Measured on Xeon 6 (24C, AVX512-VNNI), DDR5-5600, ~19.8 GB/s NVMe RAID0, gcc 13.3 -march=native: 256-tok greedy decode 1.72 -> 2.08 tok/s (+21%), expert-matmul -24%; second prompt +20%. Integer-path outputs bit-identical (associative s32 adds, same no-saturation bounds); S=1 int4 adopts the same int8-activation tradeoff IDOT already makes for S>=2. Co-authored-by: Claude Fable 5 <noreply@anthropic.com> |
||
|
RDouglas
|
cfd16b4c7e | PIN: mlock pinned experts into physical RAM (macOS: stop the compressor evicting the hot-store) (#7) | ||
|
JustVugg
|
193d2ce92d |
DSA lightning indexer: full implementation with auto-detection
- Faithful to the official modeling (transformers glm_moe_dsa): q from the q_a latent via wq_b (32 heads x 128), k = LayerNorm(wk(h)) shared across heads, interleaved RoPE on the first 64 dims, ReLU(q.k/sqrt(128)) weighted by weights_proj(h)/sqrt(32), causal top-2048 per query. - 'full' layers compute the selection (+ maintain the indexer k-cache from token 0); 'shared' layers reuse it (IndexShare, index_topk_freq=4). - Selection restricts both attention paths (absorbed decode + prefill reconstruction). MTP row stays dense. - Auto-detected like MTP: if out-idx-* weights are present for all full layers, DSA arms itself; DSA=0 disables; DSA_FORCE/DSA_TOPK for testing. - Validated on the tiny oracle (which ships indexer weights): selection machinery forced on with keep=all keys reproduces dense attention exactly (TF 32/32, gen 20/20); sparse smoke runs clean; kill switch verified. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> |
||
|
RDouglas
|
dd8c907800 | macOS/Apple Silicon port: NEON kernels + platform shim (PR #1 by RDouglasSharp) | ||
|
JustVugg
|
7a5ffd192d |
async I/O: WILLNEED readahead of the next expert block while computing the current one
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> |
||
|
JustVugg
|
757b0a8369 |
MLA weight absorption for decode (S<=4): no per-token k/v reconstruction — score via (W_K^T q)·latent, ctx via W_V(Σ a·latent). Validated exact: f32 TF 32/32, gen 20/20 with ABSORB=1 forced
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> |
||
|
JustVugg
|
741f49a1ca |
integer-dot kernels (Q8_0-style activations, AVX2 maddubs): int8 always, int4 when S>=2 — measured 1.4-2.5x
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> |
||
|
JustVugg
|
2dd6800aea |
sampling defaults for int4 reality: temp 0.7, nucleus 0.90 (official 1.0/0.95 tuned for full precision; the int4 tail is quantization noise)
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> |
||
|
JustVugg
|
1a8d2bcff3 |
prefill progress: engine reports layer N/78 on stderr, coli spinner shows it live
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> |