DeepSeek Recipe
- Multi-protocol API adaptation. Chat Completions, Responses and Messages API requests are converted into one conversation format, sharing templates and encoding logic to produce a per-model-version prompt and token sequence. Output is parsed into events (body text, reasoning, tool calls) and converted back into each protocol's streaming response. Conversion, encoding and parsing are written in Rust, with Python bindings.
- Tokenizer handling and acceleration. Two problems: (1) linear tokenization of very long contexts adds latency (hurting TTFT) and significant CPU cost; (2) special tokens that mark roles, reasoning boundaries and tool calls can also appear as plain text, and user-supplied special-token text must not affect inference. The fix is special-token boundary matching:
Input encoding: special tokens inserted explicitly while building the prompt are recorded with positions and mapped straight to their IDs; everything else uses the base vocabulary. Strings like <｜User｜> and <｜image｜> in user input, tool results or tool-call arguments are encoded as ordinary tokens, so text can't inject control tokens.
Output parsing: the end-of-thinking marker </think> and the ｜DSML｜ special token in DSML tags are recognized by token ID. The same text spelled with ordinary tokens stays as plain content and can't end reasoning or trigger a tool call.
Segmented acceleration: special-token positions are known when the prompt is built, and with BPE the plain-text segments between them can be encoded independently. So the encoder skips special-token matching inside plain text, encodes segments in parallel and concatenates in order, cutting tokenization time on very long requests.
Streaming output parsing. The backend returns chunks at token granularity, so a UTF-8 code point, a DSML tag or a tool argument can be split across chunks. A cross-chunk state machine keeps partial progress in memory and emits settled content early.
- Cross-chunk parsing: incomplete UTF-8 code points are buffered until later tokens complete them. The parser tracks phases (reasoning, body, tool arguments) and uses KMP to match delimiters incrementally; a partial delimiter at a chunk's end keeps its match state for the next chunk, with no rescanning of earlier text.
Tool argument conversion: the tool name is returned as soon as it's parsed; JSON fragments are then produced directly from DSML argument boundaries. String values are JSON-escaped piece by piece; arrays, objects and other non-string values pass through as raw JSON. No full argument object is built, avoiding repeated parse/serialize, and long arguments stream out while being generated.
- Load balancing
Both DP and EP need balanced per-GPU load to avoid idling while ranks wait on each other. Prefill/Decode Load Balancers balance attention-sublayer load via routing and scheduling; the Vision Encoder Load Balancer balances across encoders; the Expert-Parallel Load Balancer balances the MoE sublayer via redundant experts and dynamic adjustment.
- Prefill Load Balancer
Four layers, all driven by a fitted attention-cost measure called compute-quota, a pre-fitted quadratic cost model cost(q, c), where q = uncached tokens in the sequence and c = prefix length already computed:
Bucketing: 4 buckets by total input length: <16K / 16-64K / 64-256K / ≥256K. Engines per bucket are re-split each round by that bucket's share of compute-quota demand in the previous round.
Across DP groups: engines pull work according to remaining compute-quota. Since any engine can hit any prefix, there's no prefix-affinity routing.
Within a DP group: each arriving request goes to the DP rank with the lowest current compute-quota load.
Per GPU: FIFO dequeue with a variable chunk size. Each step has a time budget in nanoseconds; the scheduler solves the quadratic cost(q, c) for q and takes the largest q that fits the remaining budget.
- Decode Load Balancer
Routes across every GPU of every decode instance, balancing both request count (sets per-step batch size) and total KVCache length (sets attention cost). Heuristic: short requests go to GPUs with low concurrency but high total KVCache length; long requests go to GPUs with low total KVCache length. Based on DualPath's intra-group scheduling [1].
- Vision Encoder Load Balancer
Images are routed per machine: the prefill machine randomly picks a machine running one vision encoder per GPU, and on it the image goes to the GPU currently handling the fewest image patches ((width/patch_size) × (height/patch_size)) to produce the embedding.
- Expert-Parallel Load Balancer
Each GPU holds some redundant experts beyond its resident ones, used as replicas of hot experts. Per-expert load is tracked over the last N steps and experts are reassigned to GPUs every M steps; expert weights move within the parallel group over the scale-up network via symm_mem.
Prefill: start with 4 redundant experts per GPU (padded up so total experts is a multiple of ep_size); 100-step stats window, re-arranged every 1000 steps. Measured moe_lb (slowest rank / mean) = 1.0320 (EP2/DP2) and 1.0618 (EP4/DP4).
Decode: start with 1 redundant expert per GPU (same padding), e.g. 7 experts per GPU total at EP64 after rounding (excluding shared experts); 1000-step window, re-arranged every 10000 steps. Measured moe_lb = 1.090.
- Autoscaling
With PD disaggregation the right Prefill/Decode ratio shifts with traffic; a fixed ratio causes queuing and idling, so the cluster autoscales the ratio.
Method: Prefill instance count is driven by uncached token volume; Decode by the number of requests in flight. Changes are rate-limited and bounded. For large-EP decode, where long-tail requests waste resources, an instance told to go offline gets a 1-minute grace period; anything still running after that is handed off, with its KVCache, to other instances in one batch.
Headroom quota: (1) a pool of idle, ready machines as scale-up stock; (2) always more instances online than computed demand. Specifically, no scaling down when quota utilization is above 0.8.
Fault tolerance: restarts take minutes, so resilience comes from headroom plus request-level re-routing, not faster restarts. With large EP, all GPUs in an instance are one communication domain, so one bad GPU takes down the instance. On restart the platform re-sizes EP to the healthy GPUs, and the instance runs at whatever EP size it actually gets.
- KVCache Storage
Prefix-cache reuse greatly cuts prefill recomputation as contexts grow and agents spread. V4.1-Flash's KVCache has two parts: Block KVCache (on a 3FS cluster) and State KVCache (in local host memory).
Block KVCache: paged in 128-token blocks; per the model architecture, stores the main KV and indexer K for 4 layers.
SWA KVCache: a 134-token ring for each of the first 20 layers (window 128 + speculative rollback 6); compressed intermediate state is not stored.
- Block KVCache
Block KVCache lives only on the dedicated 3FS cluster, not in host memory. Compared with tiered storage, this lets every prefill machine hit any prefix (no clash between prefix affinity and load balancing), decouples storage from compute autoscaling, and keeps cache safe from compute-node failures (compute nodes fail an order of magnitude or more often than storage). Two levers: more capacity for hit rate, more bandwidth for read latency.
Capacity: 3FS [2] provides about 100 PiB of SSD for KVCache. TTL is 3 days for agent traffic and 7 days for chatbot traffic. Chatbot traffic has a heavier long tail (a few entries read very many times), so misses cost more; experience shows a ~3-day TTL would lower chatbot hit rate by 0.5%-1%.
Bandwidth: 3FS is tuned for random I/O at KVCache block size (about 112 KiB). Peak aggregate read bandwidth is 10 TiB/s; peak aggregate write about 850 GiB/s.
Framework side: decode instances write Block KVCache to storage, so both prefill and decode output gets reused. A hit has three stages, only the first on the critical path: (1) match: the Entrypoint does a hash-chain prefix match against the local kvcache-store in CHECK role, metadata only, but hit length must be known before routing so it can't be hidden; (2) fetch: a separate process asynchronously moves KVCache from 3FS into host shm; (3) swap-in: a thread in the engine process moves shm → HBM, aiming to overlap layer-by-layer H2D (Host-to-Device) transfer with the forward pass. CUDA copy engines handle fine-grained block copies poorly, so swap-in uses GDR, which saturates PCIe from the RDMA NIC without using SMs.
- SWA KVCache
SWA KV isn't persisted; it stays in the memory of the machine that produced it. A logically unified directory, sharded by key across Entrypoints, records where each entry lives; on a hit with no local copy it's fetched cross-machine. It's only reused within a session over minutes, so a small, fast-turnover pool covers all active sessions: in production its TTL is about 180 seconds with a hit rate above 98%. Eviction misses are recomputed by Encoder SWA Bounded Replay [3].
SWA KVCache is stored only at multiples of 128 tokens. After Prefill: take the sequence end minus one token and round down to a multiple of 128; if under 128 tokens, nothing is stored. Chunked prefill stores only at the final chunk. During Decode, a separate copy area in GPU memory receives the request's current window every time it crosses a 128-token boundary; at request end, the most recently completed copy is stored, anchored down to the 128 boundary. A copy can overshoot the boundary by at most 5 tokens, which the ring's 6 extra slots absorb.
- Hit policy
Hit length is set by Block KVCache; whether the last block needs recomputing is set by SWA KVCache, and on an SWA miss it backs off one block. Per YOCO truncation, for input length L the maximum hit is L - 134 (window 128 + rollback 6). Within that cap, find the Block KVCache hit L1 (a multiple of block size 128), then check SWA: if SWA KVCache exists between L1 - 128 and L1, the hit length is L1; otherwise L1 - 128, with 128 as the recompute window.
- 3FS distributed file system
Block KVCache is huge and lives on 3FS, which must be both highly available and fast despite node crashes and network jitter.
Data uses erasure coding (10+2 Erasure Coding); each EC group tolerates 2 simultaneous node failures.
Reads go to 11 Storage nodes in the group at once; any 10 responses are enough to decode.
A GC Worker on each Storage node periodically evicts least-recently-used KVCache (LRU) to keep disk usage under a watermark.
- SWA Bounded Replay
Not persisting SWA KVCache and recomputing it on a miss is the main KVCache-management change from V4 to V4.1 (see the V4.1 technical report [2]).
V4's approach: persist plus exact recompute. SWA KVCache was written to persistent storage at two positions; on a miss it recomputed "layers × window" tokens through a full forward pass to rebuild window state exactly. That's expensive and needed enough SSD to keep misses very rare, so SWA KVCache took nearly half of persistent KVCache capacity without matching hit gains: Block KVCache is reused over days with a long tail, while SWA KVCache is reused only within an active session for minutes and dies when the session ends or moves to the next turn. Holding minute-lifetime data on days-retention SSD is especially wasteful for short multi-turn chats. The V4 report's Zero SWA Caching (no persistence, exact recompute only) was too costly in production and was not used in V4.1.
V4.1's approach: SWA Bounded Replay recomputes only the most recent 128 tokens and truncates attention inside that span: replaying from position s, the query at position i attends only to keys from max(s, i − 127) to i, giving an approximate SWA KVCache. Recompute cost becomes a constant independent of context length and layer count. Used in two places:Encoder SWA Bounded Replay: prefix-hit length now depends only on Block KVCache. On an SWA miss, the last 128 tokens of the hit prefix are recomputed alongside the uncached suffix; this only rebuilds SWA KVCache, and Block KVCache is reused as-is, neither recomputed nor overwritten.
Decoder SWA Bounded Replay: prefill runs only in the encoder layers. In the CED architecture, decoder-layer Block KVCache is projected from the encoder's last-layer hidden state, so only decoder SWA KVCache must be computed per layer; exact rebuild would need 20 decoder layers over the last 2560 tokens. Instead each prefill recomputes only the prompt's last 128 tokens in the decoder layers, used only for decode and not added to the prefix cache. Total prefill compute is nearly halved.
- DSpark speculative decoding
DSpark is DeepSeek's speculative decoding framework. Compared with DFlash, it adds a low-rank logit bias that depends only on the previous token id on top of the base logits: a very light serial correction that restores local conditioning between adjacent tokens and improves draft accuracy (see the DSpark paper [4]).
Variable-length verification: acceptance probability differs by request and position, so fixed-length verification wastes compute on unlikely tokens. A Confidence Head predicts each draft token's acceptance rate and verification compute goes to the tokens with the highest marginal benefit. In production the draft proposes 5 tokens at a time; using the profiled hardware throughput curve and each draft token's acceptance rate, tokens are ordered by expected acceptance and the prefix subset with maximum effective throughput is verified. Versus always verifying 5 tokens, this gives +6% throughput in production.
With Asynchronous Scheduling: when schedule(N+1) runs, step N's confidence hasn't come back yet, so the CPU only fixes the total budget; per-request lengths are recomputed on the GPU by a tilelang kernel (cumulative product, global sort for top-k, fused remap scatter), with results written back to a pinned buffer.
- Asynchronous Scheduling (Zero-Overhead Batch Scheduler)
Decode step N runs schedule(N) → forward(N) → sample(N) → post-process(N). schedule and post-process are pure CPU, and post-process(N) needs a GPU-CPU sync to get sample(N), so the GPU idles during post-process(N) and schedule(N+1). Following prior work [5] [6], CPU and GPU are pipelined asynchronously: while the GPU runs forward(N) → sample(N), the CPU runs schedule(N+1) → launch_forward(N+1) → post-process(N-1) → launch_sample(N+1), so the GPU never idles.
- Structured output
Structured Output (SO) constrains output to an exact format (JSON schema, regex, EBNF, etc.).
Compile stage (dedicated SO machines): the constraint is compiled into a grammar automaton.
Runtime: at each decode step the automaton is updated, the set of legal tokens is determined, and a mask is applied in sampling to force conforming output.
Compilation is decoupled from inference: SO compiles asynchronously and the request fetches the result before entering the engine.