# DeepInfer: DeepSeek's serving engine, and a throughput chart that multiplies by 3.25

> Satyajit Ghana — Head of Engineering @ Inkers Technology
> canonical: https://ai.thesatyajit.com/articles/deepinfer
> date: 2026-10-23
> tags: inference, deepseek, kv-cache, speculative-decoding, mixture-of-experts, benchmarks, vllm

On 20 October DeepSeek posted a four-part thread: "Introducing DeepInfer, our in-house inference
framework, now fully open source! Powering DeepSeek-V4.1-Flash at scale with high throughput and low
latency." The second post lists "High-performance kernels & optimized inference engine", "Distributed
inference & speculative decoding" and "SWA Bounded Replay & structured outputs". The third adds
**deepseek-recipe**, "our production token encoding & decoding component".

An hour later a community chart went round. It took SemiAnalysis's InferenceMAX plot for
DeepSeek-V4.1-Flash, where the best configuration tops out near 6,000 output tokens a second per chip,
and drew an arrow off the top of the page to "25600 TPS/GPU" for DeepInfer. One reply said the real gain
was 1.3x to 1.5x. Another asked to see the maths.

So I read the code, ran what runs on a CPU, and traced every number in that chart to its source. Labels
throughout: **measured** means I computed it from a file or a run, **reported** is DeepSeek's or
SemiAnalysis's figure taken as given, **reasoned** is my arithmetic on the other two.

<RepoCard repo="deepseek-ai/DeepInfer" />

## What shipped, and what didn't

The README is unusually direct about this. "This open-source release provides offline inference with the
core features needed for serving at scale." Then, under "Required for a complete inference service (not
included in this release)": a KV cache system, "VE/Prefill/Decode disaggregation (most required engine
functionality is already implemented)", and load balancing across data-parallel replicas.

The code agrees. `LLMEngine`'s docstring says the `LLM` class wraps it "for offline batched inference.
This release does not include an online-serving wrapper." The hooks that would save KV to the store or
hand it to a decode machine (`init_pd_split`, `pe_prepare_save_kvcache`, `pe_do_save_kvcache`) sit under a
comment that translates as "from here on these are all virtual functions", and each body is a docstring.
What does ship is the engine that runs on either side of that split: `EngineType` has `NORMAL`,
`PROMPT` (prefill only) and `COMPLETION` (decode only), and the scheduler handles all three (measured).

The repository is 348 tracked files: 270 Python files holding 54,929 lines, and 5,908 lines of C++ in 15
`.cpp` and 18 `.hpp` files. There are no `.cu` files at all (measured). Twenty-three Python files open with
`# Derived from vLLM; modified by DeepSeek.`, and the acknowledgements say it "was originally built
primarily on vLLM, with additional inspiration from SGLang and TensorRT-LLM" (measured, reported). The
licence is Apache-2.0; deepseek-recipe is MIT (measured).

What fills the gap is `doc/inference_deployment_qa.md`, a long Q&A describing the whole service. Almost
everything below about storage and balancing comes from it, so it is reported.

## The system the engine was cut from

<Figure src="https://media.thesatyajit.com/ai/articles/deepinfer/fig1.png?v=dc74df2179d14d51" alt="Architecture diagram: user JSON enters DeepSeek Recipe, which sends a token list and sampling parameters to a stack of Entrypoints. The Entrypoint routes to a Prefill Load Balancer, which feeds a stack of Prefill Service instances each with an Expert-Parallel Load Balancer. A Vision Encoder Load Balancer and Vision Encoder Service feed the prefill side. Prefill sends KV cache to a Decode Service stack behind a Decode Load Balancer, and decode returns tokens to the Entrypoint. An External KVCache Storage (3FS) sits under both, read by prefill and written by decode." caption="The DeepSeek-V4.1-Flash serving system: recipe, entrypoint, separate vision-encoder, prefill and decode services each behind its own load balancer, and a 3FS cluster holding the KV cache. Black arrows carry requests, blue arrows KV cache (DeepInfer repository, doc/images/architecture_overview.png)." />

The shape is the one DeepSeek described for V3: separate prefill and decode, attention in data parallel,
the MoE in expert parallel with as many EP ranks as DP ranks. What changed is the numbers (reported):

