FlashMLA is DeepSeek's library of optimised attention kernels — the single most expensive operation in a transformer. "MLA" is Multi-head Latent Attention, DeepSeek's attention design that compresses keys and values into a small shared latent vector so the KV cache is a fraction of the usual size. This repository contains token-level sparse attention kernels for both prefill and decoding with FP8/FP4 KV caches, plus a famous fused kernel that performs Q-norm, Q-RoPE, the attention itself, the conjugate O-RoPE and the cast to FP8 all in one GPU launch, avoiding the memory round-trips that separate kernels would pay. On a B200 it reaches up to 1 460 TFLOPS in prefill and 950 TFLOPS in decoding; on the Huawei Ascend 950 the sparse kernels hit 410 TFLOPS in prefill (95% of theoretical peak) and 360 TFLOPS in decode (83%). The September 2026 release dropped support for NVIDIA Hopper and for DeepSeek V3/V3.2/V4.0 and changed the FP8/FP4 KV cache format, so it is deliberately not backward-compatible.

DeepGEMM is a unified tensor-core kernel library that collects the main computation primitives of a modern LLM into one CUDA codebase: GEMMs (matrix multiplications) in FP8, FP4 and BF16, fused mixture-of-experts with overlapped communication ("Mega MoE"), the MQA scoring kernel for the sparse "lightning indexer", and HyperConnection (HC) operators. Every kernel is compiled at runtime by DeepJIT, so installing the package needs no CUDA compilation and the kernel can be specialised to the exact shapes you use. It borrows ideas from NVIDIA's CUTLASS and CuTe but deliberately avoids their heavy template machinery, so the library stays small enough to read — DeepSeek positions it as a teaching resource for GPU optimisation as much as a production dependency. Despite the lightweight design it matches or beats expert-tuned libraries, reaching 1 550 TFLOPS on H800, and a fully API-compatible Ascend port exists as DeepGEMM-Ascend.

DeepEP ("DeepEveryParallel") is the communication library that makes mixture-of-experts work across many GPUs. In an MoE model each token is routed to a handful of experts, and those experts live on different GPUs, so every layer requires an all-to-all shuffle: send each token to the GPUs holding its experts (dispatch), then send the results back (combine). DeepEP provides high-throughput, low-latency GPU kernels for exactly that, including FP8 dispatch to halve the bytes on the wire, plus NVLink exchange of weights and gradients for redundant experts — which is precisely what EPLB needs when it shuffles hot experts between GPUs. V2 added a lightweight NCCL "Gin" backend, analytical SM and queue-pair sizing (no more auto-tuning), and experimental primitives for pipeline, context and data parallelism plus remote memory access for Engram. The Ascend port keeps the same buffer API and reports 313–375 GB/s dispatch across EP8–EP128 on Ascend 950DT.

DeepSelect does one job: find the k largest values in a list, extremely fast. That sounds trivial until you realise DeepInfer does it twice per token at enormous scale — once in the sparse "lightning indexer" to pick the 512 most relevant past blocks to attend to, and once in the sampler to restrict softmax to the top few thousand of a 129 280-token vocabulary. A naïve torch.topk is memory-bound and badly tuned for these shapes, so DeepSelect specialises for exactly two cases — bfloat16 indexer workloads with topk ≤ 4 096 (optimised around topk=512), and float32 sampling over a ~128K vocabulary — and achieves 2× to 20× speedups over torch.topk. It supports both CUDA and Ascend, though only the CUDA path takes float32 input. Because top-k does no floating-point arithmetic, the library is benchmarked in effective memory bandwidth rather than FLOPS.

