[Core] Release NCCL communicator memory in sleep mode by aoshen02 · Pull Request #46234 · vllm-project/vllm

aoshen02 · 2026-06-20T14:58:18Z

What

Sleep mode (CuMemAllocator) already offloads weights and discards the KV cache
while an engine sleeps, but each PyNccl communicator keeps its dynamic GPU
buffers resident the whole time. On an 8-GPU TP+EP deployment that is
~958 MiB/GPU sitting idle through a sleep/wake cycle (multi-model hot-swap,
RL weight sync, etc.).

This PR releases that memory with NCCL's native
ncclCommSuspend(NCCL_SUSPEND_MEM) / ncclCommResume (NCCL ≥ 2.29.7), which
frees a communicator's allocations while keeping its topology and connection
state, so resume is cheap — no re-init, no bootstrap rendezvous.

How

Implemented along vLLM's existing NCCL conventions — no new module:

pynccl_wrapper.py — add ncclCommSuspend/ncclCommResume to
NCCLLibrary.exported_functions + wrapper methods. They exist only in
NCCL ≥ 2.29.7, so the binding skips them gracefully (one-time warning) on
older NCCL.
pynccl.py — PyNcclCommunicator.suspend() / resume(), mirroring the
existing register_comm_window(); each guards on disabled and
nccl_version >= 22907.
parallel_state.py — suspend_nccl_comms() / resume_nccl_comms()
iterate the authoritative _groups registry and call .suspend()/.resume()
on each group's NCCL communicator. Collective (internal cross-rank barrier),
and synchronous on return (verified against NCCL mem_manager.cc), so no
extra device sync is needed.
gpu_worker.py — suspend_nccl_comms() after allocator.sleep() (so the
reported freed bytes include the NCCL release) and resume_nccl_comms() after
allocator.wake_up(), before any collective runs again.

NCCL version behavior (graceful, auto-activating)

The suspend/resume symbols are NCCL ≥ 2.29.7 only. vLLM does not pin NCCL
directly — it arrives via the torch wheel:

torch	bundled NCCL	this PR
2.11.0 (current `requirements/cuda.txt`)	2.28.9	no-op — symbols absent, skipped + warning
2.12.1	2.29.7	active — ~958 MiB/GPU freed

So on stock images today it is a safe no-op, and it activates automatically
once vLLM's CUDA torch pin moves to ≥ 2.12.1 — no code change required.

Why this is not a duplicate

feat: add suspend()/resume() for CRIU-safe snapshots #35934 (suspend/resume for CRIU) destroys and rebuilds NCCL; this keeps
the comm alive and only releases its buffers (ncclCommSuspend), avoiding the
re-init / re-bootstrap cost.
[Feat] Add fast EP scaling with NCCL communicator split #44994 splits NCCL comms for elastic EP; it does not free comm memory.
NCCL communicator memory offload is otherwise unclaimed upstream.

This is the NCCL half of #45623, split out as a focused, self-contained change:
no dependency on the CUDA-graph offload machinery, and it applies to every
sleep-mode config (the graph-offload piece has a MoE/EP coverage gap and lands
separately).

Tests

Validated end-to-end on 8×H200 (image vllm/vllm-openai:v0.24.0-ubuntu2404,
torch 2.11.0+cu130 with NCCL upgraded to 2.30.7), driving real
vllm serve --enable-sleep-mode → POST /sleep?level=1 → /wake_up, measuring
the free-memory delta around the actual ncclCommSuspend call on each rank.

Parallel-topology sweep (Qwen3-30B-A3B):

Config	GPUs	comms suspended	NCCL freed / rank	output preserved
TP2	2	2 (tp, ep)	958 MiB	✓
TP4	4	2 (tp, ep)	960 MiB	✓
TP8	8	2 (tp, ep)	960 MiB	✓
TP4 + EP	4	2 (tp, ep)	960 MiB	✓
TP8 + EP	8	2 (tp, ep)	960 MiB	✓
TP4 + DP2 + EP	8	3 (tp, ep, dp)	1438 MiB	✓ *
TP2 + PP2	4	3 (tp, ep, pp)	1438 MiB	✓