- **Prefill** runs EP2 to EP8, "the smallest EP/DP that HBM allows". **Decode** runs EP64 to EP128,
  "the largest EP/DP that the available GPUs allow". No tensor parallelism.
- Prefill and decode GPUs come in roughly a 1:1 ratio, so "the smallest efficient deployment unit is 64
  GPUs of prefill plus 64 GPUs of decode".
- The KV cache lives on a dedicated 3FS cluster of about 100 PiB, kept 3 days for agent traffic and 7 for
  chatbot traffic, and decode writes its KV back too. Any prefill machine can hit any prefix, so the
  prefill balancer does no prefix-affinity routing.
- "A single cluster currently serves 155 trillion tokens per day, with overall revenue of US\$38 per
  GPU-hour (at valley pricing)." Neither number can be checked from the repository.

Here is one request through it, with what each stage holds. Solid boxes are code you can read today;
dashed ones are prose.

<RequestLifecycle />

In words: deepseek-recipe turns the API request into tokens. The entrypoint matches the prefix against
store metadata and computes the hit before routing. The prefill balancer buckets requests into four
context tiers (under 16K, 16K to 64K, 64K to 256K, 256K and up). A prefill engine swaps the cached KV in
layer by layer over GPUDirect RDMA while its forward runs, and computes the rest. The KV and the
sliding-window state move to a decode machine. The decode balancer places the request to even out both
request count and total KV length. The decode engine drafts, verifies and streams tokens back.

## Inside the engine: C++ schedules, Python launches

The control loop is Python (`LLMEngine.step`), but the scheduler, sequence state, block manager,
input-tensor preparation and output processing are C++ behind pybind11. `Scheduler::_schedule` releases
the GIL on entry. Three design choices stand out against [vLLM](/articles/vllm) and
[SGLang](/articles/sglang).

**One step type per DP group.** vLLM's V1 scheduler mixes prefill chunks and decode tokens in one batch.
DeepInfer's `NORMAL` engine never does: each step the whole DP group votes, through an all-reduce, on
whether to run prefill or decode. The comment says only that a DP group must run one step type; the
MoE all-to-all, which every rank joins, is the likely reason (reasoned). Every `prefill_attempt_interval` decode steps it asks whether *all*
ranks have prefill work; every `prefill_any_attempt_interval` steps it asks whether *any* do, so a lone
rank's prompt waits a bounded number of steps (measured, from `scheduler.hpp`).

**Prefix caching is not the engine's job.** vLLM hashes blocks inside the engine; SGLang keeps a radix
tree. DeepInfer's `BlockManager` is a reference-counted free list with no hashing at all: the prefix
index lives in the entrypoint and the bytes in 3FS. Decode admits sequences shortest-output-first and,
when blocks run out, preempts from the longest-output end. The header comment says it preempts the
shortest; the code sorts by output length descending and evicts `running[l]`, the longest (measured,
`_schedule_decode`).

**Async scheduling.** It follows SGLang and NanoFlow, the Q&A says: while the GPU runs `forward(N)` and
`sample(N)`, the CPU runs `schedule(N+1)`, launches it, and only then post-processes step N−1. That is why
the C++ keeps two counters per sequence (`schedule_cnt`, `process_cnt`) and accepts being one step behind
on acceptance rates.

## A KV cache for three kinds of attention memory

DeepSeek-V4.1-Flash has 40 layers, a 128-token sliding window, compressed attention with a sparse
indexer, and two Engram n-gram tables at layers 1 and 14 (measured, from the Hugging Face `config.json`).
The [V4.1-Flash-1017 article](/articles/deepseek-v4-1-flash-1017) covers the weights; here is what the
serving side keeps per request.

`cache_engine.py` allocates two pools. The **paged pool** holds the compressed main KV and the indexer
keys in fixed blocks, stored in FP4: "per-16, e4m3" for main KV and MXFP4 for the indexer (reported).
The **state pool** gives each sequence one fixed `StateBlock`: an SWA ring per layer plus the compression
accumulators. The ring is the 128-token window plus `DEEPINFER_MAX_BACKTRACKING_TOKENS`, which defaults to
draft length plus one, so 134 tokens at 528 bytes each (measured from `environ.py` and `cache_engine.py`).

