What you need before you start

The code walkthroughs assume Python, some C++, and basic transformer knowledge. The introduction establishes tokens, sampling, attention, caching and experts without code. These references supply further background:

If you only have time for one, read the GPU memory-traffic post. Most of the "why" answers in Pass 4 come back to it.

The mental model

1. DeepInfer turns prompt tokens into new tokens through repeated model steps. A step batches work from independent requests. Speculative decoding can check several proposed tokens per request together.
2. The engine batches work and moves less memory per token. Its CPU, the general-purpose processor, plans the next step while the GPU runs. CPU work sets the pace if it takes longer than the GPU work it can overlap.
3. GPUs that exchange expert work must join compatible communication steps, even when one has no local requests. Supported repeatability settings also keep a request’s answer independent of its batch companions.

One picture of the whole system

The release contains only the box labelled "one engine". Everything outside it is described, not in the code doc.

A separate component, DeepSeek Recipe, turns requests into tokens doc. Load balancers choose a prefill engine and later a decode engine doc. Inside each engine, the scheduler chooses a batch, the worker runs the model, and the sampler chooses output tokens code csrc/engine/scheduler.hpp:275, deepinfer/worker/worker.py:89. The KV cache saves reusable attention state.

DP, or data parallelism, gives different requests to different GPUs. EP, or expert parallelism, spreads experts across GPUs, so tokens travel to the experts they need. A rank identifies one process in a communication group. The expert group must join matching exchanges, a form of lockstep. Separate groups can progress independently.

The six threads

1. The life of a request. From prompt in, to tokens out.
2. Memory and the KV cache. How the engine stores earlier results, and when it discards or recomputes them.
3. The model. What DeepSeek-V4.1-Flash computes per token, and where it differs from a plain transformer.
4. Spreading work across chips. How attention and experts are split over many GPUs, and how hot experts get balanced.
5. Making decode faster. Batching, recorded GPU work, overlapping the CPU with the GPU, and guessing tokens ahead.
6. Getting the right answer every time. Sampling, and making the output the same whatever else is in the batch.

Number formats such as FP8 and FP4 get no thread of their own. Thread 2 covers cache formats. Thread 3 covers model-weight formats.

Where each thread appears, and where to skip

• You know vLLM well. Skim Pass 1. In Pass 2, read the scheduler vote and async steps closely, because they differ most from vLLM. Then go to Pass 3.
• You know transformers but not serving. Read every pass in order.
• You only care about the model. Read thread 3 in each pass. You will still need the MoE parts of thread 4.
• You want design practice. Read Pass 0 and Pass 4, then try Pass 5 before you read its answers.

How each pass works

Passes 2–5 begin with recall questions. Mixed exercises ask you to choose which mechanism applies, with answers folded below. Glossary links let you revisit definitions.

The toy used in every thread

Treat “the cat sat on the mat” as six tokens at positions 0–5. Each example introduces the parts of its toy model when they are needed. These settings differ from the introduction’s browser models.

The life of a request

A request waits while the engine chooses its next batch. The engine processes the prompt in large steps, then generates tokens in smaller steps.

A restaurant kitchen with an order rail

Before each round of cooking, the head chef picks orders from a rail. A new order needs prep first, then gets one course per round.

waiting:  A = "the cat sat on the mat"   6 tokens
          B = a 5-token prompt
          C = a 4-token prompt
budget:   8 tokens per prefill step

1. Take A. It needs 6 tokens and the budget has 8. A fits whole, leaving 2.
2. Take B. It needs 5, but only 2 remain. Cut B and process its first 2 tokens now. This is chunked prefill.
3. Stop. A cut prompt ends the step, so C waits.

In this example, "a cut prompt ends the step" is the rule for the engine type this series follows. Pass 2, Thread 1 shows the code.

The next prefill step takes B's last 3 tokens and all 4 of C, 7 in all. Each prompt now has its first output token.

Say A's output is "and", "slept", ".", then the end token. Prefill gave "and", so A needs 3 decode steps at one token per step.

step 1  prefill   A[0..5]  B[0..1]          -> A gets "and"
step 2  prefill   B[2..4]  C[0..3]          -> B and C get first tokens
step 3  decode    A  B  C                   -> A gets "slept"
step 4  decode    A  B  C                   -> A gets "."
step 5  decode    A  B  C                   -> A gets the end token, A leaves
step 6  decode    B  C  (and a new request, once it has had its prefill)

The real engine does not start prefill whenever a prompt waits. It holds a vote every 70 decode steps by default, which Pass 2, Thread 1 covers.

A kitchen can cook a main course on one stove while another does prep. DeepInfer can't, because a step is all prefill or all decode code csrc/engine/prepare_sample.cpp:61.

The real engine also has many kitchens that cook in lockstep. When one does a round of prep, every kitchen does a round, even with nothing to prep. An idle kitchen cooks an empty plate, which the engine calls a placeholder step. Thread 4 explains why.

Continuous batching reforms the batch between steps: new requests join and finished ones leave. DeepInfer still keeps prefill and decode in separate steps.

Why would an engine refuse to mix prefill and decode in one step, when mixing would keep the GPU busier?

Memory and the KV cache

The engine reuses stored keys and values while they remain in the cache. Long-range records live in fixed-size blocks. A fixed ring holds recent state. Eviction, preemption or a missing window can require recomputation.

A filing cabinet and a sticky-note pad

Picture the KV cache as a cabinet of numbered drawers, all one size. A drawer holds the cards for 4 tokens in the toy, or 128 in the real engine. A clerk keeps a list of free drawers.

A new customer gets enough drawers for their whole file, and the clerk writes the drawer numbers on a slip. Each drawer is a block. The slip is the block table, and the clerk is the block manager.

Each customer also gets a sticky-note pad with a fixed number of slots, where a new note covers the oldest. It is the ring for sliding-window attention.

1. A has 6 tokens and a block holds 4. So A needs 6 / 4 = 1.5, rounded up to 2 blocks.
2. The clerk's free list says drawers 7 and 2 are free. A's block table becomes [7, 2].
3. Positions 0 to 3 go in drawer 7. Positions 4 and 5 go in drawer 2, which still has room for positions 6 and 7.

A's block table is [7, 2].

Position 8 needs a new drawer, which the clerk hands out then.

The pad has 4 slots: a 3-token window plus 1 spare. Position p goes in slot p mod 4, the remainder after dividing p by 4.

After the prompt, positions 0 to 5:

Positions 4 and 5 covered positions 0 and 1. The window for position 5 needs positions 3, 4 and 5, and all three are there.

The spare supports speculative decoding in Thread 5. A step can write a guess, then reject it. The spare stops that guess overwriting a note the window still needs. The real engine has 6 spares on a 128-token window, 134 slots in all code deepinfer/worker/cache_engine.py:88-93.

When the drawers run out

Say only 1 drawer is free and prompt B needs 2. B waits. In the real engine too, a new prompt gets every block for its whole length before it starts code csrc/engine/scheduler.hpp:386-388. So a half-done prompt never stalls waiting for its second half.

Decode is harder, because any running request may need one more drawer at any step. If none is free, the clerk takes drawers back from someone, which is preemption. Pass 3 covers who loses them.

Unlike an office drawer, a block can belong to two requests with the same prompt start. The clerk counts the slips that name each drawer and frees it when the count reaches zero.

In a plain transformer each layer would need its own cabinet. In V4.1 only 4 of the 40 layers write long-range cards, and each number takes about 4 bits. So one real block holds 128 tokens for the whole model in about 114 KB inference from the formats in deepinfer/worker/cache_engine.py:424-483.

You might think the KV cache size is "layers times tokens times the size of one key and value". It is much smaller here, for three reasons. Only 4 layers own long-range KV code deepinfer/hugging_face/config.py:62-75. The window is a fixed ring per request. The stored numbers are 4 or 8 bits. The size estimates are my calculations inference. Pass 2 calculates the real layout.

Long-range memory lives in blocks that grow with the request, and the window lives in a fixed slot. That is why the real engine keeps two pools, one for blocks and one for state slots code deepinfer/engine/llm_engine.py:265-266.

A forked request looks new, but it shares the parent's drawers and copies only the last, half-full one.

The model

The first two layers read only a 128-token window. Later layers also read up to 512 selected long-range entries. Each layer uses 6 of 384 routed expert networks, plus a shared expert.

A newsroom writer

For each word, a writer does three things.

1. They reread the last page closely. That is the window.
2. They ask a librarian, who scans a quick index, for the most useful old clippings. That is the indexer.
3. They send their sentence to a few specialist desks, such as sport, finance and science, and blend what comes back. That is the mixture of experts.

The writer keeps four versions of their working notes and blends them before each task. In the model, those notes are arrays of numbers called the hidden state. Blending four parallel versions is hyper-connections. A phrasebook of common 2-, 3- and 4-token sequences is Engram.

Follow "mat" at position 5 through the toy's two layers.

position:     0      1      2      3      4      5
token:       the    cat    sat    on     the    mat   <- the query
                                 |------ window -----|
layer 1:  window only            sees pos 3, 4, 5
layer 2:  window + 1 older pick  sees pos 3, 4, 5 + the librarian's pick
                                 the librarian scores pos 0, 1, 2 -> picks pos 1 "cat"

The toy’s second layer has four experts, E0–E3, and chooses two for “mat”. Say the router gives them these scores.

The top 2 are E2 and E1. Their scores total 0.85. After rescaling to a total of 1, E2 gets 0.60 / 0.85 = 0.71 and E1 gets 0.25 / 0.85 = 0.29. The output is 0.71 times E2's answer plus 0.29 times E1's. The real model also has one shared expert for every token. It scales the routed part by 1.5, a config value passed to the gate kernel code deepinfer/model_executor/models/v4_modules/mlp.py:378-381.

A writer has one archive. Here only layers 2, 8, 14 and 20 write long-range memory. Each layer that uses long-range memory reads one of those four code deepinfer/hugging_face/config.py:62-75. In layers 2 to 19, each entry also merges two neighbouring tokens.

The librarian uses rough 4-bit summaries to select entries, so it can miss relevant ones. The window keeps every recent position without selection or merging. It still stores numbers with limited precision.

The code calls the YOCO halves “encoder” and “decoder”, but all 40 layers are causal: a position sees only itself and earlier positions. The split describes cache ownership: the first 20 layers build their own long-range memory. The last 20 share one memory that layer 20 builds code deepinfer/model_executor/models/v4_modules/yoco.py:19-21.

Why would a model with 384 experts per layer send each token to only 6?

Spreading work across chips

Each GPU runs attention for its own requests, but experts are spread across an expert-parallel group. Every layer needs an exchange within that group.

GP practices and specialist clinics

Each GP practice keeps its own patients and files. That is data parallel attention: each GPU has its own requests and KV cache.

Each town hosts only some specialist clinics, so GPs refer patients across towns. That is expert parallel MoE.

The practices exchange referrals, and every result returns before the next round. That is the all-to-all. A long queue at one clinic can delay the group. The GPUs must join matching exchanges. They need not reach them at precisely the same time.

The toy has two GPUs. GPU0 holds E0 and E1. GPU1 holds E2 and E3. Each token in layer 2 goes to two experts.

So E0 gets 3 jobs, E1 gets 2, E2 gets 6 and E3 gets 1, 12 in all.

GPU0 does 3 + 2 = 5 jobs and GPU1 does 6 + 1 = 7. The mean is 6, but the layer takes as long as the slowest GPU. The real engine measures this with moe_lb, the busiest GPU divided by the mean. Here it is 7 / 6 = 1.17.

Fixing the hot expert

E2 is hot. Give each GPU one spare slot, so there are 6 slots for 4 experts.

Fix 1: copy E2 onto GPU0, change nothing else. E2's 6 jobs split 3 and 3.

Fix 2: copy E2, and also move E1 to GPU1.

GPU0:  E0 (3)  E2 copy (3)           = 6 jobs
GPU1:  E1 (2)  E2 (3)  E3 (1)        = 6 jobs     moe_lb = 6 / 6 = 1.00

GPU0:  E0 (3)  E1 (2)  E2 copy (3)   = 8 jobs
GPU1:  E2 (3)  E3 (1)                = 4 jobs     moe_lb = 8 / 6 = 1.33   worse!

The copy alone made it worse, and Fix 2 works because it also moves an expert. The real balancer, EPLB, does both. It decides which experts get copies, then places all the copies on GPUs code deepinfer/model_executor/eplb/policy.py:96-123 and :12-70. Pass 3 runs the real algorithm by hand.

You cannot copy a clinic overnight. Expert weights are numbers the engine can copy. The real engine re-plans and moves them every 1,000 steps for prefill and every 10,000 for decode doc.

A practice with no patients could close. A GPU with no requests can't, because the others wait for it in the all-to-all. It runs one fake token, the lockstep rule from Thread 1.

You might think more GPUs per expert group always means faster decode. More GPUs also mean more traffic and more chances for one slow GPU to delay the group. Each GPU reads fewer expert weights per step, which helps decode. So prefill uses the smallest group that fits in memory and decode uses the largest doc. Pass 4 gives the numbers.

Each GPU runs attention locally, but the expert exchange makes it wait for the slowest GPU in its group.

Making decode faster

A decode step has costs that can be shared across its rows, such as reading model weights. Small batches can add rows cheaply, but row computation and memory traffic grow with the batch. The engine tries to gain more useful tokens than those rows cost.

A waiter and a junior chef

A waiter prepares the next round’s orders while the kitchen cooks this one. The waiter represents the CPU, which prepares work for the GPU. Overlapping those tasks is zero-overhead scheduling.

A junior chef plates five dishes ahead, guessing the next orders. The head chef tastes all five at once and keeps them up to the first wrong one, which they replace. That is speculative decoding.

The waiter

Say the CPU needs 3 milliseconds, or ms, per step to pick the batch and read the results. The GPU needs 10 ms to run the model.

WITHOUT the waiter (the CPU waits for the GPU, then the GPU waits for the CPU)

CPU  |sched|                 |read+sched|                 |read+sched|
GPU        |=====10 ms=====|            |=====10 ms=====|
     one step = 3 + 10 = 13 ms  ->  about 77 steps per second

WITH the waiter (the CPU plans step N+1 while the GPU runs step N)

CPU  |sched 1|sched 2|read 1|      |sched 3|read 2|      |sched 4|read 3|
GPU          |=====step 1=====|=====step 2=====|=====step 3=====|
     one step = 10 ms  ->  100 steps per second, 30% more

The catch is that the CPU plans step N+1 before it knows step N's token. The real engine has the GPU fill in that token itself code deepinfer/worker/worker.py:191-197. Pass 2 shows how.

The junior chef

Request A has "the cat sat on the mat and". The draft model guesses "slept", "on", "the", and the main model checks all three in one step.

The check is a weighted coin flip, not a test of whether the guess is the main model's top choice. Ideal rejection sampling preserves the main model's distribution. DeepInfer implements it with small numerical approximations, described below and in Pass 2.

At the rejected position, the main model samples a replacement. It subtracts the draft’s probabilities from its own, then sets negative results to zero. Say the main model’s probabilities at the 2nd position are "on" 0.2, "in" 0.6 and "under" 0.2. The draft's were "on" 0.5, "in" 0.3 and "under" 0.2. The remaining weights are "on" 0, "in" 0.3 and "under" 0, so the replacement is "in". The real kernel scales the draft’s probabilities by a tiny factor before subtracting. Pass 2, Thread 6 shows why.

The step gave 2 tokens, "slept" and "in", where plain decode gives 1. The real engine guesses 5 at a time doc.

A head chef tastes one dish at a time. The model scores all proposed positions in one pass. At small batch sizes, sharing weight reads and launch costs can make six rows cheaper than six separate steps. It does not make the extra rows free.

The junior chef guesses for free, but the real draft model costs time every step. Each checked guess also adds a row of work. In a big batch, that growing cost becomes harder to hide.

You might think checking more guesses is always better, because a correct guess is a free token. An unlikely guess can waste the work needed to check it. The real engine can predict how likely each guess is to pass and skip the weak ones. The Q&A doc reports that this cuts guesses checked per step from 5.00 to between 3.2 and 4.4. At about 60 requests per GPU that gives 4 to 7% more throughput doc. Pass 4 shows when it loses.

Why must the GPU fill in tokens under async scheduling, rather than wait for the CPU to read them?

Getting the right answer every time

Seed and position fix each random draw. Repeatable tokens also need numerical results that do not change with the batch.

A raffle with pre-numbered tickets

Picture a raffle where your ticket numbers depend only on your name and the round, not on who else entered. Enter again and you get the same tickets.

In DeepInfer, each random draw depends only on the request's seed and the token's position code deepinfer/model_executor/layers/tilelang_ops/random_sampling.py:19-20.

Picking a token

The model produces a score, or logit, for each possible token. Higher scores mean higher probabilities. Say the vocabulary has four tokens, A, B, C and D, with logits 2.0, 1.0, 0.5 and −1.0.

1. Temperature 0.5. Divide each logit by 0.5. That gives 4, 2, 1, -2.
2. Softmax turns scores into probabilities. Exponentiate each score: e⁴, e², e¹, e⁻², where e is about 2.718. The resulting positive values, 54.6, 7.39, 2.72 and 0.14, total 64.8. Divide each by that total to get probabilities 0.842, 0.114, 0.042 and 0.002.
3. Top-p 0.9. Keep the most likely tokens until they add up to at least 0.9. A alone is 0.842, not enough. A and B make 0.956, so keep both.
4. Renormalise. A becomes 0.842 / 0.956 = 0.881. B becomes 0.114 / 0.956 = 0.119.
5. Draw. The seed and position give a number between 0 and 1. Say it is 0.93. A covers 0 to 0.881, B covers 0.881 to 1. So 0.93 picks B.

Run the same request again with the same seed, and step 5 gets 0.93 and picks B again.

Why the batch can change the numbers

Step 5 repeats only if steps 1 to 4 get exactly the same logits. Floating-point addition rounds its results. Changing the addition order can change the answer.

in 64-bit floats:      (0.1 + 0.2) + 0.3  = 0.6000000000000001
                        0.1 + (0.2 + 0.3) = 0.6

in 32-bit floats:      (100000000 + 1) + (-100000000) = 0
                       (100000000 + (-100000000)) + 1 = 1

I checked those calculations with NumPy. A GPU adds long rows in chunks, and the chunking can depend on the batch size. So the same request, alone or in a batch of 64, can get logits that differ in the last bit. Near a boundary like 0.881, one bit can flip the token, and every later token differs too.

Batch invariance keeps a request's numerical results independent of batch composition, including through stable reduction choices. DeepInfer offers it as a switch, on NVIDIA only code deepinfer/worker/runtime_config.py:62-69.

A raffle drum doesn't care what else is in the room. A GPU does, because of the chunked sums, so the logits must also repeat.

Speculative decoding adds a second break. If the engine checks 3 guesses in one run and 5 in another, the draws line up with different verification roles. Both runs can target the same distribution, subject to the numerical approximations above, while producing different tokens. Pass 3 shows why one setting must be off for exact repeats.

You might think the same seed and temperature always give the same output. Repeatability also needs identical input tokens and sampling parameters, stable numerical results and a fixed verification schedule. The README's NVIDIA recipe enables batch invariance, disables variable-length verification and supplies a fixed seq_seed code README.md:28.

All four runs below use the same model, input tokens and sampling parameters on the supported NVIDIA runtime. Variable-length verification is off, so the verification schedule stays fixed.

A repeatable output needs repeatable draws and repeatable numbers. A fixed seed and verification schedule give the first. Batch invariance supplies the second within the supported configuration.

Temperature 0 looks like a separate "always pick the top token" mode. In DeepInfer it becomes 0.00001, and the engine still draws code deepinfer/sampling_params.py:76-79. So two near-equal tokens can each win.

The request we follow

This trace follows the offline example’s long-context test. It comes from reading the source, rather than observing a run. examples/offline_inference.py --long_context stretches each sample prompt to 16,389 tokens and asks for 4,096 output tokens code examples/_prompts.py:26-27, examples/offline_inference.py:63-69. The test tests/models/test_offline_inference_long.py runs that configuration upstream.

I follow sample 1, the English one: "Question: What is the largest planet in the solar system? Answer: Jupiter ..." code examples/_prompts.py:72-76.

DP rank r gets the five sample prompts repeated r + 1 times code examples/_prompts.py:100. So rank 0 gets five prompts, r0 to r4 after their request ids. Ours is r1. Samples 2 to 4 carry images, and samples 0 and 1 are text only.

Rank 0 has one more prompt than it has running slots. A prefill step holds at most 32,768 tokens, two prompts minus 10. For the trace, assume sufficient free cache and that every rank's first prompts are ready by the vote at decode step 70. Preprocessing runs asynchronously. Reading the source cannot tell us when it finishes.

The life of a request

The request enters a C++ queue, waits for a group vote, gets prefilled in two chunks, then decodes 4,096 tokens.

Processes and who owns what

Each DP rank is its own operating-system process, with its own LLM, LLMEngine and C++ scheduler. examples/_launch.py starts one process per DP rank code examples/_launch.py:36-39.

The driver worker runs in the engine's own process code deepinfer/engine/llm_engine.py:402. Extra worker processes exist only for CP ranks 1 and up code deepinfer/engine/llm_engine.py:377. Our run has CP1, so there are none. The shared-memory broadcast in deepinfer/distributed/shm_broadcast.py only links a driver to its CP followers, so it plays no part here code deepinfer/model_executor/parallel_utils/parallel_state.py:124.

Each step the eight ranks run small all-reduces: communication operations that combine a value from each rank and return the result to every rank. One votes on the step type. Another checks whether any rank has work left code csrc/engine/scheduler.hpp:316-332, deepinfer/engine/llm_engine.py:665-679.

Eight independent schedulers serve the eight GPUs. Each step, they agree on its type and whether to continue. Requests are split across ranks before the engine sees them. Here the script splits them. In production the load balancers do, and they are described, not in the code doc.

Step 0: LLM.generate and add_request

LLM.generate loops over the prompts and calls _add_request for each code deepinfer/entrypoints/llm.py:90, :135-141. For each prompt it does four things.

1. It deep-copies the sampling parameters code deepinfer/entrypoints/llm.py:151.
2. It takes the next request id from a counter, so ours is "1" code deepinfer/entrypoints/llm.py:152.
3. It turns the token list into a 32-bit NumPy array code deepinfer/entrypoints/llm.py:157.
4. This model has a vision encoder, so it attaches an empty image dict even to text-only prompts code deepinfer/entrypoints/llm.py:158-159.

LLMEngine.add_request then checks and builds the request code deepinfer/engine/llm_engine.py:472.

• It sets the end token from the tokenizer code deepinfer/engine/llm_engine.py:527.
• It rejects an empty prompt code deepinfer/engine/llm_engine.py:554.
• It checks every token id is inside the vocabulary. For a bad id, it logs the error and returns None. The request gets no output code deepinfer/engine/llm_engine.py:580-589.
• It builds the C++ Sequence code deepinfer/engine/llm_engine.py:594.
• It hands the sequence to the scheduler and to the multimodal processor code deepinfer/engine/llm_engine.py:620-628.

The scheduler's add_sequence stamps the arrival time and reserves room for the prompt plus 4,096 output tokens. Then it pushes the sequence onto the back of the waiting queue code csrc/engine/scheduler.hpp:135-142.

The Sequence struct is at csrc/engine/sequence.hpp:275. These are the fields of r1 that this pass uses.

Sequence r1 after add_request
  input.tokens          16,389 prompt tokens, 0 output tokens
  cache.num_tokens      0          tokens whose KV is ready          (sequence.hpp:203)
  cache.block_table     []         no blocks yet                      (sequence.hpp:206)
  cache.state_id        -1         no state slot yet                  (sequence.hpp:207)
  schedule.status       WAITING                                       (sequence.hpp:236)
  schedule.schedulable  false      because multi_modal_data is set    (pybind_sequence.cpp:35)
  mtp.draft_tokens      []         no drafts yet                      (sequence.hpp:118)
  mtp.info_slot         0          slot 0 means "none"

The first engine steps are empty

_run_engine calls step() in a loop until no rank has unfinished work code deepinfer/entrypoints/llm.py:176-185.

Each step() first calls the C++ Scheduler::schedule code deepinfer/engine/llm_engine.py:885. Only then does it ask the multimodal processor for finished work code deepinfer/engine/llm_engine.py:939-945. The processor marks a request schedulable only inside that fetch() call code deepinfer/ve/multimodal_processor.py:251, :303.

Every prompt carries the empty image dict from point 4 above. None is schedulable on any rank at the first schedule(). As I read the code, the first step then goes like this.

1. The scheduler is a NORMAL engine, so it must choose prefill or decode code csrc/engine/scheduler.hpp:295.
2. decode_step is 0, so this is an "ALL point" code csrc/engine/scheduler.hpp:309-310.
3. A rank "wants" prefill only if its cache is at most 90% full and the head of waiting is schedulable code csrc/engine/scheduler.hpp:313. No head is schedulable, so no rank wants it.
4. The vote fails and the scheduler builds a decode step code csrc/engine/scheduler.hpp:318-322. running is empty, so the step is empty.
5. The worker runs a placeholder step code deepinfer/worker/worker.py:125-127. After it, decode_step is 1 code csrc/engine/scheduler.hpp:680.

The next ALL point comes when decode_step reaches 70, the default interval code deepinfer/environ.py:174. So the earliest next ALL opportunity is after 69 more decode steps inference. Whether it starts prefill depends on readiness and available resources. Each empty step still makes a full trip through the model with one fake token on every GPU in this engine group.

The translated comment on the two kinds of vote point code csrc/engine/scheduler.hpp:296-307:

NORMAL engine: decide whether the whole group does prompt or decode this step. Within DP the step type must match, and a vote unifies it. ALL points, every prefill_attempt_interval decode steps: only do it if every rank has prefill work. This avoids one rank working while the rest run empty prompt steps alongside. ANY points, every prefill_any_attempt_interval decode steps: do it together if any rank has work, as a backstop for an isolated rank.

Why is the vote an all-reduce on a CPU group, not on the GPUs?

Prefill step 1: r0 whole, r1 cut

At decode step 70 the ALL vote runs again. Under our readiness and capacity assumptions, every rank wants prefill, so this vote passes code csrc/engine/scheduler.hpp:313-319. Rank 0 runs _schedule_prefill code csrc/engine/scheduler.hpp:351. A second group vote checks whether any rank scheduled prefill. If none did, the group decodes instead code csrc/engine/scheduler.hpp:328-342.

The metadata calls the number of new positions processed in this step seqlen_q. seqlen_k includes the earlier cached positions too. In the arithmetic below, ceil means round up to the next integer.

1. r0, 16,389 tokens. Its block table is empty, so the scheduler checks the length limit and asks for all its blocks at once code csrc/engine/scheduler.hpp:374-397. It needs ceil(16,389 / 128) = 129 blocks, and they are free code csrc/engine/sequence.hpp:291-292. r0 fits within the 32,768-token budget. Its seqlen_q = 16,389, leaving 16,379 tokens. r0 becomes RUNNING and gets an MTP info slot code csrc/engine/scheduler.hpp:522-526. That slot is a row in a GPU pool where the draft model leaves its guesses for this request. MTP is the code's older name for the draft model.
2. r1, 16,389 tokens. It also gets all 129 blocks at once code csrc/engine/scheduler.hpp:518-520. Its 16,389 tokens are more than the 16,379 left, so the scheduler cuts it code csrc/engine/scheduler.hpp:470-475. seqlen_q = 16,379. r1 stays WAITING but leaves the queue, with 10 tokens still to go code csrc/engine/scheduler.hpp:533-541.
3. Your turn. Does the scheduler look at r2 in this step? Use csrc/engine/scheduler.hpp:543.

Pass 1 flagged "a cut prompt ends the step" as the rule for one engine type, NORMAL. A prefill-only PROMPT engine keeps looking, though its budget is usually spent by then.

The scheduler reserves every block of a prompt on its first chunk code csrc/engine/block_manager.hpp:206-214. So a half-done prompt never waits for blocks for its second half.

Back in schedule(): the per-step metadata

schedule() then builds one SequenceMetadata per scheduled sequence code csrc/engine/scheduler.hpp:688-729.

• For r1, seqlen_q = 16,379 and seqlen_k = num_tokens + inflight + seqlen_q = 0 + 0 + 16,379 code csrc/engine/scheduler.hpp:697.
• A guard checks the last position fits the position tables code csrc/engine/scheduler.hpp:700-705.
• The scheduler takes a metadata object from a pool code csrc/engine/scheduler.hpp:706-707.
• The scheduler marks the sequence as in flight with 6 tokens, which is nextn plus 1 code csrc/engine/scheduler.hpp:689, :726.
• Because this is a prompt step, decode_step goes back to 0 code csrc/engine/scheduler.hpp:676-677.

Python gets the metadata list and a SchedulerOutputs code deepinfer/engine/llm_engine.py:885-887. The struct is at csrc/engine/scheduler.hpp:45.

SchedulerOutputs, prefill step 1 on rank 0
  scheduled_seqs      [r0, r1]
  prompt_run          true
  num_batched_tokens  32,768
  seqlens_q           [16,389, 16,379]
  finished_seqs       []
  ignored_seqs        []

The worker round

step() hands everything to the worker code deepinfer/engine/llm_engine.py:952, :961-975. Worker.execute_model_step does four things every step code deepinfer/worker/worker.py:96-103.

1. Launch the forward pass of this step code deepinfer/worker/worker.py:113-127.
2. Post-process the previous step, if one is waiting code deepinfer/worker/worker.py:143-144.
3. Launch sampling for this step code deepinfer/worker/worker.py:146-159. Just after, the engine's pre_sampler_sync hook frees last step's pooled objects and finished requests code deepinfer/worker/worker.py:162-165, deepinfer/engine/llm_engine.py:830-846.
4. Post-process this step at once if it is a prompt step or async scheduling is off. Otherwise keep it for next time code deepinfer/worker/worker.py:182-187.

This is a prompt step, so step 4 happens at once. Thread 5 explains why decode steps wait.

After the forward pass over all 32,768 tokens, the sampler picks one row per prefill sequence, the last token of its chunk code csrc/engine/prepare_sample.cpp:197-198. Even a non-final chunk gets sampled. A translated comment: "every prefill chunk must push its last token through sampling. The results of non-final chunks are thrown away afterwards" code csrc/engine/prepare_sample.cpp:173.

Processing the output of prefill step 1

Post-processing calls Scheduler::process_outputs code deepinfer/engine/llm_engine.py:727, csrc/engine/scheduler.hpp:872. For each scheduled sequence it does the following.

• It clears the in-flight mark code csrc/engine/scheduler.hpp:903.
• It increases num_tokens by seqlen_q code csrc/engine/scheduler.hpp:927. That gives 16,389 for r0 and 16,379 for r1.
• r0 now has num_tokens >= len, so its prefill is done. It joins running, gets its first token appended and has the stop rules checked code csrc/engine/scheduler.hpp:937-990.
• r1 has num_tokens < len. It goes back to the front of waiting with nothing appended code csrc/engine/scheduler.hpp:1011-1019.

Then build_request_outputs makes one RequestOutput per token produced code deepinfer/engine/llm_engine.py:743, csrc/engine/outputs.hpp:94. r0 gets one and r1 gets none.

Prefill step 2: r1's last 10 tokens

The prompt step reset decode_step to 0, so the next step is an ALL point again. I expect the vote to pass again, because the other ranks have 10 to 40 prompts each still waiting inference.

On rank 0, running holds r0, so num_curr_seqs starts at 1 code csrc/engine/scheduler.hpp:355.

• r1 already has a block table, so the scheduler skips the allocation checks. Its work is 16,389 − 16,379 = 10 tokens, and it fits whole. r1 becomes RUNNING and gets an MTP info slot code csrc/engine/scheduler.hpp:522-526. seqlen_k = 16,379 + 10 = 16,389.
• r2, the one-image prompt, fits whole once its image is ready. It needs 16,389 of the 32,758 tokens left.
• r3, the three-image prompt, gets cut to the rest of the budget. The scheduler moves the cut back if it would split an image code csrc/engine/scheduler.hpp:488-493, :217-262.

This step samples r1's first output token, at position 16,389, from the last row of its 10-token chunk. It can't be the end token, because first_n_token_no_eos=1 bans it code csrc/engine/prepare_sample.cpp:192-195, deepinfer/model_executor/layers/sampler.py:184-185.

Continuous batching, and the fifth prompt

One more prefill step finishes r3. Rank 0 then has four running requests, the max_num_seqs=4 limit. When the scheduler reaches r4, num_curr_seqs + 1 > max_num_prompts, so it stops code csrc/engine/scheduler.hpp:440-442.

With ignore_eos and 4,096 output tokens for everyone, I expect r4 to wait until one of the first four finishes inference. This is continuous batching: a slot opens only when a request leaves.

Rank 0 still "wants" prefill at every ALL point, because r4 sits schedulable at the head of waiting. But _schedule_prefill returns an empty batch, because the slots are full. A second vote catches this code csrc/engine/scheduler.hpp:328-335. If no rank scheduled anything, the group goes back to decode. If another rank did, rank 0 runs an empty prompt step alongside code csrc/engine/scheduler.hpp:339-340.

Decode steps until the end

How Sequence r1 changes up to its first decode step, which Thread 5 covers:

The block count is ceil((len + inflight + drafts) / 128) code csrc/engine/sequence.hpp:291-292. It reaches 130 once that sum passes 16,512. With 6 in flight and 5 drafts, that is at about len = 16,502, roughly 112 output tokens in inference. _append_slot then adds one block code csrc/engine/scheduler.hpp:817-821.

After each token the scheduler runs _check_stop code csrc/engine/scheduler.hpp:986, :1035-1052. With ignore_eos there is no end-token stop. r1 finishes at 4,096 output tokens, when get_output_len() == max_tokens, with status FINISHED_LENGTH_CAPPED code csrc/engine/scheduler.hpp:1048-1050.

On the next decode schedule, free_finished_seqs moves r1 out of running code csrc/engine/scheduler.hpp:562, :752-763. Its blocks and state slot go back in pre_sampler_sync, through finalize_seq and free_seq code deepinfer/engine/llm_engine.py:844-846, :804-806, csrc/engine/scheduler.hpp:738-750. Only then does r4 get a slot.

The output

Each step() returns raw pointers to pooled RequestOutput objects. _run_engine clones each into a Python-owned copy, because the pool reuses the original next step code deepinfer/entrypoints/llm.py:179-183. Outputs are grouped by request id. When no rank has work left, the loop calls flush_inflight_reqs() once to collect the step still in flight code deepinfer/entrypoints/llm.py:178.

For r1, generate returns a list of 4,096 RequestOutput objects. The example sets decode_sequence=False, so the engine never turns tokens into text code examples/_config.py:121. The script decodes them itself code examples/offline_inference.py:78.

Each RequestOutput holds one token id with its logprobs code csrc/engine/outputs.hpp:18-40. Detokenising and streaming belong to DeepSeek Recipe, a separate repo doc.

The state machine

A chunked request goes back to the front of the queue. When the head of the queue can't be scheduled, the whole queue waits behind it.

The enum is at csrc/engine/sequence.hpp:30-43. It also lists SWAPPED, but nothing in the release sets it.

The same prompt in three engine types

The scheduler has three engine types code csrc/engine/scheduler.hpp:38-42, :288-295.

One scheduler serves all three roles, and P/D disaggregation is the same code with the vote switched off. The handoff itself is described, not in the code. The hooks send_processed_prompts and pe_do_save_kvcache are empty methods code deepinfer/engine/llm_engine.py:800-858.

A request made by fork_seq enters waiting with cached state. It shares its parent's blocks and starts with the parent's KV already counted code csrc/engine/scheduler.hpp:1061-1122.

Why is the scheduler written in C++ when the rest of the engine is Python?

Memory and the KV cache

r1 holds 129 paged blocks of 114,048 bytes for long-range KV, plus one fixed state slot for all its sliding-window rings.

The 129 blocks

A block is 128 tokens, hard-coded code deepinfer/model_executor/model_metadata_utils.py:134-135. 16,389 / 128 = 128.04, so r1's prompt needs 129 blocks. The first 128 are full. The last holds 5 tokens and has room for 123 more.

The block manager is only a free list with a reference count per block code csrc/engine/block_manager.hpp:105-116. Block 0 is never handed out. A translated comment says "GPU block 0 is unused, so 0 here means empty or null" code csrc/engine/block_manager.hpp:27-28. allocate() takes the head of the free list and sets its count to 1 code csrc/engine/block_manager.hpp:82-86. free() takes 1 off, and only at 0 does the block go back, onto the tail code csrc/engine/block_manager.hpp:95-102.

What one block holds

One block holds 128 token positions of long-range KV and indexer keys for owner layers 2, 8, 14 and 20. C2 merges pairs of positions into one entry. C1 keeps one entry per position. B means bytes. The cache engine lays them out like this code deepinfer/worker/cache_engine.py:424-483. The byte counts are my arithmetic from the formats in the code inference.

One block covers 128 token positions in 114,048 bytes.

Main KV is 512 numbers at 4 bits, 256 bytes, plus 32 one-byte scales, one per 16 numbers. That makes 288 bytes code deepinfer/model_executor/layers/tilelang_ops/cache_common.py:119-127. Indexer keys are 128 numbers at 4 bits plus 4 scales, 68 bytes code deepinfer/worker/cache_engine.py:78. The Q&A doc gives the same formats doc.

Pass 1 did a full block count. Finish this one.

1. r1 needs 129 blocks.
2. Each block is 114,048 bytes.
3. Your turn. How many bytes of paged KV does r1's prompt hold, and how many bytes is that per token?

The state slot

The sliding window is not in the blocks. Each sequence gets one fixed state slot from a second BlockManager code deepinfer/engine/llm_engine.py:265-266. A slot holds a ring for every layer, indexed by position modulo the ring size code csrc/engine/input_metadata/prepare_tensor_input_v4.cpp:103-104.

With this configuration's nextn 5, the ring has 128 + 6 = 134 positions. The extra capacity lets speculative decoding undo rejected guesses. It is at least nextn + 1. A larger draft limit can therefore enlarge the ring code deepinfer/worker/cache_engine.py:88-93, deepinfer/environ.py:132. Each position holds 528 bytes: 512 bytes of 8-bit numbers and 16 one-byte scales code deepinfer/model_executor/layers/tilelang_ops/cache_common.py:109-117.

Each sequence's active state slot has a fixed size, about 3.1 MB. This is separate from its paged blocks. Memory profiling also reserves copy capacity by budgeting the state pool twice code deepinfer/worker/base.py:241-244.

The code names the two regions code deepinfer/worker/cache_engine.py:229-234. The production store would keep the stored region across requests. The Q&A doc gives the same 1,415,040 bytes for it doc. The dropped region is never stored, because prefill can rebuild it. The byte totals are my arithmetic inference.

r1's first output token sits at position 16,389. 16,389 = 122 × 134 + 41, so it lands in slot 41 of every layer's ring code csrc/engine/input_metadata/prepare_tensor_input_v4.cpp:103-104.

Blocks hold the growing history. Rings hold the fixed recent window. The production store saves the region that would cost more to rebuild.

A near-miss and a misconception

The prefill code sets a flag called prefix_caching_hit = true code csrc/engine/scheduler.hpp:373-375. It looks like a prefix-cache hit, but it isn't. The scheduler sets it for every brand-new sequence, and its only use is to skip one assert.

The block manager has a free list with reference counts. It has no vLLM-style prefix cache where new requests can find hashed blocks. csrc/engine/block_manager.hpp has no hash, no lookup and no eviction. The only sharing in the release is fork_seq, where a caller names the parent code csrc/engine/scheduler.hpp:1061. Production prefix hits come from the external store, which the Entrypoint matches by a hash chain. That is described, not in the code doc.

The model

In r1’s first decode step, layers 0 and 1 read only their windows. Later layers also read up to 512 long-range entries selected from one of four owner layers.

What goes in

r1's first decode step has 6 rows. Row 0 is the first output token, at position 16,389. Rows 1 to 5 are draft rows, but there are no real drafts yet. After prefill the engine pads the draft list with a placeholder token code deepinfer/engine/llm_engine.py:706-711. Thread 5 explains what happens to those rows.

The entry point is DeepseekModel.forward code deepinfer/model_executor/models/deepseek.py:154. It reads the embeddings code deepinfer/model_executor/models/deepseek.py:162. It runs the Engram hash once for both Engram layers code deepinfer/model_executor/models/deepseek.py:177-188. Then it loops over the 40 layers code deepinfer/model_executor/models/deepseek.py:190. At the end it collapses the 4 residual copies and applies the final norm code deepinfer/model_executor/models/deepseek.py:225-239.

The vocabulary projection is not in forward. It runs inside the sampler, and only on the rows it samples code deepinfer/model_executor/layers/sampler.py:141-142.

Three layers, side by side

For row 0, at position 16,389:

The layer types come from the model config's compress_ratios list. The owner map comes from _source_map, which maps each layer to "the nearest source layer that is not later than that layer" code deepinfer/hugging_face/config.py:62-75. The YOCO split starts at layer 20 code deepinfer/model_executor/models/v4_modules/yoco.py:19-21.

Layer 2: why position 16,389 writes an entry

Layer 2 merges each pair of tokens into one long-range entry. The prompt's last token, at even position 16,388, had no partner yet. So during prefill it waited in the compression state, ca_state code deepinfer/worker/cache_engine.py:463-465.

Row 0 now brings position 16,389. The decode metadata builder writes a new compressed entry only when (orig_pos + 1) % 2 == 0, which holds for 16,389 code csrc/engine/input_metadata/prepare_tensor_input_v4.cpp:80-90. The new entry is number 16,389 / 2 = 8,194, rounded down. A C2 layer packs 64 entries per block. So entry 8,194 goes in block-table entry 8,194 / 64 = 128, the 129th block, at offset 8,194 − 8,192 = 2.

The merge itself is a learned, per-number weighted average of the two tokens code deepinfer/model_executor/layers/tilelang_ops/compress.py:199.

Layer 20: the shared memory for the second half

In C1 layer 20, position 16,389 becomes entry 16,389, in block-table entry 16,389 / 128 = 128, at offset 16,389 − 16,384 = 5.

Layer 20 is the only layer from 20 to 39 that writes long-range KV. That KV comes from its own input, which is the first half's output after the residual blend and the norm code deepinfer/model_executor/models/v4_modules/attn.py:465-479. Layers 21 to 39 read it but keep their own window rings.

Layer 20 also picks 2,048 candidate blocks of 8 entries each, so 16,384 entries. Layers 24, 28, 32 and 36 then score only inside those candidates code deepinfer/model_executor/models/v4_modules/index_attn.py:204-226, :308-325. For r1, at 16,390 entries, that removes almost nothing. At a million tokens it would skip about 98% inference from the numbers.

Layer 37: a tap for the draft model

Layers 37, 38 and 39 each save their attention input, averaged over the 4 residual copies code deepinfer/model_executor/models/decoder_layer.py:144-145. The draft model in Thread 5 reads those three to make its guesses.

Which layer owns the KV a layer reads?

Pass 1 did this with a toy. These are the real owner lists from the config. KV owners are layers 2, 8, 14 and 20. Indexer owners are layers 2, 8, 14, 20, 24, 28, 32 and 36. Each layer uses the nearest owner at or below it.

1. Layer 5 reads KV from layer 2, and top-k from layer 2.
2. Layer 22 reads KV from layer 20, and top-k from layer 20.
3. Your turn. Layer 27 reads KV from which layer, and top-k from which?

What attention reads, for row 0

In layer 2, row 0 attends over two sets in one softmax, plus a learned sink.

• The window: positions 16,262 to 16,389, 128 tokens, from the ring code deepinfer/model_executor/models/v4_modules/attn.py:640-647.
• 512 entries of layer 2's compressed KV, picked by its indexer from 8,195 entries.

The query has 64 heads of 512 numbers. The single stored vector serves as both key and value code deepinfer/model_executor/models/v4_modules/attn.py:662. After attention, the code rotates the last 64 numbers of each head back by the query's position code deepinfer/model_executor/models/v4_modules/attn.py:193. Then a two-step output projection maps the result back to 5,120 numbers code deepinfer/model_executor/models/v4_modules/mla_attn.py:221-229.

In V4.1, MLA has no separate key or value for each head. The 64 query heads score straight against one 512-number vector, and that same vector is the value code deepinfer/model_executor/models/v4_modules/attn.py:662, code deepinfer/hugging_face/config.py:289.

Layer 20's C1 KV looks like "no compression, so the same as the window". It isn't. It comes from a different projection, compressor.wkv, not attn.wkv, and the indexer picks from it code deepinfer/model_executor/models/v4_modules/mla_attn.py:63-97.

Why does the second half share one long-range KV instead of building its own in each layer?

Spreading work across chips

In each MoE layer, rank 0 sends each of its 24 rows to 6 experts. One fused kernel exchanges work with 7 GPUs.

One MoE layer at EP8

Rank 0's decode step has 4 requests × 6 rows = 24 rows. Each row goes to 6 routed experts plus the shared one.

The gate is at code deepinfer/model_executor/models/v4_modules/mlp.py:359-401 and the kernel call at :528-563. The gate maps each logical expert to a physical slot through the EPLB maps code deepinfer/model_executor/models/v4_modules/mlp.py:392-393. It also masks out padding rows, so they route nowhere code deepinfer/model_executor/models/v4_modules/mlp.py:387.

The engine uses Mega MoE when the whole EP group sits inside one NVLink domain code deepinfer/model_executor/models/v4_modules/mlp.py:184-188. Eight GPUs in one machine qualify. Otherwise it falls back to DeepEP's separate dispatch and combine code deepinfer/model_executor/models/v4_modules/mlp.py:565-616.

How many experts each GPU holds

The offline example asks for at least 32 redundant experts code examples/_config.py:78-79. The engine rounds the count up so the total divides by the EP size code deepinfer/model_executor/eplb/__init__.py:21-28.

• 384 + 32 = 416, and 416 / 8 = 52. No rounding needed.
• So each GPU holds 52 expert slots per layer: 48 of its own share plus 4 copies.

At start-up, experts 0, 1, 2 and so on receive copies, regardless of load code deepinfer/model_executor/eplb/__init__.py:31-39. The offline example re-plans every 30 steps from a 3-step window code examples/_config.py:80-83. Production uses 1,000 or 10,000 steps doc.

Counting the work

Per rank per layer, 24 rows × 6 experts = 144 (row, expert) pairs. If all eight ranks have 24 valid rows, that is 1,152 pairs over 416 slots, about 2.8 pairs per slot inference. Partially filled ranks and placeholder drafts reduce the routed work. Thread 4 in Pass 4 shows why this is the main argument for Wide-EP.

The receiving side counts arrivals per slot into expert_count code deepinfer/model_executor/models/v4_modules/mlp.py:474, :555. EPLB and moe_lb get their data from that count code deepinfer/model_executor/eplb/__init__.py:168-193.

The idle rank

Say rank 5 has nothing to decode this step. It still runs a step, because rank 0's Mega MoE call needs every rank's buffer. The engine marks the step as a placeholder code deepinfer/engine/llm_engine.py:972. The worker builds one fake sequence with one row and block 0 code deepinfer/worker/model_runner.py:386-403. That row has valid_token_mask = False, so the gate routes it nowhere code deepinfer/model_executor/layers/tilelang_ops/prepare_v4_decode_metadata.py:194. The rank still joins every collective, and its EPLB step counter stays in line with everyone else's code deepinfer/model_executor/models/deepseek.py:402-409.

Every rank waits for the slowest rank in each MoE layer, regardless of its own work.

At start-up, EPLB copies experts 0 to 31 in index order code deepinfer/model_executor/eplb/__init__.py:31-39. A translated comment says the first re-plan comes early, at a quarter of the interval, "to correct cold-start imbalance as early as possible" code deepinfer/model_executor/eplb/__init__.py:121-123.

The benchmark sets DEEPINFER_FORCE_LOAD_BALANCE=1 code examples/benchmark.py:165. This makes the gate route at random to distribute work evenly across experts code deepinfer/model_executor/layers/tilelang_ops/cuda/prepare_decode_metadata.py:130. It is a benchmarking tool, not EPLB.

Making decode faster

Each decode step for r1 starts before the CPU reads the previous one, fits a recorded 24-row graph and checks 5 drafts.

The first decode step

The previous step was a prompt step, which is post-processed at once. So nothing is in flight code deepinfer/worker/worker.py:182-185. The worker's correction step finds no previous state and does nothing code deepinfer/worker/worker.py:198-200.

r1's draft list holds 5 placeholder tokens from prefill code deepinfer/engine/llm_engine.py:706-711. Varlen verification is off, so every request gets 1 + 5 = 6 rows code csrc/engine/mtp_varlen.hpp:153-158. Four requests give total_q = 24.

Choosing the CUDA graph

A CUDA graph is a recorded run with fixed shapes. At start-up the engine records one per row count. The counts are powers of two up to an interval, then steps of that interval, up to the maximum code deepinfer/worker/cuda_graph_runner.py:41-51.

• The maximum is max_num_seqs × (1 + nextn) = 4 × 6 = 24 rows code deepinfer/worker/model_runner.py:76-79.
• The interval defaults to 8 code deepinfer/environ.py:155.
• So the graphs are for 1, 2, 4, 8, 16 and 24 rows.

select(total_q) picks the smallest graph at least as big as total_q code deepinfer/worker/cuda_graph_runner.py:143-148. 24 rows fit the 24 graph exactly, with no padding. The graph counts rows, not requests, because with varlen on, two batches of 4 requests can have very different row counts.

What the placeholder drafts do

Rows 1 to 5 carry the placeholder token, and the sampler checks them as drafts. Their draft probabilities come from slot 0 of the draft pool. That slot is "pre-filled so its draft log-prob puts all mass on the pad id" code deepinfer/model_executor/layers/mtp_info_pool.py:47-52. The main model gives the pad token almost no chance. With force top-k at 4,096, the pad token is almost surely outside the kept set inference. Its target logprob is then minus infinity, and the sampler rejects it code deepinfer/model_executor/layers/tilelang_ops/cuda/fused_sample.py:383-391.

So the first decode step gives r1 one token, the fresh one at row 0. Inside the same sample call, the draft model runs and writes 5 real drafts into r1's own slot of the draft pool code deepinfer/model_executor/layers/sampler.py:236-251, deepinfer/model_executor/models/mtp/dflash.py:339-344.

The second decode step, under async scheduling

This is the CPU/GPU overlap from Pass 1, now traced through the code. The diagram shows one step() call for decode step N+1, in CPU order.

In the code, the schedule call is at code deepinfer/engine/llm_engine.py:885. The worker queues the corrections at code deepinfer/worker/worker.py:209-233, launches forward at :122 and sample at :154. It waits at :144 and :263, and refreshes the metadata at :150-152.

The worker queues the next forward pass before the CPU waits for sample(N). GPU work can then overlap CPU post-processing. If preparing and launching the next step takes too long, the GPU can still wait inference from this ordering.

What the CPU doesn't know, and how the GPU fixes it

schedule(N+1) runs before process_outputs(N). So the scheduler doesn't know N's sampled token, N's new drafts, or how many drafts N accepted. It counts step N as fully accepted, by adding 6 to total_inflight_tokens code csrc/engine/scheduler.hpp:697, :726. The block count also reserves room for that code csrc/engine/sequence.hpp:291-292.

The worker builds a remap from this batch to the previous one by matching sequence ids code deepinfer/worker/worker.py:209-211, csrc/engine/zos.cpp:142-176. Two GPU kernels use it.

• One fixes each seqlen_k. It adds this step's rows and subtracts the drafts step N rejected code deepinfer/model_executor/layers/tilelang_ops/compute_seqlens_k_zos.py:18-24.
• One copies step N's sampled token and new drafts into this step's input rows code deepinfer/model_executor/layers/tilelang_ops/gather_input_ids_zos.py:21-30.

A translated docstring: "under ZOS the CPU-side seqlens_q and seqlens_k lag one step. When the batch overlaps the previous step's batch, they must be fixed on the GPU from the previous step's sampling results" code deepinfer/worker/worker.py:197.

Verifying 5 drafts

From step 2 on, r1's 6 rows are its newest token and 5 real drafts. The sampler accepts drafts by rejection sampling code deepinfer/model_executor/layers/tilelang_ops/cuda/fused_sample.py:383-394. A draft counts only if every earlier draft of the same request passed code deepinfer/model_executor/layers/tilelang_ops/sampler.py:243-247. Thread 6 shows the arithmetic.

After the step, process_outputs appends the accepted drafts and the one sampled token. Then it reduces num_tokens to the accepted count. A translated comment: "decode: backfill cache.num_tokens from the optimistic value to the number actually accepted, so unaccepted drafts don't occupy KV" code csrc/engine/scheduler.hpp:997-1000.

Three ways to run the same decode step

The CPU must fix every shape before launch. The GPU supplies any input values the CPU does not know yet.

A placeholder step and graph padding both add fake rows, but they have different jobs. A placeholder step lets an idle rank join the collectives. Padding makes a batch fit a recorded graph size code deepinfer/worker/model_runner.py:187.

Async scheduling lets the CPU run exactly one step ahead. The worker keeps a single prev_step_state code deepinfer/worker/worker.py:73, :187. Every prompt step drains it, because prompt steps are post-processed at once code deepinfer/worker/worker.py:183-185.

Why does sampling for N+1 wait until after post-processing N, when the forward pass didn't?

Getting the right answer every time

For r1, the sampler applies a fixed sequence of operations to the logits. Seed 1 and the position determine each draw.

The sampling pipeline, in order

For r1's first output token, in the prefill step:

1. Pick the rows. Only the last row of r1's chunk code deepinfer/model_executor/layers/sampler.py:140-141.
2. Compute logits. The hidden state times the output embedding, into 32-bit floats code deepinfer/model_executor/layers/sampler.py:142, deepinfer/model_executor/layers/sampler_utils.py:20-26.
3. Apply the token blacklist, empty by default code deepinfer/model_executor/layers/sampler.py:169.
4. Ban the end token for the first output, because of first_n_token_no_eos=1 code deepinfer/model_executor/layers/sampler.py:184-185.
5. Apply a grammar mask, if the request has one. r1 doesn't code deepinfer/model_executor/layers/sampler.py:187-188.
6. Force top-k 4,096. The logits shrink to their top 4,096 columns before the main kernel code deepinfer/model_executor/layers/sampler_utils.py:166-181.
7. Temperature 0.4. The kernel scales (logit − max) by 1 / 0.4 code deepinfer/model_executor/layers/tilelang_ops/cuda/fused_sample.py:180-215.
8. Top-p and top-k. r1 uses the defaults, top-p 1.0 and top-k −1, so neither filters code deepinfer/sampling_params.py:49-51, deepinfer/model_executor/layers/tilelang_ops/cuda/fused_sample.py:239-245.
9. Write the logprobs after all of the above, renormalised code deepinfer/model_executor/layers/tilelang_ops/cuda/fused_sample.py:424-434.
10. Draw the token from a stream keyed by the seed and the position code deepinfer/model_executor/layers/tilelang_ops/cuda/fused_sample.py:449-465.

The random key

Every draw is a pure function of the seed and a stream number. The stream number is 3 × position + case. Case 0 is a fresh sample, case 1 is a resample after a rejection, and case 2 is the accept-or-reject coin code deepinfer/model_executor/layers/tilelang_ops/cuda/fused_sample.py:392, :450, :481.

• For a prefill row, the position is seqlen_k, the index of the token being made code csrc/engine/prepare_sample.cpp:198. For r1's first token, seqlen_k = 16,389. So the stream is 3 × 16,389 + 0 = 49,167 with seed 1.
• For a decode row, the first row's position is seqlen_k − seqlen_q + 1 code csrc/engine/prepare_sample.cpp:206. In r1's first decode step that is 16,395 − 6 + 1 = 16,390, the index of the token row 0 makes.
• Later rows in the same request add j × 10,000,019 for row j code deepinfer/model_executor/layers/tilelang_ops/sampler.py:65, deepinfer/model_executor/layers/tilelang_ops/common.py:52.
• Draft tokens use the seed plus 1,000,000,007 code deepinfer/model_executor/layers/tilelang_ops/fused_sample.py:15, deepinfer/model_executor/models/mtp/dflash.py:290-294.

The batch row is not in the key, so r1 gets the same draw first or fourth in the batch. The code's contract: "Deterministic in (seq_seeds, sample_positions, is_nextn_predict): independent of batch size and row order" code deepinfer/model_executor/layers/tilelang_ops/random_sampling.py:19-20.

Probabilities as integers

The sampler turns each probability into an integer: the probability times 2³¹, rounded down code deepinfer/model_executor/layers/tilelang_ops/cuda/random_sampling.py:38-50. Integer sums come out the same in any order. So the threads that add up probability mass can run in any order and still agree code deepinfer/model_executor/layers/tilelang_ops/cuda/fused_sample.py:119-125. Using integers for this sum avoids the floating-point ordering problem from Pass 1.

Accepting a draft

Pass 1 did a full rejection check by hand. These are the kernel's steps, with the last one left for you code deepinfer/model_executor/layers/tilelang_ops/cuda/fused_sample.py:383-394, :436-444.

1. Take the target logprob t of the draft token and the draft's own logprob d.
2. The accept logprob is min(t − d, 0). A draft outside the kept top-k set has t = −∞, so it always fails.
3. Draw a uniform number from stream 3 × position + 2. Accept if it is below exp(min(t − d, 0)).
4. Your turn. Say t = log 0.3 and d = log 0.6, and the draw is 0.42. Does the draft pass? If not, what does the kernel sample from next?

What batch invariance changes for r1

In this run batch invariance changes four things code deepinfer/worker/runtime_config.py:67-69, deepinfer/model_executor/models/v4_modules/attn.py:87-92, :666, deepinfer/model_executor/parallel_utils/userbuffer.py:79.

1. The matrix-multiply libraries switch to deterministic algorithms.
2. FlashMLA's decode attention runs in batch-invariant mode.
3. The engine switches off a fused prefill attention kernel on some layers, so prefill takes the same path as decode.
4. With EPLB on, the MoE combine sums in one fixed order.

The kernel libraries were empty when I wrote this pass, so I could not inspect how those switches worked. I guessed they fix how GPU blocks divide a long sum. Pass 4 opens the kernels and confirms it, with one twist: for V4.1's attention shape the split was already off.

DeepInfer returns logprobs after steps 3 to 8 above, then renormalises them code deepinfer/model_executor/layers/tilelang_ops/cuda/fused_sample.py:369, :424-434. For r1's first token the end token's logprob is minus infinity, because step 4 banned it. Perplexity measures the probabilities assigned to a reference sequence. Here it measures the processed distribution. It is not raw-model perplexity unless temperature, filtering and masks leave that distribution unchanged. Matching settings makes regression comparisons meaningful, as in the offline test.

When the request doesn't go to plan

An abort only changes a request’s status. Its queue determines which resources the engine frees later.

A 50,000-token prompt aborted between its chunks

A NORMAL engine has max_num_batched_tokens=32768, chunked prefill on, and one 50,000-token prompt called P. Nothing else waits.

1. Prefill step k. P's block table is empty, so the scheduler asks for ceil(50,000 / 128) = 391 blocks code csrc/engine/scheduler.hpp:386, csrc/engine/sequence.hpp:291-292. It cuts P to 32,768 tokens but reserves all 391 blocks and one state slot code csrc/engine/scheduler.hpp:470-475, :518-520.
2. After step k. num_tokens = 32,768 < 50,000, so P goes back to the front of waiting. It is still WAITING and still holds 391 blocks code csrc/engine/scheduler.hpp:1011-1019.
3. The abort. abort_request("P") calls Scheduler::abort_sequence. It scans running and waiting, finds P, sets FINISHED_ABORTED and returns code csrc/engine/scheduler.hpp:144-163.
4. The next schedule. _schedule_prefill moves the aborted P into ignored_seqs code csrc/engine/scheduler.hpp:363-367, or the decode branch sweeps it out of waiting code csrc/engine/scheduler.hpp:569-578.
5. The output. build_request_outputs emits one output per ignored sequence, with finish reason "abort" code csrc/engine/outputs.hpp:297-303.

I cannot find a line that frees P’s 391 blocks. The scheduler only pushes into ignored_seqs, and only the output builder reads it code csrc/engine/scheduler.hpp:365, :382, :394, :408, :572, :780, csrc/engine/outputs.hpp:297. Two push paths free the blocks first. One is an over-long fork code csrc/engine/scheduler.hpp:407. The other is a preempted request in stop mode code csrc/engine/scheduler.hpp:779. The abort path doesn't, and pre_sampler_sync frees only finished_seqs code deepinfer/engine/llm_engine.py:844-846.

So I think an abort between chunks leaks the blocks and the state slot until the engine restarts inference. A comment suggests why this path may be uncommon: "We do NOT abort in PE" code csrc/engine/scheduler.hpp:150-151. The release calls abort_request only for failed images, and those requests never got blocks code deepinfer/engine/llm_engine.py:942-945. No test covers abort.

Where the request sat when the abort came

process_outputs skips an aborted running request code csrc/engine/scheduler.hpp:905-915, and free_finished_seqs moves it to finished_seqs code csrc/engine/scheduler.hpp:562. build_request_outputs walks only scheduled_seqs and ignored_seqs, so no output says "abort" code csrc/engine/outputs.hpp:94-303. An online server has to report the abort itself. Pass 5 asks you to design that.

In the first two rows the output token is get_last_token_id(). If nothing was generated, that is the last prompt token code csrc/engine/outputs.hpp:297-298. A client that ignores the finish reason would show it as model output.

An abort is only a mark, so each queue that can hold the request needs its own cleanup path.

abort_request only changes the request’s status code csrc/engine/scheduler.hpp:153. A later schedule frees memory only if the request’s queue has a cleanup path.

A request finishes while async scheduling has it in flight

Request Q is in a decode batch with ZOS on. At step N, Q samples the end token, but the CPU doesn't know that when it schedules N+1. The timeline is my reading of the code.

• At post-process(N), _check_stop sets FINISHED_STOPPED and the scheduler releases Q's draft-pool slot code csrc/engine/scheduler.hpp:986-994, :1035-1041.
• sample(N+1) keeps the slot it captured. The comment: "Keep the slot captured for an in-flight sample after ZOS releases the request" code csrc/engine/sequence.hpp:411-414.
• When the sampler emits a stop token, it zeroes Q's next draft confidences, so the wasted step takes no verify budget. Translated: "the request hit EOS, set the confidence of all its next-round draft tokens to 0" code deepinfer/model_executor/layers/sampler.py:252-254.
• At post-process(N+1), Q is finished, so process_outputs pushes 0 tokens code csrc/engine/scheduler.hpp:905-915.
• Q's blocks come back in pre_sampler_sync of step N+2, after forward(N+2) is queued code deepinfer/engine/llm_engine.py:844-846. A new owner writes them only in stream order after N+1 inference.
• DeferredRelease holds a step's pooled objects one round longer. Translated: "release a step's resources one round late, so the caller has finished reading them" code deepinfer/engine/llm_engine.py:67-71.

A finishing request can remain in one extra forward step. Fixed-5 verification uses six rows. Variable-length verification can reduce its draft allocation inference.

Why does the scheduler reserve all 391 blocks of P on the first chunk, when only 256 hold data after it?

When memory runs out

Prefill waits for free blocks. Decode can't wait, so it takes blocks from the running request with the longest output.

Four running requests and an empty free list

A rank has four running requests and no free blocks. Outputs so far: A 3,000 tokens, B 1,200, C 400, D 50. _schedule_decode does this code csrc/engine/scheduler.hpp:622-646.

1. Sort running by output length, longest first: [A, B, C, D] code csrc/engine/scheduler.hpp:624-625.
2. Set l = 0 and r = 4. Serve from the back, so D first.
3. No block is free, so preempt from the front, which is A code csrc/engine/scheduler.hpp:630-633. l becomes 1.
4. D gets its slot from A's freed blocks through _append_slot code csrc/engine/scheduler.hpp:635-637. r becomes 3.
5. C, then B, get slots the same way.
6. l == r, so the loop ends. D, C and B run. A is out.

Step 3 preempts whenever the free list is empty, even if D needs no new block code csrc/engine/scheduler.hpp:630. I think this happens in practice inference. The victim has generated the most output so far. It may still be far from finishing.

The function's header comment says the opposite. Translated: "pick this round's decode seqs by LCFS, and when blocks are short, preempt the seq with the shortest output" code csrc/engine/scheduler.hpp:556-557. The sort is descending and victims come from the front, so the code takes the longest. The finish-reason comment agrees with the code. Translated: "if the KV cache is about to overflow, we forcibly stop the task that occupies the most KV cache" code csrc/engine/sequence.hpp:55-56. This explanation follows the code.

What happens to A

_preempt has two modes code csrc/engine/scheduler.hpp:766-790.

The defaults are at deepinfer/engine/arg_utils.py:40 and deepinfer/entrypoints/llm.py:68, and the finish reason text at csrc/engine/sequence.hpp:55-56. A decode-only engine has no prefill to recompute with, so it must stop code csrc/engine/scheduler.hpp:101-104. A grammar's state can't be replayed, so those requests stop too. Translated: "SO request preemption doesn't support recompute yet, finish it directly" code csrc/engine/scheduler.hpp:769-780.

Recompute has two hidden costs. A's re-prefill waits for the next vote point and for the cache to drop to 90% code csrc/engine/scheduler.hpp:313. With YOCO truncation on, the rerun can also give a different result. A comment allows it: "with yoco truncation enabled, result changes caused by preemption are allowed" code deepinfer/engine/llm_engine.py:888.

SWAPPED is in the status enum but nothing sets it code csrc/engine/sequence.hpp:33. There is no swap to host memory, and the Q&A doc confirms decode doesn't offload KV doc.

The engine preempts the request with the most output, regardless of the header comment. I think the aim is to free the most memory per victim inference, at the cost of the most lost work.

Why stop the victim by default, instead of recomputing it?

A prefix hit with a missing window

This case is described, not in the code. The production Entrypoint matches a new prompt against the KV store before routing doc. These are the Q&A doc's hit rules for a prompt of L = 10,000 tokens.

1. The cap. The hit is at most L − 134 = 9,866 tokens. The 134 is the 128-token window plus 6 for draft rollback doc.
2. Block KV. The store holds whole 128-token blocks, so the block hit is 77 × 128 = 9,856 doc.
3. The window. Window state exists only at multiples of 128. If one exists between 9,728 and 9,856, the hit is 9,856. If not, back off one block to 9,728 doc.
4. The replay. On a miss, prefill starts at 9,728 and reuses the block KV for 9,728 to 9,855. Those 128 tokens rerun only to rebuild the window, with queries limited to keys from 9,728 on doc.

A missing window costs prefill on 272 tokens instead of 144. It adds 128 tokens, regardless of prompt length doc.

The release has the window limit in code. swa_valid_start clamps every query's window to start no earlier than it code csrc/engine/input_metadata/prepare_tensor_input_v4.cpp:150-155. The load plan zero-fills the ring below it code csrc/engine/input_metadata/prepare_tensor_input_v4.cpp:380-383. Nothing in the release sets it above 0, but Python can, for a store to use code csrc/engine/pybind_sequence.cpp:132.

Storing the window at 128-token points

Also described, not in the code. The production store keeps window state only at multiples of 128 doc. A decode step adds up to 6 tokens, so it can cross a multiple of 128 mid-step.

