Skip to content

Inkling (Thinking Machines 975B MoE): new engine, o200k tokenizer, bf16->int4 converter#312

Open
nbeerbower wants to merge 15 commits into
JustVugg:devfrom
nbeerbower:inkling
Open

Inkling (Thinking Machines 975B MoE): new engine, o200k tokenizer, bf16->int4 converter#312
nbeerbower wants to merge 15 commits into
JustVugg:devfrom
nbeerbower:inkling

Conversation

@nbeerbower

Copy link
Copy Markdown

What

Support for Thinking Machines Inkling (975B total / 41B active MoE, Apache 2.0) β€” a second full architecture for colibri, following the olmoe.c standalone-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 from tokenizer.json so the GLM/cl100k path is byte-identical to before. New tok_unicode_o200k.h range tables.
  • c/tools/convert_inkling_int4.py β€” bf16 checkpoint β†’ colibri container (int4 experts + .qs per-row scales, bf16 residents, f32 norms/convs/router). Threaded quantization, shard-resumable, --watch converts while the download is still running. Note for reviewers: TML's fused w13 tensors store gate/up rows interleaved; the converter de-interleaves exactly per transformers' conversion mapping.
  • c/tools/make_tiny_inkling.py β€” tiny random-init oracle generator (HF transformers β‰₯ 5.14).
  • Zen-4-friendly kernels in inkling.c: bf16 residents dotted via AVX512-BF16 vdpbf16ps (halves resident RAM vs f32 β€” required: f32 residents would be ~172 GB), int8 expert dot via AVX-512 VNNI (IDOT family; IDOT=0 scalar 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):

    $ SNAP=~/Models/inkling_i4 ./inkling -p "The capital of France is" -n 24
    The capital of France is Paris.
    The capital of the UK is London.
    The capital of Russia is Moscow.
    The capital of Italy
    

    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 check passes (62 tests), -Wall -Wextra clean, 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

@nbeerbower

Copy link
Copy Markdown
Author

Pushed a Stage-B performance commit (3ef6775) β€” same architecture, the expert path grew the glm.c playbook:

  • cache slots keep container rows packed (int4 stays 4-bit in RAM, ~28 MB/expert β†’ double the cache per GB); new matmul_q4 unpacks nibbles in-register, numerically identical to the int8 path
  • moe() routes everything first, then fills all missing experts in one parallel burst (prefill batches the whole sequence's misses, so the NVMe finally sees queue depth), then computes
  • usage-history pinning (SNAP/.coli_usage, top PIN_N/layer non-evictable) + RAM-sized auto cap

Measured on the 975B checkpoint (Ryzen 9 7900, NVMe, cap auto=64/layer, 1024 pinned):

Stage A Stage B cold Stage B pinned+warm
prefill (5 tok) 150.2 s 31.5 s 21.7 s
decode 0.06 tok/s 0.21 tok/s 0.25 tok/s
cache hit ~0% 72.0% 81.5%

All four oracle modes (f32, int4-container VNNI, IDOT=0 scalar, bf16 residents) remain token-exact, make check 68/68, 0 warnings.

πŸ€– Generated with Claude Code

https://claude.ai/code/session_016KLDCYJSyHxd39ChTr4T27

@nbeerbower

Copy link
Copy Markdown
Author

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):

  • CUDA tier (backend_cuda_ink.cu, CUDA=1): bf16 residents (~36 GB) live in VRAM on a warp-per-row kernel with the same numeric contract as the CPU vdpbf16ps path; freed RAM goes to the expert cache. Pure-CPU default build unchanged and dependency-free.
  • Cache warming (glm's .coli_usage idea): per-expert usage counts persist and accumulate; top PIN_N/layer pinned at startup. Fully toggleable (PIN=off, PIN_N, USAGE_SAVE=0). The -f promptfile mode exists to build the history for a binary with no serve mode yet β€” happy to drop it if it reads as scope.
  • docs/inkling.md: quickstart, env reference, benchmark methodology.

Measured on the 975B (Ryzen 9 7900 / A6000 / NVMe, commit 87981c1, greedy):

Configuration Prefill Decode Hit
Stage A 150.2 s 0.06 tok/s ~0%
CPU perf pass 21.1 s 0.25 tok/s 81.5%
+ CUDA residents 18.4 s 0.32 tok/s 83.6%
trained-prompt pins (best case) 1.9 s 2.51 tok/s 100%
steady state (novel prompt, diverse history) 35.4 s 0.25 tok/s 82.2%

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); make check 68/68, 0 warnings on both builds.

πŸ€– Generated with Claude Code

https://claude.ai/code/session_016KLDCYJSyHxd39ChTr4T27

@JustVugg

Copy link
Copy Markdown
Owner

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:

  1. A model he can actually run, or a tiny one the CI can execute end-to-end. This repo already has the pattern: make_glm_oracle.py builds a glm_tiny/ and the engine is scored token-exact against it. An engine with an oracle test is an engine that can't silently rot.
  2. A named maintainer β€” you, on the record, for that file.
  3. Isolation from glm.c's path. This is the part I'd check hardest, and it's specific per PR (see below).

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.

@JustVugg

Copy link
Copy Markdown
Owner

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.

inkling.c (+1,255) is isolated and that's the right shape β€” a separate file, like olmoe.c. backend_cuda_ink.{cu,h} likewise. If the whole thing were that, my only objection would be "nobody can run it".

The part I'd check hardest is c/tok.h (+115/-1) β€” 118 changed lines in the tokenizer glm.c depends on. That's the one place this PR can break the flagship engine, and it's the reason I'd want a real review rather than a wave-through. If the o200k support is purely additive (new tables, new branches, glm's path untouched byte-for-byte), it's defensible. If it refactors anything shared, that needs to be visible.

Concretely, that's testable without an Inkling checkpoint: if the tokenizer change is additive, the GLM oracle still scores token-exact against glm_tiny/, and tests/test_stops.c (which loads a synthetic tokenizer and checks the special flags) still passes. If you can show those green with your tok.h, you've retired the only risk I can actually name.

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 convert_inkling_int4.py, the way make_glm_oracle.py builds glm_tiny/ β€” that's the path, and it would also give the engine an oracle test so it can't rot silently.

nbeerbower and others added 15 commits July 16, 2026 14:32
…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>
@nbeerbower

Copy link
Copy Markdown
Author

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.

@JustVugg

Copy link
Copy Markdown
Owner

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!

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants