DeepInfer Explained

Tour: One step on many GPUs

AdvancedGuided toursabout 11 min read + 5 min deep dives

Code version: DeepInfer @ 9bd9cf2 · 2026-10-21

In one sentence

We follow one prefill step on 4 GPUs. Two engines vote, each driver feeds its CP follower, the ranks share keys and experts, and the drivers sample.

The idea#

The tour Tour: One engine step followed one step on one device. This tour runs the same kind of step on four devices, and watches the communication. At each stop you see which ranks work, which group talks, and the shapes of the data.

The set-up. python examples/offline_inference.py --ep 4 --cp 2 on one host with 4 GPUs. The script computes dp_size = max(4 // 2, 1) = 2 and a world of 4 ranks (examples/offline_inference.py:24-28). So there are two engines (two DP replicas), and each engine has two CP ranks:

RankEnginecp_rankProcessDevice
000 (driver)engine 0, started by _launch.py0
101 (follower)spawned by engine 01
210 (driver)engine 1, started by _launch.py2
311 (follower)spawned by engine 13

Rank = engine_rank × 2 + cp_rank (deepinfer/worker/base.py:329). The groups are: CP [0, 1] and [2, 3]; DP [0, 2] and [1, 3]; EP [0, 1, 2, 3]. The EP group is the world group, because ep_size equals the rank count (deepinfer/model_executor/parallel_utils/parallel_state.py:91-106).

The work. Engine 0 has one new request A with 8 prompt tokens. Engine 1 has no waiting request. The scheduler counter is at decode_step = 100, an ANY point. ZOS is on (the default). H is the hidden size of the checkpoint; the repository does not fix it.

StopWhat happensRanksGroup
1The engines vote for a prefill step0, 2DP (Gloo)
2Each engine schedules its batch0, 2none
3The driver sends the method call0 → 1, 2 → 3CP (shared memory)
4The driver sends the input metadata0 → 1, 2 → 3CP (shared memory + NCCL)
5Each rank takes its mirror rowsallnone
6Attention gathers K and VallCP (DeepEP BucketBuffer)
7MoE dispatch and combineallEP
8EPLB counts the tokens per expertallEP (every 128 steps)
9The last all-gather, and the drivers sampleall; sampling on 0, 2CP
10Outputs0, 2none

The picture#

Figure: the four ranks of the example (panel 1) and the ten stops of the tour, with the group that communicates at each stop (panel 2). Click to zoom.
Open full size ↗

Read panel 2 from top to bottom. Amber cards talk inside one engine (CP). Violet cards talk across both engines (EP). The green card is the only stop where the two drivers talk to each other directly.

1The engines vote for a prefill step#

We are in Scheduler::_schedule (C++), called from LLMEngine.step() on ranks 0 and 2. Only the engine processes run a scheduler; ranks 1 and 3 wait in their worker loop.

            bool is_any_point = (decode_step != 0) && (decode_step % prefill_any_attempt_interval == 0);
            bool is_all_point = (decode_step % prefill_attempt_interval == 0);
            if (!is_any_point && !is_all_point) {
                return _schedule_decode(acquire);
            }
            bool want = (get_gpu_cache_usage() <= GPU_CACHE_PREFILL_WATERMARK) && _has_prefill_waiting();
            bool do_prefill;
            if (is_any_point) {
                do_prefill = (_dp_allreduce(want ? 0 : 1, c10d::ReduceOp::MIN) == 0); // ANY:任一 rank 想做即触发
            } else {
                do_prefill = (_dp_allreduce(want ? 1 : 0, c10d::ReduceOp::MIN) == 1); // ALL:全员想做才触发
            }

decode_step = 100, so this is an ANY point. Engine 0 wants a prefill (A is schedulable and the cache has room) and sends 0. Engine 1 sends 1. MIN(0, 1) = 0: both engines prefill. With ZOS, the vote uses the Gloo copy of the DP group [0, 2] and a CPU tensor (deepinfer/engine/llm_engine.py:313, csrc/engine/scheduler.hpp:170-182).

Deep dive

Stop 1 in detail: the other two all-reduces

Before this call of step(), LLM._run_engine asked has_unfinished_requests(). Engine 0 sent 1 and engine 1 sent 0; the MAX is 1, so both engines call step() (deepinfer/engine/llm_engine.py:673-678).

After vote 1, engine 0 runs _schedule_prefill and gets a batch. Vote 2 then asks whether any engine got a batch: engine 0 sends 0, engine 1 sends 1, MIN = 0. Engine 1 returns an empty prompt step (csrc/engine/scheduler.hpp:332-340). The page Data parallelism and step votes explains the encodings.

2Each engine schedules its batch#

Still in LLMEngine.step(). Engine 0 now has one SequenceMetadata for A, with seqlen_q = 8. The scheduler pads the CP work in its token budget to a multiple of 2 × cp = 4; 8 is already a multiple (csrc/engine/scheduler.hpp:515). Engine 1 has no sequence. The engine hands both cases to the worker in the same way:

    def _execute_model_async_step(
        self, seq_metadata_list: list[SequenceMetadata], scheduler_outputs: SchedulerOutputs
    ) -> np.ndarray[np.uint64]:
        with cpu_timer.record("run_workers"):
            request_outputs = self._run_workers(
                "execute_model_step",
                driver_kwargs=dict(
                    seq_metadata_list=seq_metadata_list,
                    scheduler_outputs=scheduler_outputs,
                ),
                is_prompt=scheduler_outputs.prompt_run,
                is_placeholder=scheduler_outputs.is_empty(),
            )
        self.deferred_release.collect_request_outputs(request_outputs)
        return request_outputs

On engine 1, is_placeholder is true. Its worker runs the model on one dummy sequence with one token (deepinfer/worker/worker.py:125-127, deepinfer/worker/model_runner.py:386-403).

Deep dive

Stop 2 in detail: the data on each engine

# engine 0 (rank 0)                       # engine 1 (rank 2)
seq_metadata_list = [A]                     seq_metadata_list = []  # worker uses [placeholder]
is_prompt, is_placeholder = True, False     is_prompt, is_placeholder = True, True
A.seqlen_q = 8, A.seqlen_k = 8              placeholder: seqlen_q = 1, block_table = [0]

Both engines will run a prompt step. That is the point of the vote: the ranks of both engines must call the same collectives in the stops that follow.

3The driver sends the method call to its follower#

We are in WorkerBase.run_workers on ranks 0 and 2. Before the driver runs execute_model_step itself, it sends the method name and the keyword arguments to its CP follower:

    def run_workers(
        self,
        method: str,
        driver_kwargs: dict[str, Any] | None = None,
        **kwargs: Any,
    ) -> Any:
        with cpu_timer.record("broadcast_method", enable=False):
            assert self.is_driver_worker
            cp_group = PG.cpu_cp_group
            if cp_group and cp_group.size() > 1:
                cp_broadcast_object((method, kwargs))

        return getattr(self, method)(**(driver_kwargs or {}), **kwargs)

On ranks 1 and 3, the follower loop gets the tuple and calls the same method (deepinfer/worker/base.py:129-133). The channel is the MessageQueue of the CP group: a ring buffer in shared memory, written by CP rank 0 (deepinfer/model_executor/parallel_utils/parallel_state.py:17-23).

Deep dive

Stop 3 in detail: what the follower gets

The follower gets ("execute_model_step", {"is_prompt": True, "is_placeholder": False}) on rank 1, and the same with is_placeholder = True on rank 3. The driver_kwargs (the sequence list and the scheduler outputs) stay on the driver (deepinfer/engine/llm_engine.py:965-973). So in execute_model_step, a follower has seq_metadata_list = None, and it skips the ZOS correction, which only the driver does (deepinfer/worker/worker.py:116-118).

The pickled tuple is a few tens of bytes, far below the 1 MiB chunk size of the queue (deepinfer/model_executor/parallel_utils/parallel_state.py:124). The page distributed/shm_broadcast.py: MessageQueue explains the queue.

4The driver sends the input metadata#

We are in ModelRunner._prepare_input. The driver builds the input metadata of its batch in C++, moves it to the device, and broadcasts it in the CP group. The follower receives it:

        if PG.context_parallel_group().size() > 1:
            with cpu_timer.record("broadcast_input_metadata"):
                input_metadata_dict = {
                    "is_prompt": is_prompt,
                    "type": type(input_metadata),
                    "run_cuda_graph": run_cuda_graph,
                    "total_q": total_q,
                    "num_reqs": num_reqs,
                }
                if not run_cuda_graph:
                    assert hasattr(input_metadata, "is_prompt")
                    input_metadata_dict.update(input_metadata.to_flatten_dict())
                broadcast_tensor_dict(input_metadata_dict, src=0)

        return input_metadata, total_q, run_cuda_graph, cuda_graph_used, graph_size

    def _prepare_input_follower(self) -> tuple[InputMetadata, int, bool, CudaGraphData | None, int]:
        input_metadata_dict = broadcast_tensor_dict(src=0)
        # 只 pop 非 InputMetadata 字段的 key;is_prompt/total_q/num_reqs 本身也是字段,留给 from_flatten_dict。
        input_metadata_type: type[InputMetadata] = input_metadata_dict.pop("type")
        run_cuda_graph: bool = input_metadata_dict.pop("run_cuda_graph")
        total_q = input_metadata_dict["total_q"]
        num_reqs = input_metadata_dict["num_reqs"]
        if run_cuda_graph:
            graph_size, cuda_graph_used = self.cuda_graph_runner.select(total_q)
            input_metadata = self._graph_input_metadata(cuda_graph_used, num_reqs)
        else:
            input_metadata = input_metadata_type.from_flatten_dict(input_metadata_dict)
            cuda_graph_used = None
            graph_size = 0
        return input_metadata, total_q, run_cuda_graph, cuda_graph_used, graph_size

broadcast_tensor_dict sends a list of keys, small values and (dtype, size) pairs through the shared-memory queue. Then it sends each tensor with a device broadcast from CP rank 0 (deepinfer/model_executor/parallel_utils/communication_op.py:51-82).

Deep dive

Stop 4 in detail: shapes on engine 0

FieldValue on rank 0 and, after the broadcast, on rank 1
input_idsint32 [8], the 8 prompt tokens
positions[8]: 0 to 7
cu_seqlens_q, cu_seqlens_kint32 [2]: [0, 8]
total_q8
cp_mirror_len2 × ⌈8 / 4⌉ = 4 (csrc/engine/input_metadata/prepare_tensor_input_base.cpp:146-155)

On engine 1, the placeholder has total_q = 1 and cp_mirror_len = 2 × ⌈1 / 4⌉ = 2. A second, smaller broadcast follows in stop 9, for the sampler inputs (deepinfer/worker/model_runner.py:250-275).

The broadcast is CP-only: ranks 0 and 1 talk, and ranks 2 and 3 talk. Each CP group has its own queue and its own NCCL group.

5Each rank takes its mirror rows#

We are in DeepseekModel.forward, on all four ranks. Each rank embeds all tokens, computes the mirror indices, and keeps only its own rows:

        if hidden_states is None:
            hidden_states = self.embed_tokens(input_metadata.input_ids)

        residual = None
        # 跨层持有的运行时值 (shared_hidden / pre_rope_kv / topk): 逐层传同一引用,由各层原地写入。
        cross_layer_runtime = CrossLayerRuntime()

        if input_metadata.is_prompt and PG.context_parallel_group().size() > 1:
            input_metadata.prepare_for_mirror_cp()
            hidden_states = gather(hidden_states, input_metadata.mirror_cp_index)

A has 8 tokens and cp = 2, so it has 4 parts of 2 tokens. Rank 0 takes parts 0 and 3; rank 1 takes parts 1 and 2:

Rankmirror_cp_indexHidden states after the gather
0[0, 1, 6, 7][4, H]
1[2, 3, 4, 5][4, H]
2 (placeholder)[0, −1][2, H], row 1 is zeros
3 (placeholder)[−1, −1][2, H], all zeros

The values come from the repository's reference function for the kernel (deepinfer/model_executor/layers/tilelang_ops/prepare_cp_mirror.py:126-157). valid_token_mask is mirror_cp_index != −1 (deepinfer/model_executor/input_metadata/input_metadata.py:208).

Deep dive

Stop 5 in detail: the reverse index

The kernel also returns mirror_cp_reverse_index, one entry per token. For A it is [0, 1, 4, 5, 6, 7, 2, 3]. An all-gather makes the tensor [rank 0 rows | rank 1 rows], and token 2 is row 4 of it. Stops 6 and 9 use it. The page Context parallelism (mirror CP) derives these numbers.

Rank 3 has no real row. It still runs every layer, because its collectives must match the other ranks.

6Attention gathers K and V over the CP group#

We are in the attention layer (v4_modules/attn.py), inside each decoder layer, on all ranks. In a prefill step with CP, the layer opens a CP session and starts the all-gather of its projected K and V. It computes the queries while the data moves:

        use_cp = input_metadata.is_prompt and PG.context_parallel_group().size() > 1
        shared_hidden_bf16 = self._resolve_shared_hidden(hidden_states_bf16, shared_hidden_bf16)
        projected_kv, index_weights = self._project_current_hidden_inputs(hidden_states_bf16)
        yoco_full_kv_stored = cross_layer_runtime.consumer_full_metadata is not None
        if yoco_full_kv_stored:
            # YOCO compact prefill: the query has been trimmed, but the KV still has to cover the full context.
            # Here the full shared hidden is used to generate and write the main/index KV, and the later main path skips the KV update.
            # With CP > 1 an inner cp session is opened here (before the main session below, with no nesting).
            self._store_yoco_full_kv(cross_layer_runtime, kv_cache)
        elif self.shared_hidden_kv_proj is not None:
            # The full KV update path above was not taken and this layer needs to generate KV from the
            # shared hidden (decode / non-compact prefill). Only the KV projection is done here; afterwards
            # _compute_kv_ca_index calls _store_main_kv to do the compression and the write.
            assert projected_kv is None and shared_hidden_bf16 is not None
            projected_kv = self.shared_hidden_kv_proj(shared_hidden_bf16)

        with userbuffer.cp_session(enabled=use_cp):
            # CP: launch the all_gather of the compressor input first, overlapping it with the Q projection below.
            # Under the yoco_full_kv_stored scenario the main branch does not update kv, so this communication step is not needed
            wait_projected_kv = _make_cp_gather_getter(projected_kv, use_cp) if not yoco_full_kv_stored else None

            q_a, kv_swa = self._project_q_and_swa(
                hidden_states,
                shared_hidden_bf16,
            )
            del hidden_states, hidden_states_bf16, shared_hidden_bf16
            kv_swa, wait_kv_swa = self._prepare_swa_kv(
                kv_swa,
                kv_cache.swa_cache,
                input_metadata,
                use_cp,
            )

The all-gather runs through CPManager on a DeepEP BucketBuffer (deepinfer/model_executor/parallel_utils/cp.py:42-51). Each rank sends 4 rows and gets 8 rows: [2 × 4, ...]. The layer reorders them with mirror_cp_reverse_index back to token order (deepinfer/model_executor/models/v4_modules/attn.py:826-829). So rank 0 and rank 1 both store the K and V of all 8 tokens in their KV caches. Each rank computes attention only for its 4 query rows.

Deep dive

Stop 6 in detail: who talks to whom

7MoE: dispatch and combine across both engines#

We are in DeepseekMoE.forward (v4_modules/mlp.py), after the attention of each layer. This is the stop where the two engines meet. The EP group holds all 4 ranks. So a token of A can go to an expert on rank 2 or rank 3, which belong to engine 1:

    def forward(
        self,
        hidden_states: torch.Tensor | tuple[torch.Tensor, torch.Tensor],
        hidden_states_bf16: torch.Tensor,
        input_metadata: InputMetadata,
    ) -> torch.Tensor:
        if self.use_mega_moe:
            return self.mega_moe_forward(hidden_states, hidden_states_bf16, input_metadata)

        topk_idx, topk_weights = self.gate_forward(hidden_states_bf16, input_metadata)

        try:
            # Dispatch
            event_overlap = self.dispatch(
                hidden_states,
                topk_idx,
                topk_weights,
                input_metadata,
                None,
            )
            moe_in_hidden, _dispatched_topk_idx, dispatched_topk_weights, handle = event_overlap.current_stream_wait()

            # MoE compute
            with cuda_graph_timer.record(
                "routed_experts",
                enabled=(self.layer_idx == 1 and not input_metadata.is_prompt),
                scale=eplb.current().num_moe_layers,
            ):
                moe_out_hidden = self.routed_experts_forward(
                    moe_in_hidden,
                    dispatched_topk_weights,
                    handle.psum_num_recv_tokens_per_expert,
                    expected_m=None,
                )

            # Combine
            combine_pre_evt = self.ep_buffer.capture()
            with limit_num_sms(envs.DEEPINFER_USERBUFFER_COMM_SMS, enabled=PLATFORM.is_gpu):
                shared_out = self.shared_forward(hidden_states)
            event_overlap = self.combine(
                moe_out_hidden,
                shared_out,
                handle,
                combine_pre_evt,
            )
            moe_combined_hidden, _combined_topk_weights = event_overlap.current_stream_wait()

Each rank routes its own rows: 4 rows on ranks 0 and 1, 2 rows on ranks 2 and 3. The gate masks the rows with valid_token_mask = False (deepinfer/model_executor/models/v4_modules/mlp.py:387). After the combine, each rank again has its own rows, [4, H] or [2, H].

Deep dive

Stop 7 in detail: why engine 1 must take part

Dispatch and combine are collectives over the EP group. Ranks 2 and 3 have no real token, but the experts on their devices can receive tokens of A. Assume that engine 1 ran a decode step now. It would call dispatch with the decode token limit, and it would run other layers (the draft model) after the forward pass (deepinfer/model_executor/models/v4_modules/mlp.py:476). The vote in stop 1 prevents this.

The page shows the DeepEP path. If the EP group fits in one NVLink domain, the layer can use mega-MoE instead, with the same group (deepinfer/model_executor/models/v4_modules/mlp.py:184-188, Mega-MoE: the whole MoE layer in one call). This site does not explain DeepEP internals.

8EPLB counts the tokens per expert#

The dispatch call also passes a counter: cumulative_local_expert_recv_stats=self._eplb_layer().expert_count (deepinfer/model_executor/models/v4_modules/mlp.py:474). It adds the number of tokens that each local expert slot receives in this layer. After the forward pass, the sampler step calls the load balancer:

        expert_lb = eplb.current()
        if expert_lb is not None:  # dense 模型没有 EPLB
            expert_lb.step(input_metadata.is_prompt)

            def _moe_lb_post_func():
                sampler_output.moe_load_balance_ratio = expert_lb.take_moe_load_balance_ratio()

            sampler_output.post_funcs.append(_moe_lb_post_func)
    def step(self, is_prompt: bool) -> None:
        """由 DeepseekForCausalLM.sample 每 step 调用,更新重排与负载统计。"""
        if self._ep_group.size() <= 1:
            return

        self._update_rearrangement(is_prompt)
        self._update_load_balance_stats()
        self.expert_count.zero_()

With EP above 1, step updates the rearrangement window, but only if rearrangement is on. Then it adds the token counts of the step to the statistics and sets the counters back to zero.

Deep dive

Stop 8 in detail: the EP all-gather every 128 steps

_update_load_balance_stats stores the token sum per layer for this step. Every DEEPINFER_MOE_LB_STAT_INTERVAL steps (default 128), it all-gathers the window over the EP group and computes the mean and the maximum per rank (deepinfer/model_executor/eplb/__init__.py:167-192, deepinfer/environ.py:118). This all-gather is one more reason why all 4 ranks must run the same number of steps, also the placeholder ranks. Shape of the counter: [num_moe_layers, num_local_physical_experts] (deepinfer/model_executor/eplb/__init__.py:49). See Expert load balancing (EPLB).

9The last all-gather, and only the drivers sample#

After the last layer, each CP group gathers the output rows and puts them back in token order. Every rank of engine 0 now has the hidden states of all 8 tokens, [8, H]:

        if input_metadata.is_prompt and PG.context_parallel_group().size() > 1:
            with userbuffer.cp_session():
                gathered_hidden_states = userbuffer.cp.all_gather(hidden_states)()
                hidden_states = gather(gathered_hidden_states, input_metadata.mirror_cp_reverse_index)

Then the worker calls the sampler on all ranks. The driver sends the sampler inputs (row 7 for A: the last prompt token) with a second broadcast_tensor_dict (deepinfer/worker/model_runner.py:250-261, csrc/engine/prepare_sample.cpp:197). A follower makes SamplerInputs with no driver part (deepinfer/worker/model_runner.py:271-275). In the sampler, it computes the logits and returns an empty output:

        logits = get_logits(hidden_states, embedding, self.vocab_size)
        del hidden_states

        if driver_inputs is None:
            if need_prompt_logprobs:
                # 结果只在 driver 上有效;follower 只参与 embedding gather 集合通信。
                compute_prompt_logprobs(
                    prompt_hidden_states,
                    embedding,
                    None,
                    vocab_size=self.vocab_size,
                    force_topk=self.force_topk,
                )
            if self.max_MTP_value > 0 and not is_prompt:
                mtp_module.decode_forward(
                    main_input_metadata=main_input_metadata,
                    nextn_kv_caches=nextn_kv_caches,
                    decode_input=MTPDecodeInput(first_token_positions=decode_inputs.mtp_first_token_positions, driver=None),
                )
            return SamplerOutput(post_funcs=[])

So rank 0 samples the first output token of A. Rank 2 samples a token for the placeholder; the engine ignores it in stop 10.

Deep dive

Stop 9 in detail: shapes in the sampler

  • Hidden states: [8, H] on ranks 0 and 1. On ranks 2 and 3 they are [1, H], because the placeholder has one token (total_q = 1).
  • selected_token_indices = [7] on engine 0: one row goes to the logits, fp32 [1, vocab].
  • A prefill step runs no draft model. In a decode step, the followers would also run the draft layers (deepinfer/model_executor/layers/sampler.py:155-160). The code gives no reason; our reading is that the draft MoE collectives need every rank.

10Outputs#

We are back in Worker.execute_model_step. A prompt step is post-processed at once, also with ZOS (deepinfer/worker/worker.py:182-187). Only the driver calls the engine's callback and gets request outputs:

        with cpu_timer.record("worker.post_process.sync"):
            step_state.sample_end_evt.synchronize()
        with cpu_timer.record("worker.post_process.graph_timer_flush"):
            # flush the time of completed graph, no effect on the running graph.
            cuda_graph_timer.flush_out()
        with cpu_timer.record("worker.post_process.do_process"):
            self.model_runner.post_process(step_state.seq_metadata_list, step_state.sampler_output, step_state.is_prompt)
            request_outputs = np.empty(0, dtype=np.uint64)
            if self.is_driver_worker:
                request_outputs = self.engine_callbacks.post_process_fn(step_state)
        return request_outputs

On rank 0, LLMEngine.post_process gives the token of A to the scheduler and builds a RequestOutput. On rank 2, the step is a placeholder, so the engine uses no sampler output (deepinfer/engine/llm_engine.py:987). Ranks 1 and 3 return an empty array and wait for the next method call.

Deep dive

Stop 10 in detail: what the next step looks like

The prompt step sets decode_step = 0 on both engines (csrc/engine/scheduler.hpp:676-677). That is an ALL point. Engine 1 still has no work, so vote 1 fails, and both engines run decode steps. In a decode step:

What can go wrong#

  • A follower dies. The driver waits in the next CP broadcast or all-gather. If the driver dies, the followers get SIGKILL (deepinfer/worker/base.py:129-130).
  • One idle engine slows the other. Engine 1 runs placeholder steps with full passes through all layers. Its empty ALL votes delay the next prefill of engine 0 (Data parallelism and step votes).
  • A CP group across hosts. The shared-memory queue assumes one host per CP group (deepinfer/distributed/shm_broadcast.py:442).
  • Wrong device order. Each engine sees cp_size devices from CUDA_VISIBLE_DEVICES. A follower uses local_rank = cp_rank % ngpus (deepinfer/engine/llm_engine.py:384). If the variable lists fewer devices than cp_size, two ranks share a device.

Why it matters. This step shows the three kinds of communication in DeepInfer. The DP group only votes. The CP group moves the inputs and the keys and values inside one engine. The EP group ties all engines together in every MoE layer. Most of the rules of the multi-GPU code follow from the last point.

Check yourself#

  1. 1
    In this tour, which ranks run a scheduler, and which ranks sample?
    easy
    Show answer

    Ranks 0 and 2, the drivers of the two engines. The followers (ranks 1 and 3) have no scheduler and return before sampling (deepinfer/model_executor/layers/sampler.py:145-161).

  2. 2
    Through which channel does rank 1 learn that the next method is execute_model_step?
    easy
    Show answer

    Through the shared-memory MessageQueue of the CP group: run_workers calls cp_broadcast_object((method, kwargs)), and the follower loop dequeues it (deepinfer/worker/base.py:300, deepinfer/worker/base.py:132).

  3. 3
    Which stop is the first in which ranks of engine 0 and ranks of engine 1 exchange device data?
    medium
    Show answer

    Stop 7, the MoE dispatch. The votes in stop 1 exchange one-element CPU tensors over the Gloo group, between the drivers only. Stops 3, 4, 6 and 9 stay inside one CP group. The EP group is the world group [0, 1, 2, 3] (deepinfer/model_executor/parallel_utils/parallel_state.py:102-103).

  4. 4
    How many rows does rank 1 send in each attention all-gather of this step, and how many rows does it get back?
    medium
    Show answer

    It sends its cp_mirror_len = 4 rows and gets 2 × 4 = 8 rows: the all-gather result has shape [cp_size × rows, ...] (deepinfer/model_executor/parallel_utils/cp.py:42-51).

  5. 5
    Rank 3 has no real token in this step. Name three collectives that it must still call, and the group of each.
    hard
    Show answer

    (1) The broadcasts of the input metadata and the sampler inputs from rank 2, in the CP group [2, 3] (deepinfer/worker/model_runner.py:217, deepinfer/worker/model_runner.py:261). (2) The CP all-gathers in every attention layer and after the last layer, also in [2, 3] (deepinfer/model_executor/models/deepseek.py:241-244). (3) The MoE dispatch and combine in every layer, in the EP group [0, 1, 2, 3] (deepinfer/model_executor/models/v4_modules/mlp.py:578-610). Every 128 steps, also the EPLB all-gather in the EP group.