Say the boundary is b. The window ending at b needs positions b − 128 to b − 1. The step might reach b + 5 before the engine copies the ring. The ring then holds the latest 134 positions, b − 128 to b + 5, so the oldest needed position just survives. That fits the Q&A doc's claim that an overshoot of up to 5 tokens fits in the 6 extra slots doc. The arithmetic is mine inference.

The release has the pieces but doesn't use them. KVSTORE_GRAIN = 128, translated: "persist only where the token count is a multiple of 128" code deepinfer/worker/cache_engine.py:32-33. A helper asserts that a stored position is a multiple of 128 code deepinfer/worker/cache_engine.py:787, but only tests call it. The hook that would trigger the decode copy returns False code deepinfer/engine/llm_engine.py:811-813. Tests such as TestStateCacheOperationsWithExtraState check that store and load round-trip code tests/worker/test_state_cache_ops.py:121, :189.

What the engine does when memory is short

A new prompt can wait for memory. Decode frees memory from another request so the remaining requests can continue.

When prefill takes a shortcut

With YOCO truncation, layers 20 to 39 run only the last prompt tokens. That saves work but changes the answer slightly.

A 1,000-token prompt with compact truncation

Set DEEPINFER_YOCO_PREFILL_TRUNCATION_MODE=compact and DEEPINFER_YOCO_PREFILL_KEEP_TOKENS=140, the README's settings code README.md:25. A 1,000-token prompt arrives in one chunk.

1. Layers 0 to 19 run all 1,000 rows.
2. The keep window starts at data_len − keep_tokens = 860 code csrc/engine/input_metadata/prepare_tensor_input_v4.cpp:220-225.
3. Layer 20's KV still gets all 1,000 tokens, because the boundary step keeps the full hidden state code deepinfer/model_executor/models/v4_modules/yoco.py:93-123, deepinfer/model_executor/models/decoder_layer.py:134-142.
4. Layers 20 to 39 run only the 140 rows from 860 to 999.
5. Their windows start no earlier than 860. A query at position 900 sees keys 860 to 900, 41 of its usual 128 code csrc/engine/input_metadata/prepare_tensor_input_v4.cpp:155.

The code requires KEEP_TOKENS >= window_size + max_backtracking_tokens, which is 134 code csrc/engine/input_metadata/prepare_tensor_input_v4.cpp:354-357. The translated comment: "after rollback there must still be window-size tokens". So the Q&A doc's "the last 128 tokens" is the idea, and 134 is the real minimum.

The first decode query, at position 1,000, needs keys 873 to 1,000 in each second-half window. Those entries exist, but they came from rows that saw clipped windows. So the window KV is approximate, and so is the output. The Q&A doc calls it "an approximate SWA KV cache" doc.

Compact mode can't compute the dropped rows, so prompt logprobs are off for every request code deepinfer/engine/llm_engine.py:529-538. A chunk wholly before the keep window still keeps 1 placeholder row code csrc/engine/input_metadata/prepare_tensor_input_v4.cpp:221-225. So the first chunk of the earlier 50,000-token prompt would run 1 row through layers 20 to 39.

The three truncation modes

The modes are documented at deepinfer/environ.py:188-194. mask computes every row so a test can compare it with compact. examples/yoco/yoco_equiv.py does that, but never compares with off code examples/yoco/README.md. No test in the release measures how far compact drifts from off, which needs a GPU.

If a cheaper prefill changes what the second half sees, it changes the answer. Comparing compact and mask modes checks whether they agree. It does not check whether they match truncation off.

YOCO truncation is approximate. Prompt logprobs are therefore off, and preemption may change results. The Q&A doc says it "nearly halves total prefill computation" doc. Pass 4 asks what that is worth.

Can one token appear twice in a softmax?

This is an open question, not a known edge case. In layers 20 to 39, the long-range KV has an entry for every token, including the last 128. The window covers those 128 too. So a query could see token t twice in one softmax. It sees it through the window, and through the C1 entry if the indexer picks it.

The two vectors come from different projections, attn.wkv and compressor.wkv, so they differ code deepinfer/model_executor/models/v4_modules/mla_attn.py:63-97. I found no step that removes the overlap. The arithmetic-work estimator counts window and top-k keys separately code tests/engine/test_flops_estimator.py:82-89. My guess is that it's by design and the model learned to use both inference. Logging the indices on a real run would settle it.

Why must KEEP_TOKENS be at least 134 and not just 128?

When one expert gets too much

EPLB gives hot experts extra copies, then places the copies to reduce the busiest GPU’s load. Both stages are greedy: they choose the best next move without searching every complete layout.

A full EPLB run by hand

I ran the real rebalance_experts_global from deepinfer/model_executor/eplb/policy.py:73, which is plain PyTorch, on my CPU. Its output matches the calculations below.

Setup: 5 logical experts, 2 GPUs, 3 slots per GPU, so 6 physical slots and 1 redundant copy. The window has 2 steps.

The starting layout. Slots 0 to 4 hold experts 0 to 4. The spare slot copies expert 0, because start-up copies go to experts in index order code deepinfer/model_executor/eplb/__init__.py:31-39. GPU0 = {e0, e1, e2}, GPU1 = {e3, e4, e0}. e0's tokens split in half.

moe_lb = (115 + 125) / (100 + 100) = 1.20

Stage A: who gets the copy. For each candidate expert, calculate the result of adding one copy. Find the hottest single copy of any expert at each step, and sum over the steps. The candidate with the smallest sum wins code deepinfer/model_executor/eplb/policy.py:105-123.

• Copy e0: each e0 copy carries 45, so the hottest copy is 45 at both steps. Sum 90.
• Copy anything else: e0 still carries 90 at both steps. Sum 180.
• So e0 gets the copy, which the starting layout already gave it.

Stage B: where each copy goes. Sort the copies by load per copy, heaviest first: e0, e0, e1, e2, e3, e4. Ties keep slot order code deepinfer/model_executor/eplb/policy.py:31-35. A copy's load per step is its expert's tokens divided by its copy count.

Place each copy on the non-full GPU that gives the smallest sum, over steps, of the busiest GPU. Ties go to the lower GPU number code deepinfer/model_executor/eplb/policy.py:53-69.

For example, e2 on GPU1 leaves both GPUs at (65, 85). The peaks are 65 and 85, so the score is 150.

moe_lb = (105 + 105) / 200 = 1.05

e1 and e2 are both busy at step 1, so the policy split them. It paired e1 with e3, which is busy at step 0. Scoring each step's peak pairs experts whose busy steps don't overlap. The test file calls this "load complementarity across batches" code tests/models/test_eplb_policy.py:112.

When greedy makes it worse

Same setup, different tokens.

The starting layout {e0, e1, e2} and {e3, e4, e0} gives moe_lb 1.057. The real policy, run on my CPU, returns {e2, e4, e0} and {e3, e1, e0}, which gives 1.114. It made the balance worse on the very window it planned from.

The policy starts from scratch every time and ignores the current layout code deepinfer/model_executor/eplb/policy.py:98-103. _rearrange_experts then applies the new plan without checking that it beats the old one code deepinfer/model_executor/eplb/__init__.py:256-278. Each rearrangement also stalls every rank while weights move code deepinfer/model_executor/eplb/__init__.py:222.

The tests check that the maps are consistent, that the policy is deterministic, and some hand-built cases code tests/models/test_eplb_policy.py:14-117. None checks that the result beats the input.

EPLB’s greedy heuristic can produce a worse balance than the layout it replaces. The tests do not establish how often this happens with production loads. Nothing in the code guards against it.

A very hot expert

Stage A has no cap on copies per expert. It keeps copying whichever expert still has the hottest single copy code deepinfer/model_executor/eplb/policy.py:110-123. The tests include more copies than logical experts, and a "single spike" distribution code tests/models/test_eplb_policy.py:14-24, :65-79.

Copies only help if the gate splits the hot expert's tokens between them. The policy assumes an even split, since it divides by the copy count code deepinfer/model_executor/eplb/policy.py:31, :37. The gate kernel does the split, using the logical-to-physical map code deepinfer/model_executor/models/v4_modules/mlp.py:392-393. Pass 4, Thread 4 shows it splits by a fixed stride over the token's row in the batch. This spreads rows across copies without measuring their current loads. The split still depends on which rows choose the expert.

A GPU fails

The Q&A doc says a decode instance restarts on however many healthy GPUs it gets, and uses that as its EP size doc. The release isn't elastic, but the redundant-expert rounding handles odd sizes code deepinfer/model_executor/eplb/__init__.py:21-28. Take EP63 with a lower bound of 63 copies.

1. 384 + 63 = 447.
2. 447 / 63 = 7 remainder 6, so the padding is 63 − 6 = 57.
3. Extra copies = 63 + 57 = 120. Total slots = 504 = 8 × 63.

Each surviving GPU holds 8 slots, up from 7. It needs about 14% more expert memory inference from the formula.

Why does stage B score the busiest GPU at each step, instead of balancing the total over the window?

When the verifier says no

Rejected drafts have already written their KV. The engine moves its count back and lets later steps overwrite that KV.

The verifier rejects draft 3 of 5

Request R has 2,000 tokens, with KV for the first 1,999. Its newest token, at position 1,999, has no KV yet. The next decode step has 6 rows.

The counters move like this code csrc/engine/scheduler.hpp:927, :997-1000, csrc/engine/sequence.hpp:216-221.

1. num_tokens starts at 1,999.
2. advance_optimistically(6) sets previous_num_tokens = 1,999 and num_tokens = 2,005.
3. The step appends d1, d2 and t, so token_cnt = 3, and len becomes 2,003.
4. commit_processed_tokens(3) sets num_tokens = 1,999 + 3 = 2,002.

KV now counts as ready through position 2,001, including t1999, d1 and d2. The KV at 2,002 to 2,004 belongs to rejected drafts d3, d4 and d5. The engine leaves it in place. The next step's row 0 is t at position 2,002, which overwrites it, and later steps overwrite 2,003 and 2,004.

_append_slot also returns any block the drafts claimed but no longer need code csrc/engine/scheduler.hpp:821-825. Under async scheduling, a GPU kernel subtracts nextn − accepted = 5 − 2 = 3 from the next step's seqlen_k code deepinfer/model_executor/layers/tilelang_ops/compute_seqlens_k_zos.py:18-24.

Why the ring needs 6 spare slots

Paged blocks are safe because each position has its own address. The window ring reuses slots, slot = position mod 134 code csrc/engine/input_metadata/prepare_tensor_input_v4.cpp:103-104.

ring of 134: position 2004 writes slot 2004 mod 134 = 128, and covers position 2004 - 134 = 1870
next query:  t at 2002 needs window 1875 .. 2002          -> 1870 is not needed, safe

ring of 128: position 2004 would write slot 2004 mod 128 = 84, and cover position 2004 - 128 = 1876
next query:  t at 2002 needs window 1875 .. 2002          -> 1876 is needed, but now holds d5's KV

With a ring of 128, the rejected d5 would corrupt the very next query's window. With 134, a step can write 6 positions ahead without touching a slot a later window needs. The ring size is window_size + DEEPINFER_MAX_BACKTRACKING_TOKENS, and the second is at least nextn + 1 code deepinfer/worker/cache_engine.py:88-93, deepinfer/environ.py:132. C2 compression state gets compress_ratio + (6 − 1) slots for the same reason code deepinfer/worker/cache_engine.py:130-133.

The draft model's window cache uses the same rule. Translated: "the positions of rejected drafts are exactly where later draft writes go. They are overwritten before any read and need no masking" code deepinfer/model_executor/models/mtp/dflash.py:204-209.

examples/benchmark.py sets DEEPINFER_MAX_BACKTRACKING_TOKENS=1 code examples/benchmark.py:164. The engine clamps it to at least 6, so the setting has no effect with nextn 5 inference from deepinfer/environ.py:132.

A rejected draft’s KV can remain in place without affecting the next step. The count moves back, and the next writes cover it before anything reads it. The 6 spare ring slots make sure of that.

37 rows and a 40-row graph

With max_num_seqs=64 and nextn 5, the largest graph is 64 × 6 = 384 rows. With interval 8 the sizes are 1, 2, 4, 8, 16, 24, then every 8 up to 384 code deepinfer/worker/cuda_graph_runner.py:41-51.

A step has 37 rows, so select(37) returns the 40-row graph code deepinfer/worker/cuda_graph_runner.py:143-148. The worker adds 3 one-row placeholder sequences code deepinfer/worker/model_runner.py:187. Each has request id "placeholder", block table [0], and draft-pool slot 0, which the store kernel skips code deepinfer/worker/model_runner.py:386-403. Their valid_token_mask = False, so the gate routes them nowhere code deepinfer/model_executor/layers/tilelang_ops/prepare_v4_decode_metadata.py:194. The sampler trims them before it computes logits code deepinfer/model_executor/layers/sampler.py:140-141.

Three of the 40 rows, 7.5%, are padding. That is a share of rows, not a measured share of elapsed time. With interval 8, at most 7 rows are ever padded.

Choosing how many drafts to check, worked in full

This is the full example of varlen verification. Pass 4 leaves out a step, and Pass 5 leaves it all to you. I chose the numbers for this example. The algorithm is MTPVarlen::schedule code csrc/engine/mtp_varlen.hpp:150-267.

Four requests W, X, Y and Z, with nextn 3. The confidence head supplies these acceptance chances. Each value already includes the chance that all earlier drafts pass:

The engine measures step time per graph size at start-up. Say it gives 8.1 ms for 4 rows, 8.6 ms for 8 and 12.0 ms for 16. Add 0.4 ms for the sampler. Sizes 1 and 2 don't matter here.

1. Sort all 12 chances, highest first, with ties going to the lower request code csrc/engine/mtp_varlen.hpp:210-214.

   W .95, W .90, X .85, W .80, X .60, Y .50, X .30, Z .30, Y .20, Y .10, Z .10, Z .05
   

2. Start with 4 rows, one per request, and 4 expected tokens. A request's own row always yields a token code csrc/engine/mtp_varlen.hpp:228-231.

3. Add drafts in order. After each, work out expected tokens per millisecond on the smallest graph that fits code csrc/engine/mtp_varlen.hpp:235-258.

4. Pick the best: 8 rows, so 4 extra rows.
5. Hand them out to the top 4 drafts: W, W, X, W. W checks 3 drafts, X checks 1, Y and Z check none code csrc/engine/mtp_varlen.hpp:262-266.

Now change one number. Say the 16-row graph took 10.4 ms instead of 12.0. Then 16 rows give 9.65 / 10.8 = 0.894 tokens per ms, the best in the table. Every request would check all 3 drafts. I checked both versions with a short script.

Both the estimated cost curve and confidence determine the chosen length. Within that estimate, rows fitting the same graph share one time estimate, so the scan tends to fill graph sizes. Actual routed work can still change the runtime.

Under async scheduling, the CPU fixes only the total, 8 rows here. A GPU kernel re-ranks with the newest confidences and may give the 4 extra rows to other requests code deepinfer/model_executor/layers/tilelang_ops/cuda/schedule_mtp_varlen.py:26-81, deepinfer/worker/worker.py:218-219. The total can't change, because it chose the graph shape.

Why does the CPU use confidences from two rounds back to pick the total?

When the same prompt must give the same answer

The batch-invariance test runs a prompt alone and in a batch of 64. It forces the same continuation and compares exposed tokens and logprobs exactly.

Two batch shapes, from the batch invariance test

tests/models/test_batch_invariant.py states what the authors think must stay true code tests/models/test_batch_invariant.py:52-105.

1. Start an EP8 engine with batch invariance on, nextn 5, max_num_seqs=64, max_num_batched_tokens=16384 code tests/models/test_batch_invariant.py:57-68.
2. Make 64 prompts of random non-special tokens, with lengths around 2,500 code tests/models/test_batch_invariant.py:73-86.
3. Shape A. Run all 64 in one generate call: temperature 1.0, top-p 0.7, 50 tokens, top 5 logprobs code tests/models/test_batch_invariant.py:88-97.
4. Shape B. Run each prompt alone, forcing shape A's tokens as golden_token_ids code tests/models/test_batch_invariant.py:99-102.
5. Assert that token ids, logprobs, top token ids and top logprobs are equal with Python == code tests/models/test_batch_invariant.py:103-105. This checks those exposed values exactly, not every intermediate tensor or unforced stochastic generation.

Forcing the tokens removes the random choice, so the test compares numerical results. My reading of why: without forcing, one flipped token changes every later token. The test could not identify where the numerical results first differed inference.

The other 7 DP ranks run their own prompts at the same time. So even "alone" shares every MoE layer with other ranks' tokens, and the test covers that mixing too inference.

What makes it pass

DEEPINFER_BATCH_INVARIANT=1 flips the four switches Pass 2 listed. Two more pieces are invariant without the switch.

• The sampler. One GPU thread block per row, fixed reduction orders, and integer sums for the probability mass code deepinfer/model_executor/layers/tilelang_ops/cuda/fused_sample.py:119-125, :186-191. Its test checks that a permuted batch gives permuted results, and that each row alone matches the batch code deepinfer/model_executor/layers/tilelang_ops/random_sampling.py:310-332.
• The vision encoder. Its test checks invariance with the switch off code tests/models/test_ve_batch_invariant.py:1-14.

tests/models/test_kl.py compares step-by-step decode with a full prefill along the same tokens. Prompt logprob lists must be equal. For generated positions it checks a mean discrepancy with max_kl = 0.0 code tests/models/test_kl.py:190-206. The helper evaluates mean(exp(d) - 1 - d) in double precision, where d is the logprob difference code deepinfer/testing.py:4-11. Floating-point cancellation can make that expression zero for unequal inputs. This is strong regression evidence for the tested configuration, but it does not prove bitwise equality of all decode and prefill values. Batch-invariant mode also disables a fused attention path used by prefill on some layers code deepinfer/model_executor/models/v4_modules/attn.py:87-92.

The case that breaks it: varlen on

With DEEPINFER_ENABLE_MTP_VARLEN on, I expect the test’s logits to match but unforced tokens to differ inference.

1. The total row count comes from step times measured on the GPU every step code csrc/engine/mtp_varlen.hpp:139-148, deepinfer/engine/llm_engine.py:995-1002.
2. Each request's share comes from a ranking across the whole batch code csrc/engine/mtp_varlen.hpp:262-266.
3. A request's row count decides which random stream makes which token. Row j uses position P + j × 10,000,019, and the last row draws fresh instead of flipping an accept coin code deepinfer/model_executor/layers/tilelang_ops/sampler.py:24, :65.

Varlen can target the same distribution, within the sampler's numerical approximations, without producing the same tokens. The README's recipe for repeatable output turns it off code README.md:28. The evaluation script turns it on, so that recipe does not guarantee exact repeats for its runs code eval/run_eval.py:120-121.

A grammar that leaves no way out

Structured output uses xgrammar code deepinfer/structure_output/compiler.py:6-7. Each step, the CPU advances the grammar matcher and builds a bitmask of allowed tokens code deepinfer/model_executor/layers/sampler.py:352-376.

The grammar ends. After the matcher accepts the end token, masks allow everything, so draft rows past the end aren't wrongly blocked code deepinfer/structure_output/matcher.py:154-156, :189-190. The scheduler stops the request at the end token code csrc/engine/scheduler.hpp:986.

The grammar breaks. If the matcher can't fill a mask, the sampler marks only that request as failed. Translated: "raising would take down every request plus an NCCL barrier timeout. An SO bug kills only this request" code deepinfer/model_executor/layers/sampler.py:377-386. The scheduler appends token 0 and sets FINISHED_SO_FAILED code csrc/engine/scheduler.hpp:974-980. A client that ignores the finish reason sees a stray token 0.

Every token is blocked. If the grammar mask, the token blacklist and the first-token end ban together block every token, the row is all minus infinity. The sampling kernel then asserts "Invalid logits: contains NaN or +inf, or all entries are -inf" code deepinfer/model_executor/layers/tilelang_ops/cuda/fused_sample.py:218. A device assert usually kills the whole process inference. I found no earlier check for this mix.

The draft model ignores the grammar. The verifier rejects forbidden drafts, because their masked target probability is 0 code deepinfer/model_executor/layers/sampler.py:433-437. So I expect fewer drafts to pass under a strict grammar inference.

Sampling and verification

Pass 1 worked this in full and Pass 2 left a gap. Do this one alone.

Vocabulary of 5 tokens, logits 3.0, 2.0, 2.0, 0.0, −1.0. Temperature 1.0, top-p 0.8, top-k −1. The draw is 0.55. Which token does the sampler pick? Then: a draft token is token 2, with draft probability 0.5, and the accept coin draws 0.75. Does the draft pass?

DeepInfer turns temperature 0 into 0.00001. The sampler still draws, rather than always picking the top token code deepinfer/sampling_params.py:76-79. Two tokens whose logits differ by 0.00001 get odds of about 73% and 27%. Without a seed, each request gets a random one code deepinfer/sampling_params.py:86-90. That seed is fixed when the SamplingParams object is built, then deep-copied per request code deepinfer/entrypoints/llm.py:151. So one shared object gives every request in a generate call the same seed. Ten identical prompts then give ten identical answers inference from those lines.

Why does the sampler turn probabilities into integers before it adds them up?

One scheduler per rank, and a vote

A combined engine chooses one phase for its expert-parallel group. During prefill, requests already decoding in that group wait.

Each DP rank runs its own C++ scheduler. Per step, the ranks use tiny all-reduces to agree on the kind of step and whether to carry on code csrc/engine/scheduler.hpp:316-332, deepinfer/engine/llm_engine.py:665-679. A step is all prefill or all decode code csrc/engine/prepare_sample.cpp:61.

What it costs in a combined engine

Consider a NORMAL engine running both phases. The decode chart gives about 13 ms at EP64 with 64 requests per rank doc. The prefill chart gives 376 ms for a 32,768-token step at EP4 doc. These are different configurations. Combining them below illustrates the scheduling cost. It does not measure a combined engine. The decode estimate is 246 tokens/s/user divided by 3.25 tokens/step, about 76 steps/s inference.

1. Waiting for the vote. A ready request that misses an ALL point can wait another 70 decode steps for the next one. That is about 70 × 13 ms ≈ 0.9 s code deepinfer/environ.py:174. An ANY point every 100 decode steps offers a backstop, about 1.3 s at this step time code deepinfer/environ.py:177.
2. The prefill step. A 32,768-token prefill step at EP4 takes about 376 ms doc. If a combined engine had that prefill duration, its running requests would see a gap of hundreds of milliseconds inference.
3. The pipeline drains. Prompt steps are post-processed at once, so async scheduling restarts from empty after each one code deepinfer/worker/worker.py:183-185.

Vote opportunities do not guarantee admission: readiness, cache usage and available slots still matter. These intervals are not upper bounds on a request’s total wait.

Separating prefill and decode

Mixed prefill/decode batches let running requests continue during prompt processing, although the extra work can lengthen each step. DeepSeek's described deployment separates the phases: prefill runs on 8-GPU machines and decode on 64- to 128-GPU supernodes, with KV sent between them doc. Both roles use the same scheduler code but skip the vote code csrc/engine/scheduler.hpp:288-293.

That requires a small KV cache. Thread 2 shows a 16K-token prompt's KV is about 15 MB. The Q&A doc budgets 400 Gbps per 8 GPUs when P and D sit in different data centres doc.

I think separating P/D suits DeepSeek’s scale. Decode is the expensive phase, and separate machines keep prefill from interrupting it. On one small machine, like Pass 2’s offline example, I think mixing phases would reduce the delays caused by voting inference.

Why the scheduler is C++

Async scheduling hides CPU work only when it is shorter than the GPU step. At 13 ms steps and 60 requests per rank, the CPU has about 200 microseconds per request per step inference. That budget also covers the sampling launch, grammar masks and output building. An engine comment names the work left: "TODO(CPU): Eliminate O(num_reqs) python operations in decoding step" code deepinfer/engine/llm_engine.py:931.

Async scheduling still depends on CPU speed. Once CPU and GPU work overlap, a step takes at least as long as its CPU work.

A tiny KV cache, kept for days

At about 891 bytes per token, KV is cheap to keep on SSD for days and to send between data centres.

A 128-token block holds 114,048 bytes, exactly 891 bytes per token of paged capacity for the pinned configuration, as Pass 2 found inference from deepinfer/worker/cache_engine.py:424-483. This is the growing paged-cache component, not total GPU memory per token.

The Llama estimate uses Meta’s architecture table. The V3 estimate uses its official configuration and assumes a one-byte latent-cache format. These are layout estimates, not measured allocations inference. DeepInfer's active state slot adds about 3.1 MB per request, whatever its length. Memory profiling also reserves copy capacity by budgeting the state pool twice code deepinfer/worker/base.py:241-244. Weights, workspaces and unused capacity are separate costs.

What the small KV makes possible

One of the Q&A doc's deployment goals is to "retain KV cache for more than 72 hours" doc. At 891 bytes per token that is plausible.

• The 3FS store has about 100 PiB of SSD for KV doc. That is 1.13 × 10¹⁷ bytes, or about 1.26 × 10¹⁴ tokens, 126 trillion inference.
• One cluster serves 155 trillion tokens per day doc, so the store holds most of a day's tokens even with no reuse. Longer retention therefore depends on prefix sharing and the mix of new tokens. The Q&A describes retention periods, or time-to-live (TTL), of 3–7 days doc.
• At Llama-3-70B's 320 KB per token, it would hold about 340 billion tokens, about 3 minutes of traffic inference.

A storage server reads at least 40 GB/s doc, about 45 million tokens of KV per second inference. A prefill GPU computes about 87,000 tokens per second at a 32K chunk doc. So one storage server serves cached KV about 500 times faster than one GPU recomputes it inference.

Compared with 7 days, a 3-day TTL reduces the chatbot cache hit rate by 0.5 to 1% doc.

What it costs

Three choices make the KV small, and each may cost quality.

1. Shared KV. Only 4 of 40 layers write long-range KV. Pass 2 showed a block would be about 11.6 times bigger without the sharing inference.
2. Compression. Layers 2 to 19 merge token pairs, so their long-range memory is half resolution.
3. 4-bit numbers. Main KV is FP4 with one scale per 16 values code deepinfer/model_executor/layers/tilelang_ops/cache_common.py:119-127.

The release does not isolate these choices in controlled comparisons, or ablations. It therefore does not establish their separate effects on quality.

Bounded replay

Without bounded replay, a prefix hit with a missing window must rebuild the window from scratch. In the second half of the layers that is 2,560 tokens, 20 layers × 128 doc. With replay it costs 128 tokens at any context length doc.

The decoder version, YOCO truncation, saves more. In compact mode, layers 20 to 39 run only the last 140 rows of each prompt, which nearly halves prefill compute doc. The prefill chart's 87,000 tokens per second already includes it, labelled "Decoder Bounded Replay (140 tokens)" doc.

Pass 3 showed the cost: a slightly different answer. The release doesn't measure how different, and I'd measure that first with a GPU.

The alternative: vLLM's prefix cache

vLLM keeps a hashed prefix cache in each engine's GPU memory. DeepInfer's block manager has none code csrc/engine/block_manager.hpp:105-136. Its prefix reuse lives in an external store shared by every prefill engine doc.

With a shared store, "every prefill machine can hit any prefix" doc. The router can choose an engine by load without sending requests to whichever engine holds the prefix. Compute nodes also fail "more than an order of magnitude" more often than storage nodes, so the cache outlives them doc.

This needs small KV. A 100K-token prefix is 32 GB at 320 KB per token, but 89 MB at 891 bytes inference.

A big model with a small active part

V4.1 holds about 760 billion parameters but uses about 15 billion per token. Storing the complete model and computing the active part are different costs.

These are estimates from the model config, compared with the checkpoint’s stated 510 GB inference from research/03-model.md. B means billion parameters. GB measures storage in billions of bytes.

Per token, the model uses 6 routed experts and 1 shared expert in each of 40 layers, about 9.9B parameters. Attention adds about 5B, so about 15B in all, 2% of the total inference.

What that means for memory

One routed expert in one layer is 3 × 5,120 × 2,304 ≈ 35.4 million parameters. That is about 18.8 MB at FP4 with scales, or about 0.75 GB across 40 layers inference.

The slot counts come from the redundant-expert rounding in the GPU-failure example in Pass 3 code deepinfer/model_executor/eplb/__init__.py:21-28. I calculated the sizes inference.

By default the Engram tables stay in host memory, and the engine fetches rows across ranks code deepinfer/model_executor/models/deepseek.py:473-477. The full 203 GB tables are not resident in GPU HBM. Per-token lookups still transfer selected data. I did not measure their latency here inference.

How FP4 reduces weight reads

Decode reads every local expert that receives a token, every step. At EP64 that is up to 5.3 GB per step, against about 20 GB in BF16 inference.

The alternatives

• Dense, like Llama-3-70B. Every token touches all 70B weights. V4.1 stores about 11 times more parameters and uses about a fifth as many per token inference.
• DeepSeek-V3. 256 routed experts, top-8, and its first 3 layers dense. V4.1 has 384 FP4 experts, top-6, and MoE in every layer. These V3 settings come from its official configuration.

What the kernels confirm about the model

The kernel sources confirm some model formulas that Pass 3 left as guesses.

• Router scores are sqrt(softplus(logit)), the only scoring the gate kernel allows code vendor/DeepGEMM/deep_gemm/include/deep_gemm/epilogue/sm100_mega_gate.cuh:19-23, vendor/DeepGEMM/csrc/apis/mega_gate.hpp:99. The kernel ranks experts on score plus a learned bias. The weights it emits drop the bias, are normalised over the top 6 and are scaled by 1.5 code vendor/DeepGEMM/deep_gemm/include/deep_gemm/epilogue/sm100_mega_gate.cuh:212-213.
• No group-limited routing. V3 first picks a few expert groups, then experts inside them. V4.1's gate picks a flat top 6 of 384 code vendor/DeepGEMM/deep_gemm/include/deep_gemm/impls/sm100_bf16_mega_gate.cuh:398. Groups cap how many GPUs a token talks to. Dropping them makes sense once one NVLink domain holds every expert inference.
• The attention sink adds exp(sink) to the softmax denominator and nothing to the numerator code vendor/FlashMLA/csrc/cuda_kernels/sm100/decode/sparse/head64/kernel.cuh:327. That equals one extra key whose value is all zeros. A head that finds nothing useful can put its weight there and output close to zero.

The "Flash" checkpoint is 510 GB doc. It uses few active parameters per token, which supports fast decode per GPU.

Why decode goes wide and prefill stays narrow

Decode can spread expert weights across more GPUs to reduce each GPU’s reads. The prefill chart’s large chunks already provide many tokens per expert, so wider groups offer less benefit and add communication.

EP8 against EP64

The decode chart at 131,072 tokens of context, nextn 3 and 75% acceptance doc doc/images/wide_ep_decode_pareto.png:

The weight-read argument

Per decode step, each GPU reads these weights. All figures are my estimates inference.

1. Attention weights. Every GPU runs attention for its own requests, so it reads all of them: about 5B parameters, about 5 GB in FP8.
2. Expert weights. Each GPU reads its own expert slots, if they receive tokens. That is about 39 GB at EP8 and 5.3 GB at EP64.
3. Total. About 44 GB at EP8 against about 10 GB at EP64.

Assume a peak bandwidth of 8 TB/s, comparable to NVIDIA’s published Blackwell specification. Those reads alone would require at least about 5.5 ms at EP8 and 1.3 ms at EP64 inference. Sustained bandwidth can be lower.

At 64 requests the chart has EP8 at about 61 steps per second, or 16.4 ms per step. EP64 runs about 76, or 13.2 milliseconds per step. That 3.2 ms gap is close to my predicted 4.2 ms. The HBM estimate is rough, and it assumes every slot gets a token. I think reduced weight reads explain most of the gain inference.

At 256 requests the step is longer, so the fixed weight read matters less and the gain shrinks to 6%.

Tokens per expert

At 64 requests and nextn 3, a rank has 64 × 4 = 256 rows, so 1,536 (row, expert) pairs per step inference.

• At EP8 those spread over 52 slots per GPU, about 30 tokens per slot.
• At EP64 they spread over 7 slots per GPU, about 219 tokens per slot.

A 30-row multiply does little arithmetic per byte of weight. At 219 rows it does about 7 times more work per byte read. This is how EP "increases batch size and GPU matrix-multiplication efficiency", in the Q&A doc's words doc.

How Mega MoE hides the exchange

Mega MoE hides much of the all-to-all exchange that EP adds. It uses one long-running kernel whose thread groups do different jobs at once code vendor/DeepGEMM/deep_gemm/include/deep_gemm/impls/sm100_fp8_fp4_mega_moe.cuh:341.

1. No send and receive calls. Each GPU maps the other GPUs' buffers into its address space, so a remote write is a plain store over NVLink code vendor/DeepGEMM/deep_gemm/include/deep_gemm/layout/sym_buffer.cuh:33-39.
2. The shared expert goes first. It needs no remote tokens, so its first multiply runs while tokens arrive code vendor/DeepGEMM/deep_gemm/include/deep_gemm/scheduler/mega_moe.cuh:446.
3. The receiver pulls. Each GPU fetches the tokens for its experts, so its multiply warps start once one tile is full code vendor/DeepGEMM/deep_gemm/include/deep_gemm/impls/sm100_fp8_fp4_mega_moe.cuh:540-599.

A hot expert's tokens are split between its copies by a fixed stride, (rank + row × 23333) % copies code vendor/DeepGEMM/deep_gemm/include/deep_gemm/epilogue/sm100_mega_gate.cuh:204-206. Pass 3 guessed the split was even. It cycles through copies as row indices change, with a different starting offset per rank. Tokens selecting an expert need not occupy uniformly spaced rows, so a particular batch can still split unevenly. The policy assumes an even split code deepinfer/model_executor/eplb/policy.py:31.

