feat(vllm): pause generation during in-flight refit - #3839
Conversation
Signed-off-by: Youngeun Kwon <youngeunk@nvidia.com>
|
/ok to test 00ab4ef |
Signed-off-by: Youngeun Kwon <youngeunk@nvidia.com>
|
/ok to test abf3b2a |
GPU validationValidated exact PR head Full non-Gym performance recipe (10 steps, generation logging enabled)
Additional integration coverage
|
youngeunkwon0405
left a comment
There was a problem hiding this comment.
Team review by 5 coordinated agents (RL-code, vLLM-upstream verification, tests, independent bug scan, adversarial verification). No critical findings — 5 suggestions and 1 informational note posted inline. No merge conflicts; CI already triggered on this head.
Upstream verification (vLLM v0.25.1, pinned): AsyncLLM.pause_generation(*, mode="keep", clear_cache=...) / resume_generation() usage is exactly right — signature, PauseMode = Literal["abort","wait","keep"]. Keep-mode pause quiesces the engine before returning (idle-callback future), and weight-update collective_rpc utility calls still execute while paused — so the weight write can no longer race an in-flight forward pass. The broadcast axes (run_rank_0_only_axes=["tensor_parallel", "pipeline_parallel"]) match every sibling engine-wide op.
Nice side effect worth calling out: the old recompute_kv_cache_after_weight_updates=true in-flight path was effectively a silent no-op — reset_prefix_cache refuses to reset while requests hold KV blocks, reset_prefix_cache_async discards the returned False, and invalidate_kv_cache's all() over the filtered results then reported success. This PR is the first path that actually delivers in-flight KV recompute, via pause-time clear_cache (preempts running requests and frees their KV).
Minor non-actionable note: pause_generation/resume_generation raise on misconfig and propagate RayTaskError on worker failure, so the interface's False contract (and the collector's "Failed to pause" branch) is unreachable for the real backend — fail-loud is right per the error-handling guidelines; just a small signaling inconsistency with invalidate_kv_cache's catch-and-return-False.
Tests: all 5 new tests verified against the extracted source (simulation harness; runtime confirmation left to CI as noted in the PR description) and are correctly bucketed (Algorithms + vLLM base shards; pytest.mark.asyncio strict-mode marker correct). Linter: ruff format + check pass on all 6 changed files.
Generated by Claude Code
Signed-off-by: Youngeun Kwon <youngeunk@nvidia.com>
|
/ok to test c815f32 |
|
Claude review follow-up is in Addressed:
Validation on CW-DFW:
Scope decisions are recorded in the inline threads: matched-main performance/convergence comparisons are useful follow-up characterization rather than a merge gate for this correctness PR, and the bounded PPO critic-warmup pause inefficiency needs broader control-flow signaling. The existing full 10-step recipe runtime validation was run on |
Signed-off-by: Youngeun Kwon <youngeunk@nvidia.com>
|
/ok to test 74b9063 |
|
Follow-up design update in
CW-DFW validation: Slurm |
Signed-off-by: Youngeun Kwon <youngeunk@nvidia.com>
|
/ok to test f94fd34 |
|
Addressed the second Claude review in f94fd34:
Validation: CW-DFW Slurm job 16757596 completed 0:0 in the requested nemo_rl.0825.sqsh image. All 14 focused tests passed, including test_async_trajectory_collector_prepare_for_refit, and Ruff 0.9.9 check plus format verification passed. |
Signed-off-by: Youngeun Kwon <youngeunk@nvidia.com>
Signed-off-by: Youngeun Kwon <youngeunk@nvidia.com>
|
/ok to test cb8948d |
|
Updated the branch to current main to clear the freshness gate. The previous failure was because the PR was 12 commits behind while CI allows at most 10.\n\nCurrent main introduced a separate pause_generation(mode) / continue_generation() contract for the SGLang WeightSynchronizer. To avoid a semantic signature collision, this PR now uses the refit-specific hooks pause_generation_for_refit(clear_cache=...) and resume_generation_after_refit(), leaving the SGLang WeightSynchronizer API unchanged.\n\nValidation on current main: CW-DFW Slurm job 16764115 completed 0:0; 14 focused async-refit/vLLM tests and all 57 WeightSynchronizer tests passed; Ruff 0.9.9 check and format verification passed. |
|
Hi @terrykong, I did the self-review for this PR line by line. Can I ask for your final review for the merge? |
|
Post-refactor real-GPU validation completed successfully for both KV-cache modes. Common setup:
Preserve existing KV cache (
|
Signed-off-by: Youngeun Kwon <youngeunk@nvidia.com>
|
/ok to test 537edb0 |
terrykong
left a comment
There was a problem hiding this comment.
Third-round review by a team of six agents (RL expert, bug finder, test reviewer, design reviewer, comment reviewer, devil's advocate), focused on what changed since the last round.
One docs correction below. No blockers, and no code changes requested.
It is worth listing what we chased down and cleared, so it does not get re-raised:
- vLLM API usage is correct.
AsyncLLM.pause_generation(mode="keep", clear_cache=...)andresume_generation()match v0.25.1 exactly, including"keep"being a validPauseMode. - Skipping
invalidate_kv_cache()on the native path is right, and stronger than what it replaces.clear_cache=Truereachesreset_prefix_cache(reset_running_requests=True), which preempts running requests and drops their blocks, so preserved requests really do re-prefill. The oldreset_prefix_cache()was a no-op while running requests still held KV. - The broadcast is correct.
run_rank_0_only_axes=["tensor_parallel", "pipeline_parallel"]matches every sibling call and reaches every DP rank, which matters becauseresume_schedulerdoes an all-reduce across DP ranks. - The predicate swap changes no default.
should_use_async_rolloutsmatches the old chain for vLLM, Megatron, TRT-LLM and Dynamo. SGLang is the only backend whose result can differ, and only whenuse_async_rollouts: true, which defaults toFalseand is set tofalsein the one config that mentions it. Hence the docs note rather than a code comment. - No deadlock, no missing state reset, no threading issue. The two new flags are only touched on the collector's actor thread; rollout threads use the
threading.Event. - The trtllm assert was not lost — it lives inside
should_use_async_rollouts. - Linters pass on all changed files.
Thanks for the pause/resume approach here — routing through vLLM's native keep-mode is a real correctness improvement over writing weights into an actively decoding engine, and the 10-step 32xH100 run with 10 matched pause/resume cycles is good evidence it works end to end.
Generated by Claude Code
Co-authored-by: Terry Kong <terryk@nvidia.com> Signed-off-by: Youngeun Kwon <youngeunk@nvidia.com>
|
/ok to test 83e7e5b |
terrykong
left a comment
There was a problem hiding this comment.
nit on the docs but overall lgtm. thanks for finally
Two conflicts, both from main moving under the telemetry work: * pyproject.toml -- kept the nemo-lens base dependency and took main's tilelang entry, which broadens the marker to aarch64. taplo then de-aligned the nvidia-cudnn-cu13 trailing comment, since the inserted nemo-lens comment block ends its alignment group. * trajectory_collector.py -- import-only. main's "pause generation during in-flight refit" (#3839) added should_use_async_rollouts alongside the generation interfaces; taken as a union with the telemetry imports. The collector body auto-merged. prepare_for_refit still gates on _refit_pause_cleared, including on the new early-return path in resume_generation_after_refit, so idle/refit_event_wait still wraps the wait it was written for. Signed-off-by: Raj Singh <rajsin@nvidia.com>
PR3 (#3591) was SQUASH-merged into main as b3b6713, so none of its commits are ancestors of main while PR4 still carries all of them. Git therefore sees PR3's whole diff as independently added on both sides, which is why all 16 conflicts name b3b6713 and why the PR showed CONFLICTING despite the content being identical. That made the classification, not the content, the work. For each conflicted file: is main's version byte-identical to PR3's head (3d9ce21), and does PR4 add anything beyond it? Three groups fell out. GROUP A -- pure squash artefacts, resolved by taking OURS (10 files) fleet_health.py, collective_weight_synchronizer.py, membership.py, nccl_reshard_weight_synchronizer.py, grpo_sc_generation_shard_recovery.sh, test_watchdog_pump.py, test_membership.py, test_reconcile_communicator.py, test_reshard_rebuild.py, test_weight_synchronizer.py main == PR3 exactly and no other PR touched them, so PR4's side is main's content plus PR4's delta. Taking ours loses nothing. GROUP B -- PR4 contributes nothing, resolved by taking THEIRS (2 files) single_controller_utils/setup.py (#3480, #3727, #3821 on top of PR3) tests/unit/single_controller/test_refit_recovery.py (#3480 on top of PR3) GROUP C -- genuine merges (4 files), one per upstream PR below. The six upstream PRs that contributed real content, and what each needed: #3480 recover replay buffer from native TQ checkpoints single_controller.py: rollout_recovery imports. Kept alongside ours. setup.py, test_refit_recovery.py, L1 harness: group B / additive. #3765 log toolcall and thinktag violation rate single_controller.py: VIOLATION_TAG_KEYS. Auto-merged, verified present. #3727 support non-colocated MInf single_controller.py: MegatronGeneration import, kept alongside ours. L1 harness: grpo_megatron_generation_gym_single_controller.sh entry. #3821 warm-start the value model from a critic-pretrain checkpoint config.py: the max_num_epochs validator. Ours only adds restart_dead_shards to FleetHealthConfig, so both survive; verified the field landed in the right class and the validator is intact. #3655 nemo-lens telemetry vllm_generation.py: the @trace_fn decorator on generate. Ours adds restart_shard in a different region; both kept. #3839 pause generation during in-flight refit vllm_generation.py: pause_generation_for_refit / resume_generation_after_refit. Auto-merged, verified present -- worth knowing it exists, since it pauses engines around a refit and this PR restarts them. Verified after resolving: no conflict markers; all four lint hooks clean (the single pyrefly error is the pre-existing unrelated transfer_queue import); 1122 unit tests pass; both submodule pointers and uv.lock/pyproject byte-identical to main. Both sides' work was checked individually rather than assumed: EngineSupervisor wiring, restart_dead_shards, restart_shard, recreate_worker, desired_membership and the report_refit call on our side; the six items above on main's. Note for anyone reproducing locally: #3655 adds a nemo-lens dependency that the pre-merge container image does not carry, so tests fail at import with ModuleNotFoundError: nemo until the venv is refreshed. Plain upstream/main fails the same way in that image; it is not a merge defect. Signed-off-by: asolergibert <asolergibert@nvidia.com>
…le group
Job 6718090 ran the full lane on PR4 and split cleanly by transport: every
collective-transport recovery variant FAILED (recovery, recovery-refit, recovery-restart,
recovery-frozen, gym-router-failover-survival) and every nccl_reshard one PASSED.
The rebuild was not the problem -- it succeeded, and said so:
refit: rebuilding communicator over shards [1]; world_size 3, port 26813
and the run died 20s later, out of the recovery itself:
collective_weight_synchronizer.py:288 reconcile_communicator
self._generation.prepare_refit_info(state_dict_info)
-> vllm_generation.py ray.get(futures)
-> ActorDiedError: VllmAsyncGenerationWorker vllm_policy-0-0
prepare_refit_info fanned out with run_all_workers_single_data, which walks the whole
worker group -- including the shard the reconcile had just removed. The nccl_reshard
synchronizer passes because its reconcile makes no such call.
This call is PR4's own (a47e203), added so a RESTARTED engine gets state_dict_info it
would otherwise never have; update_weights_from_collective asserts on it. The intent is
right and the fan-out was wrong: the path runs precisely BECAUSE a shard is absent, so
addressing the whole group is wrong exactly when it executes.
Three sibling fan-outs had already been converted to _refit_leader_workers, and one of
them carries the comment that describes this failure exactly -- "after a shard is lost that
would call its dead actor and fail the refit, undoing the rebuild that just happened".
prepare_refit_info was the fourth and was missed. That is the fifth instance of the bug
class in design_vllm_fault_tolerance.md section 8.5.5, and the first to be caught by a
functional lane rather than by reading.
_refit_leader_workers is the same set the group call reached -- rank 0 over
tensor_parallel/pipeline_parallel -- restricted to the live membership, and it falls back
to every leader when no membership is recorded, which is every run that never loses a shard.
test_refit_fan_outs_address_surviving_leaders_only now covers all four by name and asserts
both halves: the surviving-leader dispatch is present AND run_all_workers_single_data is
absent. Restoring the old call fails it with the message above.
Two neighbours inspected and deliberately NOT changed:
invalidate_kv_cache still walks the whole group, but wraps everything in
try/except Exception and returns False, so a dead shard degrades it rather than
killing the run. Pre-existing, and the survivors' caches going un-invalidated after a
weight update is a real but separate correctness question.
pause_generation_for_refit / resume_generation_after_refit arrived from upstream #3839
and are called from trajectory_collector, not from the SingleController refit path.
Signed-off-by: asolergibert <asolergibert@nvidia.com>
Summary
Implementation
The trajectory collector still stops new rollout batches first. For async vLLM with in_flight_weight_updates enabled, VllmGeneration broadcasts pause_generation_async to every DP engine leader. Each worker calls AsyncLLM.pause_generation(mode="keep", clear_cache=...). After refit, the collector broadcasts resume_generation_async before admitting new rollout work.
Other backends and the drain-before-refit path retain their existing behavior.
Validation
Post-refactor real-GPU validation completed successfully for both KV-cache modes.
Common setup:
grpo-qwen3-30ba3b-4n8g-async-1off(non-Gym perf path)train_global_batch_size=2048,max_total_sequence_length=4096coreai_dlalgo_genaiPreserve existing KV cache (
recompute_kv_cache_after_weight_updates=false)16780048,COMPLETED(0:0) in00:33:54Ready for refittransitions, and 10 resume requestsAsync GRPO training complete!0.0077, average reward0.5942, validation accuracy0.5430Recompute KV cache after weight updates (
recompute_kv_cache_after_weight_updates=true)16783855,COMPLETED(0:0) in00:34:20kv_cache_management_modechanged frompersisttorecomputeReady for refittransitions, and 10 resume requestsAsync GRPO training complete!-0.0018, average reward0.5942, validation accuracy0.5430Together these runs exercise the final backend-agnostic collector contract and the vLLM native pause/resume implementation with both
clear_cache=falseandclear_cache=true.