The layout is ordered by what must persist. Only the first 20 layers' SWA goes to disk, and the Q&A
gives the size: "20 layers × 134 tokens × 528 B = 1,415,040 B per sequence". The arithmetic holds
(reasoned). The other 20 layers are the decoder half of V4.1's encoder-decoder split: their block KV is
projected from the last encoder layer, so after a cache hit only their SWA needs rebuilding.

That is what **SWA bounded replay** buys. Rebuilding a sliding-window cache exactly means running the
last 2,560 tokens through 20 decoder layers. DeepInfer recomputes only the last 128, truncating each
query's window to the replayed span, which "yields an approximate SWA KV cache". The README's default
keeps 140 tokens. The Q&A says this nearly halves total prefill compute (reported). It is a lossy
shortcut, and the release says so: approximate, used for decode only, never written to the prefix
cache.

The **Engram tables** are handled the opposite way from the community recipes. `engram.py` shards them
across an engram-parallel group into a DeepEP `EngramBuffer` and fetches rows by n-gram hash across
ranks (measured). SemiAnalysis's vLLM recipes for the same model offload the tables "to host DRAM through
UVA while the KV cache stays GPU-resident" (reported, InferenceX changelog), because those recipes run
on four GPUs. Spread over 64 decode GPUs, the 202.76 GB of tables is about 3.2 GB each (reasoned).

## Kernels: mostly reused, some new

Every GPU kernel in the tree is TileLang: 93 files and 23,331 lines under `tilelang_ops/`, with CUDA and
Ascend variants (measured). They cover the glue: Engram gathers, KV stores and loads, the DSpark
verify-length kernel, sampling, structured-output masks. The heavy lifting is DeepSeek's existing
libraries, pinned as git submodules: [DeepGEMM](/articles/deepgemm), FlashMLA, DeepEP and DeepSelect,
plus `tile-kernels` from PyPI and FlashAttention-4 for the vision encoder only (measured). TileLang
itself is covered [here](/articles/tilelang).

On whether the public kernels are the production ones: the internal versions are "broadly the same";
public FlashMLA "lacks one hand-tuned SASS optimization, which makes the decoding kernel 5%–10% slower",
and end-to-end performance is "at most 5% apart" (reported).

## Expert parallelism and EPLB

With 384 routed experts and 64 redundant slots at EP64, each GPU holds 448 / 64 = 7, matching the Q&A's
"7 experts per GPU" (reasoned). The policy in `eplb/policy.py` is pure PyTorch, so I ran it on a CPU. On
synthetic Zipf-skewed loads over 384 experts, a plain six-per-GPU placement gave a max-over-mean GPU load
of 3.58 at skew 0.6 and 10.50 at skew 1.0; after the policy placed 448 copies it gave 1.004 and 1.009
(measured, 2 threads on a shared 4-core VM). Those numbers are in-sample, planned and scored on the same
counts. The Q&A's production figure, re-planned every 10,000 steps on a 1,000-step window, is
`moe_lb` = 1.090 (reported).

## DSpark: how many drafts to verify

[DSpark](/articles/deepseek-dspark) drafts five tokens and a confidence head estimates each one's chance
of acceptance. The DSpark paper describes a load-aware verifier; DeepInfer's version is
`MTPVarlen::schedule` in `csrc/engine/mtp_varlen.hpp`, and it is not a threshold. It is a greedy:

1. For every draft of every request, take the cumulative acceptance rate: the chance it *and every draft
   before it* survive. Sort all of them, across the batch, highest first.
2. Start with one query token per request. Add drafts down the sorted list one at a time, tracking
   expected tokens and the step time from a profiled curve of step time against query tokens.
3. Keep the size where expected tokens per millisecond peaks; then give each request one extra query
   token per draft of its own that made the cut.

Because cumulative rates fall within a request, the cut always keeps a prefix of each request's drafts.
Under asynchronous scheduling the CPU fixes only the total, and a TileLang kernel redistributes it on the
GPU with the freshest confidences.

To see the trade, I rebuilt the greedy in the widget below. The step-time curve is DeepInfer's own
decode chart read back as milliseconds per step: 9.99 ms at 32 query tokens per rank, 13.20 at 256, 18.21
at 512 and 31.53 at 1,024 (reasoned, from the chart in the next section).