100B+ scale — Qwen3-235B-A22B-FP8, TP8+EP, 8×H200: 2 comms/rank,
960 MiB/rank freed

mergify · 2026-06-29T23:01:12Z

This pull request has merge conflicts that must be resolved before it can be
merged. Please rebase the PR, @aoshen02.

https://docs.github.com/en/pull-requests/collaborating-with-pull-requests/working-with-forks/syncing-a-fork

Sleep mode offloads weights and discards the KV cache, but each PyNccl communicator keeps its dynamic GPU buffers resident the whole time (~958 MiB/GPU on an 8-GPU TP+EP worker). Release them with NCCL's native ncclCommSuspend(NCCL_SUSPEND_MEM) / ncclCommResume, which frees a communicator's allocations while keeping its topology/connection state, so resume is cheap (no re-init, no bootstrap rendezvous). Implemented along vLLM's existing NCCL conventions, no new module: - pynccl_wrapper.py: bind ncclCommSuspend/ncclCommResume in NCCLLibrary (skipped gracefully with a one-time warning on NCCL < 2.29.7). - pynccl.py: PyNcclCommunicator.suspend()/resume(), mirroring register_comm_window(); gated on disabled and nccl_version >= 22907. - parallel_state.py: suspend_nccl_comms()/resume_nccl_comms() iterate the _groups registry and act on each group's communicator. Collective and synchronous on return, so no extra device sync is needed. - gpu_worker.py: suspend after allocator.sleep(), resume after wake_up(). The symbols exist only in NCCL >= 2.29.7, so this is a no-op on torch 2.11.0 (NCCL 2.28.9) and activates automatically on torch >= 2.12.1 (NCCL 2.29.7), with no code change. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: aoshen02 <aoshen@inferact.ai>

tlrmchlsmth · 2026-06-30T15:48:44Z

+    freed = abs(free_before - torch.cuda.mem_get_info()[0])
+    logger.info("NCCL %s: %d comms, %.1f MiB", label, len(comms), freed / 1024**2)


nit: I see why we're using abs here but potentially confusing and misleading especially if there's a bug

Suggested change

freed = abs(free_before - torch.cuda.mem_get_info()[0])

logger.info("NCCL %s: %d comms, %.1f MiB", label, len(comms), freed / 1024**2)

delta = torch.cuda.mem_get_info()[0] - free_before

direction = "freed" if delta > 0 else "allocated"

logger.info("NCCL %s: %d comms, %.1f MiB %s",

label, len(comms), abs(delta) / 1024**2, direction)

Nice catch!

matteso1 · 2026-06-30T23:58:54Z

@aoshen02 clean approach, and ~958 MiB/GPU on 8-GPU TP+EP is a real win. ncclCommSuspend(NCCL_SUSPEND_MEM) keeping topology so resume skips re-init is exactly right.

heads up this and #44074 both restructure the gpu_worker sleep/wake block: #44074 routes sleep()/wake_up() through a pluggable SleepModeBackend, and your suspend_nccl_comms() lands right in it. +1 to @tlrmchlsmth on moving it after torch.accelerator.synchronize(); that ordering (sync, then release NCCL mem) is exactly the guarantee the backend dispatch needs to keep, so worth pinning where the NCCL suspend sits relative to the backend's suspend/resume so both land cleanly. happy to coordinate. side note: this gives preserves_nccl a useful middle state — cumem suspends NCCL memory but keeps topology, distinct from cuda_checkpoint tearing it down.

Address review feedback: replace abs() in the sleep-mode NCCL log with a signed delta plus an explicit freed/allocated direction, so a wrong-sign result (e.g. an accidental allocation during suspend) is surfaced instead of masked. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: aoshen02 <aoshen@inferact.ai>