TileLang is a domain-specific language for writing GPU kernels in Python-like syntax that compiles to multiple hardware backends — which is how DeepInfer supports both NVIDIA GPUs and Huawei NPUs without two separate hand-written kernel sets. TileKernels (installed from PyPI as tile-kernels, not a submodule) is DeepSeek's library of dozens of pre-written TileLang kernels for MoE routing, Engram gating, FP8/FP4 quantisation and manifold hyper-connection operators, most of which run near the hardware's compute or bandwidth ceiling. Separately, DeepInfer ships about 45 of its own TileLang kernels in deepinfer/model_executor/layers/tilelang_ops/, each written twice — once under cuda/ and once under asc/ — with a dispatcher that picks the backend at import time. These in-tree kernels cover the engine-specific glue: building KV indices, loading paged caches, RMSNorm, the fused sampler, and the GPU-side draft-length scheduler.

_pkg = importlib.import_module(f"{__package__}.{PLATFORM.tilelang_ops_backend}")   # "cuda" or "asc"
for _info in pkgutil.iter_modules(_pkg.__path__):
    _mod = importlib.import_module(f"{_pkg.__name__}.{_info.name}")
    for _name, _obj in vars(_mod).items():
        if _name.endswith("_tl"):          # every kernel factory ends in _tl
            globals()[_name] = _obj

DeepInfer takes token IDs in and gives token IDs out; it knows nothing about chat formats. deepseek-recipe is the separate MIT-licensed Rust library (with Python bindings) that fills that gap: it converts Chat Completions, Responses and Anthropic Messages requests into one internal Conversation format, renders that into the exact prompt string V4/V4.1 expects, encodes it to token IDs, and parses the model's output stream back into protocol responses with reasoning content, tool calls and stop sequences. Two details are genuinely clever. First, special-token injection defence: when building the prompt it records exactly where it inserted control tokens like <|User|>, so if a user literally types that string it is encoded as ordinary text and cannot hijack the conversation. Second, segmented parallel tokenisation: because the special-token boundaries are already known, the ordinary text between them can be BPE-encoded independently and in parallel, then concatenated — which measurably cuts tokenisation latency on very long prompts. Its behaviour is aligned with the official DeepSeek API, so a third-party deployment behaves identically.

DeepInfer runs a speculative draft model; DeepSpec is the separate MIT-licensed codebase for training and evaluating one. It is a full stack: data preparation (download prompts, regenerate target-model answers, build a target cache — which can run to tens of terabytes), training code for several draft algorithms including DSpark, DFlash and Eagle3, and an evaluation harness that measures acceptance rate on GSM8K, MATH-500, AIME-25, HumanEval, MBPP, LiveCodeBench, MT-Bench, Alpaca and Arena-Hard. Released checkpoints target Qwen3-4B/8B/14B and Gemma-4-12B rather than V4.1, so it is the research vehicle, not the production artefact. You only need it if you want to train your own draft head; DeepInfer uses the DFlash head already baked into the V4.1 checkpoint.

3FS is DeepSeek's own distributed file system, built for AI workloads on modern SSDs and RDMA networks. It matters here because it is the storage layer underneath the production KV cache: roughly 100 PiB of SSD holding the attention keys and values of past requests so that a repeated prompt prefix never has to be prefilled twice. Data is stored with 10+2 erasure coding, meaning each group tolerates two simultaneous node failures, and a read fires requests to all 11 nodes in the group and decodes from whichever 10 respond first — latency insurance by design. A GC worker on each storage node evicts least-recently-used cache under a watermark, and DeepSeek sets a TTL of 3 days for agent traffic and 7 days for chatbot traffic (chatbot has a heavier long tail of repeat hits, and dropping to ~3 days costs it 0.5–1% of hit rate). Published 3FS cluster figures reach ~6.6 TiB/s aggregate read on 180 storage nodes.

In ordinary attention, every layer stores a full key and value vector for every token and every head; for a 128k-token context that cache becomes larger than the model weights. MLA instead projects keys and values down into a single shared low-rank "latent" vector per token, caches only that, and projects back up on demand — so the cache shrinks by roughly an order of magnitude while quality is preserved. In the code this shows up as a chain of low-rank projections: an A matrix down, an RMSNorm, then a B matrix up, with separate handling for the RoPE (rotary position) part of the key that cannot be compressed. DeepInfer fuses the up-projection of values with the output projection (_init_wv_o_proj) so the two never materialise separately. The V4.1 variant adds gating and a second compression stage on top, which is why you see kv_lora_rank, qk_rope_head_dim and compress_dim all over the cache-engine code.

