Bringing up Kimi-K3 on day zero

Day zero, from inside the runtime

Plow — a Packet Language for On-device Workers, where a worker is a warp on NVIDIA and a wave on AMD — is the compiled inference engine inside Infervisor. Instead of deciding, dispatching, and synchronizing work while a request waits, plow compiles a model ahead of time into one validated execution artifact, and a small Rust runtime — plowrt — runs it. At its core is a switchable SM/CU interpreter: one persistent workgroup per streaming multiprocessor (NVIDIA) or compute unit (AMD) walks a compiled instruction stream, and plowrt selects the interpreter build per target and per phase — prefill, decode, flash — behind one contract. The whole forward pass becomes a single persistent kernel driven by on-device counters, so at serve time the host does almost nothing. The mechanics are on the plow page; this post is about what happens when a genuinely new model arrives and that machinery has to absorb it.

Kimi-K3 is that kind of model. It is not a bigger version of anything we had compiled — it is a hybrid-attention Mixture-of-Experts with several structural departures, each of which produces a plausible but wrong model if you assume it away. "Day-zero support" meant getting it from cannot load to runs correctly at TP8 on eight GPUs with real weights — on AMD MI355X (CDNA4, gfx950), the same 256-CU die as MI350X clocked higher, eight of them behind dual AMD EPYC hosts. K3's weights force the parallelism: at roughly 1.56 TB the model does not fit on one GPU, so TP8 is not a tuning choice, it is the floor to load at all. This is that path — including the bugs that a nearly-right model hides, and the measured cost once it is right.

The model we hadn't run

K3 stacks 93 layers over a 7168-wide residual stream, and the layers are not all the same. Twenty-four of them are MLA — multi-head latent attention, full attention through low-rank q/kv waists — and the other sixty-nine are KDA, Kimi Delta Attention: a linear attention that carries a fixed-size recurrent state instead of a growing KV cache. The MLA layers land roughly every fourth block, with one extra in the tail, so the ratio is not a clean 3:1. Every block then feeds a LatentMoE that routes each token to 16 of 896 experts, plus two shared experts that are always on.

The traps, all of which we hit: K3's MLA uses no RoPE at all — rotary_emb is None, the 64 "rope" dimensions are carried un-rotated, and position lives entirely in the KDA layers — so applying a rotation the model does not want yields silently wrong logits. MLA also adds an output gate before o_proj that K2 did not have. The residual wiring is not the usual add-attn; add-mlp: every layer runs a softmax mix over a running prefix-sum and a stack of snapshots, and every twelfth layer (attn_res_block_size = 12) resets that prefix, so at the other layers the "wrong" plain-residual wiring is close — right up until it isn't. And the experts run at a latent width of 3584, not the model's hidden 7168; sizing them at hidden is a silent 2× overcount. The expert weights ship as mxfp4 (w4a16, one E8M0 scale per 32 values, bias 127 — a decode we confirmed against the checkpoint's own scale histogram). None of these crash. They just quietly produce a model that is almost right, which is the hardest kind of bug to gate on.

A compiler that proves its own rewrites

Before any of that reaches a GPU, the model is lowered to a shape-inferred operator graph and handed to an e-graph. Plow uses egglog equality saturation: it represents the graph and a set of 19 shape-preserving fusion rules as one e-graph, saturates it, and extracts the lowest-cost equivalent form. K3 exercised two rules written for exactly this architecture — kda-gated-norm-fuse (five e-nodes to one, emitting KdaGatedNorm) and mla-out-gate-fuse (three to one, emitting MlaOutGate) — alongside the usual norm→linear, gated-MLP, and residual→norm fusions.

A fusion that is fast is worthless if it is not legal, so every rule is gated by a machine-checked proof. Plow carries a Lean development with six checkpoints, A through F, each certifying one stage of the pipeline. Checkpoint A is the one that binds the compiler to the proofs: the exact rule names egglog reports as fired are submitted to Lean, which rejects any name not in its set of proven-sound rules. A rule cannot enter rules.egg and quietly take effect — it must also gain a soundness theorem, or checkpoint A fails the build.