<VerifyBudget />

Below a few hundred query tokens, a step costs about the same however many tokens it carries, so
verifying all five drafts is free and the greedy keeps them. Past that, each extra draft costs time and
the low-probability tail is cut first. With mean acceptance 0.75 and 60 requests per rank, my
reconstruction verifies 3.88 drafts per request and gains 2.8% over verifying all five; at 96 requests,
3.42 drafts and 8.2% (reasoned). The Q&A's own table, from EP64 benchmark data: 3.58 drafts and +5.4% at
60 requests per rank, 3.24 and +11.0% at 96 (reported). Its fixed-five throughput, 13,730 tokens/s per
rank at 60, sits within 4% of my curve's 13,271 (reasoned), which says the curve I read back is about
right.

## What the benchmark measures

<Figure src="https://media.thesatyajit.com/ai/articles/deepinfer/fig2.png?v=803683d40083d432" alt="Line chart: output throughput per GPU against per-user speed for EP64DP64, 384 plus 64 experts, context 131,072, nextn 3 fixed, per-token acceptance 75 percent giving tau 3.25. Points labelled 256, 128, 64, 32, 16 and 8 requests per DP rank run from about 26,000 tokens per second per GPU at about 100 tokens per second per user down to about 2,600 at about 325." caption="DeepInfer's decode benchmark: tokens per second per GPU against tokens per second per user, one point per batch size per rank. Note the subtitle: nextn=3 fixed, acceptance assumed at 75% (DeepInfer README, decode_pareto.png)." />

The README says these were "Measured on 4 SM100 GPUs for prefill and 64 for decode — each within one
scale-up domain". The script that makes the plot is `examples/sweep/run_sweep.sh`, and the axes come
from `plot_pareto.py`:

```python
# examples/sweep/plot_pareto.py (abridged)
def decode_tau(points, draft_acceptance):
    """τ = 1 + nextn × a"""
    assert not points[0]["mtp_varlen"]
    return 1 + points[0]["nextn"] * draft_acceptance

xs = [p["decode_avg_sps"] * tau for p in points]
ys = [p["decode_avg_sps"] * p["decode_num_requests"] * tau / p["cp"] for p in points]
```

`decode_avg_sps` is steps per second, measured over 60 steps after 50 of warm-up, with fixed-length
verification; a comment in `benchmark.py` translates as "each step always commits 1+nextn tokens, so it
does not depend on the acceptance rate". Then both axes are multiplied by τ = 1 + 3 × 0.75 = 3.25, which
is an input to the plotting script, not a measurement. Every request runs `ignore_eos` at 131,072
tokens of context with 4,096 not cached (measured, from the scripts).

Reading the six points off the PNG's pixels against its gridlines gives 26,270 tokens/s per GPU at 103.1
tokens/s per user, 22,806 at 178.5, 15,746 at 246.2, down to 2,619 at 325.2. Divide the first axis by
the second and you get the batch label every time, 254.9 for 256 and 64.0 for 64, because the y-axis is
the x-axis times requests per rank (measured). Per-user speed over 3.25 gives the step time: 31.53 ms at
256 requests and 9.99 ms at 8 (reasoned).

How fair is 3.25? Two things say it is optimistic. Speculative acceptance is a chain: a draft only counts
if every draft before it was accepted. The Q&A reports a fixed-five accepted length of 3.29 at the
production operating point; a chain with a constant 0.75 per position reproduces that (1 + 0.75 + … +
0.75⁵ = 3.29), and the same chain cut at three drafts gives 2.73, not 3.25 (reasoned). Independently,
SemiAnalysis measured V4.1-Flash's DSpark acceptance on SPEED-Bench coding prompts: 3.27 at three drafts
with thinking off, 2.90 with thinking on, and 3.51 at five drafts with thinking on (reported). DeepSeek's
own 3.29 at five drafts sits below even that. Scale 3.25 down to about 2.7 and the decode chart drops by
roughly a sixth (reasoned).

