DeepInfer: What DeepSeek Open-Sourced, and What It Kept

In April 2025, DeepSeek wrote that it would not open-source its inference engine; on 20 October 2026 it did, with Huawei's Ascend 950 listed next to NVIDIA's Blackwell. What came out is the engine and its kernels for one model, running offline; what stayed in-house is everything that turns an engine into a service. Read DeepInfer as the reference for how DeepSeek-V4.1-Flash is meant to run, not as a replacement for vLLM or SGLang.

Everything below comes from the repository at commit 9bd9cf2 (21 October 2026), its deployment Q&A, and the release notes of vLLM v0.32.0 and SGLang v0.5.22. DeepInfer needs Blackwell datacenter GPUs or Ascend 950 NPUs, and we have not run it: every number here is DeepSeek's, and says so.

Why DeepSeek released it now

DeepSeek's April 2025 note gave three reasons for keeping its engine private: it was "based on an early fork of vLLM from over a year ago", it was "tightly coupled with our internal infrastructure", and a small research team could not maintain a large open-source project. Instead, it promised to extract reusable parts as standalone libraries and to work with existing engines so that new models had support on day one.

It kept that promise for eighteen months. FlashMLA, DeepGEMM and DeepEP came out in February 2025, TileKernels with V4 in April 2026, and Ascend 950 versions of DeepGEMM and DeepEP at the end of September 2026. Meanwhile vLLM and SGLang learned to serve each new model, usually within days. DeepSeek announced a companion with it, deepseek-recipe, the Rust converter between API requests and token sequences that matches the official API's behavior.

The engine came last: by October 2026, most of what makes it fast was already public as separate libraries.

DeepInfer doesn't contradict the 2025 reasoning; it is shaped by it. The parts that are tied to DeepSeek's infrastructure, the KV cache store and the load balancers, are still private. The parts DeepSeek maintains anyway, for two chip families, are public: the engine, the model code, the kernels. The Q&A states the commitment going forward: "Inference code for future DeepSeek models, along with subsequent performance work, will continue to be open-sourced in DeepInfer." For DeepSeek, that is cheaper than porting each new architecture into two external engines, on two hardware stacks, before every launch.

What is in the box

DeepInfer runs one model, DeepSeek-V4.1-Flash-1017, on two kinds of hardware: NVIDIA SM100 GPUs (the B200 and GB200 class) and Huawei Ascend 950 NPUs. It exposes an offline batch API, LLM.generate() with SamplingParams, and no HTTP server. Prerequisites are Linux, GCC 13, Python 3.12, and CUDA 13 or CANN 9.2.

The hardware floor matters. Hopper (SM90), RTX Blackwell and the GB10 (SM12x) are not targets, so most teams that run DeepSeek models today on H100s or H200s can't run DeepInfer at all.

The model is the one described in our DeepSeek piece: 552 billion backbone parameters (763 billion in the checkpoint with the Engram tables and the drafter), arranged as a 20-layer causal encoder and a 20-layer decoder, 384 routed experts plus one shared with 6 active per token, FP8 weights with FP4 experts, a 128-token sliding window, a million-token context, two Engram memory layers, and a built-in draft model for speculative decoding (config).

The tree is about 55,000 lines of Python and 6,000 of C++:

• deepinfer/: engine, worker, model code (model_executor/models/v4_modules/ holds attention, compression, Engram, hyper-connections, MoE), and 93 files under tilelang_ops/: entry points plus 30 CUDA and 30 Ascend kernel implementations.
• csrc/: the scheduler and input preparation in C++.
• vendor/: FlashMLA, DeepGEMM, DeepEP, DeepSelect, DeepGEMM-Ascend and DeepEP-Ascend as git submodules.
• deep_cluster/: hardware and topology detection for both platforms.
• eval/: harnesses for AIME 2026, HLE and GPQA Diamond.

What remains of vLLM is the skeleton. Twenty-one files in deepinfer/ carry the header "Derived from vLLM; modified by DeepSeek", including the engine loop, the worker, the model runner and the cache engine. The scheduler is now C++ (scheduler.hpp), and the model and kernels are DeepSeek's own. One tree serves both chips through a small platform layer (PlatformGPU, PlatformNPU) and TileLang kernels compiled per backend.

What is specific about it

Several of these mechanisms already exist in vLLM or SGLang in some form; DeepInfer shows the version that runs in production, with production settings.

Wide expert parallelism for decode