Extracting the cheapest form is not free on a model this deep — a naively printed 93-layer DAG un-shares into an exponential blow-up — so extraction walks the hash-consed term graph directly rather than materializing an s-expression. By the time a plan reaches the runtime, its shape is fixed and its legality is settled. What Lean does not prove is that the emitter wired up every op the model actually needs — and that gap is exactly where this bring-up went wrong before it went right.

One kernel, no host in the loop

The compiled plan does not run as a sequence of launched kernels. Plow expands it, once, into a device instruction stream and runs the entire forward pass as a single persistent kernel whose grid is sized to the GPU's compute-unit count and which stays resident for the model's life. There is a deliberate two-format split: the on-disk wire packet stream is variable-length and decoded on the host exactly once; the on-device ISA is a fixed 64-byte instruction, so the interpreter's inner loop is an indexed load and a switch — nothing is parsed on the GPU. For K3 the interpreter is specialised to the opcodes the model uses, which cuts its .text roughly in half; it ships as 49 code objects.

Dependencies are expressed as atomic counters and resolved entirely on-device. Each workgroup owns a compute unit and walks its own stream of instruction-and-slice entries; for each one it spins until its wait-counters reach a threshold equal to the number of producing workgroups, executes only its slice, takes a single agent-scope fence, and atomically bumps its successors. "All producers are done" is not a barrier — it falls out of the counter reaching its threshold. The counters are spaced a cache line apart so signalling one doesn't ping-pong a line across the eight L2 domains.

A fair objection to compiling the whole model into one kernel is that consecutive packets jump to op bodies scattered across the binary and thrash the instruction cache. We measured it instead of arguing it: on the real K3 decode interpreter the L1 instruction-cache hit rate is 99.88%. Most of the binary is never fetched — only taken paths are, and the op bodies are loops — so the hot instruction set is small against gfx950's L1I. gfx950 has no instruction-prefetch instruction to lean on, so we bounded the loss from above instead: forcing a cold fetch before every packet, the worst case a perfect cache could ever recover is 2.2% of the token. The megakernel does not pay for its size in instruction bandwidth.

This is also where a whole class of "silently does nothing" bugs lives. One K3 op had no arm in the AMD interpreter's switch, so it fell through to the default case and computed nothing — finite, plausible, wrong. An unclaimed weight is just as invisible: if nothing declares a tensor, its absence can never surface as a missing-weight error, and a MoE whose experts never bind will still emit fluent, believable logits. Both shapes of silence bit us on K3, and the next two sections are where they were caught.

Per-XCD queues, and the counter problem underneath them

If the token is a chain of counter handoffs, the price of one handoff sets the floor for everything. So we measured it in isolation: one empty b=256 packet, first workgroup arriving to last workgroup finishing. The atomic is never the cost — the release read-modify-write is 0.07–0.14 µs in every variant we tried. What costs is the cache maintenance around it. All 256 workgroups issue a writeback and an invalidate, and on CDNA4 those are per-L2: MI355X has eight L2 partitions, one per XCD (accelerator complex die), so each partition performs the same writeback and the same invalidate 32 times over, and they serialise.

Electing one workgroup per XCD to do the maintenance on everyone's behalf is a 3.8× cheaper handoff — and the reason it wasn't already there is a scheduling fact, not a kernel one. The interpreter's default work queue is global: a workgroup claims whatever stream entry is next, so which XCD executes which slices of a packet is decided at run time. A per-XCD rendezvous needs the opposite — it must know, at emit time, how many workgroups of packet p will arrive on domain d, because that count is the rendezvous threshold. The fix is per-XCD queues: the compiler stable-sorts the stream by L2 domain and emits one window per domain; all eight windows drain concurrently inside a single launch, each workgroup serving the domain it is physically on — read from the hardware's XCD id register, not inferred from its block index — with a per-domain cursor so work-stealing still survives within a domain. Under that dispatch mode the arrival count becomes a static constant, and each domain's leader bumps the global counter by exactly its own share, so no consumer's threshold changes and the protocol's meaning is untouched.

