Tour: One step on many GPUs
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:
| Rank | Engine | cp_rank | Process | Device |
|---|---|---|---|---|
| 0 | 0 | 0 (driver) | engine 0, started by _launch.py | 0 |
| 1 | 0 | 1 (follower) | spawned by engine 0 | 1 |
| 2 | 1 | 0 (driver) | engine 1, started by _launch.py | 2 |
| 3 | 1 | 1 (follower) | spawned by engine 1 | 3 |
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.
| Stop | What happens | Ranks | Group |
|---|---|---|---|
| 1 | The engines vote for a prefill step | 0, 2 | DP (Gloo) |
| 2 | Each engine schedules its batch | 0, 2 | none |
| 3 | The driver sends the method call | 0 → 1, 2 → 3 | CP (shared memory) |
| 4 | The driver sends the input metadata | 0 → 1, 2 → 3 | CP (shared memory + NCCL) |
| 5 | Each rank takes its mirror rows | all | none |
| 6 | Attention gathers K and V | all | CP (DeepEP BucketBuffer) |
| 7 | MoE dispatch and combine | all | EP |
| 8 | EPLB counts the tokens per expert | all | EP (every 128 steps) |
| 9 | The last all-gather, and the drivers sample | all; sampling on 0, 2 | CP |
| 10 | Outputs | 0, 2 | none |
The picture#
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 diveStop 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_outputsOn 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 diveStop 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 diveStop 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_sizebroadcast_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 diveStop 4 in detail: shapes on engine 0
| Field | Value on rank 0 and, after the broadcast, on rank 1 |
|---|---|
input_ids | int32 [8], the 8 prompt tokens |
positions | [8]: 0 to 7 |
cu_seqlens_q, cu_seqlens_k | int32 [2]: [0, 8] |
total_q | 8 |
cp_mirror_len | 2 × ⌈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:
| Rank | mirror_cp_index | Hidden 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 diveStop 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 diveStop 6 in detail: who talks to whom
- Ranks 0 and 1 exchange the K and V of A. Ranks 2 and 3 each send 2 rows; of these 4 rows, only one is a real token.
- The two CP groups do not talk to each other in this stop.
- The CP buffer was sized at start-up from
max_num_prefilling_tokens = max_num_batched_tokens // cp_size(deepinfer/model_executor/model_loader.py:120,deepinfer/model_executor/parallel_utils/userbuffer.py:97-117). - Gathered tensors alias the buffer. The next session overwrites them, so the layer copies what it keeps (
deepinfer/model_executor/parallel_utils/cp.py:10-14).
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 diveStop 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 diveStop 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 diveStop 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_outputsOn 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 diveStop 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:
- No mirror layout and no CP all-gather: both checks need
is_prompt(deepinfer/model_executor/models/deepseek.py:168,deepinfer/model_executor/models/v4_modules/attn.py:766). Ranks 0 and 1 both run A's decode tokens, with the full KV cache that stop 6 stored. - With a CUDA graph, the driver sends only five small values; the graph broadcasts the metadata tensors itself (
deepinfer/worker/model_runner.py:202-209,deepinfer/worker/cuda_graph_runner.py:61-65). - Engine 1 runs placeholder decode steps until it gets work.
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_sizedevices fromCUDA_VISIBLE_DEVICES. A follower useslocal_rank = cp_rank % ngpus(deepinfer/engine/llm_engine.py:384). If the variable lists fewer devices thancp_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#
- 1In 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). - 2Through which channel does rank 1 learn that the next method iseasy
execute_model_step?Show answer
Through the shared-memory MessageQueue of the CP group:
run_workerscallscp_broadcast_object((method, kwargs)), and the follower loop dequeues it (deepinfer/worker/base.py:300,deepinfer/worker/base.py:132). - 3Which 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). - 4How 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 = 4rows and gets2 × 4 = 8rows: the all-gather result has shape[cp_size × rows, ...](deepinfer/model_executor/parallel_utils/cp.py:42-51). - 5Rank 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.