Instead of one giant feed-forward network per layer, an MoE layer holds hundreds of smaller "expert" networks plus a tiny gate that scores them. For each token the gate picks the top few experts (V4.1 uses 384 routed experts in the benchmark configs), so only a small fraction of the parameters do work on any given token. The model can therefore be enormous in total parameters while staying cheap per token. DeepInfer's gate uses sqrt-softplus scoring plus the noaux_tc top-k selection with a load-correction bias — and interestingly keeps two FP32 biases, one for text and one for image tokens. A "shared expert" runs for every token regardless of routing. The whole engine's decode strategy is built around the fact that the expert weights, not the activations, dominate memory traffic.

Full attention costs grow with the square of context length, which is unaffordable at 128k+. DSA makes attention sparse: a cheap scoring network — the "lightning indexer" — estimates how relevant each past block of tokens is to the current query, and then real attention runs only over the top few hundred blocks. The indexer is a multi-query attention (MQA) pass computing relu(q @ k) * weights logits, deliberately tiny compared to the main attention, and the selection uses DeepSelect's fast top-k. DeepInfer adds two refinements: candidate blocks, where one designated layer produces a restricted candidate set that later layers select within, and cross-layer index sharing, where consecutive layers reuse the same top-k or the same index keys instead of recomputing them. These sharing relationships are derived from the model config and validated fail-fast at startup.

Not every layer needs to see the whole context. In V4.1 a set of layers use sliding-window attention: each query attends only to a fixed window of recent keys — 128 tokens — which makes their cost constant regardless of how long the conversation is. The engine stores that window in a ring buffer per sequence rather than in the paged block cache, since it is small, fixed-size and constantly overwritten. In the code the SWA cache lives inside a StateBlock, one fixed-size slot per sequence, and a dedicated kernel (get_kv_swa_index_decode) maps logical positions to ring slots. The interaction between SWA and prefix caching is the single trickiest part of the whole system, and is what SWA bounded replay exists to solve.

YoCo ("You Only Cache Once") splits the model into a producer half and a consumer half. The first n_prefill_layer layers compute real KV caches from the prompt; the layers after that — the "consumer" or decoder layers — derive their block KV from the last producer layer's hidden state instead of computing it themselves. In DeepSeek's CED (compress-encode-decode) architecture this means prefill only has to run the encoder layers fully, and the decoder layers' only per-layer state is their sliding window. DeepInfer exposes two truncation modes: mask keeps every query row but narrows what the SWA can see, and compact actually removes the truncated query rows so the consumer layers do less work. The two are numerically equivalent inside the retained region, and examples/yoco/yoco_equiv.py exists purely to prove that — comparing SWA pivots and visible keys cheaply (L1), or comparing logprobs element-by-element across 8 GPUs (L2).

DEEPINFER_YOCO_PREFILL_TRUNCATION_MODE = EnvVar.choice(default="off", choices=("off", "mask", "compact"))
DEEPINFER_YOCO_PREFILL_KEEP_TOKENS     = EnvVar.int()
# Constraint enforced at startup:
#   KEEP_TOKENS >= window_size + DEEPINFER_MAX_BACKTRACKING_TOKENS
#   V4.1: minimum is 134 = 128 (SWA window) + 6 (backtracking, = nextn + 1).
#   The README's recommended value is 140, i.e. the minimum plus a little headroom.