The ordering is the delicate part, and it is where the saving mostly lives. Publishing is a local relaxed add per follower after its __syncthreads() has retired its stores into this XCD's L2; the workgroup whose add comes back last is the leader, and its single writeback flushes the whole partition — which is precisely the followers' stores. Observing is the mirror: the first consumer workgroup on a domain invalidates once and opens a local flag, and the others never touch the global counter at all, which is less fabric traffic, not more. Neither rendezvous needs a fence, because every participant shares one L2 partition. Measured on the real model, the sound version is −4.02 ms (−12.1%) against the unsafe ceiling's −4.91 ms, so it captures 82% of what deleting the ordering would buy — and the missing 18% is the ordering, which is the right shape for that gap to have. It is not free: three XCD-local counter arrays sized ops × 8 triple the per-token counter re-arm (315 KB → 1.26 MB).

Three traps are worth recording, because two of them were measurement bugs rather than correctness bugs, and an audit found them rather than a benchmark. Reading the XCD id register per call puts it on the gate and on every successor signal — twice per packet, across 324,940 workgroup-packets in a token — and it cost 0.43 ms until it was hoisted to kernel entry. An earlier revision gated the cross-GPU system-scope acquire by the same leader test; K3's collectives run at b=14, so the expected number of elected leaders on a collective was 0.44 — nearly every XGMI acquire was being deleted, on a TP8 run. And a blob emitted for the hierarchy, run under a plowrt built before the hierarchy's base pointer existed, faulted the GPU: that pointer lives in an alignment pad of the kernel argument block which the older host left uninitialised. The correctness gate caught it; a benchmark would have reported a number.

The wave/warp switch

Porting the interpreter between vendors comes down to one abstraction: the worker. On NVIDIA a warp is 32 lanes; on AMD a wave is 64. Plow keeps eight sub-groups either way — a 256-thread block of 8 warps on NVIDIA, a 512-thread workgroup of 8 waves on AMD — so most kernel logic is identical and a cross-lane reduction simply begins one shuffle step earlier on AMD (offset 32 rather than 16). The port is mechanical because the shape is preserved; it is the register budget and the dispatch model underneath that are not.

That segment is the mechanism that answers the vendor difference. On NVIDIA the plan runs as one cooperative launch. On AMD it runs as n_seg launches and a single drain: the stream is partitioned into wave-class segments, and the host relaunches the interpreter once per segment. The trick that keeps this from costing a host round-trip is that consecutive segments are chained by the AQL barrier bit in the dispatch header, and the counters are zeroed once per dispatch group, never per segment — so a producer that ran in an earlier launch still satisfies a consumer in a later one. Zero the counters between segments and the next segment waits forever; drop the barrier bit and the chain is silently racy. Both traps are load-bearing, and both were carried across from the reference driver with their reasons attached. (The per-op-segment variant of this, PLOW_SEG_PER_OP, still desynchronises ranks at TP8 and is off — one of the knobs listed at the end.)

Eight GPUs, one all-reduce per sublayer

K3 runs tensor-parallel across all eight MI355X. The sharding is Megatron-style: q, kv, gate, up, expert-up and the vocabulary head split column-wise; o_proj, down and expert-down split row-wise; the low-rank latent waists, the router, norms and embeddings stay replicated because they are too small to be worth cutting. Pairing a column-parallel producer with a row-parallel consumer means each sublayer needs exactly one all-reduce.

And that all-reduce is not a launched collective. In plow a collective is just another entry in the resident interpreter stream, gated by system-scope counters mapped into peer memory over the xGMI mesh — so there is no per-collective kernel launch and no host synchronization inside a token. A conventional RCCL/NCCL path pays a launch per collective; at hundreds of collectives per token that is real latency, and inlining them removes it. It is also, exactly because the reduce is a hand-placed stream entry indexed by peer slot, where the sharpest bugs on this branch hid.

A second sharding target is the vocabulary head — a 163,840 × 7,168 table that, replicated, has every rank stream the same ~2.35 GB per token, seven-eighths of it redundant at TP8. Making it column-parallel gives each rank 20,480 of the vocabulary rows, but it changes the sampler: once a rank holds only its slice of the logits, a per-rank argmax would have the ranks disagree on the token, so the on-device argmax becomes a cross-rank fold that rebases each rank's local winner by its vocabulary offset and takes the global maximum. This one (PLOW_K3_SHARD_HEAD) sat off for a while because the equivalence harness could not check it: at TP8 the dumped logits are vocab/tp wide, so the shape comparison fails by construction. That turned out to be the wrong gate — token identity does not need the logits dump at all, since each rank argmaxes its own slice and the cross-rank fold reduces to the same global winner, so the emitted token is the object to compare. It matches, and the knob now ships on.

