Commit Graph

62 Commits

Author SHA1 Message Date
JustVugg 87cd0aafed Merge dev into dual-ssd-mirror: resolve #298/#192-era conflicts
Combined resolution: mir_pread/st_prefetch_rep (this PR) now carry dev's
DISK-CLASS accounting unwind (dc_wall_exit) and O_DIRECT prefetch skip
(g_direct); direct-path keeps dc_direct=1 plus the mirror read counters.
Verified: clean build, token-exact tiny models unchanged, test_st_mirror
and test_st_pread pass.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
2026-07-20 21:24:39 +02:00
FABIOTESS 5e42e70ec3 int3-g64 (fmt=5): per-group-scale 3-bit weight format — engine, converter, tests
New weight format: int3 with ONE f32 scale per 64-input group (3.5 bits/weight
effective). Per group: 16B low plane (2 bits/val, int2 layout) + 8B high plane
(1 bit/val), values [-4,3] stored v+4. Same quantization math as
tools/quant_ablation.py _quant_last_dim(bits=3, group=64) from #132, whose
OLMoE ablation measured int3-g64 BEATING the shipped per-row int4 on quality
(-7.5 vs -9.3pp) at ~14% fewer bits.

Engine (placed per the #391 split): matmul_i3 + pack_int3_g64 + I3_* layout
helpers in quant.h next to their kernel family; fmt=5 branches in colibri.c's
qt_bytes/qt_alloc/qt_fill/matmul_qt/embed_row/qt_addrow/qt_matvec_rows.

Format detection now goes through #413's qt_resolve_fmt: fmt=5 registers its
distinct weight-byte layout O*ceil(I/64)*24 and its scale cardinality
O*ceil(I/64) there, validated against [O,I] like every other format. int3-g64
and grouped-int4-at-gs=64 carry the SAME scale count, so the weight bytes are
the int3 tag; row formats keep precedence for the small-I shapes where byte
counts coincide. The io_uring expert path still used the raw ?1:?2:3 byte
inference (it missed fmt=4 grouping entirely and never set gs) — converted to
qt_resolve_fmt like the other two expert paths.

Backends: qt_cuda_upload returns 0 for fmt=5 (tensor stays CPU-side, the
documented fallback), the dense CUDA matmul gate excludes fmt=5, and Metal's
existing fmt gates (gemm fmt<=3, moe fmt 1/2) already reject it.

Converter: quant_int3_g64 in convert_fp8_to_int4.py; --ebits 3/--xbits 3 now
emit it (previously bits=3 silently produced int4).

Tests: tests/test_int3.c (bit-exact pack/unpack vs reference, matmul_i3 vs
dequant-matmul incl. short tail groups and the real GLM I=7168, QT plumbing,
qt_resolve_fmt disambiguation incl. the same-scale-count fmt=4/fmt=5 pair,
outlier-rows RMS: int3-g64 3.3x lower error than per-row int4),
tests/test_int3_load.c (hand-rolled .safetensors fixture through st_init +
qt_from_disk: fmt=5 detected and loaded next to an int4 control tensor),
tests/test_int3_convert.py (NumPy pack round-trip vs independent decoder).

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
2026-07-20 17:31:05 +01:00
JustVugg 31526ae619 fix(build): test_pipe_block references glm.c, renamed to colibri.c by #391
The #270 rebase resolved the TEST_BINS list but missed the test_pipe_block
rule prerequisite and its #include, both still pointing at glm.c (which #391
renamed to colibri.c) — 'No rule to make target glm.c' broke the Linux C test
suite on dev. Point both at colibri.c. Build-verified locally.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
2026-07-20 18:15:39 +02:00
Vincenzo Fornaro 04ed0049a4 Merge pull request #444 from cdhdt/perf/idot-avxvnni-accumulators
idot: independent accumulators on the AVX-VNNI int8/int4 dot kernels (x86 parity with the NEON path)
2026-07-20 18:11:09 +02:00
JustVugg d34461c8d8 Merge remote-tracking branch 'origin/dev' into pr270
# Conflicts:
#	c/Makefile
2026-07-20 17:51:50 +02:00
JustVugg 48d4af691d Merge remote-tracking branch 'origin/dev' into pr330
# Conflicts:
#	c/Makefile
2026-07-20 17:41:10 +02:00
Vincenzo Fornaro 505a0f69aa Merge pull request #440 from okuvshynov/test_macpro_2019
fix: old x86 Mac scalar -> vnni
2026-07-20 17:20:31 +02:00
woolcoxm d43b54534f tests: efficiency suite — tiny-model regression gates + full-model optimization dossier
Two layers of efficiency coverage for the engine, both parsing the telemetry
glm.c already emits (REPLAY/PROFILE/[PROF]/CUDA-tier) but nothing previously
asserted on:

1. test_inefficiency.py — tiny-model asserted regression tests (8 tests, run
   in make test via test-python). Gate on: throughput floor, PROFILE phase
   accounting sanity, disk-wait not dominant on a resident model, CPU greedy
   determinism, and (when a CUDA build is present) CUDA init, dense VRAM
   upload, and CPU-vs-CUDA argmax agreement >= 70%. CUDA tests auto-skip with
   a clear build hint on CPU-only binaries.

2. test_efficiency_report.py — opt-in optimization dossier for a real model.
   Turns on every instrumentation flag (PROF, COLI_CUDA_PROFILE, CACHE_ROUTE,
   DISK_SPLIT, LOOKA) and prints 9 sections (provenance, throughput + tail
   latency, where-time-goes, attention breakdown, expert cache, disk I/O +
   phase split, routing quality + predictability, speculation, GPU tiers),
   each flagging inefficiency with the concrete knob to move tok/s. Never
   fails CI.

tools/efficiency.py is the shared harness: parse_run() captures every signal,
run_engine() wraps the subprocess. Reuses PROFILE_RE/SPEED_RE from
tools/benchmark_cuda_fixture.py and extends the tok/s regex to also catch the
run_text (parenthesized) format the full-model PROMPT path uses.

Makefile adds: efficiency / efficiency-cuda / efficiency-report targets.

Verified end-to-end on the full glm52_i4_g64 model (CPU + CUDA).
2026-07-20 10:49:07 -04:00
cdhdt 9c698b29a6 bench: microbench for AVX-VNNI idot accumulator A/B (not a gate) 2026-07-19 23:09:51 +02:00
Oleksandr Kuvshynov 5c84b25645 fix: old x86 Mac scalar -> vnni
Older macs still need -march, not -mcpu.

Before this fix, no ymm (or zmm, vnni, etc.).

Standalone repro:

```
% clang -O3  -mcpu=native glm.c -o /tmp/glm
clang: error: unsupported option '-mcpu=' for target 'x86_64-apple-darwin25.4.0'
% clang -O3  -mcpu=native glm.c -o /tmp/glm -lm
% otool -tv /tmp/glm | grep -c ymm
0
```

I'm not entirely sure why `-lm` would ignore `-mcpu=` rather than error, but either way we don't get vector instructions:

colibri itself:

```
% ulimit -l unlimited; OMP_NUM_THREADS=16 RAM_GB=700 PIN=auto PIN_GB=all PIN_FILL=1 ./coli run --ram 700 --model ~/projects/llms/glm5p2-i4 --ngen 256 "Explain the difference between concurrency and parallelism"
...
== GLM C engine (glm_moe_dsa), cache=8 experts/layer | experts@8-bit dense@8-bit | idot: scalar ==
```

After the fix, we'll get it:

```
% clang -O3  -march=native glm.c -o /tmp/glm -lm
% otool -tv /tmp/glm | grep -c ymm
2621
```

colibri run:
```
% ulimit -l unlimited; OMP_NUM_THREADS=16 RAM_GB=700 PIN=auto PIN_GB=all PIN_FILL=1 ./coli run --ram 700 --model ~/projects/llms/glm5p2-i4 --ngen 256 "Explain the difference between concurrency and parallelism"
...
== GLM C engine (glm_moe_dsa), cache=8 experts/layer | experts@8-bit dense@8-bit | idot: avx512-vnni ==
```
2026-07-19 16:40:09 -04:00
Steve Markgraf 63966ba3f8 Dual-SSD streaming: COLI_MODEL_MIRROR reads experts from two model copies
A second (read-only) copy of the model on another drive is registered as a
per-shard read replica; expert loads are split between the drives by a
deterministic (layer,eid) hash weighted by COLI_DISK_WEIGHTS=<primary>,<mirror>
or, by default, by a startup bandwidth probe using the engine's own access
pattern (parallel ~19 MB reads, O_DIRECT). Cold decode is disk-bound
(~11 GB/token), so two NVMe queues add up.