Engram is an external memory subsystem: alongside the normal transformer weights, the model carries enormous embedding tables indexed by n-gram hashes of the recent tokens. At certain layers the engine hashes the last N tokens, looks up the corresponding rows, projects them through a shared KV projection, and gates the result into the hidden state — effectively giving the model a vast, cheap associative memory that is read rather than computed. Because the tables are far too big for one GPU they are sharded across ranks and gathered by hash at runtime (DeepEP provides the remote-memory primitive). Engram also brings a compressed tokenizer: raw token IDs are mapped into a smaller compressed vocabulary before hashing, so the tables stay a manageable size. One independent write-up describes V4.1-Flash as a 552B backbone plus 196B Engram parameters, which gives a sense of the scale involved.

A standard transformer has one residual stream: each sub-layer reads it, computes something, and adds the result back. mHC replaces that single stream with hc_mult parallel streams — a "trunk" of shape [tokens, hc_mult, hidden]. Before each sub-layer the trunk is mixed into that sub-layer's input with learned coefficients, and afterwards the output is mixed back into the trunk. The mixing coefficients are normalised with a Sinkhorn iteration, which keeps the mixture well-conditioned. V4.1 also uses a "shifted" variant where the mix is offset by half a layer, so each sub-layer applies the coefficients computed by the previous one and the final mix-down happens at the end of the model instead of in a separate head. DeepInfer implements this with a lazy data structure (HCLazyPostData) so the post-mix of one sub-layer and the pre-mix of the next are fused into a single kernel call.

V4.1-Flash accepts interleaved images, so DeepInfer carries a full vision path. An image is resized to a patch-aligned "anyres" geometry on the CPU by worker processes (off the step critical path), uploaded as raw RGB, padded into patches on the GPU, run through a Vision Transformer, spatially downsampled, and projected by an MLP "aligner" into the language model's embedding space — where the resulting vectors replace the image placeholder tokens in the prompt. The token layout is strict: [IMAGE_START] then, per row, n_llm_w image tokens plus a newline token, then [IMAGE_END]. The scheduler must never cut a prefill chunk in the middle of an image, which is why scheduler.hpp carries image-boundary logic. In production the vision encoder runs on its own single-GPU instances behind their own load balancer.

Most of V4.1 does not run in 16-bit. Weights and caches are stored in 8-bit (E4M3) or 4-bit formats, with a shared scaling factor per small group of values to preserve dynamic range — this is what "per-16" or "per-32" means in the Q&A. UE8M0 is a scale format with 8 exponent bits and no mantissa, i.e. the scale is always a power of two, which makes applying it a cheap exponent add rather than a multiply; four of them pack into one int32. MXFP4 is the 4-bit format with per-32 UE8M0 scales. DeepSeek's exact production choices are documented: the main KV cache uses per-16 E4M3 (like NVFP4 but without the global scale), the indexer KV uses per-32 UE8M0 MXFP4, and the MoE routed-expert weights also use per-32 UE8M0 MXFP4. SWA KV stores 528 bytes per token — 512 bytes of FP8 data plus 16 bytes of UE8M0 scale.

A sequence's KV cache grows one token at a time and you don't know how long it will get, so reserving a contiguous slab per request wastes enormous amounts of memory. The paged approach — borrowed from operating-system virtual memory — chops GPU memory into fixed-size blocks (V4.1 uses 128 tokens per block) and gives each sequence a block table: a list saying "my tokens 0–127 are in physical block 41, tokens 128–255 in block 9…". Blocks can then be allocated and freed individually with no fragmentation, and two sequences sharing a prompt prefix can point at the same physical blocks. Slot mapping is the companion idea: a flat array telling the store kernel exactly which physical slot each newly-computed token's KV should be written into. In DeepInfer the free list is a C++ intrusive doubly-linked list with reference counts, so allocate and free are O(1) and block 0 is reserved as a null/padding sentinel.

enum class AllocStatus { OK, LATER, NEVER };   // can we fit this sequence now, later, or never?

