Commit Graph

29 Commits

Author SHA1 Message Date
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
Vincenzo 00902a2ea9 Merge pull request #224 from michael-denyer/fix-kv-alloc-double-free
glm: fix kv_alloc double-free on KV re-allocation (stale pre-Metal free block)
2026-07-15 07:47:16 +02:00
Michael Denyer 416570339f tests: remove accidentally committed Linux test binaries + gitignore them — fixes make test-c on non-Linux (#226)
c/tests/test_schema_gbnf and c/tests/test_compat_direct were committed
as Linux x86-64 ELF executables (slipped in via #111). On any other
platform make test-c considers them up to date and execs them, failing
with OSError: [Errno 8] Exec format error. Remove them and add the two
names to .gitignore alongside the other test binaries already listed,
so each platform rebuilds its own.
2026-07-15 07:40:42 +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
Kushida 8c18b7af68 openai_server: fix path traversal in web static serving (relative_to, not startswith) + empty-prompt validation, with tests (#212) 2026-07-14 22:07:41 +02:00
ZacharyZcR ec89136029 GPU resident pipeline: batch CUDA attention, head-sharded kv_b, prefill expert groups, W4A16 mixed dispatch (#111)
* Fuse CUDA expert MLP execution

* Group CUDA expert transfers by device

* Instrument grouped CUDA expert execution

* Bound grouped CUDA decode scratch

* Execute expert groups across GPUs in parallel

* Release host backing for multi-GPU experts

* Define quality-preserving memory policies

* Overlap cold expert loading with resident compute

* Adapt expert placement with session LFRU

* Fuse q4 expert gate and up dispatch

* Plan CPU work on physical cores

* Batch grouped expert CUDA kernels

* Separate VRAM and RAM expert placement

* Add ragged multi-sequence decode forward

* feat(runtime): add continuous decode scheduler

* Route concurrent API requests through batch scheduler

* Harden multiplex request lifecycle and framing

* Cancel disconnected multiplex requests

* Bind API port before starting the engine

* fix automatic KV slot allocation

* add native int4 Tensor Core grouped GEMM

* add Tensor Core throughput benchmark

* optimize packed int4 low-row kernels

* add asynchronous CUDA staging streams

* document validated six-GPU dense acceleration

* tune six-GPU expert hot set

* raise validated expert hot-set target

* add CUDA MLA absorption core

* fuse grouped expert gate and up projections

* Warn for explicit lossy routing flags

* Add full-resident expert placement mode

* Adapt VRAM expert slots to live routes

* Accelerate int4 matvec on AVX-512

* Reduce AVX-512 and RoPE decode overhead

* Seed every GPU expert layer after prefill

* Limit live GPU swaps during decode

* CUDA batch MLA attention, kv_b head-sharding, fused o_proj, expert-group dispatch, W4A16 kernels

Lab-qualified on the 6x RTX 5090 machine (914-token request benchmark):
- batch MLA absorption kernel (COLI_CUDA_ATTN=1): whole-batch attention on
  device, 154.8s -> 102.4s
- attention -> o_proj fusion on the layer device: -> 97.4s
- kv_b head-sharding across cards (COLI_CUDA_ATTN_SHARD=1), no weight
  duplication: -> 94.05s
- per-device expert-group dispatch with pinned-buffer async transfers,
  W4A16 tensor-core kernels for the shared expert, OMP hot-thread tuning

Negative results (reverted, kept out): GPU-side weighted scatter-add
(atomics + per-layer D2H lose 43.8%), shared-expert fused small-batch
kernel (-38.8%), W4A4 grouped tensor cores (int4 activations corrupt
output). Details in the lab research log.

* GPU resident pipeline: device-resident prefill attention chain, GPU expert groups in prefill, batched router, W4A16 mixed dispatch

COLI_CUDA_PIPE=1 keeps the prefill data plane on the layer home device;
control flow (routing, cache/pin management) stays on CPU. Any CUDA
failure falls back to the unchanged CPU path.

- Device primitives + unit tests (tests/test_pipe_cuda.cu): rmsnorm
  (strided), interleaved RoPE, silu-mul, residual add, fixed-order row
  merge (no atomics), device-input GEMM, persistent per-device scratch.
  All verified against the engine's CPU math on SM120 (worst 1.2e-5).
- attn_pipe_prefill: q_a -> norm -> q_b -> rope -> kv_a -> norm -> rope ->
  batch attention -> o_proj in one device chain (q_a/q_b/kv_a colocated
  with kv_b); only the final [S,D] and the new KV rows return to host.
  Attention 41.2s -> 30.8s on the 1571-token benchmark.
- Prefill batch-union now uses the GPU expert groups (previously gated to
  S<=64, leaving all VRAM-resident experts idle during prefill - measured
  21ms of GPU expert time in a 148s prefill). Expert phase 78.9s -> 69.0s.
- Router computed as one batched matmul instead of S sequential rows
  (bit-identical math).
- W4A16 tensor-core path for expert groups (COLI_CUDA_TC_W4A16=1) with
  row-count mixed dispatch: >=16 rows per expert use tensor cores, smaller
  batches keep the naive kernel (tensor cores measured negative below
  ~16 rows). Expert phase 69.0s -> 64.3s, decode unaffected.

Net on the 1571-token prefill benchmark: 148.8s -> 114.3-126.8s
(component timings stable across runs; wall drifts +-3-5s because
.coli_usage placement learning shifts the expert tiers between runs).
PROFILO now also prints the prefill-phase breakdown.

* Skip OMP hot-thread tuning when CUDA is enabled

The active-spin worker team measured 66.9s->20.9s on the CPU-only Zen5
build, but on the six-GPU full-residency workload the spinning workers
contend with the CUDA dispatch threads: ~4x slower prefill with the
process stuck near 1.8 cores. Gate the tuning on COLI_CUDA so each
configuration keeps the behavior it was measured to prefer.

* Inc.2a: sparse layers fully resident on the layer device, residual hops cards at layer boundaries

COLI_CUDA_PIPE=2 keeps the residual stream on the layer home device for
consecutive sparse layers (cudaMemcpyPeer at boundaries): in/post norms,
attention chain, both residual adds and the shared-expert MLP run on
device. Per layer only the post-norm activations (router + CPU-tier
experts + group gather), the new KV rows and, on DSA indexer layers, the
pre-attention norm leave the card. Per-layer transfers drop from ~130MB
to ~70MB. A device-side snapshot at layer entry makes any mid-layer CUDA
failure fall back to the unchanged CPU path idempotently.

1571-token prefill: 127.1s (PIPE=1 control) -> 117.6/118.9s, components
attention 30.8->26.1, other 31.8->22.5-24.5; output verified coherent
against the control.

* Head-sharded attention inside the pipe: negative on PCIe star topology, gated opt-in

Slicing q per card from the home device and collecting ctx back
serializes ~95MB/layer through the home card's PCIe link: attention
26.1s -> 41.4/44.4s on the 1571-token benchmark (two repeats), wall
117.6 -> 135-138s. The standalone host-path sharding won because six
cards uploaded from host RAM in parallel; a home-device star has no
such parallelism without NVLink. Kept behind COLI_CUDA_PIPE_SHARD=1
for interconnects where peer bandwidth does not share one root port.

* Inc.3: device-resident KV shadow for decode attention

Decode re-uploaded the whole latent+rope window per layer per token
(~300MB/token at 1571 context). Each layer now keeps a device shadow of
the compressed KV on its kv_b card, bulk-synced when behind and appended
incrementally; the host cache stays canonical. Invalidation on kv_bind
(slot switch), kv_alloc (resize) and on any overwrite of mirrored rows,
with the legacy full-upload path as fallback.

Measured (COLI_CUDA_PIPE gate): short-context decode 5.48 -> 5.59/5.87
tok/s, 1571-context decode 4.14 -> 4.22 tok/s. Decode remains CPU-expert
bound; the shadow removes the transfer tax, not the compute.

* tools: unified user-experience benchmark (bench_ux.sh)

Two fixed scenarios (short chat, long-document QA), TTFT + decode tok/s
+ first-line drift check, TEMP=0 DRAFT=0 enforced, medians over REPS
runs. Encodes the measurement discipline from the lab record: same
binary per comparison, judge medians because .coli_usage placement
learning drifts wall times between runs.

* tools: bench_ux.sh executable bit

* gitignore compiled test binaries

* tools: expert_atlas.py — measure per-expert topic affinity (#175)

Diffs .coli_usage across 10 themed probe batches (code/math/chinese/
prose/science/law/poetry/structured/translation/casual, 3 prompts each)
driven through a running API server — one engine load total. Every
touched expert gets a topic-affinity vector, entropy, and a specialist/
generalist label; output experts.json feeds the Brain page hover.

* serve: persist .coli_usage after every turn in mux mode, not only at exit

run_serve_mux saved the learning cache once at shutdown; a crash lost
the whole session's routing history, and live consumers of the file
(expert_atlas.py diffs it between probe batches) saw a frozen snapshot.
Now saved per turn like the interactive path (165KB write, negligible).

* web: Brain hover shows measured expert atlas when published

If /experts.json (from tools/expert_atlas.py, #175) is served next to
the app, the tooltip upgrades from the depth heuristic to measured
data: specialist/generalist label, entropy, and the top-3 topic
affinities. Row index maps to real layer (row+3, last row = MTP 78).
Falls back to the heuristic when no atlas is published.

---------

Co-authored-by: JustVugg <JustVugg@users.noreply.github.com>
2026-07-14 18:18:05 +02:00
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>
2026-07-14 16:35:39 +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
Ronit Baldaniya 7216992c84 test: cover native-Windows platform detection without uname (#145) 2026-07-14 14:11:32 +02:00
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>
2026-07-14 13:49:17 +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
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).
2026-07-13 16:11:31 +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
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>
2026-07-13 08:57:10 +02:00
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>
2026-07-12 22:44:56 +02:00
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
2026-07-12 13:38:40 +02:00
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>
2026-07-12 01:35:39 +02:00
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
2026-07-11 19:39:13 +02:00
nalepy 89d95fc73b Windows 11 native port, phase 1: MinGW-w64 static build, compat shims, setup + docs (#40)
* fix: Windows port audit fixes — _FILE_OFFSET_BITS guard, O_BINARY st.h, getrusage peak, oracle diagnostic, setup.sh wmic

Audit remediation (all MEDIUM issues fixed):
- compat.h: compile-time #error if _FILE_OFFSET_BITS < 64 on _WIN32
- compat.h: COMPAT_O_RDONLY macro (O_RDONLY|O_BINARY on Windows, belt-and-braces)
- st.h: use COMPAT_O_RDONLY in both open() call sites (plan §1 row 1)
- compat.h: getrusage shim now uses PeakWorkingSetSize (ru_maxrss = peak, not current)
- glm.c: oracle mismatch diagnostic — prints position/expected/got on TF failures
- setup.sh: replace deprecated wmic with /proc/meminfo (MSYS2 provides it)
- .gitignore: *.exe, glm_tiny/, olmoe_hf/, olmoe_i4/

LOW issues addressed:
- _FILE_OFFSET_BITS guard prevents silent 32-bit off_t wrap at >4GB offsets
- coli Windows venv path (Scripts/python.exe) fixed earlier
- posix_fadvise do{}while(0) kept intentionally — no caller checks return code

Verification: oracle 32/32, 27/27 Python tests, rename EEXIST, >4GB pread offset.

* docs: add Windows 11 native port section to README

- Toolchain: MinGW-w64 (winlibs or MSYS2), GCC 16.1 tested
- Build instructions: make glm.exe, tiny oracle verification
- Runtime: SNAP=..., coli chat, coli serve all work
- Status: Phase 1 complete (compiles, correct, static-linked)
- Update platform requirements to include Windows 11 natively
2026-07-11 12:59:49 +02:00
ZacharyZcR 9e6e1ac327 Isolated sequence KV contexts: KVState extraction, KV_SLOTS serve contexts with per-slot persistence, budget-aware pool accounting (#29)
* Add isolated sequence KV contexts

* Log projected multi-slot KV footprint
2026-07-10 16:19:43 +02:00
ZacharyZcR 76e858a80f Bounded FIFO inference scheduling: admission queue, OpenAI-shaped 429/503, queued-disconnect detection (#28) 2026-07-10 12:52:07 +02:00
ZacharyZcR 3cbc52ab00 coli plan: read-only Disk/RAM/VRAM placement planner (mirrors engine budget math, versioned --json, --auto-tier) (#27)
* Add tiered resource planner

* Apply resource plans to runner configuration

---------

Co-authored-by: JustVugg <JustVugg@users.noreply.github.com>
2026-07-10 12:51:23 +02:00
ZacharyZcR 1453dab7ae REPIN follows live session heat (decaying heat map; .coli_usage stays the persistent signal) + disk→RAM→VRAM promotion (#26) 2026-07-10 12:48:53 +02:00
ZacharyZcR 3e4d08b6bf OpenAI-compatible HTTP API: stdlib-only gateway over SERVE with KV prefix reuse across stateless requests (#21)
* Add OpenAI-compatible HTTP API

* Support browser API clients

* Handle missing KV cache during rewind
2026-07-10 11:04:56 +02:00
ZacharyZcR 13e8f09ffc Organize project tools and local workflows: c/tools, c/scripts, c/tests, root Makefile (flat C core untouched) (#22) 2026-07-10 10:08:39 +02:00
ZacharyZcR 8a2e4439ba Local project checks and contribution templates: make check, dependency-free C/python tests, issue forms + blank issues, CONTRIBUTING (#20)
* Add lightweight project checks and templates

* Ignore generated test binaries

* Keep project checks local
2026-07-10 08:56:41 +02:00