The DeepGEMM README gives no Mega MoE numbers, only a link to an external pull request code vendor/DeepGEMM/README.md:18. Its test demands bit-for-bit equality with the separate DeepEP dispatch, multiplies and combine code vendor/DeepGEMM/tests/test_mega_moe.py:345-346. This checks numerical agreement on the tested inputs. The release has no benchmark that isolates the speedup from fusion.

Why prefill stays narrow

At EP4, a 32,768-token chunk makes 32,768 × 6 = 196,608 pairs per rank. Over 104 slots per GPU that is about 1,890 tokens per slot inference. The larger token count makes arithmetic more important relative to weight reads. Going wider then adds communication for less weight-read benefit.

Hence the Q&A doc's rule: prefill uses "the smallest EP/DP that HBM allows", decode "the largest EP/DP that the available GPUs allow" doc. Balance also worsens as EP grows. For prefill, moe_lb is 1.045 at EP2 and 1.074 at EP4 doc. At EP4, 78 GB of experts per GPU is about what fits beside the KV cache. So I think EP4 is the low end in practice inference.

Why EP and not tensor parallel

Tensor parallel splits every weight matrix across GPUs. I see three reasons to avoid it here. All three are inference.

1. The experts are small. At 5,120 × 2,304, an 8-way split leaves very thin multiplies.
2. Attention has one KV head. The model stores one 512-number vector per token code deepinfer/hugging_face/config.py:55. Splitting heads would leave every GPU holding a full copy of that KV.
3. Every layer would need an all-reduce of the hidden state. EP sends only routed tokens to their experts.

The Q&A doc gives the outcome without reasons: "TP and CP are not used for now" doc.

What load balancing is worth

Production decode measures moe_lb at 1.090 with EPLB on, against 1.0 for perfect balance doc. So the busiest GPU does 9% more than average, and every rank waits for it.

The Q&A doc also reports roughly 80% of a perfectly balanced offline benchmark when assessing overall load balancing doc doc/inference_deployment_qa.md:125-127. The charts above use random expert routing to remove much of the imbalance code examples/benchmark.py:165. The 80% report does not establish a conversion factor for every chart: production context lengths, concurrency and acceptance rates must also match before comparing throughput.

Complete the expert placement

Setup: 4 logical experts, 2 GPUs, 3 slots per GPU, so 2 redundant copies. Two steps.

Start-up copies go to e0 and e1, so GPU0 = {e0, e1, e2}, GPU1 = {e3, e0, e1}. That gives GPU1 140 and 130, GPU0 60 and 50, so moe_lb = 270 / 190 = 1.42.

Stage A. Round 1: copying e3 cuts the hottest copy to 60 and 50, sum 110. Anything else leaves 220. So e3. Round 2: copying e3 again cuts its copies to 40 and 33.3. The hottest copy is now 40 at step 0, from e2 or an e3 copy, and 40 at step 1, from e1. Sum 80, against 110 for anything else. So e3 has 3 copies.

Stage B. Copies in order of load per copy: e3, e3, e3, e1, e2, e0.

Your turn. Where does e0 go, and what is the final moe_lb?

EPLB also considers which steps make each expert busy, as Pass 3 showed. It can only plan from past load. Production re-plans every 10,000 decode steps from a 1,000-step window doc. A burst lasting a few seconds may finish before the next layout change inference.

What each speed trick is worth

Drafts help when their extra committed tokens outweigh the time spent proposing and checking them. The balance changes with batch size.

The cost of a decode step

The Q&A doc's table compares fixed-5 and dynamic verification at EP64 doc. For each row I worked out the step time as concurrency × accepted length / throughput. The row count is concurrency × (1 + verified length). A line fitted to all 12 points gives this inference.

step time  ≈  8.8 ms  +  0.0168 ms per row

The fit groups costs into a roughly fixed term and a term that grows with rows. Weight reads, communication, launches and arithmetic contribute, but this fit does not measure their separate shares.

Why dynamic verification loses at 16 and wins at 96

Dynamic verification drops the drafts least likely to pass. That saves rows but loses their tokens.

At concurrency 16 the step is mostly fixed cost. Dropping 9 rows saves about 0.15 ms of a 10.4 ms step, about 1.5%. Accepted length falls about 2.4%, from 3.32 to 3.24 doc. The Q&A doc reports a net −1.5% doc.

At concurrency 96, the fitted per-row term is more than half the step time. Dropping 169 rows saves about 2.8 ms of an 18.5 ms step and loses about 4.5% of accepted tokens. The net is the reported +11% doc.

For small batches, the scheduler’s own estimate should favour checking every draft. I see three possible reasons it does not. All are inference, listed from strongest to weakest code evidence.

1. The total uses confidences from two rounds back, which are stale code csrc/engine/mtp_varlen.hpp:160-168.
2. The confidence head may be badly calibrated on later drafts. The engine logs a calibration error per step for exactly this code deepinfer/model_executor/layers/tilelang_ops/sampler.py:589.
3. The start-up cost curve is measured at full load with random routing and short context code deepinfer/worker/model_runner.py:581-584. So it may be wrong for small batches.

Choose how many drafts to check

Three requests U, V and W check up to 2 drafts each. Each confidence already includes the chance that all earlier drafts pass.

The cost curve gives 7.0 ms for the 4-row graph, 7.6 ms for 8 rows and 9.5 ms for 9 rows. Add 0.4 ms for the sampler.

Your turn. Fill in rows 8 and 9, pick the best, and say how many drafts each request checks.

Async scheduling

Pass 1’s toy overlapped 3 ms of CPU work with a 10 ms GPU step. That gave 30% more steps per second. The real gain depends on the CPU time, which the release doesn't publish. The engine's cpu_timer does time the schedule, the post-process wait and a "bubble" segment code deepinfer/engine/llm_engine.py:881-887, :991-992. A GPU run with async scheduling on and off would settle it. The benchmark already puts the setting in its trace names code examples/benchmark.py:567.

The costs are clearer.

• Complexity. The GPU must fix lengths and copy tokens the CPU didn't know code deepinfer/worker/worker.py:191-245. The finishing-request example shows the resulting lifetime problem.
• Wasted rows. A finishing request can remain in the next forward step. Fixed-5 verification gives it six rows. Variable verification can reduce its draft allocation.
• No tests. I found no engine-level test of the async edge cases. Safety rests on assertions, such as the varlen scheduler's counter check code csrc/engine/mtp_varlen.hpp:195-196.

CUDA graphs and padding

Graphs remove launch cost, which matters when a 13 ms step is a few hundred kernel launches. They cost padding. At the default interval of 8, a step wastes at most 7 rows code deepinfer/environ.py:155. At 0.0168 ms per row that is at most about 0.12 ms, under 1% of a step inference from my fit. The benchmark uses an interval of 32 code examples/benchmark.py:169. That means fewer graphs, shorter start-up and less graph memory, but up to 31 padded rows, about 0.5 ms.

Speculative decoding helps only when the extra accepted tokens outweigh drafting and verification costs. Per-user speed and aggregate tokens per GPU can respond differently. At high concurrency each checked draft takes a row another request could have used. That trade changes with load, which is why the varlen scheduler exists.

What determinism costs

Batch invariance makes numerical choices stable across batch compositions. The tests enforce exact equality for some exposed results and aggregate regression checks for others.

What the switch changes inside the kernels

Pass 2 guessed that the switch fixes how kernels divide long sums. The kernel sources confirm it.

FlashMLA, decode attention. The switch changes one decision in the dispatcher code vendor/FlashMLA/csrc/api/sparse_decode.cpp:219-223:

// Split-KV only pays off when a request has enough work
bool enable_split_kv = !enable_batch_invariant && !(topk + extra_topk <= 640);

Split-KV cuts one request's keys into pieces for different GPU cores, then merges the partial softmaxes. A scheduler sets the cut points by dividing the whole batch's key blocks evenly between cores code vendor/FlashMLA/csrc/cuda_kernels/smxx/decode/get_decoding_sched_meta/get_decoding_sched_meta.cu:54, :65. So a long request joining the batch moves your request's cut points and changes its sum order. With the switch on, one thread block owns each request from start to end, and the merge kernel never runs code vendor/FlashMLA/csrc/cuda_kernels/sm100/decode/sparse/head64/kernel.cuh:732, vendor/FlashMLA/csrc/api/sparse_decode.cpp:444-448.

V4.1's decode attention reads a 128-token window code deepinfer/model_executor/models/v4_modules/attn.py:857 plus 512 top-k entries, exactly 640. So split-KV is off for V4.1 decode either way, and I read batch invariance as free in this model's decode attention inference. No README or comment gives a cost figure.

DeepGEMM, matrix multiplies. The switch is host-side only, read by no GPU kernel code vendor/DeepGEMM/csrc/jit_kernels/heuristics/runtime.hpp:15-26. It changes four choices.

1. Dense BF16 multiplies skip cuBLASLt, which picks algorithms by shape outside DeepGEMM's control code vendor/DeepGEMM/csrc/apis/gemm.hpp:519-523.
2. The same for einsum code vendor/DeepGEMM/csrc/apis/einsum.hpp:81, :115.
3. The hyper-connection kernel splits its sum into a fixed 16 pieces, not a number based on batch size code vendor/DeepGEMM/csrc/apis/mega_mhc.hpp:15-17.
4. The MoE gate doesn't split its sum at all: "deterministic mode is also batch invariant, so K is never split" code vendor/DeepGEMM/csrc/jit_kernels/heuristics/mega_gate.hpp:110-111.

Point 3 shows that determinism doesn't need split-K turned off. It needs the split count fixed, so the summation tree is the same for any batch.

DeepEP and Mega MoE, the expert exchange. Mega MoE sums each token's expert outputs in 32-bit floats in a fixed order. The kernel says: "Shared slot goes first, then routed slots in top-k order" code vendor/DeepGEMM/deep_gemm/include/deep_gemm/impls/sm100_fp8_fp4_mega_moe.cuh:1449-1450. That order doesn't depend on the batch, so it needs no switch.

DeepEP's separate combine, used across NVLink domains, is different. By default it sums partials on each expert's GPU first, rounding to 16-bit floats at each stage code vendor/DeepEP/deep_ep/include/deep_ep/impls/ep/combine_utils.cuh:67. The grouping depends on which experts share a GPU, so the rounding changes when EPLB moves experts. DeepInfer turns early summing off only when batch invariance and EPLB are both on code deepinfer/model_executor/parallel_utils/userbuffer.py:79. That costs bytes on the wire: one slot per top-k expert instead of one per GPU code vendor/DeepEP/csrc/buffers/ep.hpp:222.

One gap remains. The gate picks a hot expert's copy with (rank + row × 23333) % copies code vendor/DeepGEMM/deep_gemm/include/deep_gemm/epilogue/sm100_mega_gate.cuh:204-206. The row is the token's position in this batch, so the batch picks the token's GPU. With Mega MoE that doesn't matter, because copies share weights and the sum order is fixed. With DeepEP and redundant experts but EPLB off, the row could change the early-sum grouping, and the switch wouldn't stop it. I think that gap is real but narrow inference.

I therefore expect little batch-invariance overhead with one NVLink domain, Mega MoE and V4.1’s attention shape. Other configurations can lose faster cuBLASLt paths or small-batch split-K.

Why the tests compare floats exactly

The offline PPL test compares each sample's decode perplexity with a stored value using == code tests/models/test_offline_inference.py:152-158. This checks the final scalar exactly. Perplexity aggregates token logprobs, so matching it does not prove that every token, logprob or intermediate value matched.

The golden file also tests across batches. DP rank 0 runs the five sample prompts with seeds 0 to 4. Rank 1 runs the same five with the same seeds, then five more, in a batch twice the size code examples/_prompts.py:100, :113. Rank 1's first five stored perplexities equal rank 0's exactly, on all four platform and workload pairs code tests/models/offline_decode_ppl.json:1-90. Our Pass 2 request, sample 1 in the long workload on NVIDIA, stores 1.000645962988749.

The two ranks share every MoE layer but run different batch sizes. Numerical changes that alter a stored perplexity fail the test. Aggregation or rounding can hide other changes. The golden file can be refreshed with examples/refresh_offline_decode_ppl.py, so a failure needs investigation before accepting a new baseline inference about intent.

The Ascend perplexities match across ranks too, although Ascend does not support the batch-invariance switch code deepinfer/worker/runtime_config.py:62-63. The engine always enables deterministic kernel mode there code deepinfer/worker/runtime_config.py:67. Matching these aggregates does not establish batch invariance for every exposed value or configuration.

What it buys

• Debugging. Stable numerical results make a forced continuation easier to compare alone and in a batch. Scheduling or lifetime bugs can still depend on batch composition.
• Prefill/decode regression checks. test_kl.py compares exposed logprobs along the same tokens, with an aggregate discrepancy limit of zero code tests/models/test_kl.py:190-206. As Pass 3 explains, this does not prove whole-pipeline bitwise equality or validate every P/D handoff and recomputation path.
• Repeatable evaluation, with the supported NVIDIA recipe, identical inputs and parameters, a fixed seed and varlen off. The evaluation script enables varlen, so that recipe does not guarantee its runs repeat code eval/run_eval.py:120-121.

The alternative

Allowing batch-dependent numerical choices gives kernels more freedom to optimise each shape. The cost of fixing those choices needs measurement for each workload.

DeepInfer combines position-keyed random draws, a sampler designed to be invariant without the switch, and regression checks across batching and execution paths. I think these checks help a team shipping a model and engine together. Each check still has limits.

Determinism depends on every kernel, as Pass 1 showed. A fixed seed cannot prevent one changed bit in the logits from changing the sampled token.

The online server and its aborts

You must put an online server on top of LLMEngine. Clients stream tokens and can cancel at any time. The release has no online wrapper code deepinfer/engine/llm_engine.py:141. From Pass 3:

• An aborted running request gets no final output from the engine. It just stops appearing code csrc/engine/outputs.hpp:94-303.
• An aborted waiting request that holds blocks, such as a prompt between chunks, seems never to free them code csrc/engine/scheduler.hpp:363-367, :569-578.

Design the abort path. Say what the server tracks, what it sends the client, and what you change in the engine.

One 8-GPU box, prefill and decode together

A customer has one 8-GPU SM100 machine and no budget for P/D disaggregation. They run DeepInfer as a NORMAL engine at EP8. Their users complain that answers stall for a third of a second every few seconds.

A decode step at EP8 with 64 requests per GPU takes about 16 ms, from the EP8 curve in the decode chart doc. The prefill chart, at EP4, shows about 68 ms for a 4,096-token step and about 376 ms for 32,768 tokens doc. Assume EP8 prefill is similar per GPU.

Explain the stalls. Then decide: change the settings, or make the engine mix prefill and decode in one step?

Size the KV store

You run a cluster like DeepSeek’s. Use these invented traffic assumptions.

• 155 trillion tokens per day, 90% input and 10% output.
• 80% of input tokens hit the cache, so 20% are new.
• Decode KV is also written to the store, as the Q&A doc says production does doc.
• 60% of stored tokens need 3-day retention, 40% need 7 days.
• Each storage server has 20 × 30.73 TB SSDs and reads at 40 GB/s doc. It uses 10 data units plus 2 redundant units, called 10+2 erasure coding doc.
• Requests average 20,000 tokens.

How many bytes does the store hold, how many servers do you need, and is capacity or read bandwidth the limit? Can you keep each request’s window state as long as its blocks?

Put a prefix cache inside the engine

You want vLLM-style prefix caching inside one engine, for a deployment without a shared store. What do you add to the block manager and the scheduler? What is hard about this model in particular?

A KV budget of 1 KB per token

You are designing the next model's cache layout. Your budget is 1,024 bytes per token for long-range KV and indexer keys, in V4.1's formats. Per raw token, a C2 owner layer costs 178 bytes and a C1 owner layer 356 bytes, from Pass 2's block layout. V4.1 uses 3 C2 owners and 1 C1 owner, 890 bytes.

Propose a layout and argue for it. One rule from the code: a KV-sharing group must not cross the YOCO boundary code deepinfer/model_executor/models/v4_modules/cross_layer.py:115-116.

64 GPUs, and one expert gets 3 times the traffic

A decode instance runs EP64 with 384 experts and 1 redundant copy per GPU, so 7 slots per GPU. Every expert gets L tokens per step, except one, which gets 3L. The redundant copies sit where start-up put them, on experts 0 to 63.

What does this do to step time? What do you do about it? What if the hot expert is hot only for a few seconds at a time?

Find a better expert placement

Pass 3 did a full run, and Pass 4 left out the last step. Here you do every step: 6 experts, 2 GPUs, 4 slots per GPU, so 2 copies. Both steps have the same tokens.

Work out stage A and stage B by hand. Then ask whether a better layout exists.

Pick the verify lengths, and the graph sizes

Two requests, A and B, each with 3 drafts. Their confidences, already multiplied along each request:

The cost curve is 6.0 ms for 2 rows, 6.1 ms for 4 rows and 9.0 ms for 8 rows. Add 0.2 ms for the sampler. Run the varlen scan alone, without a worked step.