struct Blocks {  // NOTE: GPU block 0 is unused, so 0 means "empty"
    Blocks(int _n) : _blocks(_n), _head(1), _tail(_n - 1), _size(_n - 1) {
        TORCH_CHECK(_n > 1, "GPU blocks too few: ", _n);
        for (int i = 1; i < _n; ++i) { _blocks[i].prev = i - 1; _blocks[i].next = i + 1; }
        _blocks[_n - 1].next = 0;
    }
    int _pop_front() { int r = _head; _erase(_head); return r; }   // O(1) allocate
    int num_free_blocks() { return _size; }
};

DeepInfer keeps two caches. The paged block cache holds the compressed main KV, which grows with context length. The state cache holds everything that is fixed-size per sequence: the sliding-window KV ring buffers and the compression-boundary accumulators. Each sequence gets one StateBlock slot, and that slot is deliberately partitioned into a stored region (the SWA layers that must survive across requests — the first 20 layers in V4.1) followed by a dropped region (the YoCo consumer SWA, recoverable by recomputation, plus compression state that has no cross-boundary dependency). The split exists so that persisting a sequence to disk is just "write the first state_size_to_disk bytes" — a length truncation, not a gather. The number in the Q&A: 20 layers × 134 tokens × 528 bytes = 1 415 040 bytes per sequence.

# Layout, low address → high:
#     [ stored region:  SWA                                  ]
#     [ dropped region: YOCO consumer SWA | compression state ]
#
# SWA partitioning depends only on n_prefill_layer / n_layer, never on the
# truncation mode. Consumer SWA recovery relies on prefill recomputation;
# compression state has no dependency across the persistence boundary.
self.state_size_to_disk = self.swa_store_total_bytes_inference_pad

Here is the problem. Prefix caching lets you skip prefill when someone sends a prompt you've already seen — but only if you saved all the state, and the sliding-window KV is expensive to persist (it is per-layer and constantly rewritten). If you don't save it, a cache hit still forces you to recompute the window exactly, which for V4.1's decoder would mean pushing the last 2 560 tokens through 20 layers. Bounded replay says: don't rebuild it exactly — rebuild it approximately, by replaying only the last 128 tokens and truncating the attention window to within the replayed segment. Replaying from position s, the query at position i sees only keys in [max(s, i−127), i]. The cost becomes a constant, independent of context length and layer count. DeepSeek reports the quality loss as negligible, and the win is large: because misses are now cheap, the persistent cache can store only the global KV and skip SWA KV entirely, shrinking the stored footprint dramatically.

A naïve decode step is schedule(N) → forward(N) → sample(N) → post-process(N). Two of those four stages are pure CPU work, and you need a GPU→CPU synchronisation before post-processing to learn which token was sampled. So the GPU sits idle during post-process(N) and schedule(N+1) — a bubble on every single token. Asynchronous scheduling removes it by pipelining: while the GPU runs forward(N) and sample(N), the CPU concurrently runs schedule(N+1), launches forward(N+1), post-processes step N−1, and launches sample(N+1). The GPU queue never drains. DeepSeek credits SGLang and NanoFlow for the idea and enables it by default (DEEPINFER_ZERO_OVERHEAD_SCHEDULING, on for NORMAL and COMPLETION engines).

A decode step launches thousands of tiny kernels, and at a few microseconds of CPU overhead each, the launching alone can exceed the compute. A CUDA graph captures the entire sequence of launches once and replays it as a single driver call. The catch is that a graph is frozen: fixed shapes, fixed memory addresses. DeepInfer therefore captures one graph per batch size on a ladder (default step of 8, via DEEPINFER_CUDAGRAPH_BATCH_SIZE_INTERVAL), pads the real batch up to the next captured size, and keeps stable input buffers that it overwrites each step. There is a second, separate runner for the DFlash draft model. During startup the engine also times every captured graph to build a decode performance curve — (batch size, milliseconds) pairs — which is cached to disk and later fed to the speculative-decoding scheduler so it can predict how long a candidate batch size would take.

