[rl] Stateful Trainer Send: IPC [2/N] - #48981
Conversation
Signed-off-by: haoaaron <ahao@anyscale.com>
|
Documentation preview: https://vllm--48981.org.readthedocs.build/en/48981/ |
Signed-off-by: haoaaron <ahao@anyscale.com>
Signed-off-by: haoaaron <ahao@anyscale.com>
|
@codex review |
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 7181dd4a10
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
| ipc_handles = self._all_gather_and_merge_handles(ipc_handles) | ||
| self._do_send( |
There was a problem hiding this comment.
Merge handles on the designated sender
When IPCTrainerInitInfo.rank differs from the default process-group rank—for example, rank 0 is local to a trainer subgroup—the designated sender reaches _do_send, but _all_gather_and_merge_handles returns merged handles only on torch.distributed.get_rank() == 0. The explicit sender therefore receives empty handle dictionaries while the global rank-0 process does not send, causing workers to fail to find their GPU UUID. The merge destination must follow self.is_sender/the supplied trainer rank, or the API must use and document the corresponding process group.
Useful? React with 👍 / 👎.
|
Find a codex comment make sense to me, could you check: The old static path ran _post_send_sync() inside _send_unpacked, while weight_refs was still in scope: beforeweight_refs.append(weight) The new send_weights() moves the barrier outside the scope that owns the tensors: afterself._send(source) # _send_unpacked returns -> weight_refs dropped Under FSDP, materialize_full_tensor calls full_tensor(), so each yielded tensor is a temporary whose only reference is weight_refs. None of the current examples hit it (the two single-process examples send live module params, and the FSDP example uses packed), but the engine docstring explicitly supports multi-rank + unpacked, so it's reachable through the public API. Suggested fix: move _post_send_sync() back inside _send_unpacked (within weight_refs' scope), and drop it from the tail of send_weights(): def _send_unpacked(self, source: WeightSource) -> None: |
|
Maybe we should add backward compatibility? |
| """Update info for IPC weight transfer backend.""" | ||
| """Per-round update info for the IPC weight transfer backend. | ||
|
|
||
| Whether the transfer is packed is a must-agree wire param carried on the |
… vLLM version compat vLLM main (post-0.26.0) removed IPCTrainerSendWeightsArgs and the static trainer_send_weights path via: vllm-project/vllm#48042 — Stateful Trainer Send: New Abstractions [1/N] vllm-project/vllm#48981 — Stateful Trainer Send: IPC [2/N] Add a vllm_version_is conditional branch: - 0.26.0: preserves the existing static NPUIPCTrainerSendWeightsArgs + NPUIPCWeightTransferEngine.trainer_send_weights path unchanged. - main: introduces NPUIPCTrainerInitInfo (backend='npu_ipc'), NPUIPCTrainerWeightTransferEngine (subclass of upstream's IPCTrainerWeightTransferEngine), and delegates HTTP transport to HTTPVLLMWeightSyncClient. Register the trainer-side engine via WeightTransferTrainerFactory only on main. Signed-off-by: hfadzxy <starmoon_zhang@163.com>
Signed-off-by: haoaaron <ahao@anyscale.com> Signed-off-by: aoshen02 <aoshen@inferact.ai>
… vLLM version compat vLLM main (post-0.26.0) removed IPCTrainerSendWeightsArgs and the static trainer_send_weights path via: vllm-project/vllm#48042 — Stateful Trainer Send: New Abstractions [1/N] vllm-project/vllm#48981 — Stateful Trainer Send: IPC [2/N] Add a vllm_version_is conditional branch: - 0.26.0: preserves the existing static NPUIPCTrainerSendWeightsArgs + NPUIPCWeightTransferEngine.trainer_send_weights path unchanged. - main: introduces NPUIPCTrainerInitInfo (backend='npu_ipc'), NPUIPCTrainerWeightTransferEngine (subclass of upstream's IPCTrainerWeightTransferEngine), and delegates HTTP transport to HTTPVLLMWeightSyncClient. Register the trainer-side engine via WeightTransferTrainerFactory only on main. Signed-off-by: hfadzxy <starmoon_zhang@163.com>
… vLLM version compat vLLM main (post-0.26.0) removed IPCTrainerSendWeightsArgs and the static trainer_send_weights path via: vllm-project/vllm#48042 — Stateful Trainer Send: New Abstractions [1/N] vllm-project/vllm#48981 — Stateful Trainer Send: IPC [2/N] Add a vllm_version_is conditional branch: - 0.26.0: preserves the existing static NPUIPCTrainerSendWeightsArgs + NPUIPCWeightTransferEngine.trainer_send_weights path unchanged. - main: introduces NPUIPCTrainerInitInfo (backend='npu_ipc'), NPUIPCTrainerWeightTransferEngine (subclass of upstream's IPCTrainerWeightTransferEngine), and delegates HTTP transport to HTTPVLLMWeightSyncClient. Register the trainer-side engine via WeightTransferTrainerFactory only on main. Signed-off-by: hfadzxy <starmoon_zhang@163.com>
… vLLM version compat vLLM main (post-0.26.0) removed IPCTrainerSendWeightsArgs and the static trainer_send_weights path via: vllm-project/vllm#48042 — Stateful Trainer Send: New Abstractions [1/N] vllm-project/vllm#48981 — Stateful Trainer Send: IPC [2/N] Add a vllm_version_is conditional branch: - 0.26.0: preserves the existing static NPUIPCTrainerSendWeightsArgs + NPUIPCWeightTransferEngine.trainer_send_weights path unchanged. - main: introduces NPUIPCTrainerInitInfo (backend='npu_ipc'), NPUIPCTrainerWeightTransferEngine (subclass of upstream's IPCTrainerWeightTransferEngine), and delegates HTTP transport to HTTPVLLMWeightSyncClient. Register the trainer-side engine via WeightTransferTrainerFactory only on main. Signed-off-by: hfadzxy <starmoon_zhang@163.com>
… vLLM version compat vLLM main (post-0.26.0) removed IPCTrainerSendWeightsArgs and the static trainer_send_weights path via: vllm-project/vllm#48042 — Stateful Trainer Send: New Abstractions [1/N] vllm-project/vllm#48981 — Stateful Trainer Send: IPC [2/N] Add a vllm_version_is conditional branch: - 0.26.0: preserves the existing static NPUIPCTrainerSendWeightsArgs + NPUIPCWeightTransferEngine.trainer_send_weights path unchanged. - main: introduces NPUIPCTrainerInitInfo (backend='npu_ipc'), NPUIPCTrainerWeightTransferEngine (subclass of upstream's IPCTrainerWeightTransferEngine), and delegates HTTP transport to HTTPVLLMWeightSyncClient. Register the trainer-side engine via WeightTransferTrainerFactory only on main. Signed-off-by: hfadzxy <starmoon_zhang@163.com>
… vLLM version compat vLLM main (post-0.26.0) removed IPCTrainerSendWeightsArgs and the static trainer_send_weights path via: vllm-project/vllm#48042 — Stateful Trainer Send: New Abstractions [1/N] vllm-project/vllm#48981 — Stateful Trainer Send: IPC [2/N] Add a vllm_version_is conditional branch: - 0.26.0: preserves the existing static NPUIPCTrainerSendWeightsArgs + NPUIPCWeightTransferEngine.trainer_send_weights path unchanged. - main: introduces NPUIPCTrainerInitInfo (backend='npu_ipc'), NPUIPCTrainerWeightTransferEngine (subclass of upstream's IPCTrainerWeightTransferEngine), and delegates HTTP transport to HTTPVLLMWeightSyncClient. Register the trainer-side engine via WeightTransferTrainerFactory only on main. Signed-off-by: hfadzxy <starmoon_zhang@163.com>
### What this PR does / why we need it? | Files | Upstream vLLM change | vllm-ascend adaptation | |-------|---------------------|------------------------| | `vllm_ascend/__init__.py` | [vllm#48841](vllm-project/vllm#48841) — added `from triton.experimental import gluon`, `from triton.experimental.gluon import language as gl` and `from triton.language.core import _aggregate` to `triton_utils/__init__.py`, requiring Triton 3.6+ API absent from `triton-ascend 3.2.1` | Pre-register `triton.experimental.gluon` / `.language` as `sys.modules` stubs; stub `triton.language.core._aggregate`. Gated on `os.getenv("VLLM_VERSION") != "0.26.0"`, mirroring `vllm_version_is` env var path. | | `tests/ut/patch/platform/test_patch_structured_output.py` | [vllm#49665](vllm-project/vllm#49665) — changed `VLLMValidationError` base class from `ValueError` to `VLLMClientError`, breaking `pytest.raises(ValueError)` assertions | Updated 3 assertions from `pytest.raises(ValueError, ...)` to `pytest.raises(VLLMValidationError, ...)` | | `tests/e2e/pull_request/one_card/test_guided_decoding.py` | [vllm#49665](vllm-project/vllm#49665) — same `VLLMValidationError` base-class change; on 0.26.0 the upstream validation may still raise `ValueError` | Version-gated assertion: `pytest.raises(ValueError, ...)` on 0.26.0 vs `pytest.raises(VLLMValidationError, ...)` on main | | `vllm_ascend/ops/vocab_parallel_embedding.py`<br>`vllm_ascend/_310p/ops/vocab_parallel_embedding.py` | [vllm#49731](vllm-project/vllm#49731) — added `*, disable_tp: bool = False` keyword arg to `ParallelLMHead.__init__()` and `VocabParallelEmbedding.__init__()` | Version-gated `__init__` with `vllm_version_is("0.26.0")`: else branch accepts `disable_tp` kwarg and forwards it (Ascend manages TP via `lmhead_tp_enable()`) | | `vllm_ascend/distributed/kv_transfer/kv_p2p/mooncake_connector.py`<br>`vllm_ascend/distributed/kv_transfer/kv_p2p/mooncake_hybrid_connector.py` | [vllm#26593](vllm-project/vllm#26593) — `KVConnectorBase_V1.__init__()` sets `self._kv_transfer_config`, required by `requires_kv_delivery` property | Version-gated `__init__`: else branch sets `self._kv_transfer_config` on vllm main | | `vllm_ascend/ops/fused_moe/fused_moe.py` | [vllm#50089](vllm-project/vllm#50089) — added `fused_output_is_reduced` kwarg to `_maybe_reduce_shared_expert_output()` and `output_is_reduced` kwarg to `_maybe_reduce_final_output()` | Version-gated both methods: else branch accepts new kwargs but ignores them (Ascend handles reduction independently); `trunc_size` may now be `None` | | `vllm_ascend/patch/worker/patch_distributed.py` | [vllm#47288](vllm-project/vllm#47288) — added `use_all2all: bool = False` param to `GroupCoordinator.__init__()` | Version-gated via `_IS_VLLM_026` env var check (avoids vllm-ascend import cycle). Else branch stores `self.use_all2all` (unused, no all2all on Ascend) | | `vllm_ascend/distributed/device_communicators/npu_communicator.py` | [vllm#47288](vllm-project/vllm#47288) — added `use_all2all: bool = False` param to `DeviceCommunicatorBase.__init__()` | Version-gated `NPUCommunicator.__init__` with `vllm_version_is("0.26.0")`: else branch accepts and forwards `use_all2all` (NPU keeps the no-op `_NpuAll2AllManager`) | | `vllm_ascend/worker/npu_input_batch.py` | [vllm#48018](vllm-project/vllm#48018) — added `use_replayssm: bool` kwarg (Mamba replay-SSM); [vllm#40996](vllm-project/vllm#40996) — added `slot_mapping_modes` kwarg (DCP hybrid attention) | Accepts both kwargs unconditionally; stores `self.use_replayssm` and `self.slot_mapping_modes` only on main (`not vllm_version_is("0.26.0")`) for interface alignment | | `vllm_ascend/worker/v2/spec_decode/eagle/aclgraph.py` | [vllm#49364](vllm-project/vllm#49364) — renamed `skip_attn` parameter to `full_cudagraph` in cudagraph mode logic | Version-gated call: 0.26.0 keeps `skip_attn=(cg_mode != PIECEWISE)`; main passes `full_cudagraph=(cg_mode != PIECEWISE)` | | `vllm_ascend/distributed/weight_transfer/npu_ipc_engine.py`<br>`vllm_ascend/distributed/weight_transfer/__init__.py` | [vllm#48042](vllm-project/vllm#48042) + [vllm#48981](vllm-project/vllm#48981) — replaced static `IPCTrainerSendWeightsArgs`/`trainer_send_weights` with a stateful `IPCTrainerWeightTransferEngine` driven by `WeightTransferTrainerFactory.trainer_init(...).send_weights()`; `packed` moved onto trainer init-info; per-round `update_info` slimmed | Whole module version-gated: 0.26.0 keeps the static `NPUIPCWeightTransferEngine`; main defines stateful `NPUIPCTrainerWeightTransferEngine` (`_send_unpacked` instance method, `packed` on `NPUIPCWeightTransferInitInfo`, no-arg `npu_generate_uuid()`, `is_sender`/`_all_gather_and_merge_handles`). `register_engine()` also registers the trainer engine in `WeightTransferTrainerFactory` on main only | | `examples/rl/rlhf_http_npu_ipc.py` | [vllm#48981](vllm-project/vllm#48981) — ported RL examples to `WeightTransferTrainerFactory.trainer_init(...)` + `engine.send_weights()` with `HTTPVLLMWeightSyncClient`/`ModuleSource` | Version-gated example: 0.26.0 uses static `NPUIPCWeightTransferEngine.trainer_send_weights`; main uses the stateful trainer engine path | | `tests/e2e/pull_request/one_card/test_npu_ipc_weight_transfer.py` | [vllm#48981](vllm-project/vllm#48981) — IPC weight-transfer E2E now exercised via the stateful trainer engine | Version-gated test: 0.26.0 keeps the manual `_post(start/finish)` + static `trainer_send_weights`; main drives `WeightTransferTrainerFactory.trainer_init(...).send_weights()` (engine owns the lifecycle) | | `tests/ut/distributed/weight_transfer/test_npu_ipc_engine.py` | [vllm#48981](vllm-project/vllm#48981) — trainer-side IPC APIs changed from static methods to instance methods on the stateful engine | Version-gated tests (`IS_VLLM_026`): `_send_unpacked`, `packed` placement, `npu_generate_uuid` arity, and `start/finish_weight_update` no-op behavior asserted per version | | `vllm_ascend/worker/v2/spec_decode/dflash/speculator.py` | [vllm#50000](vllm-project/vllm#50000) — `_prepare_dflash_inputs_kernel` gained `temperature`/`seeds` sampling params for probabilistic draft sampling | Version-gated `_prepare_dflash_inputs_kernel_ascend`: 0.26.0 branch keeps the old signature; main branch adds the four `temperature`/`seeds` pointers and the corresponding stores to stay aligned with upstream while keeping Ascend's own kernel | | `.github/vllm-main-verified.commit` | — | Updated verified main commit hash to `0351e9aa1fdf1a51329d1906881528dfe61fc88e` | ### Does this PR introduce _any_ user-facing change? ### How was this patch tested? - vLLM version: v0.26.0 - vLLM main: vllm-project/vllm@d02df74 --------- Signed-off-by: hfadzxy <starmoon_zhang@163.com>
Migrate the dense NCCL and sparse NCCL backends onto the stateful trainer engine, completing the trainer-side weight-transfer rework started in vllm-project#48042 (abstractions) and vllm-project#48981 (IPC). - `NCCLTrainerWeightTransferEngine`: rank 0 holds the `PyNcclCommunicator` and owns the concurrency, running the inference-side `update_weights` on a side thread while it broadcasts (both rendezvous inside the same NCCL calls), with a `future.done()` early-error check so a request rejected before any NCCL call surfaces instead of hanging the broadcast. - Wire params ride the init info, mirroring IPC: `NCCLTrainerInitInfo` carries the rendezvous fields plus the must-agree `packed` / `packed_buffer_size_bytes` / `packed_num_buffers`, and the sender ships them to the worker at `trainer_init`, so the two sides cannot disagree. `NCCLWeightTransferUpdateInfo` is slimmed to per-round `names` / `dtype_names` / `shapes`, and the worker reads `self.packed` recorded at the init handshake. - `SparseNCCLTrainerWeightTransferEngine` is modeled as a delta backend: sparse patches differ every round, so they are not a stable `WeightSource`. The engine takes no `source`; each round's patches go to `send_weights(patches)`. `SparseWeightPatch` gains `full_shape`, required on this path. - `source` is now optional on `TrainerWeightTransferEngine` / the factory to support delta backends; full-resync backends (NCCL, IPC) reject `None` themselves. - Multi-rank trainers: every rank builds the engine and calls `send_weights()`. Non-senders hold no communicator and skip the client RPCs and the broadcast, but still iterate the `WeightSource` and run `metadata()` to stay in the trainer-side collectives (e.g. the FSDP `full_tensor()` all-gather). - Removes the static NCCL trainer API (`trainer_send_weights`, `NCCLTrainerSendWeightsArgs`, and the `NCCLWeightTransferEngine.trainer_init` re-export); ports the NCCL / HTTP / FSDP-EP / sparse examples to `WeightTransferTrainerFactory.trainer_init(...).send_weights()`. Signed-off-by: haoaaron <ahao@anyscale.com> Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Signed-off-by: haoaaron <ahao@anyscale.com>
Signed-off-by: haoaaron <ahao@anyscale.com> Signed-off-by: Tej Kiran <kiran.tej@amd.com>
Signed-off-by: haoaaron <ahao@anyscale.com> Signed-off-by: root <root@smci355-ccs-aus-m02-09.cs-aus.dcgpu>
The stateful IPCTrainerWeightTransferEngine (vllm-project/vllm#48981, #48042) is present at v0.27.1 and main, and the static trainer_send_weights path's IPCTrainerSendWeightsArgs was removed upstream before v0.27.1. Remove the vllm_version_is("0.26.0") gates and keep only the stateful path for both refs. Signed-off-by: liaoqidan <1107297340@qq.com>
Context
Second of the planned three-PR split of the trainer-side weight-transfer rework.
(
WeightSource/ModuleSource,VLLMWeightSyncClient,TrainerWeightTransferEngine,WeightTransferTrainerFactory). Purelyadditive; no backend migrated. [rl] Stateful Trainer Send: New Abstractions [1/N] #48042
end-to-end — trainer engine + wire params on the init info + update-info
slimming + IPC examples/tests. First PR with an end-to-end correctness test of
the new trainer path.
trainer_send_weightsfrom the worker ABC, and port the remaining examples/tests/docs.
NCCL and sparse NCCL are untouched here and keep their existing static trainer
path; the worker ABC still declares
trainer_send_weights(removed in PR 3). Docs remain unchanged and are deferred to PR 4.What changed
IPC engine (
ipc_engine.py)IPCTrainerWeightTransferEngine— stateful trainer engine driven by aparameter-free
send_weights(). Straight-line (no concurrent broadcast likeNCCL):
update_weightsis the transfer and rides the client. For multi-ranktrainers all ranks iterate the source and join the IPC-handle all-gather; only
the sender (rank 0) ships the merged handles and drives the inference-side RPCs,
with a per-chunk barrier in packed mode so producers don't overwrite a buffer a
consumer is still reading. It reads
packed/packed_buffer_size_bytesfromits
IPCTrainerInitInfo.IPCWeightTransferUpdateInfoslimmed: thepackedfield is gone; per-roundpayloads carry only per-round metadata + IPC handles.
IPCWeightTransferEnginereadspackedfromself.packed(set atthe init handshake), not from a per-round field or the config.
IPCWeightTransferEngine.trainer_send_weightsandIPCTrainerSendWeightsArgsare removed; a transitional stub remains on theworker engine (raises
NotImplementedError) only to satisfy the still-abstractworker ABC member, which PR 3 deletes.
Factory (
factory.py)WeightTransferTrainerFactoryregisters theipctrainer engine (lazy import).NCCL / sparse register in PR 3.
TrainerInitInfoto define backend str.Examples & tests
examples/rl/rlhf_ipc.py,rlhf_http_ipc.py,rlhf_ipc_fsdp_ep.pytoWeightTransferTrainerFactory.trainer_init(...).send_weights()with aVLLMWeightSyncClient(Ray / HTTP) andModuleSource. The manualinit/start/finish juggling collapses to a single
send_weights()call; theinference side takes a plain
WeightTransferConfig(backend="ipc")and packingis set on
IPCTrainerInitInfo.tests/distributed/test_weight_transfer.py: IPC worker tests use astring-only config and learn
packedfrom the init handshake; thetrainer-factory registry test asserts
ipcis registered; addedtest_ipc_trainer_send_weights_drives_client_in_order(start→update→finishordering,
packedoff the per-round update-info) andtest_ipc_trainer_init_ships_packed_to_worker(asserts the init handshakecarries
{"packed": ...}).docs/training/weight_transfer/*rewrite lands with theNCCL migration (PR 3), when the trainer-side story is described once for all
backends.
Breaking changes (IPC only)
IPCWeightTransferEngine.trainer_send_weights/
IPCTrainerSendWeightsArgs. UseWeightTransferTrainerFactory.trainer_init(backend="ipc", ...).send_weights().WeightSource(e.g.ModuleSource(model)).packed/packed_buffer_size_bytesmove ontoIPCTrainerInitInfoandare no longer per-round
update_infofields; the trainer propagatespackedtothe worker at init.
Migration (IPC)
Before:
After:
The inference side is constructed with the same plain config
(
LLM(..., weight_transfer_config=WeightTransferConfig(backend="ipc"))orvllm serve ... --weight-transfer-config '{"backend": "ipc"}'); it learnspackedfrom the trainer at the init handshake.Testing
pytest tests/distributed/test_weight_transfer.py # 58 passed (8xH100)Covers IPC update-info validation, IPC handle (de)serialization, the worker
learning
packedfrom the init handshake, and the new IPC trainer engine(
send_weightsordering +trainer_initpropagatingpacked), plusmulti-process IPC integration (Ray + HTTP).
All three examples were run end-to-end on 8xH100 (facebook/opt-125m for
rlhf_ipc.py/rlhf_http_ipc.py; Qwen3-30B-A3B FSDP2 + expert-parallel,packed, for
rlhf_ipc_fsdp_ep.py) — each produces garbage from thedummy-initialized server before the sync and coherent text after.