Conversation
The V4 core and its quantization kernels are portable C; the only x86 dependencies were two unconditional immintrin.h includes (whose intrinsics were already gated behind __AVX512F__) and the platform gates themselves. Guard the includes with __AVX512F__, accept aarch64/arm64 in the Makefile.deepseek-v4 uname gate, and add an aarch64-linux branch to the parent COLI_V4_SUPPORTED triplet gate so the amalgam unit tests and the tiny oracle run in make check. Without AVX-512 the rows16 hot-expert packing stays disabled and the engine uses the portable reference paths, exactly as on pre-AVX-512 x86. Validated on NVIDIA DGX Spark GB10 (Cortex-X925/A725, gcc 13.3.0): make check green (79 Python tests, 19 C test binaries, tiny oracle 11/11 token-exact), plus tiny fixture regenerated on-device with transformers 5.14.1 (identical reference tokens) and re-passed. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Add aarch64 NEON implementations of coli_fp4_matvec_rows16_v10 and coli_fp4_dual_matvec_rows16_v10 using the same 16-row interleaved layout as the AVX-512 kernels. The 16-entry E2M1 table is exactly 64 bytes, so a four-register TBL gathers whole decoded floats; UE8M0 block scales decode through the table once per 32-column block. Arithmetic is the same (activation * value) * scale multiply-multiply-add per column with no fused multiply-add, columns ascending, one row per vector lane — so rows are bit-identical to the scalar reference and to the AVX-512 kernel. A new COLI_FP4_ROWS16_KERNEL macro in native_quant_fp4_rows16.h names targets that have a rows16 kernel; the four engine gates that packed hot experts only under __AVX512F__ now key on it, enabling hot-pin packing on aarch64 under the same conditions as x86. x86 preprocessing is unchanged: with __AVX512F__ the macro is defined and the same branches are selected. Validation on GB10 (Cortex-X925/A725, gcc 13.3.0): standalone harness compares NEON kernels against the flat scalar reference over five shapes up to 7168x2048 (single and dual, both nibble paths) — all outputs bit-identical with -ffp-contract=off and with default -O3; make check green including the token-exact tiny oracle. Note the tiny fixture leaves pin_slots_per_layer=0 (4 experts, all resident) on every platform, so rows16 execution coverage comes from the harness, not the oracle. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Enable the DeepSeek V4 engine on aarch64 Linux + NEON rows16 kernels
Inkling's audio tower is one embedding table plus one RMSNorm (~10 MB bf16 on Inkling-Small), so there is no reason for the int4 container to drop it. --keep-audio passes model.audio.* through under its original names; the engine loads them directly. Vision and MTP stay skipped. Default off, so existing conversions are byte-identical. make_tiny_inkling_audio.py is the audio counterpart of make_tiny_inkling.py: a random-init InklingForConditionalGeneration with an audio tower, saved in the native TML layout, plus a ref_inkling.json that adds "dmel" (flattened [n_frames, n_mel_bins] levels). The C engine must reproduce tf_pred and the greedy continuation token-for-token. The vision tower is stubbed out: it is irrelevant to the audio oracle and its constructor trips on tensor-typed dims with tiny configs.
input_audio content parts on /v1/chat/completions now work when the engine is Inkling: base64 WAV in, spoken answer out. The gateway decodes the WAV (PCM16 or float32, mono after mixdown, 16 kHz enforced - resampling belongs at the capture edge), runs the PCM through a numpy port of the DMel DSP, renders TMLv0 audio framing, and ships the frames to the engine as a binary extension of the SUBMIT header. The DSP is the part that has to be exactly right, so it is tested against the canonical implementation rather than the paper description: byte-identical output vs tinkernel-audio, which is itself byte-golden against the official tml-renderers 0.1.0 wheel. That covers the 100 ms periodic-Hann window centered on i*hop, zero edge padding, slaney mel with 2/(upper-lower) normalization, magnitude-domain projection, the turn-level RMS boost for quiet audio (rms < 0.01), and round-half-up quantization to 16 levels in log10 [-7, 2]. Note the HF processor skips the RMS boost; tml-renderers is what the model saw in training, so that is what we match. Framing: a multipart user message becomes one TMLv0 message per content run, in part order - text runs under <|content_text|>, each audio clip as <|content_audio_input|> + one <|audio|> placeholder per frame + <|audio_end|>. A TMLv0 message carries a single content type, so this split is the trained distribution, not a workaround. Wire: the optional 7th SUBMIT field was grammar bytes for glm; for inkling it is DMel bytes. The two engines never see each other's extension, and inkling rejects grammars upstream, so the framing cannot desync. numpy is imported lazily and only for audio requests; text-only serving still runs on the stdlib.
Two additions that together put Inkling-Small on a 128 GB Apple Silicon box with audio in: the engine consumes DMel frames at <|audio|> placeholders, and the routed-expert math runs on the Apple GPU. Audio. Inkling's audio tower is a [mel_bins*mel_vocab, D] embedding table plus one RMSNorm: a frame's embedding is rmsnorm(sum of 80 rows, eps 1e-6), and it REPLACES the placeholder's text embedding (embed_norm does not apply), matching masked_scatter in HF modeling_inkling.py. The tensors load from model.audio.* when present - either a --keep-audio container or an audio.safetensors sidecar dropped in the snapshot dir (st_init already indexes every *.safetensors). Absent tensors = text-only engine, unchanged. Entry points: --audio file.dmel on the CLI (implies --chat, emits the TMLv0 audio framing), an optional 7th SUBMIT field carrying raw DMel bytes on the serve protocol, and a "dmel" array in ref_inkling.json for the oracle harness. Validated token-exact against the HF multimodal oracle: 38/38 teacher-forced, 24/24 greedy. Frame counts are checked against placeholder counts up front; a mismatch is an error, not a silent skip. Metal (COLI_METAL=1, opt-in). Reuses colibri's batched coli_metal_moe_block: the container's nibble-packed int4 with per-row scales is bit-identical to the fmt=2 kernel, int8 to fmt=1, and the fused p13 tensor needs no reshaping because resolve() accepts interior pointers (gate = p13, up = p13 + I rows). Expert slots move from per-slot mallocs to one page-aligned slab per sparse layer, carved into slot regions up front and registered with the backend once - eviction reuses the region, so there is no register/unregister churn and unified memory stays zero-copy. Each cache round groups its (token, expert) pairs by expert and submits one command buffer; a failed submit falls back to the CPU loop for that round only. Verified token-identical to the CPU path on an int4 container with the audio oracle active (294 blocks on GPU, 0 fallbacks), and a [metal] counts line makes GPU execution provable rather than assumed. Attention, dense MLP, and shared experts stay on the CPU: at high hit rates ~90% of decode is routed-expert matmul, so that is the lever that pays first. Also: mem_avail_bytes gets a macOS branch (free + inactive + purgeable), so the auto expert-cache cap sizes to the machine instead of falling back to 16 experts/layer on a 128 GB Mac.
At S==1 a MoE round is topk pairs of 6-row kernels; the dispatch+sync latency (~135 ms/block measured) swamps the math. Inkling-Small decode: 0.14 tok/s on GPU vs 0.63 on CPU. Prefill rounds carry up to cap pairs and amortize the launch, so the GPU stays on for S >= 2. INK_METAL_MIN_S=1 restores GPU-always for A/Bs.
Three findings from a 10-config knob sweep on Inkling-Small (M-series 128 GB, frozen probe, 2 reps each, all reproducible): - OMP active-spin tuning off on Apple: the spinning team steals the SoC power budget and loses even for pure-CPU decode (0.50 -> 0.84 tok/s) and prefill (49.7 -> 24.4s). Same mechanism the M5 Max Metal report documented for GLM. - Metal residency set on by default: without it every MoE block pays per-buffer useResource churn, ~135 ms per decode block. With it the same block is ~3 ms: 0.17 -> 1.76 tok/s decode. - GPU decode back on (INK_METAL_MIN_S default 2 -> 1): with the residency set the GPU wins decode 2x over CPU, so the prefill-only gate now costs performance instead of saving it. Net on the frozen probe: 1.75 tok/s decode + 10.3s prefill, from 0.50 and 49.7 at the old defaults. Output tokens unchanged; tiny-oracle validation stays exact on both CPU and GPU paths. The phase profile now puts shared experts and attention (both CPU) above expert matmul.
kimi_k3 and inkling never called _setmode(..., _O_BINARY). On Windows the CRT
opens stdout in TEXT mode and rewrites '\n' as '\r\n', so the engine emits
\x01\x01READY\x01\x01\r\n
and coli -- which matches the sentinel byte-exactly -- waits for a byte that was
mangled on the way out. No error is printed, because nothing failed: one side is
waiting, the other already sent something else.
@brad-evony hit it in #748 on Windows 10: 93 layers loaded over 42 minutes, the
tokenizer read, and then silence. Ctrl-C landed in read_engine_turn's
stream.read(1), which is exactly where the handshake blocks.
colibri.c has done this since #195 and its comment describes the symptom word
for word ('the READY sentinel never matches and chat hangs'). The two newer
engines were written without it and nothing noticed for months. Not a v1.4.0
regression -- the bug predates it. v1.4.0 only made it reachable, because before
that the archives contained no kimi_k3.exe at all (#720).
It goes in compat.h rather than being pasted into each engine: pasting is how it
went missing, and this is the third Windows-specific behaviour to be found
living in colibri.c alone (after select()/fd_set in #736).
Adds tests/test_serve_sentinel.c, which writes the handshake through a real
stream and asserts the bytes contain no CR. Nothing in the tree looked at those
bytes before, which is why this reached a user rather than CI. It can only pass
on Linux/macOS -- that is the point of a regression guard, and CI builds every
engine on Windows since #736, so it runs where it can fail.
olmoe.c has no serve mode and is untouched. colibri.c is untouched: it is
already correct, and 14 open PRs touch that file.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
fix(win): put stdout in binary mode before the serve handshake (closes #748)
inkling: audio input, Metal expert MoE, and tuned Apple defaults
@centralware and two colleagues downloaded the 400 GB model, compiled the engine, and gave up at step 3 of the quickstart -- not on anything technical, on our prose. Their exact objections, and what each one now says: 'multi-turn wire' read as wire-wrapping, a physical technique. It meant the bytes sent to the model across a several-turn conversation. '77/77 token ids' meant nothing without knowing what encoding_k3.py is. It now says we compared our text handling to Moonshot's own and got identical numbers on 77 test conversations. 'gated on a host' meant 'we cannot test this until someone has 1.6 TB'. That paragraph is gone with the v1.4.0 README rewrite. The contradiction they reported is real and was ours: the Kimi K3 note (1.6 TB, serious machine) sits near the GLM-5.2 quickstart (372 GB, no GPU). Both true, about different models, printed together with nothing saying so. Reading them as one requirement produces exactly 'needs 1.6 TB' next to 'runs on a laptop'. Adds a per-model requirements table to the roster -- disk, RAM, GPU, one row each -- so the numbers are attached to the model they belong to, and says plainly that none of the four needs a GPU. Three electronics engineers with the hardware, the download and a working compiler are not the difficult case. If they bounce off the words, the words are wrong. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
docs: plain language for the words that stopped three engineers (#191)
Fourth and final engine for #700, after colibri.c (#716), kimi_k3.c (#719) and inkling.c (#742). Every engine that routes experts now writes the same format. m->freq moves from one flat [n_layers * n_experts] block to route_trace.h's row per layer, so the three read sites -- the hot-pin ranking, the bump in moe() and the LFRU eviction guard -- change shape but not meaning. last_access stays flat; it is a separate array route_trace.h does not own. Every olmoe layer routes and there is no MTP layer, so the only row dropped is the spare one rt_init leaves. The history is opt-in behind COLI_USAGE, exactly as kimi_k3 does it: with the variable unset nothing is loaded and nothing is written, so the default path is unchanged. PPL=1 deliberately does not save, so a loss sweep cannot fold its own tokens into the persisted ranking. Verified against the real OLMoE-1B-7B-0125-Instruct converted with tools/convert_olmoe_merged.py. Output is byte-identical to dev on the same run, down to hit=1216 miss=832 -- the eviction path decides exactly as before, which is the property the layout change could have broken. The COLI_USAGE round trip writes the -1/-2 header with the olmoe identity, leaves no row for the dropped layer, and accumulates 2048 -> 4096 across two runs. make check green. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Fourth and final engine for #700, after colibri.c (#716), kimi_k3.c (#719) and inkling.c (#742). Every engine that routes experts now writes the same format. m->freq moves from one flat [n_layers * n_experts] block to route_trace.h's row per layer, so the three read sites -- the hot-pin ranking, the bump in moe() and the LFRU eviction guard -- change shape but not meaning. last_access stays flat; it is a separate array route_trace.h does not own. Every olmoe layer routes and there is no MTP layer, so the only row dropped is the spare one rt_init leaves. The history is opt-in behind COLI_USAGE, exactly as kimi_k3 does it: with the variable unset nothing is loaded and nothing is written, so the default path is unchanged. PPL=1 deliberately does not save, so a loss sweep cannot fold its own tokens into the persisted ranking. Verified against the real OLMoE-1B-7B-0125-Instruct converted with tools/convert_olmoe_merged.py. Output is byte-identical to dev on the same run, down to hit=1216 miss=832 -- the eviction path decides exactly as before, which is the property the layout change could have broken. The COLI_USAGE round trip writes the -1/-2 header with the olmoe identity, leaves no row for the dropped layer, and accumulates 2048 -> 4096 across two runs. make check green. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…'t trust native (#631) On Linux ARM the Makefile fell through to the x86-64 else: -march=native on a gcc that doesn't know the core (gcc 13 vs Cortex-X925) silently degrades to plain armv8-a, the SDOT/SMMLA int8-int4 kernels compile out, and int4 S=1 decode lands on the f32 fallback (g_i4s=2 without DOTPROD). New branch, default ARCH=native: keep -mcpu=native if it already defines __ARM_FEATURE_MATMUL_INT8; otherwise re-add the features /proc/cpuinfo reports (asimddp -> +dotprod, i8mm -> +i8mm) as -march=armv8-a+..., kept only if the compiler accepts the modifiers (gcc < 10 has no +i8mm), else fall back to -mcpu=native with a warning. Explicit ARCH= respected: armv* -> -march, core names -> -mcpu, same convention as the Darwin branch. GB10 default build: 0 -> 22 smmla, banner idot: neon-i8mm, int4 S=1 decode kernel ~2.5x. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…s-current docs: acknowledge inference research foundations
docs(experiments): GLM-5.2 on 4x RTX A6000 (Ampere) + EPYC 7402P
Fix stateful KV tail at the NGEN limit
inkling: shared experts to the GPU, 1.75 to 2.46 tok/s on Apple Silicon
…n-r1 docs: register int4-rans256-g0 in FORMATS.md; registry parity for ordinal-less rows
cuda: fmt=8 (fp8-e4m3) — dense matmul + expert-group kernels with 128x128 block scales
…-runtime cuda: add paged ragged KV research runtime
…-closeout docs: publish GLM-5.2 inference research closeout
Build the container image in CI
Run the structural efficiency tests in CI
One conflict, two rows of the same CUDA table. This branch enriched
CUDA_EXPERT_GB ("Also accepts `auto`") and added CUDA_RESERVE_GB; dev
added CUDA_EXPERT_LOAD_BALANCE (#795). Kept the fuller description and
both new rows -- neither side is contested, they were simply added at
the same line.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Document the 79 undocumented environment variables (and fix the scan that missed them)
feat(windows): enable native HIP with deterministic runtime binding
deepseek_v4: reuse the KV prefix instead of re-prefilling every turn
`All engines (windows)` has been red on dev since #788 landed: FAIL: test_abi_is_derived_from_the_loader_source self.assertEqual(len(f.mandatory), 46) AssertionError: 47 != 46 Nobody's change is at fault. Three PRs landed the same day and each was green on its own branch: #795 RESOLVE(attention_project_ragged) -> mandatory 46 -> 47 #817 RESOLVE_OPT(fp8_set_lut) -> optional 2 -> 3 #788 the test that asserts 46 and 2 #788's CI ran before the other two were in its base, so the collision could only appear on dev. The hardcoded counts are worth keeping. The ABI is what every Windows coli_cuda.dll must export, and widening it should be a conscious act, not something a user discovers as a missing symbol at load time. So this updates the numbers rather than removing the assertion -- and adds the two new symbols by name, plus exports == mandatory + optional, so the next failure carries its reason instead of only a different integer. The parser itself was already right: it derives the ABI from backend_loader.c rather than restating it, which is why the drift showed up as a count and not as a silently wrong stub. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
`coli chat`, `run`, `info`, `plan` and `tune` all printed
GLM-5.2 · 744B MoE · int4 · streaming CPU
regardless of what was loaded. Load Inkling, Kimi K3, OLMoE or the new
DeepSeek V4 engine and the third banner line still named GLM.
## Why it was wrong, and where the wrongness lives
Not in the banner. In what the banner would have had to ask:
def model_arch(model):
...
return "glm" # <- every unrecognised model_type
That default is correct for its job -- colibri.c is the general engine, so
an unknown checkpoint should be offered to it. It is the wrong answer to a
different question, "what did the user load", and the banner needs that
one. So `model_banner_line()` reads the raw `model_type` instead, and the
dispatch in `model_arch()` is untouched.
## What it prints
GLM-5.2 · 744B MoE · 372 GB on disk
DeepSeek V4 Flash · 284B MoE · 167 GB on disk
OLMoE · 7B MoE · 4.2 GB on disk
qwen3_moe · 48L x 128E MoE · 61 GB on disk <- not in the roster
Parameter counts come from the README roster, so the two cannot drift apart
without someone noticing. A model_type that is not in the table is NOT
forced into a name: it prints its own type and the geometry measured from
its config, which is honest and still useful. No config.json at all, or an
unreadable path, keeps the original tagline -- `coli info` banners before it
validates the model directory, so this has to be safe on a bad path.
Size is measured by stat-ing the shards; no safetensors headers are parsed.
A banner runs before every command and may not cost a scan of a 400 GB
checkpoint.
## Note for the open PRs that also touch banner()
#825 and #670 both change `banner("run")` to pass a CUDA flag positionally.
`model=` is therefore **keyword-only**, so a second positional argument
raises TypeError instead of being silently read as a path. There is a test
for exactly that.
Eight tests in tests/test_cli_output.py, including the regression itself: a
`deepseek_v4` config must not produce a line containing "GLM" or "744B".
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
coli: the banner said GLM-5.2 over every model
Starting `coli chat` on DeepSeek V4 prints this while the model loads:
Exception occurred during processing of request from ('127.0.0.1', 45576)
Traceback (most recent call last):
...
File "openai_server.py", line 1771, in send_json
self.wfile.write(data)
BrokenPipeError: [Errno 32] Broken pipe
Nothing is wrong. `coli` polls /health on a 2 s timeout while waiting for
the engine and drops each connection the moment it has an answer. The
handler is still writing; socketserver catches the escaped exception and
logs a full traceback for it. On a 24-second DeepSeek load that buries the
loading spinner under stack traces, and a first-time user reasonably reads
it as a crash.
The same thing happens on Ctrl-C during a stream -- which the banner tells
the user to do: "Ctrl-C stops the answer". Cancelling an answer should not
look like a fault.
Caught in handle_one_request rather than in send_json so it also covers the
SSE writes in the streaming path, which is where Ctrl-C lands. The
connection is marked closed and the handler returns; no response is
half-drained back into a keep-alive socket.
Two tests, in a class that records handle_error instead of printing it:
a client that sends a request and closes with SO_LINGER 0 (RST, not a clean
FIN) must produce no server error, and the server must still answer the
next request -- the real damage would be a handler thread lost to the
exception rather than the noise itself.
Verified the tests fail without the fix:
AssertionError: Lists differ: [ConnectionResetError(104, ...)] != []
Full suite: 110 tests, green.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
serve: a client hanging up is not a server error
`check` has been red on dev since #832: ERROR: test_size_is_reported_without_rounding_to_zero OSError: [Errno 28] No space left on device Ran 394 tests in 549.596s FAILED (errors=1, skipped=51) My fault, and the comment on the line was the mistake: shard.truncate(shard_bytes) # sparse: costs no disk True on ext4 and APFS. NOT true on NTFS, where ftruncate allocates. The test asked a Windows runner to materialise 372 GB and filled its disk. Linux and macOS passed, so it only ever failed on one of three platforms -- and only after landing, since #832's own CI ran before the merge. The size is now supplied by patching os.path.getsize, and the shard on disk is an empty file that exists only so listdir() finds it. Same assertions, no allocation anywhere: measured, the whole class now costs 4 KB instead of 372 GB. Verified it still fails when the formatting is wrong -- changing the sub-10 GB branch from .1f to .0f makes it red -- so this is not the kind of repair that quietly turns a test into a no-op. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
c/tests/bench_omp_grain is a 16,880-byte ELF executable, mode 100755, added by 00129d2 (#808). It is the only tracked build artifact under c/tests/ and it is mine. .gitignore already guards c/tests/test_* with source-file exceptions, but the pattern was never extended to bench_* or fuzz_*, which build to extensionless binaries in the same directory. That gap is how it got in, so the fix is the gap rather than the one file: same shape, same exceptions, for both prefixes. This also matters for #801. That PR fixes tools/clean.py, which today globs only tests/test_*.exe and therefore deletes nothing on Unix; once it correctly removes extensionless test binaries, `make clean` would delete a tracked file and leave every contributor with a dirty tree. Removing the binary is what keeps that fix harmless.
…fting
Measured on an i7-1355U (Raptor Lake: AVX2, no AVX-512), DeepSeek V4 Flash,
RAM_GB=18. Everything here is token-exact against the transformers oracle.
## Why this engine was 75 s/token on consumer Intel
The rows16 fp4 fast path was gated on
#if defined(__AVX512F__) || defined(__aarch64__)
Intel dropped AVX-512 from consumer parts at Alder Lake, so on most laptops
and desktops sold since 2021 the gate was false, coli_fp4_matvec_rows16_v10
returned -1, and every routed expert fell back to the flat scalar path.
601 s for 8 tokens.
The AVX2 backend mirrors the AVX-512 one rather than reinventing it: two
__m256 accumulators where AVX-512 has one __m512, the same per-column
(activation * value) * scale, ascending, no FMA. That ordering is what keeps
all three backends bit-identical to the scalar reference.
The same argument applies to the fp8 dense path, which is roughly half the
per-token compute: an fp8 rows8 layout with an in-place repack at load, and
an AVX2 bf16 dot for the LM head. The head deliberately stores its eight
products and adds them in order rather than reducing horizontally -- one ULP
there flips a near-tie argmax into a different token.
## Bytes, not just arithmetic
- O_DIRECT expert reads through the twin descriptors st.h already opens,
with page-aligned slabs. A FLOCK-packed checkpoint becomes one request;
an ordinary HF layout uses direct I/O for weights and a small buffered
read for scales. Every failure falls back to the buffered path, so an
unsupported filesystem cannot become an unusable one. V4_DIRECT=0 opts out.
- .coli_usage warm start. The pin ramp used to rediscover the same hot
experts from cold disk every session; a persisted history is already the
evidence it was waiting to collect, so pins rank on it immediately. The
file is the same bytes the other engines read and write.
## Speculative drafting: built, measured, left off
DSpark's markov table drafts candidates and the target verifies them, so
accepted tokens are always the target's own argmax -- the draft can save
forwards, never change output. Rejection restores an exact attention
snapshot and replays only verified inputs at unchanged absolute positions.
Measured on real multi-turn chat: **1 accepted candidate in 15**. On this
engine a rejection costs a recurrent-state replay that dominated the visible
decode, so V4_DRAFT defaults to 0 and the whole path is opt-in. A partial
n-gram rejection now disables it for the rest of the session unless
V4_NGRAM_PARTIAL_KEEP says otherwise.
Keeping the code with the measurement recorded beside it is the point: the
next person to try this on faster storage will find both.
## --no-dspark changed meaning, so its test did too
It was a no-op that printed a notice, because a target-only engine had
nothing to disable. It now disables drafting for real. The tiny-oracle check
asserted the old notice; asserting it still would require the engine to keep
claiming it does nothing. It now asserts the new contract instead -- with
drafting off, the engine's own attempt counter must be zero.
## Verification
- make deepseek-v4-tiny-check: 15/15, including greedy and teacher-forcing
token-exactness against DeepseekV4ForCausalLM
- make test-c: full native suite
- make test-python: 352 passed, 27 skipped
- colibri, inkling, kimi_k3, olmoe and deepseek-v4 all build clean
c/deepseek_v4_dspark.inc is added here. It was untracked while this was
developed, so the tree built only where that file already existed on disk --
a fresh clone would have failed at COLI_V4_UNIT_KV_CACHE.o.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Windows UCRT64 died in tests/test_deepseek_v4.exe before it printed a single line -- no assertion, no message, just a non-zero exit. Linux and macOS were green. compat.h maps posix_memalign onto _aligned_malloc on Windows and says at its own definition that those blocks must go to compat_aligned_free, because passing one to free() corrupts the CRT heap. Three releases of the O_DIRECT bounce buffer used plain free(), and that buffer is allocated on every direct expert read -- so the heap was gone long before any test output reached the console. The slab in the hot rows16 store had the same problem in its destroy path. That store allocates both ways (malloc for the buffered path, posix_memalign for the aligned one) and already carries an aligned_slab flag to tell them apart; the destroy loop simply did not consult it. The second expert store's destroy is left on plain free() deliberately: that one allocates slabs with malloc only, so routing it through compat_aligned_free would be wrong in the other direction. Applied the same edit to both at first and the compiler caught it -- V4ExpertSlot there has no aligned_slab member, because it never needed one. POSIX is unaffected either way: compat_aligned_free is plain free() there. make test-c and make deepseek-v4-tiny-check both pass again on Linux; the Windows path is what this restores. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Untrack the bench binary my own PR committed
deepseek_v4: AVX2 kernels, direct expert I/O, warm pins, verified drafting
Fixes GHSA-pmq2-6f2p-hjvf, GHSA-c84r-hmpc-qg8h, GHSA-2h73-p2rc-4796, GHSA-5xpg-vw35-2687, GHSA-rfqv-4g4x-j4vr and GHSA-25w8-8c74-g9c8. Reported by @zh-Processor with exact source locations, ASAN traces and working reproducers. Every one was confirmed against this tree before it was touched. ## inkling: load_scalar() wrote a file-sized tensor into a 4-byte stack slot st_read_f32 validates the header against itself -- numel*esz == nbytes -- and cannot know the destination size. `float v` is four bytes; a snapshot declaring gate.global_scale as F32 [4096] therefore dropped 16 KiB of attacker bytes over the frame and its return address, with RSP fully controlled in the reported crash. st_read_f32_cap already existed for exactly this and takes the caller's capacity. A scalar's capacity is 1. ## inkling: load_cfg() indexed layer_types[] by layer count, not array length Both loops ran to num_hidden_layers while indexing a JSON array of an unrelated, equally attacker-chosen length. 66 layers with a one-element layer_types read past a malloc'd array that starts at capacity 8, then dereferenced whatever followed as a string. config.json alone was enough. A short array now means "unspecified", falling through to the same default an absent key uses; element types are checked, because kids[i]->str is only a valid pointer when t == J_STR. ## SERVE: max_tok was the one submit field nobody validated prompt_reject() asks `np + want > ctx_max`. A negative want makes that sum smaller than np, so the context check passed for any prompt, kv_alloc came out shorter than the prompt, and prefill wrote past the K/V cache. The official gateway forces a positive integer, so this is unreachable through openai_server.py -- but the SERVE protocol is public and anything bridging it exposes it, so the check belongs beside the one plen already had. ## Router: all-NaN logits left the top-k pick at -1, and -1 was used as an index NaN > bv is false for every bv, so best stayed -1 and became score[-1], usage[-1]++ and a pread at a negative offset. Release builds did not fault: they finished and returned wrong numbers over quietly corrupted memory. colibri.c has guarded this since test_logit_nan.c was written. inkling, kimi_k3 and olmoe never received it -- the recurring shape of defects here, a fix that lands in one engine and not its siblings. So the guard goes in route_trace.h, the one header all four already include, instead of being pasted a third and fourth time. ## /profile served telemetry to anyone, with --api-key set /health and /experts gained an _is_authed() gate; /profile, served in the same block before require_auth(), did not. It carries prompt and completion token counts and per-phase timings for the last 120 turns. A test asserted the old behaviour, so this was deliberate once. It is not a client requirement: the dashboard sends Authorization: Bearer to /profile exactly as it does to the other two (web/src/lib/api.ts). The test now asserts the gate instead, and says why. ## Serve: unbounded threads, and a timeout a drip could renew forever One thread per connection with no ceiling, 8 MiB of stack each, and a `timeout` that is per socket operation -- so a byte every 29 s renewed it indefinitely. The comment claimed it stopped exactly that. - MAX_CONNECTIONS (64): the accept loop is bounded. Over the cap we close rather than queue, so a flood costs the attacker's socket, not our address space. - MAX_CONNECTIONS_PER_IP (8): a global cap alone converts exhaustion into starvation. Measured that while testing -- one source held every slot and a legitimate client was refused -- so one address is bounded well under the server cap. - READ_DEADLINE (30 s): cumulative, accept to end of body. Every read shrinks the socket timeout to the time left, so a drip runs the clock down instead of resetting it. It covers the read phase only: send_response hands the socket back to the ordinary timeout, because a 600-second generation is normal and must not inherit a header clock. All three are env-overridable (COLI_MAX_CONNECTIONS, COLI_MAX_CONNECTIONS_PER_IP, COLI_READ_DEADLINE). Verified under a real slowloris: 40 dripping connections from one address hold 4 of 16 slots, a client from another address still gets 200 OK during the attack, and every connection is reclaimed when the deadline expires. ## Not fixed here, because they already are GHSA-4gw4-j89j-4c8r and GHSA-wc4x-3786-cxh7 describe primitives that #413 closed: st.h now validates data_offsets ordering, file bounds and numel*esz == nbytes with a shape-overflow guard, and tok.h rejects negative and implausible token ids. Both re-checked against this tree. 113 server tests pass, including three new ones for the connection limits; all four engines build clean. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
security: six advisories — loader bounds, router NaN, serve limits
## The roster said four families; there are five DeepSeek V4 Flash landed in #165 and was tuned in #839, but the README still opened with "Four families run today" and its table stopped at OLMoE. Someone scanning the front page had no way to learn the engine exists. It is now in the opening line, in the roster table, and in the hardware table above it -- ~167 GB on disk, 16 GB of RAM minimum and 22 comfortable, measured on the reference box rather than estimated. ## The DeepSeek section described a version that no longer exists It called the path "experimental" and said "DSpark is intentionally kept for a separate stacked follow-up". DSpark is in, and the honest state is more interesting than either claim: - the checkpoint streams with no conversion -- routed experts stay native fp4, dense stays fp8-e4m3 with UE8M0 block scales - greedy, one KV slot, no tools or grammar yet: said plainly, because finding that out from a rejected request is worse - --ram is the knob that matters. 43 x 256 routed experts are ~137 GiB and a token touches 301 of them, so the cache hit rate sets tok/s. It changes speed only, never output. - speculative drafting is implemented, verified, and OFF, with the numbers that made that call: 1 accepted in 15 for the markov drafter, 10 in 24 for full MTP, and a 14-token answer that took 495 seconds to replay its rejected suffixes That last one is the point of documenting it at all. The code stays, the measurement stays beside it, and whoever retries this on faster storage starts from evidence instead of from scratch. ## Repo layout described a tree that has not existed since July It listed `glm.c`, renamed to `colibri.c` in #391 three weeks ago, and no other engine -- so the file that runs GLM was wrong and the four files that run everything else were missing. Also absent: quant.h, compat.h, expert_store.h, route_trace.h, kv_prefix.h, the Metal and Vulkan backends, resource_plan.py, and docker/. Every path and every make target in the new listing was checked to exist on this branch before it was written down. The rule behind the layout is now stated, because it is the one that keeps being violated: one .c per model family, over shared single headers. An engine owns its architecture and nothing else. The recurring defects in this tree -- the OpenMP thread count, the KV prefix reuse, the NaN router guard -- are all the same shape: a mechanism that landed in one engine and never reached its siblings. ## Also `#### Other supported models` now sits where the roster table is, so `[Full roster ↓](#other-supported-models)` in the opening paragraph resolves to the table instead of to prose four sections earlier. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
docs: DeepSeek V4 joins the roster, and the repo layout matches the repo
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
14 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.
Release v1.5.0 — 51 pull requests from 15 contributors.
A fifth engine: DeepSeek V4 Flash
@DrewZt's #165 brought the target-only CPU engine — MLA + DSA sparse attention, 43 layers, 256 routed experts plus one shared, top-6. The official checkpoint streams with no conversion: routed experts stay native fp4, dense stays fp8-e4m3 with UE8M0 block scales.
#839 then made it usable on the hardware most people own. The rows16 fp4 fast path was gated on
__AVX512F__ || __aarch64__, and Intel dropped AVX-512 from consumer parts at Alder Lake — so on most laptops and desktops sold since 2021 every routed expert fell back to the scalar path. Measured: 601 s for 8 tokens. The release adds the AVX2 backend, an fp8 rows8 dense path, O_DIRECT expert reads,.coli_usagewarm pins, and KV prefix reuse so a second chat turn stops re-prefilling the first.Speculative drafting is implemented, verified, and off: 1 accepted candidate in 15 for the markov drafter, 10 in 24 for full MTP, and one 14-token answer that took 495 seconds replaying its rejected suffixes. The code stays with the numbers beside it.
Eight security advisories, published with this release
Fixed in #841 and disclosed today rather than patched quietly — the fix commits were already public, and silence would only have kept users on 1.4.0 from learning they should upgrade.
load_scalar()wrote a file-sized tensor into a 4-byte stack slot; RSP was fully controlled in the reported crashload_cfg()indexedlayer_types[]by layer count instead of array length — a craftedconfig.jsonalone, no weightsmax_tokaccepted negatives, which inverted the context check and overflowed the KV cache-1and used it as an index; release builds did not fault, they returned wrong numbers over corrupted memory.colibri.chad been guarded sincetest_logit_nan.c; the three sibling engines never received it/profileserved per-turn telemetry to unauthenticated callers even with--api-keysetReported by @zh-Processor (six), @aeonframework and @ajmeese7 (two, already closed by #413 and re-verified here). All credited on the advisories.
Everywhere else
Windows gained native HIP with fail-closed runtime binding (#788, @Kenneth-Javier), validated on physical gfx1151. Apple Silicon:
METAL=1builds on a stock macOS again (#807), and Inkling's shared experts reach the GPU for 1.75 → 2.46 tok/s (#757, @rgbkrk). CUDA got fmt=8 dense and expert-group kernels (#817, @kreuzzelg) and a paged ragged KV runtime (#795).OMP_NUM_THREADSis now sized from physical cores on every platform (#805, @ThefloorMiner) — and, after this release exposed that it reached only GLM, for all five engines. A one-iteration OpenMP region is gone (#806), measured on a 744B model across three runs per side by the reporter of the counter-evidence.Kimi K3 learns from use (
.coli_usage, autopin, decay), the adaptive LRU survives autopin (#815), the stateful KV tail at the NGEN limit is fixed (#567), and the CLI banner finally names the model you actually loaded instead of GLM-5.2 (#832).Thanks
@ZacharyZcR (10 PRs), @ThefloorMiner (7), @terrizoaguimor, @rgbkrk, @mgua, @kreuzzelg, @bherald, @DrewZt, @winklemad, @monotophic, @dyKiU, @anrasi, @RDouglasSharp, @Kenneth-Javier — and @zh-Processor, @aeonframework, @ajmeese7 for the security work.
Special mention to @ThefloorMiner, who answered a performance claim by running the real 744B model with
.coli_usagesnapshotted between runs, three runs per side and a byte-for-byte output diff — and reported numbers that contradicted the original framing. That is the standard this project wants.