aoshen02 · 2026-07-01T00:50:56Z

@aoshen02 clean approach, and ~958 MiB/GPU on 8-GPU TP+EP is a real win. ncclCommSuspend(NCCL_SUSPEND_MEM) keeping topology so resume skips re-init is exactly right.

heads up this and #44074 both restructure the gpu_worker sleep/wake block: #44074 routes sleep()/wake_up() through a pluggable SleepModeBackend, and your suspend_nccl_comms() lands right in it. +1 to @tlrmchlsmth on moving it after torch.accelerator.synchronize(); that ordering (sync, then release NCCL mem) is exactly the guarantee the backend dispatch needs to keep, so worth pinning where the NCCL suspend sits relative to the backend's suspend/resume so both land cleanly. happy to coordinate. side note: this gives preserves_nccl a useful middle state — cumem suspends NCCL memory but keeps topology, distinct from cuda_checkpoint tearing it down.

Thanks, I missed some context, will take a look right now.

tlrmchlsmth

pre-commit is failing:

vllm/distributed/parallel_state.py:152: error: Found torch.cuda API call. Please refer RFC https://github.com/vllm-project/vllm/issues/30679, use torch.accelerator API instead.

galletas1712 · 2026-07-01T09:11:05Z

+        # Release idle NCCL communicator memory (NCCL >= 2.29.7). Collective
+        # across ranks; no-op at world_size 1 or on older NCCL. Done after the
+        # cumem unmap so the reported freed bytes include the NCCL release.
+        suspend_nccl_comms()


we actually have some ongoing work that will allow flashinfer's comms to be suspended as well (see flashinfer-ai/flashinfer#3727), (flashinfer-ai/flashinfer#3745). Could we have a more generic suspend/resume surface in device_communicators, that each communicator can override?

Nice catch, will do.

aoshen02 · 2026-07-05T09:52:40Z

pre-commit is failing:

vllm/distributed/parallel_state.py:152: error: Found torch.cuda API call. Please refer RFC https://github.com/vllm-project/vllm/issues/30679, use torch.accelerator API instead.

fix

Per review on vllm-project#46234: instead of parallel_state reaching into device_communicator.pynccl_comm, add suspend()/resume() to DeviceCommunicatorBase (default no-op) and override in CudaCommunicator to forward to its pynccl_comm. parallel_state now walks device communicators polymorphically (suspend_device_comms/resume_device_comms), so FlashInfer / all2all comms can add their own suspend later without touching parallel_state. Non-CUDA backends inherit the base no-op. The freed-bytes probe also moves to the torch.accelerator memory API. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> Signed-off-by: aoshen02 <aoshen@inferact.ai>

Signed-off-by: aoshen02 <aoshen@inferact.ai>

aoshen02 · 2026-07-05T11:59:12Z

This pr rely on #45321

aoshen02 requested a review from njhill as a code owner June 20, 2026 14:58

mergify Bot added the v1 label Jun 20, 2026

mergify Bot added the needs-rebase label Jun 29, 2026

aoshen02 force-pushed the nccl-comm-offload-sleep branch from 666ecc1 to 73998ef Compare June 30, 2026 07:25

mergify Bot removed the needs-rebase label Jun 30, 2026

aoshen02 closed this Jun 30, 2026

aoshen02 force-pushed the nccl-comm-offload-sleep branch from 73998ef to ba22cb6 Compare June 30, 2026 13:39

aoshen02 reopened this Jun 30, 2026

aoshen02 added 2 commits June 30, 2026 21:57

Merge branch 'main' into nccl-comm-offload-sleep

0dc82d9

Merge branch 'main' into nccl-comm-offload-sleep

cfd92a1

tlrmchlsmth reviewed Jun 30, 2026

View reviewed changes

Merge branch 'main' into nccl-comm-offload-sleep

919adc4

aoshen02 mentioned this pull request Jul 1, 2026

