Skip to content

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

Merged
JustVugg merged 15 commits into
JustVugg:devfrom
nbeerbower:inkling
Jul 23, 2026
Merged

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

Conversation

@nbeerbower

Copy link
Copy Markdown
Contributor

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
Contributor 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
Contributor 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

Copy link
Copy Markdown
Contributor 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!

@JustVugg

Copy link
Copy Markdown
Owner

@nbeerbower this is a serious piece of work and a second full engine is exactly where the project wants to go β€” thank you. One thing before merge, and it's the same bar the GLM engine holds itself to: a tiny test model + oracle check we can run in CI and locally.

For GLM we keep glm_tiny (a few-MB randomly-initialized model in the real container format) and the engine's TF=1 teacher-forcing mode prints N/32 vs oracle β€” token-exactness against a transformers reference. That's what lets us merge refactors all day without fear. Since none of the maintainers can run the real Inkling 975B, without a tiny model your engine would be effectively untestable by anyone but you β€” and the first refactor that touches shared code (tok.h, Makefile) could silently break it.

I see you already ship tools/make_tiny_inkling.py β€” could you: (1) generate + commit (or attach) a tiny Inkling snapshot, (2) wire a TF=1-style oracle check (32-token teacher-forcing vs HF transformers on the same tiny model), and (3) add the one-liner to CI? With that in place this merges as-is β€” the code is additive and CI is already green. Happy to help with the CI wiring.

@JustVugg

Copy link
Copy Markdown
Owner

Status check, and sorry for the wait. This conflicts with dev now β€” the engine was renamed and split in #391 (glm.c β†’ colibri.c plus header modules) and c/Makefile and c/tok.h have both moved since. Are you still working on it?

On the substance, the structural choice here is right and worth saying out loud: inkling.c as a separate engine alongside colibri.c, sharing the container and tokenizer infrastructure but not entangling the GLM path, is exactly how a second model family should arrive. It's the same shape as the DeepSeek V4 work in #165, and it's why both are reviewable at all at their size.

Two things before I can take it seriously as a merge candidate:

  1. A tiny-model oracle. make_tiny_inkling.py is there, but what settles a new engine in this project is a token-exact check against a transformers reference on a tiny checkpoint that lives in the repo β€” the way glm_tiny does. If tools/make_tiny_inkling.py already produces that and there's a teacher-forcing comparison, point me at the command; if not, that's the piece I'd want most.
  2. CI must build it. A second engine that isn't compiled by make check rots within weeks. Add DeepSeek V4 Flash CPU inference with NVMe expert streamingΒ #165 wires its build and tiny-oracle into the check target; the same approach would work here.

If you're still on it, rebase and tell me about the oracle. If you've moved on, say so and I'll close it as groundwork with credit rather than leaving it open β€” no hard feelings either way, and the design note above stands regardless.

nbeerbower and others added 15 commits July 22, 2026 20:22
…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
@JustVugg
JustVugg merged commit e4b5bf3 into JustVugg:dev Jul 23, 2026
10 checks passed
@nbeerbower

Copy link
Copy Markdown
Contributor Author

Still on it — rebased onto dev and pushed, so it's mergeable again. The glm.c→colibri.c split was a clean rebase: inkling.c / backend_cuda_ink.* are isolated new files, and the tokenizer commit dropped out entirely since o200k is already upstream (that was #330), so tok.h is now byte-identical to dev. The only real re-apply was the Makefile inkling target onto the new $(GPUCC) build.

On the two asks β€” both are in the PR now and green in CI:

1. Tiny-model oracle. It's the glm_tiny shape exactly. tools/make_tiny_inkling.py builds a small random-init InklingForCausalLM via HF transformers (8 layers, hitting every branch: sliding + global attention, log-length tau, dense + sparse MLP, 2 shared experts, unpadded vocab) and writes ref_inkling.json. The engine reproduces it token-for-token:

python3 tools/make_tiny_inkling.py tiny_inkling
SNAP=tiny_inkling ./inkling 8 0 tiny_inkling/ref_inkling.json
#  teacher-forced argmax: 36/36 match
#  Matching tokens: 24/24

bits=0 keeps experts f32 so the forward is bit-exact; the greedy pass exits non-zero on any mismatch, so it's a real gate, not a print. Green in f32, int4-container+VNNI, IDOT=0 scalar, and bf16-resident modes.

2. CI builds it. The Engine (Linux, CPU) job now make colibri inkling, the CUDA-syntax job also nvcc -c backend_cuda_ink.cu, and there's a new inkling-oracle job that pip-installs CPU torch + transformers>=5.14 and runs exactly the two commands above. All three pass on this push. I kept the oracle out of the dependency-free make check gate on purpose β€” that one runs on Windows/macOS where I haven't confirmed the engine builds yet, so it's a Linux CI job instead. Happy to fold it into make check cross-OS as a follow-up once I've verified the Windows/macOS builds.

One honest scope note: the engine has grown well past this PR in my fork (Schneewolf-Labs/sabrewing) — GPU expert tier, int8 residents, MTP draft/verify, LoRA serving, per-phase profiling. I'd suggest landing this Stage-A→CUDA-warming version + oracle + CI first (reviewable at size, the shape you described), and I'll bring the perf/CUDA/MTP work as follow-up PRs rather than dropping another 10k lines at once. Sound right?

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