<Figure src="https://media.thesatyajit.com/ai/articles/deepinfer/fig3.png?v=4798c07f940c4b16" alt="Two decode curves on the same axes. The EP64DP64 curve with 384 plus 64 experts lies above and to the right of the EP8DP8 curve with 384 plus 32 experts at every batch size; at 256 requests per rank they are close in throughput, at 64 requests EP64 is both faster per user and higher per GPU." caption="EP8 against EP64 for decode, same benchmark and same assumed acceptance. Wide-EP mostly buys per-user speed at the same batch (DeepInfer Q&A, wide_ep_decode_pareto.png)." />

The Wide-EP chart makes the case for 64-GPU decode. At 64 requests per rank, EP8 reaches 12,646 tokens/s
per GPU at 197.7 per user; EP64 reaches 15,746 at 246.2 (measured from pixels). Fewer experts per GPU, fewer weight bytes read per step.

<Figure src="https://media.thesatyajit.com/ai/articles/deepinfer/fig4.png?v=33c7b1d5f167405e" alt="Prefill chart: non-cached input throughput per GPU against latency per prefill step for EP4DP4 with decoder bounded replay of 140 tokens. Throughput rises from about 60,000 tokens per second per GPU at 4,096 tokens per step and 68 ms to about 89,000 at 65,536 tokens per step and 730 ms." caption="Prefill: non-cached input tokens per second per GPU against step latency, EP4DP4 with decoder bounded replay on (DeepInfer README, prefill_pareto.png)." />

Prefill on four GPUs runs 60,051 tokens/s per GPU at 4,096 tokens a step in 68.3 ms, and 89,474 at 65,536
tokens in 731.7 ms (measured from pixels). This curve has no assumed factor in it.

## The chart that ran off the page

<Figure src="https://media.thesatyajit.com/ai/articles/deepinfer/fig5.jpg?v=dd27600790b8f967" alt="The community chart. At the top, red hand-written annotations in Chinese: 25,600 output TPS per GPU at 100 TPS per user P90; 23,500 at 175; 15,000 at 250; all data B300 DeepInfer on 64 cards, DP64 EP64. A red arrow runs down to the SemiAnalysis InferenceMAX chart of DeepSeek V4.1 Flash agentic output token throughput per chip against P90 interactivity, updated 2026-10-20, where the highest curve, GB300 NVL72 Dynamo vLLM, peaks near 6,100 tokens per second per chip, and a hand-drawn B300 DeepInfer label sits above it." caption="@SpXMerlin1D's chart on X: DeepInfer's points annotated above SemiAnalysis's InferenceMAX plot, with the note 'too far away to fit on the chart'. Two empty stretches of the arrow are cut out at the dashed lines; nothing else is changed (@SpXMerlin1D, X, 20 October, built on SemiAnalysis InferenceMAX)." />

The red text reads: "Output 25600 TPS/GPU, 100 TPS/User P90 output speed"; "23500 … 175"; "15000 … 250";
"All data in the figure is B300 DeepInfer on 64 cards. Parallelism: DP64 EP64." The post says "Data comes
from SemiAnalysis's InferenceMAX". Four things about that:

1. **The DeepInfer numbers are not InferenceMAX's.** They are DeepInfer's README chart, read by eye:
   25,600 at 100 against 26,270 at 103.1, and so on. SemiAnalysis's benchmark repository contains no
   DeepInfer configuration at all (measured, by search of a fresh clone).
2. **"P90" is wrong.** InferenceMAX's x-axis is P90 interactivity across a trace replay. DeepInfer's
   is the mean step rate of a fixed batch, times 3.25.
3. **"B300" is unverified.** The README says SM100 and "one scale-up domain" of 64 GPUs. An eight-GPU
   B300 box is not a 64-GPU NVLink domain; that needs a rack-scale system such as GB200 or GB300
   NVL72 (reasoned). DeepSeek doesn't say which.
4. **Per GPU means different GPUs.** For a disaggregated row, InferenceMAX's API divides output by
   prefill plus decode GPUs: GB300 Dynamo vLLM at concurrency 7,480 reports 250,443 tokens/s on 18 + 16
   GPUs, 7,366 per GPU (measured, API rows fetched 23 October). DeepInfer divides by decode GPUs only,
   and its Q&A says prefill and decode come 1:1.