DeepSeek runs decode at EP64 to EP128: each GPU of a 64- to 128-GPU supernode holds a few of the 384 experts, and attention runs data-parallel with as many replicas as there are expert ranks. Prefill uses EP2 to EP8 on ordinary 8-GPU machines, "the smallest EP/DP that HBM allows". Wider EP means fewer expert weights read per GPU per step, so a larger batch fits in the same step time; the Q&A's chart shows EP64 ahead of EP8 at every batch size: at 64 requests per rank, about 15,700 tokens per second per GPU at 246 per user, against 12,600 at 198 for EP8 (131,072-token context, DeepSeek's measurement; both curves use the same assumed draft acceptance, so the gap holds even if the absolute numbers don't, see the gotchas).

Experts are not equally popular, so each GPU also holds redundant copies of hot experts, and the expert-parallel load balancer (EPLB) reshuffles them from recent load statistics. The production settings are in the Q&A: for decode, one redundant expert per GPU, a 1,000-step statistics window, a reshuffle every 10,000 steps, and a measured imbalance (slowest rank over mean) of 1.090. EP64 sets the decode floor at 64 GPUs, and DeepSeek runs prefill and decode at roughly one GPU to one, so the smallest deployment it calls efficient is 64 GPUs of prefill plus 64 of decode.

DSpark: verify only the drafts likely to survive

V4.1-Flash ships with a DSpark draft model (arXiv 2607.05147): it proposes a block of 5 tokens in one pass, with a small sequential head for dependencies inside the block. The novelty for the engine is a confidence head (ConfidenceHead) that predicts, for each draft token, the probability that the target model will accept it.

In speculative decoding, a rejected draft token costs a verification slot for nothing, and the fifth token of a draft is far less likely to survive than the first. DeepInfer uses the predicted probabilities to decide how many of each request's drafts to verify: it sorts all draft tokens in the batch by expected acceptance, and picks the subset that maximizes useful tokens per second on a throughput curve profiled for the hardware.

Verification shrinks by up to a third while accepted tokens barely move; the fuller the batch, the more the saved slots are worth.

DeepSeek's estimates, from the Q&A, computed from EP64/DP64 benchmark data:

The gain grows with the batch, because a fuller batch makes each verification slot more expensive. At small batches it is slightly negative. At the production operating point, around 60 sequences per rank, DeepSeek puts it at +4 to 7% (+5.4% in the table).

With asynchronous scheduling, the CPU plans step N+1 before step N's confidences exist. So the CPU fixes only the total token budget, and a TileLang kernel on the GPU (schedule_mtp_varlen.py) splits it across requests once the confidences arrive.

Bounded replay: prefill that skips half the network

V4.1-Flash keeps a sliding-window (SWA) cache of the last 128 tokens in every layer, besides its compressed long-range cache. Its decoder layers compute their long-range KV from the encoder's last hidden state, so for a prompt, only the encoder runs over every token; the decoder needs nothing but its sliding-window cache. Rebuilding that cache exactly is not 128 tokens of work: each layer's window reaches 128 tokens further back through the layer below, so an exact cache needs the last 128 × 20 = 2,560 tokens run through all 20 decoder layers.

Decoder SWA bounded replay recomputes only the last few tokens in the decoder and truncates attention to the replayed segment, which gives an approximate SWA cache. The Q&A says this "nearly halv[es] total prefill computation". The replayed suffix must cover the 128-token window plus the speculative-decoding backtrack (6 tokens at 5 drafts), so at least 134 tokens; the README default is 140:

# README @ 9bd9cf2: recommended V4.1 settings (environ.py default: off)
export DEEPINFER_YOCO_PREFILL_TRUNCATION_MODE=compact
export DEEPINFER_YOCO_PREFILL_KEEP_TOKENS=140

compact removes the truncated queries from the decoder layers; mask keeps them but narrows their attention, and saves less. The same trick, on the encoder side, handles prefix-cache hits whose sliding-window cache is missing: recompute the last 128 tokens of the hit instead of the whole prefix. DeepSeek's prefill chart uses replay: about 60,000 non-cached input tokens per second per GPU at 4,096 tokens per rank (70 ms per step), rising to about 89,000 at 65,536 (730 ms), on 4 SM100 GPUs at EP4, with 131,072-token contexts.

Replay is an approximation, and what it costs in quality depends on the model. The public check so far is SGLang's: its team checked it on AIME 2026 and found the same pass@1 with replay on and off (453 of 480 samples each).

The rest of the engine