Second question. You own the CUDA graph sizes. The default is powers of two up to 8, then steps of 8. Would you change that for an engine with max_num_seqs=128 and nextn 5?

Make speculative decoding deterministic, with varlen on

Pass 3 showed why varlen breaks repeatable output. The number of drafts checked per request depends on timing and on other requests, and that changes which random stream makes each token. Can you keep varlen's speed and still get the same tokens every run, for the same prompt and seed?

What the code showed that the doc didn't

• There is no quadratic chunk-size solver in the release. The Q&A's cost(q, c) and compute_quota_limit live in the unreleased load balancer. Chunk size is set only by max_num_batched_tokens (csrc/engine/scheduler.hpp:470). flops_estimator.hpp only feeds arithmetic-work statistics.
• Prefill and decode never share a step. A NORMAL engine votes across DP ranks every 70 or 100 decode steps.
• There is no prefix cache in the block manager. The only in-release sharing is fork_seq.
• ZOS is on by default (deepinfer/engine/arg_utils.py:51-55). MTP varlen and EPLB rearrangement are off by default.
• The preemption comment says "shortest output", the code preempts the longest output (csrc/engine/scheduler.hpp:556 vs :622-639).
• The Q&A says decoder replay recomputes 128 tokens. The code requires at least 134 (csrc/engine/input_metadata/prepare_tensor_input_v4.cpp:354).
• CP is in the code and README, but the Q&A says production does not use it.
• Aborting a prompt between prefill chunks seems to leak its blocks and state slot (csrc/engine/scheduler.hpp:363-367, :569-578). Nothing frees an aborted sequence in ignored_seqs.
• FlashMLA never splits V4.1's decode attention, because 128 window + 512 top-k = 640 is at its threshold (vendor/FlashMLA/csrc/api/sparse_decode.cpp:223). So batch invariance is close to free there.
• The gate picks a hot expert's copy by (rank + row × 23333) % copies (vendor/DeepGEMM/deep_gemm/include/deep_gemm/epilogue/sm100_mega_gate.cuh:204-206). The choice depends on the batch row and ignores current load. Selected rows can split unevenly between copies.

Open questions and first runs

Questions the code can't answer without a GPU

These questions come from all seven research notes. Reading the code could not answer them.

1. Does the abort-between-chunks leak really happen? Abort a 50,000-token prompt after its first chunk, then check that get_num_free_gpu_blocks() returns to its starting value (csrc/engine/pybind_scheduler.cpp:73).
2. How long does the offline engine run empty steps at start-up? I expect about 70 placeholder decode steps before the first prefill, because nothing is schedulable at the first schedule() (Pass 2, Thread 1).
3. How big is the CPU gap that async scheduling hides? The engine records schedule, worker.post_process.sync and bubble in cpu_timer (deepinfer/engine/llm_engine.py:881-887, :991-992).
4. Does a real block measure 114,048 bytes? The cache engine logs the block count and size at start-up (deepinfer/worker/cache_engine.py:935-946).
5. What are the measured decode costs and chosen draft counts? Do they reproduce the Q&A table's 5.00 to 3.2-4.4 drafts per step? Why does varlen lose at concurrency 16: stale confidences, a badly calibrated confidence head, or a cost curve measured at full load?
6. How far does YOCO compact truncation drift from off? No test compares them. examples/yoco/yoco_equiv.py only compares mask with compact.
7. What does batch invariance cost in throughput across EP sizes, and does it hold across different CP or EP sizes at all?
8. Is there a batch-invariance gap on the DeepEP path with redundant experts and EPLB off, because the gate's replica choice depends on the batch row?
9. How often does an EPLB re-plan make balance worse on real traffic, and how long does a re-plan stall the ranks at EP64?
10. Can one token appear twice in a C1-layer softmax, once through the window and once through the indexer's pick? Logging the indices would settle it.
11. How long does the Engram fetch take from host memory across ranks, and is it hidden behind layers 0 and 1?
12. How often do the sampling rules block every token in a row? That is a row where the grammar mask, the token blacklist and the first-token end ban together block every token. Does the device assert then take the whole engine down?
13. What does the internal compute_quota_limit limiter do, and what are its fitted cost(q, c) numbers? That code is not in the release.

What to run first if you get a GPU

The offline example needs 8 SM100 GPUs for EP8. The decode benchmark curves need 64. Start small, and in this order.

1. Build the engine and check its stored results. export DEEPINFER_ACCELERATOR=nvidia && ./build.sh, then python -m pytest -v -s tests/models/test_offline_inference.py. This runs the 16,389-token request family from Pass 2 and compares stored perplexity scalars exactly. A pass gives a regression baseline. It does not prove that every intermediate value or execution path matches.
2. Read the start-up logs. Check the block size against 114,048 bytes and count the empty decode steps before the first prefill. The logs let you check two of my main inferences.
3. Write and run the abort test from question 1. Use LLMEngine to check whether the suspected leak occurs.
4. Measure async scheduling. Run python examples/benchmark.py with DEEPINFER_ZERO_OVERHEAD_SCHEDULING=0 and =1 and compare decode avg_SPS and the bubble timer. The benchmark already puts the setting in its trace names (examples/benchmark.py:567).
5. Measure varlen at small and large batches. DEEPINFER_ENABLE_MTP_VARLEN=1 python examples/benchmark.py --decode_num_requests 16, then 64. Log batch_confidence_error to check whether inaccurate confidence predictions explain the loss at 16.
6. Measure batch invariance's cost. The same benchmark with DEEPINFER_BATCH_INVARIANT=0 and =1. Pass 4 predicts almost no change in decode attention.
7. Then check the remaining cases: YOCO compact against off on long prompts, EPLB re-plans with --enable_eplb, and the C1 overlap logging.

Research and plan

• research/05-decode-speed.md: corrected the claim that ZOS defaults to off. EngineArgs.__post_init__ turns it on for NORMAL and COMPLETION engines (deepinfer/engine/arg_utils.py:51-55). The reader had looked only at the EnvVar.bool() default.

After Pass 2

• Pass 1, Thread 1: "a cut prompt always ends the step" became "ends the step", for NORMAL engines (csrc/engine/scheduler.hpp:543). A PROMPT engine keeps looking.
• Pass 1, Thread 1: added a "simplified here" line before the toy step table. The real engine only starts prefill at vote points, every 70 decode steps by default, not whenever a prompt waits. Pass 2 showed the vote.
• Pass 1, Thread 4: "every few thousand steps" became the Q&A doc's real numbers, 1,000 steps for prefill and 10,000 for decode. The vague figure hid that prefill and decode differ by 10x.
• Pass 2: fixed six line references that were one or two lines off. For example, the force top-k flag is examples/offline_inference.py:17, not :18. I built tools/check_refs.py to catch these.
• Pass 2: every offline prompt carries an empty image dict. This prevents scheduling on the first step. I now expect about 70 empty decode steps before the first prefill. Pass 1 had no way to show this, because it comes from the order of two calls inside step().

After Pass 3

• Plan, Thread 3: moved the pending C2 token in ca_state from Pass 3 to Pass 2. The request at position 16,389 provides a concrete example. Pass 3 uses a chunk wholly before the YOCO keep window instead, so no example repeats.
• Pass 1, Thread 5: added "simplified here" after the leftover-odds example. The real kernel scales the draft's odds by e^(-0.00005) before subtracting, so the leftover never rounds to all zeros (deepinfer/model_executor/layers/tilelang_ops/common.py:53).
• Glossary, Abort: the old text said the blocks come back on the next free pass. Pass 3 found that an abort between prefill chunks seems to leak them, so the entry now says it depends on the queue.
• Glossary, Rejection sampling: added the documented bias. "Keeps the output on the main model's distribution" was true only to within (1 + 5e-5).
• Glossary, Preemption: added that the victim is the longest-output request. Pass 3 showed the scheduler's own comment says the opposite.
• Pass 3, top-p exercise: my first draft credited the tie rule for keeping two tied tokens. The mass alone needed both. I rewrote the answer to say when the tie rule matters.
• Understanding that grew in Pass 3: I ran the real rebalance_experts_global on CPU. It matched my calculations and found layouts where the greedy plan worsens the starting balance. Before that, I had only the reader's port to go on.

After Pass 4

• Checked out vendor/FlashMLA, vendor/DeepGEMM and vendor/DeepEP with git submodule update --init --depth 1, as the prompt allowed for Pass 4. A reader wrote research/07-kernels.md from them.
• Pass 2, Thread 6: my guess that batch-invariant mode "pins how a long sum is split" is now confirmed. I kept the guess, marked it as a guess at the time, and pointed to Pass 4. FlashMLA already skips split-KV when the window plus top-k is at most 640. V4.1’s 128 + 512 is exactly 640. The switch therefore changes nothing in this model’s decode attention.
• Pass 3, Thread 4, Case J: the even split of a hot expert's tokens across its copies was a guess. The gate kernel shows a fixed stride over the token's batch row, (rank + row × 23333) % copies. This does not measure current load, and selected rows can split unevenly across copies. I updated the case and pointed to Pass 4.
• Pass 4, Thread 2: my first draft said Llama-3-70B-sized KV would fill the 3FS store with "a few hours" of traffic. The sum gives about 3 minutes. Fixed.
• Pass 4, Thread 3: added what the kernels confirm. Router scores are sqrt(softplus(logit)), there is no group-limited routing, and the attention sink only adds to the softmax denominator.
• Pass 4: choosing replicas by batch row may leave a narrow batch-invariance gap. It affects the DeepEP path with redundant experts and EPLB off. I couldn't confirm it without a GPU, so it is tagged as inference and listed in "What's next".
• tools/check_refs.py skipped .cu and .cuh files at first. I added the extensions and rechecked everything.

Final read, Pass 0 to Pass 5

• Pass 0: the thread table still listed "odd positions" under Pass 3. That example moved to Pass 2 after Pass 3, so the table now says so.
• Pass 0: the system diagram used EP, DP and "drafts" before anything defined them. Added one paragraph after the diagram that names all three and points to Threads 4 and 5.
• Pass 1: "they send their draft sentence" in the newsroom analogy clashed with draft tokens. Now "their sentence".
• Pass 1: "in a real office each layer would need its own cabinet" mixed the analogy with the model. Now "in a plain transformer".
• Pass 2: "MTP info slot" appeared without a definition. Added one sentence where it first appears, and a glossary entry for MTP.
• Pass 2: EP, DP and CP were not linked at first use. Linked them in the setup table.
• Pass 2: the line quoting Pass 1's "a cut prompt always ends the step" no longer matched Pass 1's revised wording. Rewrote it to match.
• Pass 3: four path:line references ended in ] instead of a backtick, so the checker never saw them. All four had a digit in the path, for example prepare_tensor_input_v4.cpp, which my earlier fix pattern missed. Fixed, and rechecked with a digit-aware pattern across all files.
• All passes: Passes 3 to 5 had no glossary links. tools/link_glossary.py now links the first use of each term in every pass, 137 links in all.
• Glossary: added MTP and Wide-EP, both used in the passes without entries.

Diagrams instead of ASCII

• Pass 0, Pass 1, Pass 2 and Pass 3: six box-and-arrow ASCII pictures are now Mermaid diagrams, which Pele draws at build time. They are the whole system, one token through one layer, request states, one MoE layer at EP8, and the two async-step timelines.
• Pass 1 to Pass 5: ASCII grids for block contents, the state slot, expert token counts, truncated prefill and draft verdicts are now tables.
• Pass 2: the ASCII held nine short code references, such as [worker.py:122], that the checker couldn't see. They are now full references in the prose, and all of them resolve.
• Pass 2, request states: the six finished states sit in one box with their causes, not six boxes. Six boxes side by side shrank the text until it was unreadable.
• Pass 1, Thread 2: new figure that steps through the drawer toy. Blocks come from the head of the free list. A fork shares full drawers. The clerk returns a drawer to the tail only when its count reaches 0.
• Pass 2, Thread 1: new figure of the prefill vote on four ranks. It follows the scheduler's rule: ALL points every 70 decode steps, ANY points every 100, and a prompt step resets the count. The arrivals are made up.
• Pass 2, Thread 3: new figure of the 40 layers. It shows which layer owns the long-range KV and the top-k picks each layer uses, which half each layer is in, and the draft-model taps. It sits after the layer-27 exercise, so it doesn't give away the answer.
• All figures: titles now say what the figure shows, not just its topic.

Independent factual and learning review, 21 October 2026

• Pass 0: pinned the source revision and separated engine-local cache management from the production service. Scoped collective coordination to one expert-parallel group. Qualified CPU/GPU overlap.
• Pass 1: replaced “never computes them twice” with cache reuse while state remains available. Named recomputation after preemption and missing windows.
• Pass 1: acknowledged the first two window-only layers and the shared expert. Decode rows now have an explicit growing cost instead of a universal “six costs little more than one” claim.
• Pass 1: carried fixed verification length, identical inputs and the supported NVIDIA configuration into the repeatability table. Qualified ideal rejection sampling by the implementation's numerical approximations.
• Pass 2: labelled the request as a worked trace from upstream configuration. Stated readiness and memory assumptions. Added the scheduler’s second prefill vote. Preserved all interactive-figure build markers.
• Passes 2 and 4: separated the 891-byte paged-cache slope from active state, reserved copy capacity and other GPU memory. Made the ring size conditional on the draft limit.
• Pass 2: distinguished perplexity of the processed sampling distribution from raw-model perplexity. Matching settings alone does not turn one into the other.
• Pass 3: stated that the batch-invariance test forces the continuation and compares exposed values exactly. Removed the claim that zero aggregate discrepancy proves bitwise equality. The helper promotes to double, but cancellation can still hide distinct inputs.
• Pass 4: limited the PPL golden test to the aggregate it asserts. Removed claims that any numerical change must fail it or that matching Ascend aggregates prove broader invariance.
• Pass 4: kept the 80% load-balancing comparison within its documented scope. Removed an unsupported conversion of every chart into production throughput. Removed an unchecked comparison with other engines.
• research/factual-review.md: corrected the numerical counterexample to account for the helper's promotion to double. The revised example uses distinct float32 inputs near zero and still produces a zero double-precision discrepancy.
• research/06-correctness.md: aligned the original test-reading notes with the corrected limits of exact exposed-value comparisons and aggregate KL/PPL checks.
• Glossary and Pass 5: separated numerical batch invariance from repeated sampled text. Stated the deterministic-drafting proposal’s equal-probability assumption rather than treating the aggregate KL test as proof. Kept its ideal-arithmetic and independent-draw limits.

Guided learning experience and verification

• Added an introduction to the next token that needs no specialist background. Kept all six original technical passes for deeper reading.
• Added tested batch, cache, expert, verification, draft-cost and arithmetic experiments plus assessment and a free-response check. Prompt completion samples its first answer token. Each scene states its toy limits.
• Improved navigation, mobile section access, glossary focus/history, numeric alternatives, local fonts and no-JavaScript reading. Pre-rendered the initial index to prevent a measured layout shift.
• Added source/model ledgers, independent reviews, reproducible builds, artifact validation, model/browser checks, local performance/screenshot evidence and a CI workflow. The review states that GPU performance and human learning remain untested.

What the release contains

The official README describes an offline inference release supporting DeepSeek-V4.1-Flash, NVIDIA SM100 and Huawei Ascend 950. It contains engine-local KV management, model execution, scheduling and sampling. A complete inference service also requires an external KV system, disaggregation and load balancing across replicas. The production Q&A describes these parts. They are outside the engine release.

Central claim ledger

The notes in research/00 through research/07 give code locations, worked calculations and questions that need accelerator tests. A separate agent checked the main claims and source identity. This is another model’s review. It does not replace reproduced results or human expert review.

What the experiments can establish

The experiments show how separate requests share work and how history differs from fixed window capacity. They also show placement bottlenecks, probability correction and floating-point ordering. They do not estimate deployment speed, GPU utilisation, cache capacity or output quality. The page does not run a language model.

The batch toy’s phases are deliberately simpler than the engine’s policy. The memory toy uses eight layers, two owners, four-position blocks and an eight-position ring. Real state includes more components and spare capacity for rejected drafts. The expert toy assumes perfect splitting and identical work per job. The probability toy has three tokens, ideal arithmetic rules and no grammar or top-p. The draft-cost toy uses constant conditional acceptance and invented milliseconds. The sum-order example is an intentionally sensitive boundary.

All central models share calculation code with their rendered numbers. The repository’s quality review records automated checks and browser results. No human participant study has tested learning.

Open questions

Engine timing and cache allocation during abort still need accelerator tests. So do exact repeats, replay drift and the placement/invariance cases in the advanced passes. These are open questions, not verified bugs or measurements. The open-questions page suggests experiments. The revisions page records corrections to the explanation.

Attribution and access

DeepInfer and the kernel repositories I inspected are DeepSeek projects. Their Apache 2.0 licences apply. I wrote the diagrams and introductory models for this explanation. The speculative-decoding algorithm is attributed above. Font licences accompany bundled font files. Accessibility work targets WCAG 2.2. Automated checks alone do not prove conformance. The page uses no analytics, accounts or remote inference calls. It stores your Day/Night preference locally.