372f8fd633081acb1407412e7aab21c9b32441b2
89 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
 Tom Olorin
|
372f8fd633 |
run_tests: normpath test binaries — forward-slash relative paths fail CreateProcess on Windows (#196)
'make test-c' passes TEST_BINS as 'tests/test_json.exe' etc.; Python's subprocess on Windows hands that to CreateProcess, which rejects forward-slash relative paths for the executable (WinError 2), so the runner this script exists for (running tests from any Windows shell) failed on every test. os.path.normpath makes it 'tests\test_json.exe' on Windows and is a no-op elsewhere. Verified: all 8 suites run via 'make test-c' on Windows 11 / MinGW GCC 16.1 and paths are unchanged on POSIX. Co-authored-by: olorin <io@zyphyr.co> Co-authored-by: Claude Fable 5 <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> |
||
|
Louie
|
ec0aadf91a |
olmoe: int8 container support (fixes SIGBUS), NEON int8 dot kernels (2.7x on ARM), macOS RSS fix, converter fallback (#187)
- load_expert_w: detect I8/U8 tensors and read them raw with their .qs scales. Before this, pointing SNAP at a convert_olmoe.py container crashed with SIGBUS (st_read_f32 walks I8 data as 2-byte elements). Container misses now cost half the I/O and skip quantize_rows entirely: 2.08 -> 3.69 tok/s on the miss-heavy 16-slot ref.json run (M5). Raw bf16 checkpoints keep working. - matmul_q: Q8_0-style integer path on ARM (per-16 activation blocks, sdot on dotprod CPUs, vmull fallback). Same IDOT family as glm.c, same semantics: IDOT=0 stays byte-exact vs the oracle (12/12); default integer path can flip an argmax tie (documented in glm.c README). End-to-end on M5: 4.5 -> 12 tok/s warm-cache decode. - rss_gb: ru_maxrss is bytes on macOS, KB on Linux. RSS lines were reading '2029 GB' on Macs. - convert_olmoe.py: snapshot_download(local_files_only=True) raises LocalEntryNotFoundError when the repo is not cached; the download fallback was unreachable. Measured on M5 MacBook (10 cores, 24 GB, macOS 26.5), OLMoE-1B-7B-0125: ref.json greedy 12/12 with IDOT=0 on both container and raw paths. Co-authored-by: x <x@Mac.fritz.box> Co-authored-by: Claude Fable 5 <noreply@anthropic.com> |
||
|
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>
|
||
|
woolcoxm
|
35ed9bc921 |
config: fix UTF-8 decode crash + null config-value TypeError on Windows (doctor/resource_plan) (#184)
Two related crash bugs in the Python planning layer:
P1: Path.read_text() without encoding= defaults to locale.getpreferredencoding()
(cp1252 on most Windows installs). HuggingFace config.json is UTF-8 — if it
contains any non-ASCII byte (Chinese fields, accented chars, emoji in
metadata), read_text() raises UnicodeDecodeError. In doctor.py this is
caught and mis-reported as 'config.json is missing or invalid' (false
negative); in build_plan() called directly it is an uncaught crash.
Fixed: read_text(encoding='utf-8') in both resource_plan.py and doctor.py.
P2: int(cfg.get('kv_lora_rank', 0)) crashes with TypeError if the key is
present but null (JSON null). dict.get() returns the default only when
the key is ABSENT; a null value returns None, and int(None) raises
TypeError. The engine validates against malformed configs in C but the
Python planner did not.
Fixed: int(cfg.get(key) or 0) — coerces both missing and null to 0.
Applied to all 8 int(cfg.get(...)) calls in build_plan().
Co-authored-by: woolcoxm <13604288+woolcoxm@users.noreply.github.com>
|
||
|
woolcoxm
|
39dcebf39a |
openai_server: constant-time API key check (hmac.compare_digest) + resource_plan: csv-parse GPU names (fixes comma-in-name drop) (#186)
P6: discover_gpus() used line.split(',', 3) to parse nvidia-smi CSV output.
If the GPU name contains a comma (e.g. 'Tesla, Inc. V100'), split produces
more than 4 fields, the int() parse fails on the wrong field, and the GPU
is silently dropped — doctor reports 'no NVIDIA device detected' with no
clue why. Fixed by using the csv module (handles quoted fields correctly).
P7: require_auth() used plain string != comparison for the API key
('Authorization' header vs expected 'Bearer <key>'). This enables a
timing side-channel that could leak the key byte-by-byte. Low impact when
bound to localhost (default), but serve() only warns when host is
non-localhost without a key, and users do expose on 0.0.0.0. Fixed by
using hmac.compare_digest (constant-time comparison).
Co-authored-by: woolcoxm <13604288+woolcoxm@users.noreply.github.com>
|
||
|
Rodolfo Hansen
|
8395435322 |
convert: NVFP4 (modelopt e2m1) dequant path for convert_fp8_to_int4.py (#83)
convert_fp8_to_int4.py could ingest FP8-blockscale, bf16 and f32
checkpoints, but not NVFP4 — the format NVIDIA modelopt emits for
REAP-pruned GLM-5.2 (quant_algo=NVFP4). Those expert weights are U8
(two e2m1 nibbles/byte) with a per-16-block FP8 scale sidecar
(.weight_scale) plus a per-tensor F32 global (.weight_scale_2).
Adds dequant_nvfp4(): e2m1 LUT (verified 1:1 against
ml_dtypes.float4_e2m1fn), low-nibble=even / high-nibble=odd unpacking,
per-block scale repeat-interleaved over group_size=16, times the global.
classify() now consumes the .weight_scale/.weight_scale_2/.input_scale
sidecars, and dequant() routes U8 tensors with a scale sidecar to the
NVFP4 path (keys is now required, so a stray U8 can't fall through to it).
Existing FP8/bf16/f32 paths are untouched.
Guards against silent corruption rather than trusting the input:
* group_size is fixed at 16 and the block-scale column count is
asserted == ceil(I/16); the old code inferred it from the data
(I // ncols), which misaligns silently on a padded/swizzled scale
layout and hard-crashes on a non-multiple-of-16 I — after a multi-GB
shard download. A partial trailing block is handled by slicing to I.
* .weight_scale_2 is asserted < 1: modelopt stores the small global and
MULTIPLIES; llm-compressor/compressed-tensors stores the reciprocal
(>= 1) and DIVIDES. The two are dtype-identical, so without this a
compressed-tensors checkpoint would corrupt every tensor by ~gscale^2.
--selftest-nvfp4 (no network) asserts the 16-code LUT, an encode->dequant
round-trip to <1e-9, and a dequant->colibri-int4 requant bound (<0.30;
the inherent ~0.17 is informational, not a precision claim).
Claude-Session: https://claude.ai/code/session_01DS7oc65c5RdA9V99otRCwt
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
|
||
|
woolcoxm
|
789169f8f9 |
convert: fix converter crashes on Windows (fcntl import, null config) (#185)
Three crash bugs in convert_fp8_to_int4.py, all on the download path:
P3: 'import fcntl' at line 211 is Unix-only — ModuleNotFoundError on Windows.
The --indir test path returns before reaching it so tests pass, but
'--repo' on Windows hard-crashes. Guarded the import: try fcntl (Unix),
fall back to msvcrt.locking (Windows), skip if neither available.
P4: repo_info retry loop had range(999) — up to ~16 hours of retries on a
bad network, then fell through to line 395 where 'info' was unbound
(NameError). Capped at 10 retries and added an explicit error + return
when exhausted. Also added an early return if no safetensors shards are
found in the repo.
P5: if the shards list was empty (wrong repo, all filtered out), the
'for i, sh in enumerate(shards)' loop never executed and 'i' was unbound
at line 460 (NameError). Now caught by the early return from P4.
Co-authored-by: woolcoxm <13604288+woolcoxm@users.noreply.github.com>
|
||
|
Harvad Li
|
bc688a6211 |
coli: raise RLIMIT_NOFILE at startup — fixes 'Too many open files' (#147)
The engine opens/mmaps every safetensors shard of the model (144+ files for GLM-5.2). On macOS the default soft fd limit is 256, so a stock terminal session fails while loading with: <model>/out-00125.safetensors: Too many open files the engine exited while loading Raise the soft limit toward the hard limit (capped at 65536) in the coli launcher before spawning the engine, so users don't need a manual 'ulimit -n' in every shell. No-op on Windows and on shells whose limit is already sufficient. Co-authored-by: Harvad Lee <hongyanab@gmail.com> Co-authored-by: Claude Fable 5 <noreply@anthropic.com> |
||
|
Ronit Baldaniya
|
7216992c84 | test: cover native-Windows platform detection without uname (#145) | ||
|
bopof
|
f29cca396c |
Makefile: detect target via $(CC) -dumpmachine (#171)
rofl0r noted on #129 that uname describes the host shell, not the build target, and that a gcc/clang toolchain should be queried with -dumpmachine. Do that: derive the target triple from $(CC) -dumpmachine (e.g. x86_64-w64-mingw32 / x86_64-pc-cygwin / x86_64-unknown-linux-gnu / arm64-apple-darwin / powerpc64le-unknown-linux-gnu) and match mingw/cygwin/darwin/powerpc64 in it. The triple follows the toolchain rather than the shell, so detection is correct under a native-Windows shell (no uname on PATH) and when cross-compiling (make CC=x86_64-w64-mingw32-gcc on Linux), and it now also distinguishes cygwin from mingw. #129's OS=Windows_NT check and uname are kept as ordered fallbacks for the rare toolchain that does not answer -dumpmachine, so no host regresses. The CUDA/METAL macOS guards now use the derived DARWIN flag. Validated on native Windows (WinLibs GCC 16.1.0 x86_64-w64-mingw32, PowerShell, uname absent): make selects the Windows branch, glm.exe links -static (no libgcc/libwinpthread/libgomp DLL deps), and all 7 dependency-free C test binaries build and pass. Co-authored-by: bopof <285767350+bopof@users.noreply.github.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com> |
||
|
Dennis Paul
|
540781528e |
quant_ablation: condition GLM contexts on [gMASK]<sop> for in-distribution scoring (#155)
The tool tokenized every context with add_special_tokens=False, so pointing it at a GLM snapshot scored the model out-of-distribution — the same bug @bokiko found in eval_glm.py (#108). Measured on this box (GLM-5.2 int4, engine SCORE mode, same 1023 target tokens, only the conditioning prefix changed): corpus no prefix +[gMASK]<sop> natural prose PPL 29.2 PPL 9.4 markdown + code PPL 131.0 PPL 24.5 It does not just depress scores, it distorts sensitivity to numerical changes: an exact-vs-IDOT kernel A/B measured without the prefix reported a penalty that halved AND flipped sign on one of two corpora once the prefix was restored (#153). A quantization-ablation tool that scores OOD is therefore worse than useless — it produces confident deltas that are artifacts. The GLM tokenizer does not add the prefix itself (add_special_tokens=True is a no-op there), so it must be prepended explicitly. Auto-detected from the vocab rather than opt-in, so it cannot be lost by omission; --prefix overrides. Models with no such prefix are unaffected: OLMoE (the tool's default) has no BOS at all, add_special_tokens True/False give identical ids, so the numbers in #108 measured on OLMoE still stand. |
||
|
RDouglas
|
101b8b736e | openai_server: COLI_DEBUG levels (1=output, 2=both sides) (#149) | ||
|
RDouglas
|
bf3e890a53 | resource_plan: macOS memory_available() via vm_stat — fixes 0-on-macOS budget (#150) | ||
|
Sai Saharsh Chittimilla
|
6e9affba2a |
Makefile: add install/uninstall targets (PREFIX/bin + libexec) (#167, fixes #164)
* Add make install/uninstall targets Neither the root Makefile nor c/Makefile had an install target, so 'build+install' never actually installed anything (fixes #164). Adds standard PREFIX/DESTDIR/BINDIR-respecting install and uninstall targets that place glm and coli in $(BINDIR). * Split install: coli in bin/, engine + support files in libexec/ Addresses feedback on #164 from JustVugg and yurivict: coli goes to $(BINDIR), glm/olmoe and their Python support modules (resource_plan.py, doctor.py, openai_server.py, tools/) go to $(LIBEXECDIR) (default $(PREFIX)/libexec/colibri), matching typical Unix/FreeBSD-port conventions for a wrapper vs. its internals. coli now resolves the engine path in this order: 1. $COLI_ENGINE if set (explicit override) 2. glm next to itself (run-in-place from a source checkout, unchanged) 3. $(LIBEXECDIR)/glm (installed layout), also added to sys.path so the Python support modules still import correctly Also adds a 'bench' target (builds iobench) since only cuda-bench existed before. Tested locally (WSL2/Ubuntu): - run-in-place: cd c && python3 coli info -> engine ready - installed: make install PREFIX=$HOME/.local && ~/.local/bin/coli info -> engine ready (found via libexec) - make uninstall cleans both bin/ and libexec/colibri/ fully |
||
|
woolcoxm
|
c333840baa |
Makefile: portable clean/test-c via python helpers — works without sh.exe on native Windows (#179, fixes #172)
Problem: 'make clean', 'make test-c', and 'make check' use POSIX shell constructs (for loop, rm -f, rm -rf) that require sh.exe. On native Windows with WinLibs MinGW (no MSYS2, no Git Bash), there is no sh.exe on PATH. GNU Make falls back to cmd.exe, which can't parse 'for test in ...; do' or find 'rm', so these targets fail with 'test was unexpected' or 'CreateProcess error'. Root cause: the Makefile's recipe lines assumed a POSIX shell is always available. The IS_WIN detection (from #129) catches the platform but the shell-dependent targets were never made portable. Fix: replace the shell-dependent constructs with small Python helper scripts (Python is already a project dependency for test-python, convert, bench). This works from cmd.exe, PowerShell, Git Bash, and MSYS2 alike. Changes: - tools/run_tests.py (new): runs each C test binary, exits non-zero on the first failure. Replaces the 'for test in ...; do ./$test || exit 1; done' shell loop in test-c. - tools/clean.py (new): removes build artifacts and test binaries. Replaces 'rm -f' and 'rm -rf' in clean. Only removes executables (.exe) and known artifact names — never .c or .py source files. - Makefile: PYTHON defaults to 'python' on Windows (not 'python3'); test-c and clean now call the Python helpers instead of shell constructs. Verified from native PowerShell (no sh.exe): make clean removes 8-19 files/dirs, make test-c runs all 7 C test suites, source files survive. Also verified from Git Bash (sh.exe present): behavior unchanged. Co-authored-by: woolcoxm <13604288+woolcoxm@users.noreply.github.com> |
||
|
woolcoxm
|
a0e9c8dbe8 |
doctor: platform-aware engine.binary check — Windows has no X_OK bit (#178, fixes #141)
On Windows, os.access(path, os.X_OK) always returns True for any existing file (NTFS has no execute bit; executability is governed by file extension, not mode bits). So the test_non_executable_engine test — which chmods the engine to 0o644 and expects 'fail' — could never pass on Windows. Fixed doctor.py to use a platform-aware check: on Windows, any existing file is executable; on POSIX, honor the mode bits via os.access(X_OK). Fixed the test to assert the correct per-platform expectation: 'pass' on Windows (chmod is a no-op for executability), 'fail' on POSIX. Co-authored-by: woolcoxm <13604288+woolcoxm@users.noreply.github.com> |
||
|
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. |
||
|
bopof
|
afc259c599 |
Makefile: detect Windows from $(OS)=Windows_NT before uname so make works from native PowerShell/CMD (#129)
On Windows $(OS) is Windows_NT in every shell; check it first so native PowerShell/CMD (no uname on PATH) doesn't fall through to the Linux branch. Non-Windows unchanged (else branch still uses uname). Linux build verified green. |
||
|
ZacharyZcR
|
f8c0552c6d |
quant_ablation: add rotation preconditioning (-rot, QuaRot/QuIP#) and int3 schemes (#132, #81)
Extends the ablation grammar to int{2,3,4,8}[-g<N>][-rot][-nohead]. -rot round-trips weights through an orthogonal Hadamard Q=diag(±1)·H/√n on the input dim, measuring the exact weight error of a deployed rotate-activations scheme. Engine-free tool (c/tools/quant_ablation.py only). Verified: syntax clean, scheme parser correct (int3/-rot/-nohead), no unsafe constructs. Findings feed the v2 int3-g64 direction (#81/#108).
|
||
|
JustVugg
|
1f0e1b7076 |
olmoe: reject quant_bits outside 2..8 (fixes degenerate output, #134) + correct ref.json to OLMoE-0125-Instruct oracle (#133)
#134: olmoe.c stored experts as int8_t but silently accepted any bits argv; bits=16 (falsely documented as f32) truncated in quantize_rows -> wraparound garbage experts. Guard bits to 2..8 with a clear error (int8 is token-exact; f32 experts are not implemented here, unlike glm.c's fmt=0 path). Doc corrected. #133: shipped ref.json was from a different checkpoint (continued 'The capital of the United States is Washington'); the intended target is allenai/OLMoE-1B-7B-0125-Instruct (per convert_olmoe.py), whose greedy oracle continues 'The official language of France is French'. full_ids/text updated to the reporter's verified oracle (engine is already token-exact 12/12 vs it). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> |
||
|
ZacharyZcR
|
5f49a5adca |
coli: honor --gpu/--vram without --auto-tier, fail-fast on CPU-only binary (#125, fixes #121)
env_for() mapped --gpu/--vram only inside the --auto-tier branch, so the standalone form started a CPU-only engine with no warning (#121 — nearly published as a GPU benchmark). Now: --gpu none disables CUDA; --gpu list/auto sets COLI_CUDA/COLI_GPUS; --vram sets CUDA_EXPERT_GB; and --gpu/--vram on a CPU-only binary exit with a clear 'make glm CUDA=1' message instead of falling back silently. Tested on a CPU-only build: fail-fast fires for --gpu/--vram, --gpu none and default still start the CPU engine (no regression). |
||
|
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. |
||
|
Dennis Paul
|
15bf8cac6b |
openai server: type tool arguments from schema + implement tool_choice, 9 new tool-calling tests (#118)
Two OpenAI-compat tool-calling bugs found against the real GLM-5.2 (dnnspaul): (1) string-typed args coerced to numbers — declared schema type now decides, string kept verbatim, bool rejected as number, schema-less params keep permissive decoding; (2) tool_choice was ignored — none/auto/required/{function} now honored, invalid returns 400. Python-only (openai_server.py + tests), engine untouched. 36/36 tests pass (verified independently in a clean worktree).
|
||
|
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. |
||
|
Dennis Paul
|
97c756a064 |
tools: quant_ablation.py — engine-free A/B of any quantization scheme vs fp16 (fake-quant, isolates weight error) (#108, #115)
Measuring "what does int4 cost?" by comparing colibri's score to a published model-card number does not work: this harness scores 0-shot log-likelihood while published numbers are few-shot/CoT, and that protocol gap can swamp the quantization effect entirely (#108). This removes the confound by construction: take an fp16 model, push its weights through colibri's own quantizer (quantize -> dequantize, in place), and score both with the SAME harness on the SAME questions. The only variable is the quantizer, so the delta IS the quantization cost. Runs on OLMoE in minutes, so a scheme can be ranked BEFORE committing to a multi-hour GLM conversion. Quantizer math is replicated from tools/convert_fp8_to_int4.py (symmetric absmax, per-row scales) and generalised with an optional group size, so grouped/finer schemes can be compared directly against what ships today. Measured on OLMoE-1B-7B, n=200/task (#108): scheme hellaswag arc_c mmlu mean delta fp16 77.0% 47.0% 47.0% 57.0% -- int4 (shipped) 74.0% 41.0% 31.5% 48.8% -8.2pp int4-nohead 73.5% 40.5% 37.5% 50.5% -6.5pp int4-g128 78.5% 45.5% 38.0% 54.0% -3.0pp int4-g128-nohead 78.5% 46.5% 38.0% 54.3% -2.7pp The per-row int4 container costs ~8pp, concentrated on the HARD task: MMLU falls to 31.5% against a 25% random baseline while easy HellaSwag barely moves -- per-row scales eat the small logit margins that hard questions depend on (the same margin erosion that flips near-tie tokens in #100). group=128 recovers ~63% of the loss. Keeping lm_head/embed in fp16 is NOT the fix (+1.7pp alone, +0.3pp atop grouping). Includes a coverage assert: transformers fuses MoE experts into 3D tensors, so a ndim==2 filter silently skips every expert and leaves ~85% of the model in fp16 while appearing to work. The tool fails loudly instead of reporting fiction. Dev-only tool (torch + transformers); the engine's dependency-free path is untouched. |
||
|
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>
|
||
|
RDouglas
|
0f9f99cc29 |
OpenAI tool calling for GLM-5.2: gated tool declaration + parse_tool_calls, non-tool path preserved, opt-in salvage (#96)
* openai_server: parse GLM-5.2 tool calls into OpenAI tool_calls
Fills the 'tools not supported yet' stub. Upstream declared nothing and
rejected tools/functions; this makes them real end to end:
render_chat: emit a tools-declaration <|system|> block; replay assistant
tool_calls as <tool_call>NAME<arg_key>K</arg_key><arg_value>V
</arg_value></tool_call>; render tool-result messages as
<|observation|><tool_response>...</tool_response> (one
observation per consecutive tool run).
parse_tool_calls: turn the model's <tool_call> markers back into OpenAI
tool_calls; strip <think> from surfaced content.
generation: non-stream returns message.tool_calls + finish_reason
'tool_calls'; stream suppresses the markers from content
deltas (marker-length hold-back so a split <tool_call> is
caught) and emits tool_calls deltas after generation.
Default-safe: with no tools in the request, render and streaming take the
exact upstream path (byte-identical); tool handling is gated on tools present.
Marker format is authoritative (GLM-5.2 chat_template.jinja).
Optional de-mangler (COLI_TOOL_SALVAGE=1, default OFF) recovers malformed
calls from heavily-quantized models by mapping a lone payload onto the tool's
primary parameter; never rewrites well-formed output, never on by default.
A [api] tool-calls telemetry line reports strict vs de-mangled.
Pure stdlib; no new deps.
* openai_server: COLI_THINK global thinking default (opt-in)
COLI_THINK=1 makes thinking the default when a request sends neither
reasoning_effort nor enable_thinking (a launch-time global switch equivalent
to the old server's --think, useful because reasoningEffort propagation from
openai-compatible front-ends is unreliable). An explicit client value always
wins. Default off => exact OpenAI-standard behavior, so this is inert unless
enabled.
* openai_server: keepalive during long prefill (SSE reasoning pings)
engine.generate() blocks silently through the cold prefill (minutes for a
large prompt), and OpenCode/undici drop the socket after their idle timeout
-> the stream is 'terminated' before the first token. Upstream's SSE path
emits nothing until generation produces text, so it has no heartbeat.
A background pump emits a reasoning_content '.' delta whenever no event has
been written for KA_GAP (10s): the channel that reliably resets the client
timer and lands in the thinking panel, so answer content stays clean. All
wfile writes share a lock so the pump and event() never interleave; a
last-write timestamp gates the pump so it's silent while real tokens flow
(decode). The pump stops as soon as generation returns.
Inert for fast responses (nothing sent if events flow within 10s). No new
deps.
* openai_server: COLI_DEBUG echoes decoded tokens to stderr
Upstream sends decoded tokens only to the client (stdout->SSE), so the
console shows no generation text -- painful when the client has disconnected
or for local debugging. COLI_DEBUG=1 tees each decoded chunk to stderr at the
engine-callback boundary (both the plain and tool-call streaming paths).
Default off; no effect unless enabled.
* openai_server: use GLM-5.2 authoritative tool-declaration block
The tools were declared with a hand-written preamble; GLM-5.2 was trained on
the '# Tools' + <tools></tools> XML structure from chat_template.jinja. With
the wrong preamble the model doesn't recognize the tools as native and
hallucinates other frameworks' syntax (observed: repetitive 'end_action' and
key+value concatenated into one arg). Switch render_chat to the byte-exact
trained format so the model emits well-formed <tool_call> blocks.
|
||
|
AutoJanitor
|
6aaa8fc37a |
Makefile: Linux PowerPC (ppc64le) build support — -mcpu, scalar fallback, x86 path untouched (#97)
PowerPC GCC uses -mcpu instead of -march, so the Linux branch failed with unrecognized option -march=native on ppc64le. Detect ppc64le and ppc64 via uname -m and use -mcpu=$(ARCH) there. The x86-64 path is unchanged. Validated on an IBM POWER8 S824 (Ubuntu 20.04, gcc 9.4): make test-c passes, teacher forcing 32/32 positions and greedy 20/20 tokens against the transformers oracle, engine reports the scalar idot fallback. Signed-off-by: Scott <scottbphone12@gmail.com> Co-authored-by: Scott <scottbphone12@gmail.com> |
||
|
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
|
1bab3243f7 |
iobench: Windows correctness — aligned free, 64-bit totals, full-range random offsets (also improves Linux offset coverage) (#64)
* 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
* fix: iobench Windows correctness — aligned free, 64-bit totals, full-range offsets
Three bugs found running iobench.exe on Windows 11 (MinGW-w64):
1. Heap corruption (0xC0000374): free() on posix_memalign'd buffer.
On Windows posix_memalign maps to _aligned_malloc, whose memory must
be released with _aligned_free. Use compat_aligned_free (plain free
on POSIX, _aligned_free on Windows).
2. Negative GB totals: 'long tot' is 32-bit on Windows (LLP64), so any
benchmark reading more than 2 GB overflowed. Now int64_t.
3. Inflated speeds: RAND_MAX is 32767 on Windows, so rand()*4096 only
generated offsets in the first 134 MB of the file — which the page
cache holds entirely after one pass, reporting RAM speed instead of
disk speed (measured 6 GB/s fake vs 2.1 GB/s real on a laptop NVMe).
Combine two rand() calls into 30 bits so offsets span the whole file.
Portable: same code path on Linux/macOS, same seeded sequence shape.
|
||
|
ZacharyZcR
|
8f5f3e3a2b |
coli doctor: read-only setup/health diagnostics (path, config, shards, disk, RAM budget, placement) (#33)
* Add read-only coli doctor diagnostics * Fix doctor JSON output assertion |