• Zero-overhead scheduling. While the GPU runs step N, the CPU schedules step N+1 and post-processes step N−1, so the GPU never waits on Python. DeepSeek credits SGLang and NanoFlow for the design.
• Batch invariance and deterministic sampling. DEEPINFER_BATCH_INVARIANT=1 makes a request's logits independent of what else is in the batch; add a fixed seq_seed for identical outputs. NVIDIA only: the engine asserts it is off on Ascend.
• Structured output through XGrammar, compiled asynchronously on a dedicated machine in production.
• Kernels. DeepGEMM's Mega MoE, which fuses dispatch, both expert matmuls and combine into one kernel, is "significantly faster than unfused MoE" within a scale-up domain, at any EP size (DeepGEMM). The public kernels are close to the internal ones: within 5% end to end, with FlashMLA's decode kernel 5 to 10% slower for lack of one hand-tuned SASS optimization.
• FP4 where it pays: the main KV cache in blocks of 16 values with E4M3 scales (NVFP4 without the global scale), the indexer cache and the routed experts in MXFP4. Our microscaling piece explains both formats.
• Engram tables in host memory. On NVIDIA, the two Engram tables (about 768 million rows) stay in CPU memory and are read over RDMA through DeepEP's EngramBuffer; on Ascend, the engine forces them into device memory.

What DeepSeek kept

The README lists what a complete service needs and this release lacks: "a KV cache system", "VE/Prefill/Decode disaggregation (most required engine functionality is already implemented)", and "load balancing across data-parallel replicas". The engine side of disaggregation is indeed there: the C++ scheduler already has prompt-only and completion-only engine types. What's missing is everything between engines.

DeepSeek released the boxes and kept the arrows.

The Q&A describes the private parts in enough detail to show what you would have to build:

• KV cache store. A dedicated 3FS cluster, about 100 PiB of SSD (about 14.4 TB of logical capacity per GPU) with 10+2 erasure coding; KV blocks live 3 days for agent traffic and 7 days for chat. A shared store lets any prefill machine hit any prefix, so routing doesn't need prefix affinity. Sliding-window caches stay in the host memory of the machine that produced them, with cross-machine pulls over a 400 Gb/s VLAN.
• Prefill balancer. Requests are bucketed by context length (<16K, 16-64K, 64-256K, ≥256K); each engine pulls work against a time budget, and a fitted quadratic cost model picks the chunk size that fills each step.
• Decode balancer. Routes across all decode GPUs, balancing request count (the batch size) against total KV length (the attention cost).
• Autoscaling and failover. A decode engine being shut down gets a one-minute grace period, then hands its unfinished requests and their KV cache to other engines.

According to the Q&A, this layer gets production to about 80% of a perfectly balanced offline benchmark. If you want it, you assemble it from open parts: NIXL or Mooncake for the prefill-to-decode transfer, LMCache, Mooncake or SGLang's HiCache for the KV tier, llm-d or NVIDIA Dynamo for routing. Our disaggregated serving piece covers how they fit.

What is at stake

Ascend becomes a first-class target

Huawei's chips get the same engine, the same feature list except batch invariance and deterministic sampling, and kernel libraries with the same APIs as the NVIDIA ones. When V4 launched in April, Reuters reported a rush on Ascend 950 orders from ByteDance, Tencent and Alibaba, with H200 shipments to China still blocked; Jensen Huang has called DeepSeek moving to Huawei chips "a horrible outcome" for the United States. With DeepInfer, an operator on Ascend 950 gets DeepSeek's own production engine and kernels; the server, the prefill/decode split and the KV store are still theirs to build.

Two caveats. DeepEP-Ascend's published bandwidth figures come from a proof-of-concept driver and firmware build (HDK) supplied to DeepSeek, not a public release. And "same API" doesn't mean the same features or performance: batch invariance and deterministic sampling are already missing on Ascend.

The advantage is in the model, and the model is already open

Most of what makes V4.1-Flash cheap to serve is a model decision: a decoder that takes its KV from the encoder, compressed and sparse attention, FP4 experts, Engram memory, a drafter trained with a confidence head. The engine exploits those decisions; it doesn't create them. That is why DeepSeek can publish the engine at little cost. Its advantage lives in the next model, which the engine will support on day one.

vLLM and SGLang already run most of this