- st.h: mirror accepted per file only if size + safetensors header are
  byte-identical to the primary (identical data_offsets by construction, so
  every pread is valid on either copy); partial mirrors work (smaller second
  SSD holding only some shards); the mirror is never written — .coli_usage /
  .coli_kv stay on the primary.
- glm.c: routing covers the coalesced slab pread, O_DIRECT, mmap/Metal and
  scale reads, plus the OMP-parallel pin/autopin warmup (streams from both
  drives). Deterministic routing keeps readahead/PILOT WILLNEED on the same
  drive as the demand read and avoids caching an expert twice. A mirror read
  error falls back to the primary (one warning, no crash). Per-drive bytes
  are reported in a MIRROR: stats line.
- tests: test_st_mirror covers validation, read equality on both replicas,
  and the rejection paths (divergent header, size mismatch, missing file).

Measured on 2x NVMe (GLM-5.2 int4, greedy, DRAFT=0, DIRECT=1, cold-ish
cache): 0.42 -> 0.57 tok/s (+36%), expert-disk service 15.2s -> 10.3s,
byte-identical output; probe chose a 48/52 split.
2026-07-19 15:52:06 +02:00
ZacharyZcR 083fda5b0a fix: update test_logit_nan to include colibri.c instead of glm.c 2026-07-19 21:18:33 +08:00
ZacharyZcR f853ea8a0b refactor: split glm.c into colibri.c + 4 header modules
Rename glm.c → colibri.c and extract four self-contained modules
into header-only files (same pattern as st.h/tier.h/grammar.h):

  quant.h      (672 lines) — SIMD matmul kernels, quantization
  sample.h     (143 lines) — RNG, top-p sampling, stop-set
  kv_persist.h (121 lines) — .coli_kv disk persistence
  telemetry.h  (189 lines) — dashboard protocol, stats, usage

Main engine file shrinks from 6588 to 5396 lines (−18%).