Every iteration, the scheduler answers one question: what goes on the GPU this step? It maintains waiting (not started), prompted (prefilled, awaiting decode) and running queues, and a step is homogeneous — either a prefill step or a decode step, never both. Key behaviours: continuous batching, where finished sequences leave and new ones join every single iteration rather than waiting for a whole batch to complete; chunked prefill, slicing a 128k prompt into pieces so it cannot monopolise the engine (with careful logic to never cut inside an image); and preemption, where if KV blocks run out the scheduler evicts sequences — choosing victims by shortest output length (LCFS, last-come-first-served) so the most-invested requests survive. A NORMAL engine running both phases must stay in lockstep with its data-parallel peers, which it achieves by voting.

// Two kinds of attempt point:
//   ALL point (every prefill_attempt_interval = 70 decode steps):
//       only prefill if EVERY rank has prefill work — otherwise one rank works
//       while the others run an empty prompt step.
//   ANY point (every prefill_any_attempt_interval = 100 decode steps):
//       prefill if ANY rank has work — a starvation backstop for an isolated rank.
bool is_any_point = (decode_step != 0) && (decode_step % prefill_any_attempt_interval == 0);
bool is_all_point = (decode_step % prefill_attempt_interval == 0);
if (!is_any_point && !is_all_point) return _schedule_decode(acquire);

bool want = (get_gpu_cache_usage() <= GPU_CACHE_PREFILL_WATERMARK) && _has_prefill_waiting();
bool do_prefill = is_any_point
    ? (_dp_allreduce(want ? 0 : 1, ReduceOp::MIN) == 0)    // ANY
    : (_dp_allreduce(want ? 1 : 0, ReduceOp::MIN) == 1);   // ALL

// Second vote: "want" only checked that the queue was non-empty — _schedule_prefill
// can still come up empty (no slots, no KV blocks). If EVERY rank came up empty we
// must fall back to decode, or the group spins forever re-voting and re-failing.

Three names for layers of the same idea. MTP (multi-token prediction) is the general technique: a small draft model proposes the next few tokens, and the big model verifies all of them in one forward pass, so a single expensive memory read can yield several accepted tokens instead of one. DFlash is DeepSeek's specific draft implementation, a lightweight head that shares the main model's KV and runs inside its own CUDA graph. DSpark is the scheduling layer on top, and it is the genuinely novel bit: a confidence head predicts the acceptance probability of each individual draft token, and the scheduler then allocates verification compute only to the tokens with the best expected return. Verification is therefore variable-length — a request whose drafts look doubtful gets fewer slots this step. Acceptance correctness is preserved by standard rejection sampling against the target distribution. Full detail and the measured gains are in section 10.

The model outputs a score for every token in the vocabulary; the sampler picks one. In DeepInfer that is far more than a softmax. The pipeline is: gather the hidden states that need sampling → matrix-multiply against the embedding to get logits → apply the structured-output bitmask → apply temperature, top-k and top-p → sample → and, if speculation is on, verify the previous round's drafts by comparing their logprobs under the old and new distributions, accept a prefix, resample at the first rejection from the residual distribution, then run the draft head to propose the next round. A notable optimisation is force_topk: rather than softmaxing 129 280 values, restrict to the top few thousand first (default 4 096 in the offline example) and maintain an index mapping. There is also golden_token_ids for teacher-forcing, used by the evaluation and equivalence harnesses.

Normally, running the same prompt twice on the same GPU can give different results, because floating-point addition is not associative and batching changes the order in which partial sums are reduced. Who else happened to be in your batch silently changes your logits in the last decimal places, which occasionally flips a token. For evaluation, auditing, debugging and reproducible agent runs this is maddening. Batch-invariant inference (DEEPINFER_BATCH_INVARIANT=1) forces deterministic reduction orders in the GEMM and attention kernels so results do not depend on batch composition. Deterministic sampling builds on it: with batch invariance on, variable-length MTP off (DEEPINFER_ENABLE_MTP_VARLEN=0) and a fixed seq_seed in SamplingParams, identical input tokens and parameters produce byte-identical output. Both are NVIDIA-only in this release.