Both engines also have Ascend paths (the vllm-ascend plugin, SGLang's NPU backend), but neither lists V4.1 on Ascend 950 in these releases; DeepInfer does. The convergence is not a coincidence: DeepSeek's stated policy since 2025 has been to work with these engines before each launch, and SGLang's V4.1 support shipped with bounded replay and Engram host tables in September. For vLLM and SGLang developers, DeepInfer is the reference to diff against: production EPLB settings, the scheduling of verification, the exact replay semantics.

The economics, in DeepSeek's words

The Q&A gives two production figures: one cluster serves 155 trillion tokens a day, and revenue is US$38 per GPU-hour "at valley pricing", DeepSeek's off-peak API rates. The 82% gross margin circulating since then is an outside estimate built on assumed GPU costs and utilization; DeepSeek didn't publish it. Whatever the margin, $38 per GPU-hour at off-peak prices takes the throughput of this engine running this model; publishing the engine doesn't hand that to anyone serving a different model.

Maturity

DeepInfer is a code drop with a production pedigree, not yet an open-source project:

• History: four commits, the first on 20 October: the initial commit, two Q&A updates, and a DeepEP bump. No tags, no releases.
• Community (as of 24 October 2026): about 21,400 stars in four days; issues are disabled and Discussions open; 194 pull requests, one merged: the DeepEP bump. Open PRs add AWS Trainium, Apple MLX, and a Mooncake KV tier: the community is already building what was left out.
• Tests: 16 test modules and per-sample perplexity baselines that must match exactly, on both platforms; no public CI.
• Docs: the README, the Q&A, the examples README. Comments in core files are often in Chinese.

The quality of the engineering is visible in the code; the maintenance model isn't. DeepSeek's own 2025 concern, a small team that can't maintain a large open-source project, hasn't gone away.

Gotchas

Prefill-only engines don't overlap scheduling. Asynchronous scheduling is on for normal and decode engines and off for prompt-only (prefill) engines: EngineArgs sets DEEPINFER_ZERO_OVERHEAD_SCHEDULING from the engine type unless you set it, and environ.py's false is only the fallback. If you split prefill and decode, expect the prefill side to run without the overlap unless you turn it on.

The decode chart assumes its acceptance rate. The benchmark counts every draft token as accepted, then plot_pareto.py rescales by τ = 1 + drafts × acceptance, with an acceptance you choose. The README chart uses 3 drafts at 75%, so τ = 3.25 tokens per step, and shows about 26,000 tokens per second per GPU at 256 requests per rank (EP64, 131,072-token context). That is an assumption, not a measurement. For comparison, the Q&A's EP64 benchmark data shows an average accepted length of about 3.1 to 3.3 with 5 drafts (its definition isn't stated); the same drafter with only 3 drafts would accept fewer, so the chart likely flatters absolute throughput. Use the chart to compare configurations, which share the assumption, and your measured acceptance for absolute numbers.

Deterministic mode turns off confidence scheduling. Deterministic sampling requires DEEPINFER_ENABLE_MTP_VARLEN=0. Speculative decoding still runs, at a fixed draft length; what you lose is the verification scheduling above, about 5% at production batch sizes, whenever you need reproducible outputs, for example in evals.

Replay truncation is off unless you set it. environ.py defaults the mode to off; the README's compact and 140 are a recommendation. With truncation on, KEEP_TOKENS must cover the window plus the backtrack, 134 for V4.1, or the engine refuses to start.

Memory planning differs per chip. Engram tables sit in host memory on NVIDIA and in device memory on Ascend. Budget HBM separately for each platform.

Verdict

DeepInfer is the most complete public description of how a frontier lab serves its own model, written in code, and the clearest sign yet that Ascend is a first-class inference target. It is not a server you deploy: one model, offline, on Blackwell or Ascend 950, without the KV store and routing that make up most of a service. Three things would change that call: the disaggregation and KV pieces landing, a second model, and independent benchmarks against vLLM and SGLang on the same hardware.

• If you operate V4.1-Flash, stay on vLLM or SGLang. Take DeepInfer's settings instead: EP sizes, EPLB windows, replay length, drafts, and the Q&A's capacity figures (about 14.4 TB of logical KV storage per GPU, the TTLs, roughly one prefill GPU per decode GPU).
• If you build engines, read the C++ scheduler, the GPU-side verification scheduling, the replay path and the platform layer. It is rare to see production choices with their numbers.
• If you run Ascend, this is the first production-grade engine for Ascend 950 with the same features as on NVIDIA. Try it, knowing it supports one model and CANN 9.2.

The libraries that carry most of DeepInfer's speed have their own pieces: DeepGEMM for the matrix multiplications, expert parallelism and DeepEP for the communication. Disaggregated serving covers the layer DeepSeek kept.