Build system: primary target is now colibri$(EXE); `make glm`
remains as a phony alias for backward compat. CI, setup.sh,
coli CLI, and all 10 test files that include the engine are
updated. make check passes (C + Python, 73 tests, zero warnings).
2026-07-19 21:12:04 +08:00
JustVugg 4fcbea9056 Merge branch 'p370' into t370
# Conflicts:
#	c/Makefile
#	c/glm.c
2026-07-18 12:41:07 +02:00
JustVugg 4fb5b4e975 sampling: survive non-finite logits instead of emitting token 0 forever (#369)
A single NaN or +Inf logit silently broke the default sampling path. +Inf
became `mx`, then `expf((Inf-mx))`/`expf((NaN-mx))` is NaN, the softmax sum went
NaN, every probability went NaN — and dist_sample's fallback loop
`if(g_pbuf[i]>0)` is false for NaN at every index, so it returned 0. Every
subsequent token: 0. No error, no warning. @KingIcyCreamProjects found it.

dist_build now takes `mx` over finite logits only, gives a non-finite logit
probability 0, and when the distribution is unusable (no finite logit, or a
non-finite/zero sum) collapses to a delta on the finite argmax and warns ONCE
on stderr — degraded, but a valid token and a visible cause, never a silent
stream of zeros. The finite argmax uses the index found during the mx pass
(robust even when lo[0] itself is NaN, where argmax_v would wrongly return 0).

tests/test_sample_nan.c: healthy logits still sample correctly; NaN/+Inf
injected at lo[0], the middle, and the last position all pick the finite
argmax; an all-non-finite vocab leaves no NaN in the buffer and doesn't crash.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
2026-07-18 01:10:33 +02:00
KingIcyCreamProjects 7f70a8db5d sampling: guard against non-finite logits (was silent token-0 spew)
On the default serve path (TEMP>0, 0<NUCLEUS<1) a single NaN or +Inf in
the logits — a bad streamed expert tile, or an fp overflow in the matmul
at a low-RAM eviction boundary — poisoned softmax: g_pbuf became all-NaN,
dist_sample never satisfied cum>=u, and the fallback returned token 0. The
engine then emitted an unbroken run of token 0 with NO error. The greedy
path was equally blind: argmax_v started bv=lo[0] and `lo[i]>NaN` is always
false, so a NaN at index 0 pinned the argmax to 0.

- argmax_v: skip NaN (x==x) and seed from -inf, so it returns the max
  finite/+Inf entry instead of being NaN-pinned to 0. Covers greedy decode
  and the speculative-verify argmax path.
- dist_build: after the softmax sum, if s is non-finite or <=0, collapse
  g_pbuf to a one-hot over the finite argmax and warn once, instead of
  dividing every entry into NaN. Covers the nucleus and verify paths.

Both are O(1)/free on the happy path (one branch after the existing loop;
one extra comparison inside the existing argmax loop). Degrade + diagnose,
never silently corrupt.

test_logit_nan (wired into TEST_BINS): asserts argmax_v skips NaN/picks
+Inf, dist_build yields a finite normalized one-hot on the max finite
logit, dist_sample emits that token (not 0), and clean logits still give a
valid distribution. Fails on stock dev, passes with this change.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
2026-07-17 14:37:30 -05:00
JustVugg 8b36736e5d Merge branch 'p357' into trial357
# Conflicts:
#	c/Makefile
2026-07-17 20:47:21 +02:00
JustVugg 70a58799d6 Merge branch 'p354' into trialsamp
# Conflicts:
#	c/Makefile
2026-07-17 20:45:53 +02:00
woolcoxm 5d16368817 sampling: partial top-keep select in attention_rows DSA — O(nk) quickselect, not O(nk log nk) qsort (#356)
The DSA lightning indexer selects the top-index_topk (2048) context keys to
attend to by finding the threshold = keep-th largest attention score. It
previously full-qsorted all nk scores per layer per token (O(nk log nk)) just
to read one pivot value, then scanned the original array in position order to
build the kept set.

Replace the qsort with partial_select_desc (Hoare quickselect, median-of-three,
descending): O(nk) average to partition the keep largest into a[0..k), then the
threshold is min of that block. The two position-order scans (>thr then ==thr)
are UNCHANGED, so the kept-position set is bit-identical -- a stronger contract
than #335's sampling heap (which was multiset-only because the heap was unstable
and changed accumulation order). The quickselect pivot IS by definition the
keep-th largest, so the new threshold equals old tmp[keep-1] exactly.

Measured (bench_dsa_select, keep=2048, median of 2000 reps):
  nk=2049:   119us -> 5.7us   (21x)
  nk=8192:   626us -> 43us    (15x)
  nk=32768:  2.8ms -> 0.28ms  (10x)
  nk=65536:  6.6ms -> 0.47ms  (14x)
The gap widens with context (linear vs n-log-n). DSA only fires past index_topk,
so this is precisely the long-conversation regime where decode latency matters.

Adds test_dsa_select (in TEST_BINS): 129 cases asserting element-wise identical
kept-set vs an independent qsort reference across shapes (random, peaked,
sorted, reverse-sorted, tie-plateau, all-equal) and edges (keep==1, keep==nk).
Also directly checks the partition invariant.

Adds bench_dsa_select (on-demand, NOT a gate): reproduces the table above.
2026-07-17 09:54:40 -04:00
woolcoxm 37c96ee07a bench: add bench_topp -- head-to-head old-qsort vs new-heap timing on V=151936 (#335)
test_topp proves correctness; bench_topp measures the latency claim. It re-implements
the OLD dist_build (full-vocab qsort) inline on a private buffer and times it against
the REAL new dist_build over identical inputs in one process: V=151936, temp=0.7,
3 shapes (realistic / uniform / plateau) x 4 nuc values (0.5/0.9/0.95/0.99), 2000
timed reps each, median reported. Deliberately NOT in TEST_BINS -- it's a microbench,
not a gate. Build on demand: make tests/bench_topp && ./tests/bench_topp
2026-07-17 09:22:26 -04:00
woolcoxm ba00889fb9 sampling: partial top-p select in dist_build — O(V) heapify + k pops, not O(V log V) qsort (#335)
dist_build() sorted the entire 151936-entry vocab by probability (qsort) on every
sampled token whenever 0 < g_nuc < 1 — the serve default — and again per draft
position under rejection sampling. Measured cost: 5.6-8.0 ms/call; the actual work
is finding the few-hundred-token head whose cumulative mass reaches g_nuc.

Replace the full qsort + linear scan with a max-heap partial select:
  - Floyd heapify g_pidx over V by descending g_pbuf prob  (O(V), cache-friendly)
  - pop winners to the array's high end until cum >= g_nuc  (k * O(log V))
  - the remaining heap prefix IS the tail -> zero it, renormalize the head

Winners land in g_pidx[out..V-1] in descending order, so s2 accumulates in the
same order as before -> head is unchanged on tie-free shapes (ties were already
unspecified under the unstable qsort). All four dist_build/dist_sample contract
properties hold: g_pbuf stays id-indexed, g_pidx stays internal, the tail is
fully zeroed, the head renormalizes to 1.

No API change, no caller change, no new globals.

c/tests/test_topp.c (new): drives the REAL dist_build via the include-glm.c
pattern against an independent double-precision reimplementation of the OLD
algorithm. 123 cases: 6 sizes (1..1519) x 5 shapes (uniform/peaked/geometric/
plateau/sharptail) x 4 nuc values, plus the g_nuc>=1 guard-off paths and V=1.
Tie-free shapes compare head values within 1e-6 rel (float vs double renorm
noise); tie shapes compare value-multisets. No scratch files -> builds clean on
Windows MinGW without the unmerged mkdtemp shim (#352).
2026-07-17 09:01:23 -04:00
ebootheee dc196633f1 pipe: blocking pipe_wait (COLI_PIPE_BLOCK=1) + PIPE_WORKERS implies PIPE=1
pipe_wait's sched_yield spin storms the scheduler for the full 0.5-3ms of
each in-flight expert read; behind COLI_PIPE_BLOCK=1 it parks on a condvar
instead (~5us wake, no lost-wakeup: workers store ready with release
BEFORE taking the mutex to broadcast, and the waiter re-checks under the
lock). Default OFF = byte-identical spin. Pthread pool only: the URING
backend has no waiter spin to replace.

Setting PIPE_WORKERS>0 in the env without PIPE now implies PIPE=1 with a
stderr note: sizing the pool declares the intent to use it (a full
benchmark campaign ran with PIPE_WORKERS=16 and the pipe silently off).
The implication fires only when the platform default left the pipe off
(no-op on _WIN32 where PIPE defaults to 1), only on a positive value
(PIPE_WORKERS=0/empty does not enable a clamped 1-worker pipe), and an
explicit PIPE=0 always wins. The rule lives in pipe_workers_imply_pipe()
so the table is unit-testable.

tests/test_pipe_block.c (in TEST_BINS, all platforms): pins the
implication table, and drives the pool through 200 generations under each
waiter against an on-disk expert fixture — spin as control, condvar arm
alternating parked and fast-path waits — verifying identical slot bytes.

Measured on the spin side (2x5090 + Gen5 NVMe, GLM-5.2 744B int4, CPU
decode): 1.98 -> 2.16 tok/s at 192 tokens (expert-disk service 44.6s ->
33.5s). On Metal/GPU decode an M5 Pro A/B (PR thread) measured no change,
consistent with the mechanism: the win is freeing the core the spinner
was stealing from the CPU matmul team.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
2026-07-16 23:20:16 -06:00
Nicholas Beerbower 6507817d9f st: chunked pread with EINTR retry and honest short-read errors
Two latent bugs in every st.h reader, both hit in the field:

 - a single pread caps at ~2^31 bytes on Linux, so any tensor past
   2.1 GB (bf16 embed/unembed tensors of large models qualify) came
   back silently truncated with perror printing '... : Success'
   (errno untouched by a short read) — the same misleading-error
   symptom glm.c fixed for its own reads in #236;
 - no EINTR retry.

st_pread_full loops in ST_PREAD_CHUNK pieces (1 GB default, override
for tests), retries EINTR, and reports offset + progress on failure.
All five read sites converted; behavior on well-formed files is
byte-identical (GLM oracle re-verified on this branch: 32/32).

tests/test_st_pread builds with -DST_PREAD_CHUNK=7 so a 96-byte tensor
exercises the multi-chunk loop, and forks a child against a shard
truncated after st_init (init's static bounds check means the pread
path only fires when a file shrinks under a live handle) asserting
exit(1) with a 'short read' message and no 'Success'.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
2026-07-16 15:20:43 -04:00
Nicholas Beerbower 9e2d41567c tests: o200k tokenizer coverage, no model download
tests/tok_o200k_tiny.json is a synthetic byte-level BPE (274 vocab, a
few KB) whose Split regex is the o200k pattern; expected ids in
tok_o200k_cases.txt were generated by HF tokenizers on that same file.
test_tok_o200k (in TEST_BINS) scores 40/40 encode + 40/40 round-trip:
case-transition splits, contractions, digit groups, the [\r\n/]* tail,
whitespace branches, CJK/Greek/Cyrillic, added-token atomicity.

The cl100k path is untouched by construction — dispatch requires
\p{Lu} in the tokenizer's own Split pattern, which cl100k lacks — and
stays covered by the GLM oracle (verified on this branch: 32/32).

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
2026-07-16 15:17:21 -04:00
Vincenzo 2dca36a8a0 Merge pull request #314 from mohamedmastouri2000-boop/fix/windows-cuda-dll-build
Add build-config stamp: flag changes (e.g. CUDA_DLL=1) force a relink instead of a silent stale binary
2026-07-16 19:49:14 +02:00
mohamedmastouri2000-boop c70d94368e Merge branch 'main' into fix/windows-cuda-dll-build 2026-07-16 19:05:43 +03:00
JustVugg 6de32c55f6 Don't trust a converted checkpoint's config for the stop set
The engine armed its stop tokens from config.json's eos_token_id and nothing
else. That trusts metadata written by third-party conversion tooling, which is
a thing we already know goes wrong: the README documents a mirror shipping
int4 MTP heads that silently give 0% draft acceptance. GLM-5.2 declares THREE
eos ids (<|endoftext|>, <|user|>, <|observation|>); a converter that rewrites
config.json with a reduced list leaves the engine stopping on fewer tokens
than the model emits, and the missed ones get detokenized and printed into the
chat as literal text while generation runs past the end of the turn.

Two independent defenses:

  - eos_token_id is now unioned with generation_config.json, which is
    HuggingFace's authority for generation (config.json often carries a
    partial legacy copy). An extra stop is harmless; a missing one is not.

  - every added-token the TOKENIZER marks "special":true is armed as a stop,
    whatever the configs say. Those are control tokens (<|user|>, <|assistant|>,
    <sop>, [gMASK], the image/video/audio markers) and are never legitimate
    content in a reply -- GLM itself lists three of them as official eos.
    <think>/<tool_call>/<arg_key> are "special":false and are deliberately NOT
    swept up: they are real output. tok.h was parsing added_tokens but throwing
    the "special" flag away, so the distinction wasn't available to anyone.

On the real per-row checkpoint this takes the armed set from 3 to 18:
  [stop] 18 stop tokens: 154820 154827 154829 154821 ... (15 from the
  tokenizer's special set)

Honesty about scope: this is hygiene for a class of bug, NOT a fix for the
trailing-junk report on #298 that prompted it. I hypothesised @woolcoxm's g64
checkpoint had lost eos ids in conversion; he checked, and it hadn't -- his
config arms all three correctly. The emit path is also innocent: is_stop() is
checked BEFORE emit() at every one of the four call sites (4215, 4256, 4908,
4987), so a correctly-armed stop cannot be printed. His trailing junk is still
unexplained and is more likely quantization noise. What this commit buys is
that a checkpoint we don't control cannot leak control tokens into a reply,
which was true before and is not now.

tests/test_stops.c covers both defenses: the union, a missing
generation_config.json, BOTH configs mutilated (the tokenizer still stops all
five control tokens while leaving <think> alone), and T=NULL (the validation
path keeps config-only behaviour).

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
2026-07-16 17:59:34 +02:00
JustVugg 2a5961a01b tests: exactness oracle for the grouped-int4 kernel (fmt=4)
matmul_i4_grouped is the reference the CUDA fmt=4 port (#298) is expected to
reproduce, and it had no test of its own. @woolcoxm is currently debugging a
CUDA backend against an oracle nobody had verified, which is two moving
targets at once -- and he can't cross-check on CPU, since a 5-prompt run
takes 8 hours on the 744B model.

This checks matmul_i4_grouped against a plain-C reference that dequantizes
nibble -> (v-8)*scale[i/gs] and accumulates in double, over 11 shapes: I a
clean multiple of gs, a partial last group (the glen clamp), odd I (the
scalar nibble tail), gs > I, gs=16/64/128, S>1, and the nibble extremes
0x00/0xFF -- which decode to -8/+7 because the format is offset-encoded, not
two's complement. Reading that backwards turns 15 into -1 and looks like
data-dependent noise rather than a bug.

All 11 shapes match to ~1e-8 relative, so the CPU kernel is exact and can be
trusted as the reference.

One note on the tolerance, because the first draft of this test got it wrong
and "found" a bug that wasn't there: the error is compared against the sum of
|terms|, not against |result|. A dot product of signed terms can land near
zero through cancellation, and then a 1e-6 absolute error -- ordinary f32
accumulator precision -- reads as a 1e-3 relative one. A wrong scale index or
a wrong group boundary shifts the result by a fraction of the terms, so it is
still caught at 1e-6.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
2026-07-16 16:37:39 +02:00
Mohamed Mastouri 9cf97a4f0c Fix Windows cuda-dll build: quote NVCC, gate MSVC warnings, add config stamp (#306) 2026-07-16 15:27:08 +03:00
fermionoid 85c7857fe2 Makefile: add -std=gnu++17 to METALXX for Apple clang 16
Apple clang 16 (clang-1600.0.26.6) defaults objective-c++ to a pre-C++11
dialect, so the raw string literal holding the Metal shader in
backend_metal.mm fails to parse:

  backend_metal.mm:12:29: error: use of undeclared identifier 'R'
  backend_metal.mm:13:10: fatal error: 'metal_stdlib' file not found

Pinning gnu++17 on METALXX fixes 'make glm METAL=1' and 'make metal-test'.
Verified on macOS 15 / M4 Max: both targets build and all metal backend
tests pass.
2026-07-16 18:11:45 +08:00
RonitBStudent 53324fdcf0 fix(build): handle portable detection edge cases 2026-07-15 22:19:25 -05:00
RonitBStudent c82593ad87 fix(build): make portable checks target-aware
Select a portable architecture from the compiler target instead of forcing x86-64-v3 on every platform. On macOS, only enable Homebrew OpenMP when its header and library actually exist, preserving the dependency-free fallback.
2026-07-15 21:13:05 -05:00
JustVugg a2391587c7 Makefile(win): link -lpsapi so the MinGW build resolves GetProcessMemoryInfo
compat.h's rss_gb() calls GetProcessMemoryInfo and links psapi via
#pragma comment(lib,"psapi.lib") — an MSVC-ism. MinGW gcc ignores that pragma
(emits -Wunknown-pragmas), and the Windows LDFLAGS never linked psapi, so on a
gcc that doesn't honor the pragma (e.g. 16.1.0 UCRT) the build fails with
`undefined reference to GetProcessMemoryInfo`. Add -lpsapi to the Windows
LDFLAGS; harmless on toolchains where the pragma also resolves it. Found while
building on native Windows 11 with winlibs GCC 16.1.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
2026-07-15 22:52:17 +02:00
Vincenzo 704b4d7eeb Merge pull request #269 from michael-denyer/arm-i8mm-smmla
ARM i8mm: SMMLA fast path for the int8/int4 IDOT drivers (opt-in via ARCH=native)
2026-07-15 16:20:32 +02:00
Michael Denyer d0c9941a67 ARM i8mm: SMMLA fast path for the int8/int4 IDOT drivers
vmmlaq_s32 computes a 2x2 int32 tile (2 weight rows x 2 activation
rows) per instruction on 8-deep segments. Tile o and s in pairs,
halving weight traffic and doubling per-instruction work at S>=2.
Four independent accumulators over a 64-deep unroll keep the loop
throughput-bound (a single chained accumulator measures no better
than SDOT: latency-bound). S=1 and all tails (odd o, odd s, I not a
multiple of 16/32) keep the existing SDOT/scalar code, and scales
apply in the same order, so results are bit-identical.

Compile-time gated on __ARM_FEATURE_MATMUL_INT8. The default Darwin
build passes no -mcpu and is byte-identical (still SDOT, IDOT_KERNEL
"neon"). Opt in with ARCH=native (new Darwin Makefile knob, appends
-mcpu=<arch>), which reports IDOT_KERNEL "neon-i8mm". The same gate
lights up on any aarch64 with i8mm (Graviton3+, Grace).

test_idot grows a driver-level exactness check through matmul_qt_ex:
fmt 1 and 2, S in {2,3,4,5,8}, O in {1,2,3,64,65}, I in {16,17,100,
1408}, bitwise float equality against a plain-C reference. Green on
both build flavors.

Measured on an M5 Pro (18 threads, matmul_qt_ex microbenchmark at
GLM-5.2 expert shapes, best of 3 process runs, vs the SDOT baseline):
  gateup int4 S=8   499.7 -> 1076.6 GF/s  (+115%)
  gateup int8 S=8   512.9 -> 1166.3 GF/s  (+127%)
  down   int4 S=8   696.9 -> 1186.0 GF/s  (+70%)
  S=1 decode rows unchanged (SDOT path untouched)
2026-07-15 14:59:06 +01:00
Robert Landers 5d1eb142ab port uring implementation
Signed-off-by: Robert Landers <landers.robert@gmail.com>
2026-07-15 12:42:10 +02:00
JustVugg c71e5d1cf8 glm: honest short-read reporting in expert pread; Makefile: -pthread for *BSD
Two small correctness fixes surfaced by community reports on non-Linux hosts:

#236 — expert_load's buffered pread paths (slab + qs scales) used
perror("pread expert") on a short read. Since pread returned a short count
(not -1), errno stays 0 and perror prints "Success" — a confusing message
right before exit(1) in the score/bench path. New pread_full() helper loops
over short reads and EINTR and reports actual/expected bytes and offset, so a
truncated shard reads as such instead of "Success".

#219 — Linux pulls -pthread in via -fopenmp; the *BSDs do not, so pthread_*
fail to link there. Added -pthread to the generic (Linux/*BSD x86-64) and
aarch64 CFLAGS/LDFLAGS. No-op on Linux, required on FreeBSD (complements #206).

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
2026-07-15 07:55:09 +02:00
Michael Denyer 69004f73c2 tests: kv_alloc re-allocation regression test
kv_alloc guards every KVState free with if(k->Lc) precisely so it can be
called again on the same KVState (context resize, slot re-init). Exercise
that path: allocate, touch the cache, allocate again at a larger size.

Fails at this commit with a double-free (fixed in the next commit):
  malloc: *** error for object 0x7: pointer being freed was not allocated
2026-07-15 00:32:58 +01:00
bopof bfd49c41b9 Makefile: select $(PYTHON) from the host, not the target triple (#205)
clean/test-c/test-python run python on the build host, but $(PYTHON) was
chosen from $(IS_WIN), which is derived from the target triple
($(CC) -dumpmachine). A Linux->mingw cross build (make CC=x86_64-w64-mingw32-gcc
...) therefore sets IS_WIN and picks `python`, which fails on hosts where only
`python3` exists (e.g. Debian/Ubuntu) — breaking the cross-compile path #171
added.

Key PYTHON off the host instead: $(OS)=Windows_NT (the #129 signal) is set in
every Windows shell and empty on Linux/macOS. EXE stays driven by the target
triple, as it should. Addresses @rofl0r's host/target note on #171.

Co-authored-by: bopof <285767350+bopof@users.noreply.github.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
2026-07-14 22:07:56 +02:00
JustVugg bad64d1b06 Makefile: opt-in NVCC_CCBIN=<compiler> to override nvcc's host compiler (fixes #211)
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
2026-07-14 20:45:11 +02:00
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>
2026-07-14 16:31:09 +02:00
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>
2026-07-14 16:28:36 +02:00
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>
2026-07-14 16:12:15 +02:00
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>
2026-07-14 14:11:28 +02:00
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
2026-07-14 13:49:29 +02:00
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>
2026-07-14 13:49:26 +02:00
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.
2026-07-13 20:54:30 +02:00
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.
2026-07-13 20:47:23 +02:00
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.
2026-07-13 15:24:25 +02:00
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
2026-07-13 14:30:36 +02:00