Conversation
…fill Large dense GEMMs (NT=S*O elements) were dispatched as a single grid; past the device threadgroup-grid limit the tail output rows were silently never computed, leaving stale scratch -> nondeterministic long-context corruption under COLI_METAL=1 (identical prompt correct on CPU / at short lengths). Chunk the dispatch to <=2^25 elements via g_gx/g_gy buffer offsets: no kernel change, disjoint output ranges, negligible overhead, no-op for small S and decode. Gated COLI_GEMM_CHUNK (default on; =0 restores old path for A/B). Adds tests/test_gemm_largebatch.mm (make gemm-test): kv_b S=7478 nerr=1.0 before, clean after.
…grid NT=S*O, not S; S=4376 is clean)
…tion Metal: fix grid-size truncation in coli_metal_gemm at large prefill (long-context corruption)
…racle-validated New standalone engine (olmoe.c pattern) for the Inkling text stack: GQA with interleaved sliding-window (512, 16 KV heads) / global (8 KV heads) attention, learned relative-position bias banks (no RoPE), tau log-length scaling on global layers, four depthwise causal short convs per layer (K/V, post-attn, post-MLP; fp32, residual inside, cached conv state for decode), dense + MoE layers, sigmoid router with loss-free bias top-k and jointly normalized routed+shared weights, mup logit scaling with unpadded-vocab slice. Routed experts stream per-expert from the fused [E,2I,D]/[E,D,I] tensors with an LRU cache; -p "prompt" mode does streaming greedy generation with eos stop. tools/make_tiny_inkling.py builds a tiny random-init oracle via HF transformers (>=5.14): the engine reproduces it exactly — 36/36 teacher-forced argmax, 24/24 greedy tokens. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
convert_inkling_int4.py maps the TML-native checkpoint onto the HF names inkling.c loads, following transformers' conversion_mapping.py verbatim — including the critical detail that every fused w13 tensor stores gate/up rows INTERLEAVED (g0,u0,g1,...), so de-interleave = even rows then odd rows. Routed experts (95% of params) quantize to int4 + .qs per-row scales with packing bit-identical to convert_fp8_to_int4.py; norms/sconvs/rel-banks/router stay f32; everything else passes through as bf16. Resumable per shard (atomic rename), --watch converts while the download is still running, --xbits 0/4/8, audio/vision/MTP tensors skipped. inkling.c grew the matching container reader: U8 expert tensors are detected at load and streamed as packed int4/int8 rows + scales with no runtime requantization. --selftest-e2e fabricates a TML-named checkpoint from the tiny HF oracle snapshot (inverse mapping, re-fused and re-interleaved), converts it both ways, and the engine reproduces the oracle exactly in both passthrough and int4-container modes (36/36 tf, 24/24 greedy). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
f32 residents were fine for the tiny oracle but the real checkpoint needs ~172 GB that way — over sabre's 187 GB before the expert cache. Large matmul weights now stay in their on-disk dtype (Wt): raw bf16 for the real model (~86 GB resident), f32 for the oracle so validation stays bit-exact. Dispatch via matmul_w; embedding rows and lm_head included. Kernels for Zen 4: - matmul_h: bf16 dot via vdpbf16ps (activations rounded to bf16 per row, matching the HF bf16 reference numerics); shift-to-f32 scalar fallback on machines without AVX512-BF16 - matmul_q: Q8 per-32-block activation quantization + VNNI _mm256_dpbusd_epi32 int8 dot (maddubs fallback on plain AVX2), same IDOT family as glm.c; IDOT=0 keeps the byte-exact scalar path Oracle regressions all pass: f32 36/36+24/24, int4 container with VNNI and with IDOT=0 36/36+24/24, and a bf16-cast snapshot through the vdpbf16ps path also reproduces the oracle exactly. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Single-threaded numpy quantization paced the whole conversion at ~15 min per expert shard (~10+ hours for the checkpoint). Slice reads stay sequential for the HDD; the quantize work fans out on a thread pool (numpy ufuncs release the GIL on large arrays), in waves of 16 experts to bound peak RAM. Measured on the real checkpoint: 20 GB expert shard -> 5.5 GB int4 in 196 s, ~4.6x faster end to end. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The checkpoint index also references mtp.safetensors (the separate MTP-head file we skip entirely), so the watch loop never saw itself as finished. Filter the completion set to the shards we convert. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Linux caps a single pread at ~2.147 GB; the bf16 embed and lm_head tensors are 2.47 GB each, so load_w died on a short read. Loop in 1 GB chunks. First real-checkpoint smoke test passes with this: coherent greedy completion from the full 975B model, RSS 90.5 GB. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…, auto cap Stage-B of the expert path, all glm.c-playbook: - cache slots keep the container rows PACKED (int4 stays 4-bit in RAM: ~28 MB/expert instead of ~57 unpacked, double the experts per GB); new matmul_q4 unpacks nibbles in-register (AVX2 interleave + VNNI dot), numerically identical to the unpack+int8 path - moe() is now three passes: route everything, fill ALL missing experts in one parallel burst (prefill batches the whole sequence's misses -> NVMe finally sees queue depth), then compute - usage-history pinning: selection counts persist to SNAP/.coli_usage (glm convention, IKU1 header); next run pins the top PIN_N per layer (default cap/4) as non-evictable slots, filled in parallel at startup - cap 0 = auto-sized from MemAvailable after residents load; generate mode defaults to auto Oracle regressions all green: f32, int4 container (VNNI + IDOT=0), bf16 residents — 36/36 teacher-forced + 24/24 greedy each. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…ode) Gate+up now run as one matmul over the fused 2I rows (half the parallel regions per expert, one buffer), plus glm.c's OMP_WAIT_POLICY re-exec. Honest result: decode unchanged at 0.25 tok/s, prefill 21.7->21.1s. The measurement that matters: decode is RAM-bandwidth-bound on the ~35 GB of bf16 residents read per token (~1.2s/token floor on dual channel DDR5), not OMP-overhead-bound like glm's int4 residents were. Keeping the changes for the structure; the bandwidth wall is the GPU tier's job. All oracle modes still token-exact. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…t prefill New minimal backend (backend_cuda_ink.cu, CUDA=1 + existing flags): bf16 resident tensors upload to VRAM at load (~36 GB on the 975B; embed stays host for row lookup) and matmul_w dispatches to a one-warp-per-row bf16 kernel with the same numeric contract as the CPU vdpbf16ps path (both operands bf16-rounded, f32 accumulate). Freed RAM goes straight to the expert cache: cap auto 64->79/layer, hit 81.5->83.6%, RSS 116->82 GB. Two overheads found the hard way (bench4 regressed before these): - pinned host staging for activations — pageable cudaMemcpyAsync silently degrades to synchronous bounce copies - the OMP hot-thread re-exec is now #ifndef COLI_CUDA, the same exception glm.c makes: a spinning 24-thread team starves the CUDA driver during every stream sync 975B decode 0.25 -> 0.32 tok/s, prefill 21.1 -> 18.4 s. GPU path reproduces the tiny oracle exactly (36/36 + 24/24); CPU-only build unchanged and still oracle-exact in all modes. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
cudaDeviceScheduleSpin (+~1%: sync latency was not the bottleneck) and lightweight phase instrumentation in generate mode. First real breakdown on the 975B (24-token decode): fill 75.9s / expert-mm 10.4s / shared 0.9s / attn 1.2s — expert I/O is 85% of wall, and at 83.6% hit decode averages ~1 miss/layer, so fills run at queue depth 1: latency-bound, not bandwidth-bound. The GPU resident tier did its job; the remaining time belongs to cache hit rate and cross-layer prefetch. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…ame) 975B decode by pin depth, same auto cap 79/layer: 19 pinned = 0.32 tok/s (83.6% hit), 40 = 0.80 (95.6%), 64 = 2.51 tok/s at 100% hit on a usage-trained prompt — decode fills run at queue depth ~1, so every pinned expert removes a ~35ms stall. Phase profile inverted: expert I/O fell from 85% of wall to ~2%, CPU expert matmul is now 90%. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
One prompt per line (# comments skipped), model loaded once, KV/conv state reset between prompts, usage history accumulated across all of them and saved at the end. Purpose: build a routing ranking that generalizes — pins trained on one prompt hit 100% on it but 79.8% on novel routing (and crowd out LRU capacity), so the history needs diverse coverage before deep pinning pays off. kv_alloc now reuses its buffers across prompts instead of leaking per call. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
PIN=off (or PIN=0) disables warming entirely — no usage seeding, no startup pinning, cold LRU start. PIN_N=0 keeps the seeded ranking but pins nothing. USAGE_SAVE=0 skips the end-of-run history rewrite so benchmark loops don't skew the accumulated ranking; PIN=off implies it. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Quickstart, modes, cache-warming env reference, the full benchmark table (Stage A 0.06 -> trained-pins 2.51 tok/s, steady state 0.25 at 82.2% hit on a novel prompt with an 11-prompt warmup history), and the validation story. usage_save now reports whether it actually wrote — the message printed 'saved' even under USAGE_SAVE=0. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
A second engine that CI doesn't compile rots within weeks, so: - the Engine (Linux, CPU) job now `make colibri inkling` — inkling is built on every push/PR alongside the flagship; - the CUDA syntax-check job also `nvcc -c backend_cuda_ink.cu` (no-GPU syntax gate, same no-pipeline rule so it can actually fail); - a new `inkling-oracle` job builds a tiny random-init InklingForCausalLM via HF transformers (tools/make_tiny_inkling.py) and has the C engine reproduce it token-for-token (bits=0, bit-exact experts). The greedy pass exits non-zero on any mismatch, so the gate fails the moment a shared-code change (tok.h, Makefile) regresses the engine — the same guarantee glm_tiny gives colibri. oracle-requirements.txt pins CPU torch + transformers>=5.14 (the first release with Inkling support); setup-python hashes it for the pip cache. Left out of the dependency-free `make check` gate on purpose: that runs on Windows/macOS too, and this oracle needs torch/transformers. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_016KLDCYJSyHxd39ChTr4T27
Inkling (Thinking Machines 975B MoE): new engine, o200k tokenizer, bf16->int4 converter
This change follows up on c5b9d5a. (#504) It removes an unnecessary `https://www.google.com/search?q=` from links in `docker/README.md`. This fixes one more anchor (on line 395), outside the TOC anchors addressed in PR 504. And even for external links it seems better to just point directly at the relevant page rather than a Google result.
Stage-A inkling.c only had greedy -p / oracle modes, so the shared openai_server.py gateway had nothing to drive. Add the serve loop it speaks, byte-identical to colibri.c's protocol: - SERVE=1 -> serve_loop: READY sentinel, STAT handshake, then per request SUBMIT <id> <slot> <len> <max_tok> <temp> <top_p> -> DATA frames -> DONE <id> STAT ... + PROF (per-turn phase timings). CANCEL honored between tokens; requests run one at a time (v1). - temperature + top-p nucleus sampling (sample_logits; temp<=0 = greedy, so the oracle path is unchanged) + a light REP_PEN repeat guard. - HWINFO / TIERS / EMAP dashboard lines (same as colibri.c) so the web Brain/Profiling pages render live expert-tier state. - prompt_reject bounds the served context (CTX_MAX, default 8192). Additive: the oracle still matches 36/36 teacher-forced + 24/24 greedy, 0 warnings. POSIX serve loop (inkling targets Linux/CUDA); the select() include is guarded like colibri.c's. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_016KLDCYJSyHxd39ChTr4T27
…nel)
With inkling.c now speaking the serve protocol, teach the gateway to
drive it. Additive — the glm path is byte-unchanged (ARCH defaults to
"glm"):
- --arch {auto,glm,inkling}; auto reads model_type from the model's
config.json. Picks the chat-template renderer (render_chat_inkling)
and the model-id.
- inkling's thinking is reasoning, not answer: split the <|content_thinking|>
/ <|content_text|> channels (split_inkling / strip_inkling_markers) into
reasoning_content on both the streaming (InklingStreamSplit) and
non-streaming chat paths.
- response_format grammars are refused for inkling (its serve loop speaks
the 6-field SUBMIT header only) with a clear 400.
51 openai_server tests pass; end-to-end serve smoke against a tiny model
returns DATA/DONE/PROF and coherent reasoning/text split.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_016KLDCYJSyHxd39ChTr4T27
- Update binary path from glm to colibri (binary was renamed on dev) - Rename _HAS_GLM → _HAS_BINARY, _run_glm → _run for clarity - Update all docstrings and comments referencing glm.c → colibri.c - All 3 CPU-only tests pass, 6 GPU tests skip cleanly - Fixes test_cuda_zero_with_gpu_env_fails_guard: already had correct assertions for both CUDA and CPU-only builds
The fast path in matmul_q was ARM NEON only — every x86 build silently fell back to a scalar, byte-at-a-time loop even on CPUs with AVX2 available. Adds an AVX2 counterpart: sign-extend both int8 operands to int16 (exact, no precision loss), multiply+horizontal-pair-sum via _mm256_madd_epi16, then reduce to a scalar. Pure integer arithmetic, so it's bit-for-bit identical to the scalar reference. Correctness (verified): IDOT=0 (scalar) and IDOT=1 (AVX2) produce the exact same 100/100 tokens on OLMoE-1B-7B-0125-Instruct, same prompt. SNAP=<olmoe_merged> IDOT=0 ./olmoe 64 8 ref.json SNAP=<olmoe_merged> IDOT=1 ./olmoe 64 8 ref.json Performance: measured on an 8-core AMD Zen (Lucienne) CPU-only box, same 100-token prompt, cache=64 (all experts resident). This machine is NOT an isolated benchmark box — it was running a full kube-apiserver/etcd/kubelet stack, clickhouse-server, and several Chrome renderers throughout (load average ~4.0 on 8 cores), and that background load dominates the signal: IDOT=0 (scalar): 1.90 / 2.37 / 2.73 / 2.84 tok/s (median 2.55, n=4) IDOT=1 (AVX2): 2.54 / 2.64 / 3.90 tok/s (median 2.64, n=3) ~4% difference at the median — within the run-to-run noise on this machine, not a confirmed speedup. Not claiming a performance win here; this patch's value is x86/ARM parity and verified correctness. Re-measure on a quieter or dedicated box for a real throughput number.
The old --repo path called snapshot_download (pulls the whole checkpoint to local disk first) then loaded every shard fully into RAM via load_file before processing -- for OLMoE-1B-7B's ~14GB bf16 checkpoint, that's a real problem on a RAM-constrained machine (this was hit directly: 21GB total RAM, most already in use by other processes). Rewrites both paths to be lazy: - --model (local dir): opens shards with safetensors.safe_open and reads tensors on demand instead of safetensors.torch.load_file's eager full-shard load. - --repo: downloads and processes ONE source shard at a time via hf_hub_download, deleting it immediately after extraction -- peak extra disk usage is one shard, not the full checkpoint. An expert's three projections can land in different shards; incomplete triples are held in a small in-RAM dict until their last projection arrives, then merged, quantized, and flushed. - Resumable: existing output shards are scanned (header-only, via safe_open) and already-converted tensors are skipped on a rerun, so an interrupted conversion picks up where it left off instead of restarting. Same output format and CLI as before (--repo/--model/--out), no behavior change for anyone already using it successfully -- this only changes how much RAM/disk it needs while doing so. Verified end-to-end: converted allenai/OLMoE-1B-7B-0125-Instruct on a 21GB-RAM machine (previously not attempted with the old eager-load version), all 1024 experts (16 layers x 64) present and complete, olmoe.c loads and runs correctly against the result.
…x TTFT 50.1s -> 1.7s
…memory to chunk size
…nch and no up[] round-trip per group call
cuda: keep async packed-int4 expert groups token exact
docs: tell ARM64 users the prebuilt Linux archive will not run
…uild docs: CUDA-backend build recipe for glibc >= 2.41 (trixie/F40+)
doctor: report shared libraries the engine cannot load
The dashboard defaulted maxTokens to 512 and capped the input at 4096, so a GLM-5.2 reasoning trace on any non-trivial prompt was cut mid-trace with no indication to the user (reported in #634). - default max output tokens 512 -> 4096 - raise the input ceiling 4096 -> 32768 so long reasoning is reachable - surface a 'Truncated' badge (with a tooltip pointing at the setting) when the stream finishes with finish_reason=length, so a cutoff is never silent Gateway/engine paths unchanged; this is dashboard-only. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
With --kv-slots > 1 and no client-supplied cache_slot, the scheduler admitted each request to min(free_slots), blind to which slot already held that conversation's KV cache. Under any interleaving of clients a turn could land on another conversation's slot and pay a full re-prefill (reported as Defect 1 in #634). Derive a stable slot from a key that does not change across a conversation's turns -- the leading system messages plus the first user message -- so every turn of one conversation routes to the same slot and reuses its warm prefix. Distinct conversations spread across slots; with more live conversations than slots, colliding ones degrade to the old re-prefill rather than to anything worse. Raw /v1/completions (no messages) and any client that pins cache_slot are unchanged. Gateway-only; the engine's per-slot KV reuse is untouched. Unit-tested in test_openai_server.py (stable across turns, in range, deterministic, spreads). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
fix(web): stop truncating reasoning at 512 tokens by default
fix(serve): pin each conversation to a stable KV slot (#634 Defect 1)
The README leads with a dashboard screenshot and the site's whole argument is the measured Expert Atlas, but no release archive has ever contained either. `release.yml` packages the engine and four .py files; web/dist is a build artifact and is not tracked, so a release user has no web/ to build from -- and `coli web` told them to "cd web && npm install" in a directory the archive does not have. #578 shipped web/public/experts.json, which is correct and does land in web/dist at build time. It just never reached anyone who did not build the front end themselves. Packaging alone would not have been enough: WEB_DIST was a single path, __file__/../../web/dist, which only resolves in a source checkout, where this file sits in c/. In an archive or an installed tree the server sits NEXT TO web/dist, one level nearer. Both openai_server.py and coli now probe for index.html across the two layouts instead of assuming one. The dashboard is built in CI rather than committed, so dist/ stays an artifact. On Windows the step forces shell: bash -- the job default is msys2, whose minimal PATH does not see the node that setup-node installs. `make install` installs web/dist next to openai_server.py when it exists, and says how to build it when it does not. The archive verification gains the dashboard: not just that the files are present, but that the packaged server actually RESOLVES them -- an archive with the right files under an unresolvable path is precisely the failure that step already exists to catch. Verified by replaying all three layouts locally: source checkout, an unpacked release tree, and make install DESTDIR -- each resolves to its own web/dist with index.html and experts.json present, and an unbuilt tree still falls through to the "web UI not built" hint. make check 211. Refs #578, #175
release: ship the web dashboard, so the Expert Atlas actually reaches users (#578)
The README still sent newcomers to the older per-row int4 mirror (mateogrgic/GLM-5.2-colibri-int4-with-int8-mtp), which measures ~9pp worse on quality and causes the think-mode loops / never-terminating generations in #455. Point to the group-scaled (gs64) container that docs/quickstart.md already recommends, and explain why gs64 over per-row. Updates the English README plus the zh-CN / zh-TW / it variants and both docker READMEs. The historical experiment log keeps its original link (it records what was actually run). New: https://huggingface.co/mastouri/GLM-5.2-colibri-int4-g64-with-int8-mtp Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
docs: point the README to the gs64 model container (#642)
Replace the project website with a cleaner, editorial landing: warm cream + tan cards + near-black blocks, bold-sans headings with a serif for supporting copy, left-aligned hero, model cards with metadata rows, a hardware ladder, and a monochrome rotating expert-atlas galaxy drawn on canvas. Pixel wordmark in Bytesized to match the pixel hummingbird mark. Self-contained (fonts from Google Fonts, no images/assets), ~28 KB. Deployed to Pages by site.yml on main. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
site: new Anthropic-style landing page
…rch (3x)
iq3_pack.encode picks a 4-bit sub-scale by trying all 16 candidates and keeping
the best, and each candidate ran its own nearest-grid search — 16 [R*8,4]x[4,256]
products per 32-weight sub-block. But the candidates differ only in the scalar
db, and
argmin_g ||m/db - g||^2 = argmin_g [ db*||g||^2 - 2*(m.g) ] (db > 0)
so m.g does not depend on the code at all. Computing it once serves all 16. The
squared error falls out of the same product,
sum (db*g - m)^2 = db^2*||g||^2 - 2*db*(m.g) + ||m||^2
whose ||m||^2 term is identical across codes and drops out of the comparison —
which also removes the per-code gather of g[idx] and the reconstruction buffer.
Rows are then encoded in blocks (IQ3_ROW_CHUNK, default 128): the search keeps a
[rows*8,256] score array live per sub-block, and letting that grow to a whole
expert tensor turns the argmin memory-bound — the win falls from 5.6x at 256 rows
to 2.4x at 2048 without blocking.
Output is BYTE-IDENTICAL to the previous encoder on every shape tested, so this
is not a quality trade: the algebra is exact and the argmin is over the same
values scaled by a positive constant. Measured on the GLM-5.2 expert shapes
(Ryzen, 1 thread): 0.27 -> 0.82 Mw/s, i.e. 747 h -> 245 h for the model's 725e9
routed-expert weights, which is what makes an E8 container of it practical.
…al pass Two things that only bite on a big model converted on a small host. RAM: embed/lm_head at [154880, 6144] is 3.8 GB once dequantized to f32, and abs/divide/rint/clip each materialise another full-size copy — ~15 GB peak on a box with 13 GB free. Every non-E8 format here carries per-row scales (or per-group scales along I), so rows never interact and quantizing a row block at a time is BIT-IDENTICAL to quantizing the whole tensor; it just caps the peak. Verified equal for int8/int4/int4-g64/int3-g64/int2. E8 is excluded: it is already blocked inside iq3_pack.encode, and its ".qs" companion is a single format tag rather than a per-row scale, so it must not be concatenated. Progress: the --indir loop printed nothing per shard, so a local run that takes days was silent until it finished (stdout is a redirected file, so it is block-buffered too). It now prints the shard, free space and an ETA, flushed — the same feedback the download path already gave.
fmt=6 shipped with only its scalar reference kernel, and on a real model that
dominated everything else: on GLM-5.2 (744B, E8 container) a 32-token decode
spent 206.6s of 224.1s — 92% — inside matmul_e8, with disk service fully
overlapped (1.1s of actual wait). The format's own arithmetic was the wall, not
I/O and not the GPU. This is the same shape of problem AVX2 matmul_i3 fixed for
int3-g64, where the scalar path cost more than half the achievable rate.
The format suggests the vectorisation: one 8-weight lane is two 4-dim grid rows
(8 contiguous codebook bytes) plus 8 signs, which is exactly one AVX2 register.
So a lane decodes straight into a register and FMAs against x, instead of being
expanded into a stack buffer and read back:
- the 8 grid bytes widen with a single vpmovzxbd;
- the sign byte (the 7 stored bits plus the parity-derived 8th) expands to 8
lane masks with AND/CMPEQ against a bit-select vector and applies as an XOR
of the float sign bit. That removes the per-weight branch, which is what made
the scalar expansion expensive;
- the grid's half-unit convention folds into the sub-scale, so the magnitudes
need no separate scaling;
- two accumulators over a super-block's 32 FMAs keep it off one dependency
chain, and one horizontal add per 256 weights replaces one per 32.
The scalar path stays for the ragged tail and for non-AVX2 hosts; the converter
enforces whole 256-blocks, so the fast path is the one real containers take.
Kernel: 280.8 -> 4141.3 Mw/s single-threaded, 14.7x (44.8 -> 3.0 ms per
[2048,6144] expert tensor). tests/test_e8_kernel green against its fixture
(worst rel 1.46e-06); AVX2 vs scalar agree to 2.2e-07 relative L2, which is
accumulation order.
End to end on GLM-5.2 (RX 9070 box, 32-token greedy decode, cold cache, single
drive), same text produced in every configuration:
prefill decode expert-matmul
CPU before 91.3s 224.1s (0.14 tok/s) 206.6s
CPU after 16.7s 38.6s (0.83 tok/s) 18.7s
VK before 91.6s 177.3s (0.18 tok/s) 161.0s
VK after 17.2s 33.6s (0.95 tok/s) 16.8s
5.3-5.9x decode, 5.5x prefill. Decode is now disk-bound again (34.9s service,
3.4s wait) rather than compute-bound, which is the regime the streaming design
targets, and the VK tier's share rose from 22.7% to 31.0% of expert hits.
E8/IQ3 is a codebook quantizer: for every 4 weights it searches a 256-entry
lattice, and repeats that for all 16 candidate sub-scales. That is ~1280 ops per
weight against ~5 for int4, so producing an E8 container of a 725e9-routed-weight
model ran to ~27 days of CPU. This makes it hours, and it is not a one-off: the
cost is in the format, so every future E8 conversion pays it.
iq3_encode.c reimplements the numpy search exactly, and measuring the two halves
separately is what shaped it (2048x6144, one thread):
numpy 0.95 Mw/s
C scalar 1.58 Mw/s 1.7x plain fusion
C AVX2+FMA 12.71 Mw/s 8.0x over scalar
So the earlier "numpy is memory-bound" reading was wrong: a scalar C port with no
DRAM traffic at all is barely faster, which puts the cost in the search itself.
SIMD is where the win is, and it needed two things beyond widening:
- Four independent min chains. A single running min over the 256 entries is a
32-deep latency chain and left the vector units mostly idle — the same shape
of problem as the fmt=5 shader's FMA chain.
- Choosing a code needs only the SUM of the 8 group minima; the indices matter
for the winner alone. The 16 candidate passes now track a bare _mm256_min_ps
with no compare/blend and no index bookkeeping, and one extra pass recovers
the indices for the winning code: 17 cheap passes instead of 16 expensive
ones.
The scalar path is kept (4 accumulators, same structure) and is what builds on
anything without AVX2, so the tool stays portable; the library itself is OPTIONAL
and iq3_pack falls back to numpy when it is absent, so no build step becomes
mandatory. IQ3_NATIVE=0 forces the numpy path, which is how the two were A/B'd.
Output is BYTE-IDENTICAL to the numpy encoder across every shape tested, on two
different hosts. That is the whole acceptance criterion: the fp16 super-scale
conversion is a hand-written round-to-nearest-even to match numpy's astype rather
than relying on the hardware path, and the error expression keeps numpy's
association so code selection cannot diverge on a near-tie.
Threads scale near-linearly to 4 (12.9 -> 47.2 Mw/s). Measured end to end on the
GLM-5.2 conversion: ~68 h remaining -> ~3-4 h.
On Grace-Blackwell GB10 (DGX Spark) and other integrated GPUs the hot expert tier and the host RAM cache share ONE physical memory pool. The auto RAM budget reads g_mem_avail_boot, a MemAvailable snapshot captured at startup BEFORE the expert tier is placed, so on unified memory it over-counts free RAM by exactly the tier size. The budget then over-commits and the process is OOM-killed during prefill. Add coli_cuda_device_integrated() (cudaDeviceProp::integrated) and, after the tier is placed, subtract the placed bytes from g_mem_avail_boot when the primary device is integrated. Both consumers -- cap_for_ram and the overcommit guard -- read the corrected value, so the budget shrinks by the tier size and no longer over-commits. Discrete GPUs have a separate VRAM pool: integrated is 0 and this is a no-op, so their behaviour is unchanged. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The MSVC CUDA_DLL build links the host against backend_loader (not cudart), so every coli_cuda_* symbol must be resolved via GetProcAddress there too. The new coli_cuda_device_integrated was missing from the loader -> undefined reference at link. Add its typedef, struct slot, RESOLVE_OPT, and wrapper. RESOLVE_OPT (not RESOLVE): a coli_cuda.dll predating #653 leaves the pointer NULL and the wrapper returns 0 ("not integrated"), so the RAM-budget correction simply doesn't apply instead of failing the whole GPU backend. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
fix(cuda): correct RAM budget on integrated/unified-memory GPUs (#653)
quant: AVX2 matmul_e8 — fmt=6 decode was 92% of the time on a real model
convert: make fmt=6 (E8/IQ3) conversion practical — 15x encoder, RAM blocking, progress
…ask64 Full 64-value groups decode in one 512-bit iteration (dot_i3g64_avx512): the 16B low plane goes through an extended variant of matmul_i2's SSE unpack (adds the unpackhi stages + lane inserts) and the 8B high plane loads directly into a mask register driving a masked +4 on the low-plane bytes. Per-group partial accumulation with post-group scaling is preserved (fma reordered within a group only, like the NEON arm); partial tail groups and non-AVX-512 builds keep the scalar loop. Compile-time gate mirrors the existing dispatch (__AVX512F__ && __AVX512BW__, no build flag changes), and the arm carries the same runtime layer as the file's existing AVX-512 path: I3_AVX512 env kill-switch (g_i3_avx512, default on) and an I3_AVX512_TEST startup selftest against the scalar decode on a fixed all-lanes-asymmetric group, wired at the I4_ACC512 site. test_int3: add I=33 and I=7 shapes (single sub-group rows exercise the scalar-only fallback under every vector arm); CHECK failures now report the failing (I,S,O) tuple.
quant: AVX-512 arm for matmul_i3 (fmt=5 int3-g64)
…#663) The C engine has a HIP backend, but the Python tooling around it only knew NVIDIA, so a working AMD/ROCm setup was reported CPU-only and `--gpu N` failed. - doctor.py cuda_linkage(): a HIP build links libamdhip64, never libcudart, so match both vendors (mirrors the vendor-aware cuda_binary() already in c/coli). Fixes #663. - resource_plan.discover_gpus(): split into NVIDIA + ROCm paths; when nvidia-smi finds nothing, fall back to `rocm-smi --showmeminfo vram --showproductname --csv`. rocm-smi reports bytes (no MiB scaling) and its column names drift across ROCm versions, so columns are matched by substring, not position. Fixes #662. Both degrade to [] / not-linked on non-AMD hosts (verified: no rocm-smi -> FileNotFoundError -> []). The ROCm path is authored without a ROCm host to test against (issues labelled hardware-owner-needed) -> reporter to verify on AMD. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Brings main's 6 direct commits (site redesign #647, favicon, #642 gs64 doc links) back into dev. Conflicts in docs/quickstart.md and docker/README.md resolved in favour of dev, which carries the fuller gs64 guidance and the correct HF link (main had an older/shorter version and a broken google.com/search redirect URL). Makes the dev->main release PR #677 clean. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
20 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Promotes 158 commits from
devtomainas the next release. Cuts a known-good baseline before the next large batch (arch-seam, autotuning, Kimi K3, entropy) lands in dev.Highlights since v1.1.1
Correctness / hardware
doctor+resource_plan(HIP no longer reported CPU-only)Performance
matmul_e8(fmt=6 was 92% of decode on the scalar kernel)matmul_i3(fmt=5 int3-g64)Web / docs
Version
c/version.pyis the single source of truth (currently1.1.1). Proposing v1.2.0 (feature-level: AMD tooling, fmt=6/fmt=5 SIMD, dashboard). I'll add therelease: v1.2.0bump commit once you confirm the number — then this PR is ready for you to merge + tagv1.2.0.Nothing from the incoming batch (#667/#670 arch seam, #673 autotune, #651 CUDA tier, #676 Kimi, #671 entropy) is included here.