if envs.DEEPINFER_BATCH_INVARIANT or PLATFORM.is_npu:
    deep_gemm.use_deterministic_algorithms(True)
    tile_kernels.use_deterministic_algorithms(True)

Structured output constrains generation so the result cannot violate a format — a regex, a JSON Schema, or an EBNF grammar. The format is compiled into a grammar automaton (DeepInfer uses xgrammar); at each decode step the automaton is advanced by the token just emitted, the set of legal next tokens is computed, and a bitmask is applied to the logits so illegal tokens get −∞ before sampling. Compilation is expensive and inference is latency-critical, so the two are decoupled: compilation runs asynchronously on a thread pool pinned to idle CPUs (in production, on a dedicated machine) and a request fetches its compiled matcher before entering the engine. DeepInfer adds two extras beyond the basics: trigger_token_ids, a token sequence that must appear before the grammar activates (so the model can think freely first), and bound_token_id, which stops free generation and forces the trigger sequence through.

cases = [
    ("…write a long story…",  {"schema_type": "regex", "schema": r"[A-Z]{2}-[0-9]{3}"}),
    ("…ramble in natural language, no JSON…",
                           {"schema_type": "json",  "schema": json.dumps(json_schema)}),
    ("…write a long poem…",    {"schema_type": "ebnf",
                            "schema": 'root ::= "answer=" ("yes" | "no") ";level=" [1-3]'}),
]
outputs = llm.generate(prompts=[p for p, _ in cases],
    sampling_params=[SamplingParams(temperature=0, max_tokens=128, so_args=so) for _, so in cases])
# The prompts deliberately ask for the OPPOSITE of the schema — if the output still
# matches the regex/schema/grammar, the constraint is genuinely binding.

At 200 decode steps per second with 60 sequences each, allocating a Python object per token per request would bury you in garbage collection pauses — and a GC pause at the wrong moment stalls the GPU. DeepInfer's answer is blunt: the worker calls gc.disable() at startup and runs a manual collection every DEEPINFER_GC_INTERVAL_STEP (default 1 000, ×10 for completion engines). Everything hot is pooled — SequenceMetadataPool, RequestOutputPool, MTPInfoPool — and the outputs handed back to Python are raw uint64 pointers into C++ storage rather than objects. Because those pointers are recycled on the next step, there is a DeferredRelease class that holds resources for exactly one extra round so the caller is guaranteed to have read them, and releases them during the next step's pre_sampler_sync — i.e. overlapped with GPU compute.

DeepInfer ships three timing systems. CPUTimer records wall time around named scopes. CUDATimer records GPU time with event pairs. And CUDAGraphTimer solves a subtle problem: once a step is captured in a CUDA graph you can't insert timing calls, so it embeds event-record nodes into the graph itself and ping-pongs between two event buffers so that reading last replay's timing never stalls this replay. On top of this sits a StatsLogger with sliding windows that periodically prints throughput, cache usage, preemptions, MoE load balance and MFU — model FLOPs utilisation, computed from the C++ FlopsEstimator, which deliberately excludes draft-model FLOPs because DFlash is extra compute spent to accelerate the target model, not useful work in itself. There is even an optional HTTP endpoint per engine process for capturing a torch profiler trace on demand.

EP (expert parallelism) splits the MoE experts across GPUs; each GPU holds a slice of the experts and tokens are shuffled to wherever their experts live. DP (data parallelism) runs independent replicas, each handling different requests — in V4.1 the attention part is data-parallel with #DP = #EP. CP (context parallelism) splits a single long sequence across GPUs, which helps latency on very long prompts; DeepInfer implements a "mirror CP" variant with its own index-preparation kernels. TP (tensor parallelism), the classic approach of splitting each weight matrix, is not used in DeepSeek's production deployment at all. The Q&A's recommended settings: prefill uses EP2–8/DP2–8 (the smallest that HBM allows); decode uses EP64–128/DP64–128, the largest the GPUs allow — this is what "Wide-EP" means.