[Core] Pluggable sleep-mode backend abstraction (RFC #34303) #44074

Merged

tlrmchlsmth reviewed Jul 1, 2026

View reviewed changes

galletas1712 reviewed Jul 1, 2026

View reviewed changes

aoshen02 mentioned this pull request Jul 1, 2026

[Roadmap] 2026 Q2 vLLM × RL Roadmap #41733

Open

18 tasks

Merge branch 'main' into nccl-comm-offload-sleep

78f5b0e

mergify Bot added the nvidia label Jul 5, 2026

github-project-automation Bot added this to NVIDIA Jul 5, 2026

aoshen02 added 2 commits July 5, 2026 19:48

Update base_device_communicator.py

a3444d1

Signed-off-by: aoshen02 <aoshen@inferact.ai>

Update pynccl_wrapper.py

ad16cec

Signed-off-by: aoshen02 <aoshen@inferact.ai>

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

Uh oh!

[Core] Release NCCL communicator memory in sleep mode#46234

[Core] Release NCCL communicator memory in sleep mode#46234
aoshen02 wants to merge 9 commits into
vllm-project:mainfrom
aoshen02:nccl-comm-offload-sleep

aoshen02 commented Jun 20, 2026 •

edited

Loading

Uh oh!

mergify Bot commented Jun 29, 2026

Uh oh!

Uh oh!

tlrmchlsmth Jun 30, 2026

Uh oh!

aoshen02 Jul 1, 2026

Uh oh!

matteso1 commented Jun 30, 2026 •

edited

Loading

Uh oh!

aoshen02 commented Jul 1, 2026

Uh oh!

tlrmchlsmth left a comment

Uh oh!

galletas1712 Jul 1, 2026 •

edited

Loading

Uh oh!

aoshen02 Jul 1, 2026

Uh oh!

aoshen02 Jul 5, 2026

Uh oh!

aoshen02 commented Jul 5, 2026

Uh oh!

aoshen02 commented Jul 5, 2026 •

edited

Loading

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

		freed = abs(free_before - torch.cuda.mem_get_info()[0])
		logger.info("NCCL %s: %d comms, %.1f MiB", label, len(comms), freed / 1024**2)

-    freed = abs(free_before - torch.cuda.mem_get_info()[0])
-    logger.info("NCCL %s: %d comms, %.1f MiB", label, len(comms), freed / 1024**2)
+    delta = torch.cuda.mem_get_info()[0] - free_before
+    direction = "freed" if delta > 0 else "allocated"
+    logger.info("NCCL %s: %d comms, %.1f MiB %s",
+                label, len(comms), abs(delta) / 1024**2, direction)

Uh oh!

Uh oh!

Conversation

aoshen02 commented Jun 20, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

What

How

NCCL version behavior (graceful, auto-activating)

Why this is not a duplicate

Tests

Uh oh!

mergify Bot commented Jun 29, 2026

Uh oh!

Uh oh!

tlrmchlsmth Jun 30, 2026

Choose a reason for hiding this comment

Uh oh!

aoshen02 Jul 1, 2026

Choose a reason for hiding this comment

Uh oh!

matteso1 commented Jun 30, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

aoshen02 commented Jul 1, 2026

Uh oh!

tlrmchlsmth left a comment

Choose a reason for hiding this comment

Uh oh!

galletas1712 Jul 1, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

aoshen02 Jul 1, 2026

Choose a reason for hiding this comment

Uh oh!

aoshen02 Jul 5, 2026

Choose a reason for hiding this comment

Uh oh!

aoshen02 commented Jul 5, 2026

Uh oh!

aoshen02 commented Jul 5, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

aoshen02 commented Jun 20, 2026 •

edited

Loading

matteso1 commented Jun 30, 2026 •

edited

Loading

galletas1712 Jul 1, 2026 •

edited

Loading

aoshen02 commented Jul 5, 2026 •

edited

Loading