Put both on one denominator (decode GPUs) at matched interactivity, and interpolate DeepInfer's curve:

| Run | Per-user speed | Output tok/s per decode GPU | DeepInfer at that speed | Ratio |
|---|---:|---:|---:|---:|
| GB300 Dynamo vLLM, DEP16 decode, c7,480 | 145.6 (P90) | 15,653 | 24,317 | 1.55x |
| GB300 Dynamo vLLM, DEP32 decode, c10,888 | 196.5 (P90) | 13,732 | 20,929 | 1.52x |
| GB300 Dynamo SGLang, DEP16 decode, c4,096 | 146.8 (P90) | 9,456 | 24,262 | 2.57x |
| GB300 Dynamo SGLang, DEP16 decode, c2,884 | 196.5 (P90) | 6,978 | 20,929 | 3.00x |

The InferenceMAX columns are measured from its API; DeepInfer's column and the ratios are reasoned.
Against the best vLLM run the gap is about 1.5x, which is what the "1.3x-1.5x" reply guessed, and before
taking 3.25 down toward 2.7. It is still not like for like. DeepInfer's side is offline decode with no
arrivals, no prefill interference and a mean; InferenceMAX's is an agentic trace replay at P90 with
mean inputs near 140K tokens, on EP16 or EP32 rather than EP64, and with its own simulated acceptance.
The honest summary is that DeepInfer's decode looks clearly ahead of the open engines' best published
V4.1 runs, by less than the chart drew. Nobody has yet run DeepInfer under InferenceMAX's harness. That
needs 64 GPUs and the store, router and transfer code DeepSeek kept back.

For a single-user number on similar hardware, see [Kimi K3 on GB300](/articles/kimi-k3-gb300-megakernel);
for why a Pareto chart is not a deployment, see [Qwen3.8 on vLLM](/articles/qwen38-pd-serving).

## What I ran

In a Python 3.12 venv with CPU PyTorch 2.14.1, on a shared 4-core cloud VM, no GPU:

- `tests/models/test_eplb_policy.py`: 26 passed, 6 skipped (CUDA cases), with an empty `deep_ep`
  stand-in so the package's imports resolve.
- `tests/test_cuda_sync_guards.py`: 6 passed, 1 skipped. `tests/distributed/test_shm_broadcast.py`: 3
  passed.
- The EPLB probe above.

What did not run: the C++ extension `deepinfer._C`. `setup.py` builds it as a `CUDAExtension` linked
against `libcuda`, so the scheduler, block manager and DSpark greedy tests (`test_sequence.py`,
`test_flops_estimator.py`) cannot load without a CUDA toolkit. Hence the widget re-implements the greedy
rather than calling it. I did not build or run deepseek-recipe's Rust crates.

## Where this leaves things

- **The release is the engine, not the service** (measured). KV storage, P/D transfer and the four load
  balancers are described, not shipped, and the Q&A says they are "tightly coupled to our internal
  infrastructure".
- **The scheduler is C++**, with no in-engine prefix cache and one step type per DP group (measured).
- **DSpark's verify length is a greedy over a profiled step-time curve**, worth a few percent at
  production concurrency and more under load (reported, and reproduced in shape, reasoned).
- **The headline decode chart multiplies a measured step rate by an assumed 3.25**, which DeepSeek's own
  acceptance numbers suggest is high by about a sixth (reasoned).
- **The viral comparison mixed sources and denominators**; on a common denominator DeepInfer is about
  1.5x vLLM's best GB300 V4.1 run (reasoned).
- **Not verified:** the GPU model, the 155-trillion-tokens-a-day and \$38-per-GPU-hour figures, and
  anything about the unreleased store and balancers.

For the request side, deepseek-recipe is the more immediately usable piece: an MIT Rust library that
renders V4 and V4.1 prompts and parses DSML tool calls the way the official API does. Combine it with
[DeepSeek Harness](/articles/deepseek-harness) and the client end of DeepSeek's stack is now open. The
serving end is open as far as one engine, which is further than [most labs go](/articles/glm-built-its-own-inference-infra).
If [prefill and decode](/articles/how-llm-inference-works) are new to you, start there.

<RepoCard repo="deepseek-ai/deepseek-recipe" />