Skip to main content
Back to timeline
vLLM BlogSource publication:

vLLM serves Thinking Machines' Inkling 2 Small on day one: 594B total parameters, 17B active, about 630 tok/s for a single user

Synopsis

Thinking Machines' Inkling 2 Small gained vLLM support on its release day. It is a text-and-image-in, text-out mixture-of-experts model with 594B total and 17B active parameters, up to 1M tokens of context, and a bundled DSpark speculative-decoding drafter. vLLM fits it with kernel fusion, programmatic dependent launch, sequence parallelism and a Gluon attention kernel rewritten for the 6:1 GQA ratio, reaching about 630 tok/s for a single user at TP4 on GB200 with acceptance length 3.5, and also supports LoRA and reinforcement learning via SkyRL.

Source-provided article image: Inkling 2 Small on vLLM: Day-0 Support
vLLM

Interpretation

vLLM can serve Inkling 2 Small on the day it ships, including the DSpark speculative-decoding drafter that comes with the checkpoint. Inkling 2 Small changes the backbone relative to the earlier Inkling: layers drop from 66 to 48, the sliding window grows from 512 to 1,024, the MoE moves from 256 routed experts (top-6) plus 2 shared experts to 384 routed experts (top-8) with every layer an MoE layer, and the image encoder changes from a 4-layer hMLP to a 23-layer ViT. Served at TP4 on GB200, a single user decodes at about 630 tok/s with acceptance length 3.5; on the TP8 agentic benchmark at 32K session context, each GPU delivers about 1,070 tok/s at roughly 84 concurrent sessions and a 91% prefix-cache hit rate.

The adaptation hinges on merging the repeated, scattered memory-bound steps in decode and prefill and overlapping the launch of adjacent kernels. Block attention residuals make every sublayer boundary perform an all-reduce, a short convolution, a residual update, the attention-residual mix and an RMSNorm, five steps that all read and write the same hidden states. The fused allreduce_conv_mix kernel runs the all-reduce over NVLink multicast and keeps everything after it in registers; the GEMM, attention, routing and activation kernels use programmatic dependent launch (PDL) so each kernel's prologue overlaps the tail of the previous one, which matters at small batch sizes where a decode step is a long chain of short kernels.

Prefill uses sequence parallelism to remove work that plain tensor parallelism repeats on every GPU. Under plain tensor parallelism, after each all-reduce every GPU runs the short convolution, attention-residual mix and RMSNorm on all tokens, which is cheap in decode but costly in prefill. Each GPU reduce-scatters, runs the short convolution, residual update, attention-residual mix and RMSNorm in one fused conv_mix kernel on only its slice of tokens, then all-gathers the result for the next GEMM; an 8K-token prefill step shrinks by 22% at TP4 and 26% at TP8, freeing 1.3–1.9 GiB per GPU for KV cache.

The decode path gets a separate Gluon kernel for relative attention because the existing FlashAttention-4 decode path is slow for this model. Inkling 2 Small's 6:1 GQA ratio silently disables GQA packing, which slows Inkling's FlashAttention-4 decode path. The new warp-specialized Gluon kernel packs the 6 heads into one 64-row tile and reads KV once: a TMA warp streams KV, a tcgen05 warp runs the MMAs and four softmax warps apply the relative bias; a batch-1 decode step at 100K context drops from 15.1 ms to 4.8 ms, while prefill keeps Inkling's sheared-bias FlashAttention-4 path.

Perspective

The demonstrated scope is vLLM serving Inkling 2 Small on GB200 at TP4 and TP8, with performance numbers from an agentic workload using DSpark probabilistic drafting and prefix caching, where each turn adds about 3K new tokens, generates 1,024 tokens at temperature 1.0 with reasoning on. Accuracy checks cover NIAH up to 1M tokens, MMMU-Pro and HLE, compared against runs without speculative decoding while DSpark is enabled.

The accuracy check only shows that scores with DSpark enabled match runs without speculative decoding; it does not give absolute scores aligned item by item with the reference implementation. The 22% and 26% prefill speedups apply to an 8K-token prefill step at TP4 and TP8 respectively, and no numbers are given for other parallelism degrees or sequence lengths.