Its size is the lesson. Sharding the head removes 2.055 of 14.021 GB per rank per token — 14.7% of the GEMV bytes — and bought 0.379 ms, 1.14%. We had predicted 2.4 ms. The reason for the 6× miss is that lm_head was already running at ~94% of peak bandwidth, on the very same GEMV kernel that reaches 23% on the projections: its bytes were the cheapest bytes in the model, so deleting a seventh of them deleted a hundredth of the time. A byte-reduction lever is worth bytes ÷ the achieved bandwidth of those bytes, never bytes ÷ peak — and, as the last section shows, only where those bytes already dominate the per-packet floor.

Correctness before latency

This architecture specialises in failures that read as success: stable, believable logits from a model quietly computing the wrong thing. The unbound-expert MoE, the op that fell through the switch, the RoPE the model does not want, the collective reading the wrong peer slot — every one of them produces finite output. A full-model harness will confidently report a latency for such a model, which is exactly why block-level "rung" gates are not enough: block-correct is not model-correct.

Two controls caught them. The first is an all-eight-ranks-token-identical gate on the real asset: with a bug, ranks disagree; without one, they are bit-identical at every prefill and decode step — verified on both the decode-only build and the full 93-layer asset. (K3 cannot be run at TP1 for a tp1-vs-tp8 control — it does not fit on one GPU — so the cross-rank identity is the model-level control, backed by the same TP path proven token-identical across TP1/2/4 on models that do fit.) The second is a block oracle: the MLA gated block checked against a reference implementation came back NoPE bit-exact, with MoE routing selecting the same expert set (the within-set order is an unresolvable bf16 tie, and the combine is order-invariant) and stage RMS relative errors around 1e-3. A structural coverage gate now also fails the build when an emitted op or a checkpoint weight goes unread — the direct defence against the two silences.

And beyond bit-level checks, we tested for accuracy, because every gate above proves the runtime is self-consistent and none of them proves the model is right — a blob that is wrong the same way on all eight ranks passes all of them. We ran a task-level accuracy evaluation end-to-end through plowrt's served OpenAI endpoint — exercising the chat template, K3's stop channel, the tokenizer, chunked prefill over ~1k-token prompts, on-device sampling and SSE — and the model answered at the accuracy expected of K3.

Only with those gates green did we let the clock start. Everything below is measured on that model.

The 28.9 ms, taken apart

Here is the honest headline. Batch-1 decode on the correct model, TP8, ctx 32k: 28.876 ms/token — 34.6 tok/s, token-identical to the control, with the counter hierarchy, per-XCD queues and a column-parallel lm_head all on. The bring-up baseline was 40.6 ms and the pre-hierarchy stack 33.1 ms. The interesting part is not the number, it is its anatomy, because the number is not where a decode-is-memory-bound intuition would put it.

The kernel-level figure and the served figure are independent instruments that share no code above the engine, and they agree to 3% — 28.876 ms from the bench at ctx 32k against 29.82 ms TPOT from the OpenAI endpoint at 30k prompt tokens. That agreement is the reason we quote either.

The decode dependency graph is nearly a chain: the critical path is 1,739 of the 2,459 packets in a token, with an average of 1.41 packets per level and never more than five counters live at once, so there is almost nothing to overlap. An empty b=256 packet cost about 5.72 µs to hand off, so 1,739 hops floored the protocol at 9.95 ms/token — the entire 10 ms target, spent before a weight is touched. The hierarchy takes that floor to 6.02 ms; it does not remove it. K3 decode is memory-bound in principle and nowhere near memory-bound in practice: its wall is the depth of the on-device dependency chain, and the per-packet work now sits close enough to the floor that chain length is the term that matters.

