Inkling (Thinking Machines 975B MoE): new engine, o200k tokenizer, bf16->int4 converter#312
Inkling (Thinking Machines 975B MoE): new engine, o200k tokenizer, bf16->int4 converter#312nbeerbower wants to merge 15 commits into
Conversation
|
Pushed a Stage-B performance commit (3ef6775) β same architecture, the expert path grew the glm.c playbook:
Measured on the 975B checkpoint (Ryzen 9 7900, NVMe, cap auto=64/layer, 1024 pinned):
All four oracle modes (f32, int4-container VNNI, IDOT=0 scalar, bf16 residents) remain token-exact, π€ Generated with Claude Code |
|
Final push for this PR β it now includes the CUDA resident tier and cache warming, bringing inkling to feature parity with what glm already has (no concepts new to colibri):
Measured on the 975B (Ryzen 9 7900 / A6000 / NVMe, commit 87981c1, greedy):
The honest number is the last row; it grows toward the best case as usage history accumulates (counts persist across runs). Phase timers show the remaining time is ~90% CPU expert matmul at high hit rates β GPU expert compute is the natural follow-up PR, not included here. All modes remain token-exact vs the transformers oracle (f32, int4-container VNNI + IDOT=0, bf16 CPU, bf16 CUDA); π€ Generated with Claude Code |
|
Thank you for this β real work, and I want to be straight with you about the bar rather than let it sit open for another week. The maintainer's rule, and it's the right one: a new engine lands only if he can run it on his own machine first. Not "the CI compiles it" β run it. Right now neither of us can: there's no checkpoint here, so the only thing that could be verified is that gcc doesn't complain. That's not verification, that's a spell-check. Why the bar is that high, concretely. This week, on the one engine we can run:
That's the state of the engine that gets daily attention. A second engine that nobody here can execute doesn't get daily attention β it gets zero, until the month it breaks and the person holding it is the maintainer, who never had the weights. What would change the answer, and I think it's reachable:
Until then the honest home is a fork, linked from the README. That's not a brush-off: it means people can find and use your work, and it doesn't hand a maintainer 10,000 lines he can't test. If it gets traction and someone can run it in CI, it comes back. I'd rather tell you this now than leave it open for a month and close it then. |
|
Specific to this one β and it's a more favourable read than #165. You have a track record here: #232 landed today after I merged and tested it (PROF=1 works, the instrumentation costs ~48Β΅s/token against tokens that take 3-30 seconds, so it's free in practice). That matters. This isn't a first PR from someone we've never seen.
The part I'd check hardest is Concretely, that's testable without an Inkling checkpoint: if the tokenizer change is additive, the GLM oracle still scores token-exact against The remaining one is the maintainer's rule: he needs to run Inkling here before it lands. 975B is not a small ask, so if there's a smaller Thinking Machines checkpoint β or a way to build a tiny one with |
β¦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>
Inkling ships an o200k-family tokenizer (case-aware Split regex, GPT-4o
lineage) rather than cl100k. tok_load now detects the family from the
pattern itself (\p{Lu} appears only in the o200k regex) so GLM behavior
is untouched, and encode dispatches to a new pretok_chunk_o200k that
replays the regex engine's backtracking order exactly: greedy optional
prefix, maximally-greedy uppercase run given back until the lowercase
run can match, contractions attached to letter runs, \p{N}{1,3}, and
the [\r\n/]* punctuation tail.
tok_unicode_o200k.h adds the two range tables the new classes need
(Lu+Lt and Lm+Lo+M), generated from Python unicodedata.
Validated against HF tokenizers on 357 adversarial strings (case
transitions, contractions, CJK, combining marks, emoji + modifiers,
zero-width chars, 300 mixed-charset fuzz cases): 357/357 identical.
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>
|
Totally get wanting to try it yourself before merging. Just wanted to get pen on paper mostly. Made a separate PR #330 for just the tokenizer bit. |
|
The model is interesting and there are many things we can try to bring it into the hummingbird catalog. Now with kimi 3 it would also be interesting. The open world is going wild! |
What
Support for Thinking Machines Inkling (975B total / 41B active MoE, Apache 2.0) β a second full architecture for colibri, following the
olmoe.cstandalone-engine precedent:c/inkling.cβ Stage-A engine: GQA with interleaved sliding-window (512, 16 KV heads) / global (8 KV heads) attention at 5:1, learned relative-position bias banks (no RoPE), log-length tau scaling past 128k, four depthwise causal short-convs per layer (fp32, residual inside, cached conv state for decode), dense + MoE layers, sigmoid router with loss-free bias and jointly-normalized routed+shared weights, mup logit scaling. Routed experts stream per-expert from the fused 3D tensors (LRU cache, colibri int4/int8 container or f32).-p "prompt"streaming greedy mode; ref-oracle harness mode.c/tok.hβ o200k pre-tokenizer support (case-aware Split regex), auto-detected fromtokenizer.jsonso the GLM/cl100k path is byte-identical to before. Newtok_unicode_o200k.hrange tables.c/tools/convert_inkling_int4.pyβ bf16 checkpoint β colibri container (int4 experts +.qsper-row scales, bf16 residents, f32 norms/convs/router). Threaded quantization, shard-resumable,--watchconverts while the download is still running. Note for reviewers: TML's fusedw13tensors store gate/up rows interleaved; the converter de-interleaves exactly pertransformers' conversion mapping.c/tools/make_tiny_inkling.pyβ tiny random-init oracle generator (HF transformers β₯ 5.14).inkling.c: bf16 residents dotted via AVX512-BF16vdpbf16ps(halves resident RAM vs f32 β required: f32 residents would be ~172 GB), int8 expert dot via AVX-512 VNNI (IDOT family;IDOT=0scalar path preserved). Scalar fallbacks keep the dependency-free default CPU path intact.Validation
Token-exact oracle (the colibri standard): tiny random-init Inkling via HF transformers β 36/36 teacher-forced argmax, 24/24 greedy, in f32, int4-container+VNNI, IDOT=0 scalar, and bf16-resident modes.
Tokenizer: 357/357 strings encode identically to HF
tokenizers(case transitions, contractions, CJK, combining marks, emoji+modifiers, 300 mixed-charset fuzz cases).Converter e2e (
--selftest-e2e): fabricates a TML-named checkpoint from the tiny oracle (inverse mapping, re-fused, re-interleaved), converts back, engine reproduces the oracle exactly in both passthrough and int4 modes.Real checkpoint (975B, int4 experts + bf16 residents, 469 GiB snapshot):
Converted weights published at nbeerbower/Inkling-colibri-int4.
Perf status (honest numbers)
Ryzen 9 7900 (24t), 187 GB RAM, NVMe snapshot, commit b60f881,
make CUDA=0, single run, cold cache: RSS 90.5 GB, prefill 150 s (5 tok), decode 0.06 tok/s at cache 16/layer β expert streaming bound, as expected for Stage A with a plain LRU. The architecture is unusually kind to long context (55/66 layers have a fixed 512-token window; 1M-token KV β 46 GB bf16, RAM-resident), so the usual roadmap applies and should pay off more than it did for GLM: hot-store/pinning + prefetch (tier.h), then a GQA CUDA tier β the attention here needs no MLA-style custom kernels.make checkpasses (62 tests),-Wall -Wextraclean, GLM oracle path untouched (tok.h dispatch is additive).Happy to split this into smaller PRs if you prefer (tokenizer / engine / converter).
π€ Generated with Claude Code
https://claude.ai/code/session_016KLDCYJSyHxd39ChTr4T27