MoE routing is not uniform: certain experts attract far more tokens than others. Since every GPU must finish before the layer completes, one overloaded GPU sets the pace for all 64 — and the imbalance is dynamic, drifting with the traffic mix. EPLB fixes this with two mechanisms. First, each GPU holds a few redundant experts: extra copies of hot experts, so their load can be split. Second, the engine counts tokens per expert over a sliding window and, every N steps, recomputes the whole placement and physically shuffles expert weights between GPUs over NVLink symmetric memory. DeepSeek's measured results use moe_lb — the slowest rank's load divided by the mean, where 1.0 is perfect:

The counts are rounded up so total experts divide evenly by ep_size — at EP64 that works out to 7 experts per GPU after rounding, excluding the shared expert.

Production runs four separate balancers, and the Q&A describes each precisely even though none is in the release. The Prefill LB buckets requests into four context-length tiers (<16K, 16–64K, 64–256K, ≥256K) and redistributes engines between tiers based on last round's demand; within a GPU it doesn't use a fixed chunk size but solves a fitted quadratic cost model cost(q, c) for the largest chunk q that exactly consumes the step's time budget, capped at 32 768 tokens. The Decode LB balances two things at once: request count (which sets batch size) and total KV length (which sets attention cost), so it heuristically sends short requests to GPUs with low concurrency but long KV, and long requests to GPUs with short KV. The Vision Encoder LB routes at machine granularity, picking the GPU currently processing the fewest image patches, computed directly from image dimensions. The Expert-Parallel LB is EPLB, above.

If a user sends a 100k-token document and asks five questions about it, prefilling it five times is pure waste. Prefix caching keeps the KV from the first pass and reuses it. DeepSeek runs a three-tier hierarchy: in-flight KV in GPU HBM; the SWA KV in local host memory with an in-memory directory sharded across entrypoint machines, so a machine that doesn't hold it can pull it over a 400G VLAN; and the persistent paged block KV on the dedicated 3FS cluster. Hiding the cost of a hit is its own engineering problem: matching is a hash-chain prefix match on metadata only but can't be hidden because the hit length determines routing; fetching from 3FS into host shared memory runs in a separate process; and swap-in from host to GPU runs in a thread, overlapping per-layer host-to-device transfers with the forward pass — using GDR/RDMA rather than the CUDA copy engine, because the copy engine performs poorly on the fine-grained block copies and RDMA saturates PCIe without consuming SMs.

In a Wide-EP deployment all GPUs of an instance form one communication domain, so a single failed GPU makes the whole instance unusable — a brutal failure mode at EP64. DeepSeek's answer is not to make restarts fast (they take minutes) but to keep redundancy: ready machines held unoccupied as scale-out inventory, and a rule forbidding scale-in whenever quota utilisation is above 0.8. On restart the platform re-derives the EP size from whatever GPUs are currently healthy, so the instance adapts to the GPUs it actually receives. There is also a neat detail for long-tail requests: when a decode engine gets a shutdown notice it enters a one-minute grace period and forwards any still-running requests — along with their KV cache — to other engines in one shot, rather than killing them or waiting for the longest generation to finish. Scaling signals are simple: prefill sizes itself from the number of non-cached tokens, decode from the number of running requests, both rate-limited and bounded.

Serving performance is a trade-off curve, not a number. Increase the batch size and you serve more tokens per second per GPU (good for your costs) but each individual user waits longer between tokens (bad for them). A Pareto chart plots both axes so you can pick the operating point your service level agreement allows — which is why DeepSeek published curves rather than headline numbers. The metrics: TTFT (time to first token) is dominated by prefill; TBT/TPOT (time between output tokens) is dominated by decode; throughput is tokens/s/GPU. They are not interchangeable, and comparing one project's tokens/s/user against another's tokens/s/GPU is meaningless. DeepInfer ships the harness to generate these curves yourself: examples/sweep/run_sweep.sh sweeps one dimension and plot_pareto.py joins the points.