A per-opcode rate attribution says the same thing from the other side, and it overturned our own ranking. GEMVs are 30% of the body time at 26% of peak bandwidth, which reads like a kernel problem — until you compare each shape's bytes against the packet floor: 14 of 17 GEMV shapes move bytes that take less time than an empty packet costs. Their low percent-of-peak is the protocol showing through, not a bandwidth failure, and only 0.683 ms of the whole GEMV term is recoverable by moving fewer or faster bytes. The body half was protocol wearing a roofline disguise. Which is why the next structural lever is batching: at batch B the same weight bytes and the same packet floor both serve B tokens, so it divides both terms of a protocol-bound token — and on K3 it is blocked on exactly one thing, a KDA recurrent state with no batch axis, which is in progress.

The roofline: what MI355X allows, and why it isn't the wall

It is still worth drawing the floor the bytes impose, because it explains why the answer is not in the matrix cores. Batch-1 decode is memory-bandwidth-bound: every generated token streams the active weights (plus KV and recurrent state) through the GPU once, and the arithmetic per byte is far too low to hide behind compute. So the bandwidth floor is bytes-touched-per-token ÷ memory bandwidth, and the compute peak barely enters.

A decode GEMV reads one weight and does one multiply-add — two FLOPs — per parameter, so its arithmetic intensity is just 2 ÷ bytes-per-weight: 1 FLOP/byte in bf16, 2 in fp8, 4 in mxfp4. The ridge points of MI355X — where a kernel tips from memory-bound to compute-bound — sit at roughly 400, 800 and 1,600 FLOP/byte for those precisions. Decode is two to three orders of magnitude to the left of that: the matrix cores run at about a quarter of one percent of peak.

So the roofline is the honest yardstick for the weight-streaming part of the token, and it says that part is ~2.76 ms. It is deliberately not a promise about the whole token: it assumes every byte streams at full bandwidth and nothing else costs time, which no real decode achieves — least of all one whose dependency graph is a near-chain. The gap between 2.76 ms and 28.9 ms is the protocol, and closing it is measured work, not a roofline argument.

What day-zero does — and doesn't — mean

• The numbers are measured, and unflattering.
• Prefill is real but partial. Time-to-first-token is linear in prompt length — ~0.63 s per 1,000 tokens, flat across a 234× range, 19.3 s at a 30k prompt — so the chunker is doing its job and nothing is quadratic. But tiling the MLA prefill onto the matrix cores is 2.25–2.79× on that kernel and only ~20% end-to-end, because it covers 24 layers and the other 69 are KDA, whose chunked prefill is not written yet and which currently run prefill as N decodes. Anyone quoting the kernel figure as a model-level one is quoting the wrong number — including us, until we measured it.
• Batch 1 throughout, structurally. The KDA recurrent state has no batch axis, so the engine refuses batched decode at load rather than aliasing one sequence's state onto another's. A slot-indexed, snapshottable state pool is the one piece of work that opens both multi-user batching and speculative decoding; it is in progress and inert behind a gate until it is correct. Multi-user serving numbers do not exist yet, and a concurrency column today would measure queueing, not batching.
• K3 is currently a gfx950-only model in plow. The NVIDIA path does not yet carry the state-stepping kernel the KDA layers need.
• On comparisons. On the same 8× MI355X, AMD's own ATOM stack measures 53.35 ms/token TPOT (18.7 tok/s) at concurrency 1 on the vendor's published recipe; plow's 28.9 ms is faster at that point, but it is a like-for-like conc-1 latency point and not a throughput claim, and neither stack is near where this model should be.

The reason a new architecture could be absorbed in this shape is the shape itself: a validated, fixed plan; one persistent kernel; counters instead of a host in the loop; and a compiler that will not let a rewrite fire without a proof. When the model changed, most of the work was teaching the compiler the new blocks — the runtime contract underneath did not move. And the first week of optimisation on top of it went where the instrument pointed rather than where intuition did: the counter hierarchy over per-XCD queues was worth 4 ms, sharding the largest weight in the model was worth 0.4, and the whole "bodies are bandwidth-bound" premise turned out to be the protocol in disguise. The rest of the honesty is the measurement: 34.6 tok/s is not a result to be proud of yet — it is a correct, instrumented, token-identical starting line, with the wall named (a 1,739-deep chain) and the next move known (shorten it, then divide it by batching). That is what we mean by day-zero support.