Skip to content

feat: vmm slot for local cuda graph offload - #3544

Open
Xdydy wants to merge 2 commits into
NVIDIA:release_v2.13from
Xdydy:release_v2.13
Open

Xdydy wants to merge 2 commits into
NVIDIA:release_v2.13from
Xdydy:release_v2.13

Conversation

@Xdydy

@Xdydy Xdydy commented Sep 18, 2026

Copy link
Copy Markdown

Description

Add a fixed-VA CUDA VMM activation slot so a Graph-captured tensor address can keep its virtual address while physical backing is released and remapped outside Graph replay. This is the TE primitive behind local whole-layer CUDA Graph + fine-grained CPU activation offload.

Fixes # (issue)

Local CUDA Graph and CPU offload currently conflict on pointer semantics:

Graph replay  requires: captured tensor addresses never change
CPU offload   requires: physical device memory is released after forward

A naive empty() reload allocates a new pointer. Capture recorded src.data_ptr(), so replay either reads stale storage or silently aliases an allocator reuse of that VA. Related: NVIDIA/Megatron-LM#3697.

This PR does not implement Megatron scheduling. It exposes the VMM slot, resident workers, and deferred H2D so a caller (Megatron fine_grained_activation_offload) can:

  1. Reserve a VA and map physical backing before capture.
  2. Capture / replay Graphs against slot.tensor.data_ptr() (stable for the run).
  3. After forward D2H, unmap/release physical pages without freeing the VA.
  4. Before backward, map new backing at the same VA and H2D.
  5. Optionally prefetch remap while the current Graph runs, then launch H2D at replay time.

CUDA driver names stay CUDA-shaped in source. On MUSA they map at compile time through musify.h (cuMemMapmuMemMap, cudaMemcpyAsyncmusaMemcpyAsync, …). Python keeps CUDA names (CUDAActivationVMMAllocation, cuda_stream); a MUSA runtime still works when torch.cuda is aliased, and device.type in ("cuda", "musa") covers TE tests that import this module without musa_patch.

Type of change

  • Documentation change (change only to the documentation, either a fix or a new content)
  • Bug fix (non-breaking change which fixes an issue)
  • New feature (non-breaking change which adds functionality)
  • Breaking change (fix or feature that would cause existing functionality to not work as expected)
  • Infra/Build change
  • Code refactoring

Changes

  • Add transformer_engine/pytorch/csrc/extensions/vmm_activation.cpp and bind it from pybind.cpp / extensions.h.
  • Add Python wrapper transformer_engine/pytorch/vmm_activation.py.
  • Add tests/pytorch/test_vmm_activation.py (async reload / host-callback lifetime).
  • Add musify.h VMM aliases used by this file (cuMemAddressReserve, cuMemCreate, cuMemMap, cuMemSetAccess, cuMemUnmap, cuMemRelease, cuMemGetInfo, …).

Not in this PR’s intended review surface: GEMM / DPA / cpu_offload.py Megatron handler changes that happen to sit on the same working branch.


Why a TE primitive (not a Megatron allocator wrapper)

Graph capture records a device pointer. CPU offload must drop physical pages. Those two facts only compose if the Graph sees a VA that outlives the pages mapped under it:

Naive:  act = torch.empty(...)     # address comes from the caching allocator
VMM:    slot = VMMActivationSlot(bytes, device)
        act  = slot.tensor(shape, stride, dtype)   # act.data_ptr() == reserved VA
cuMemAddressReserve  →  fixed VA (kept for the run)
cuMemCreate          →  physical allocation handle
cuMemMap             →  map handle at that VA
cuMemSetAccess       →  RW access for the device

offload path (after D2H):
  cuMemUnmap + cuMemRelease     →  pages gone, VA still reserved

reload path (before backward):
  cuMemCreate + cuMemMap + cuMemSetAccess at the same VA
  cudaMemcpyAsync(host → VA)

The captured Graph never sees a new pointer. Physical handles may change; the VA must not.

CUDAActivationVMMAllocation constructs the slot, wraps a non-owning from_blob view, and asserts tensor.data_ptr() == info["address"] after construct / remap / adopt. Using that view while the slot is unmapped is a hard page fault — that is intentional; it turns the silent allocator-reuse bug into a crash.


Slot lifecycle

construct
  cuMemAddressReserve + first cuMemCreate/Map/SetAccess
  tensor view at the reserved VA

forward offload (caller D2H, then)
  release_hooks_after(slots, d2h_stream)
    host func on d2h_stream after the D2H burst
    ReleaseWorker: cuMemUnmap + cuMemRelease
    VA reservation kept

backward restore
  remap_only_slot_after(slot, pinned_host, h2d_stream)   # remap now, no H2D
  … later, at Graph replay …
  launch_remap_slot_h2d(context, slot, h2d_stream)       # H2D now
  wait_remap_slot_on_stream(context, slot, compute_stream)
  adopt_async_remap()                                    # slot owns the new handle

close
  drain in-flight remap/release, unmap, release, cuMemAddressFree

Synchronous unmap_and_release() / create_and_remap() remain for tests and fallback. Training uses the async path.

Address stability check after every remap / adopt:

info["address"] == original_va and tensor.data_ptr() == original_va

Python API

transformer_engine.pytorch.vmm_activation:

Call Role
CUDAActivationVMMAllocation(shape, stride, dtype, device) reserve VA, map first backing, non-owning tensor view
release_hooks_after(allocs, stream) one host func after D2H; worker unmaps the batch
remap_hooks_after(allocs, stream) remap only, no copy
remap_and_copy_after / remap_and_copy_slot_after worker remaps and submits H2D
remap_only_slot_after worker remaps; H2D deferred
launch_remap_slot_h2d(ctx, i, stream) wait remapped, submit H2D on stream
wait_remap_slot_on_stream(ctx, i, stream) wait event_recorded, stream waits on slot event
wait_until_remap_slot_submitted host wait until that slot’s H2D event is recorded
wait_remap_copy_on_stream / enqueue_remap_copy_wait whole-context copy event (non-deferred path)
vmm_driver_memory_info() cuMemGetInfo free/total (caching-allocator stats do not cover cuMemCreate pages)
vmm_enable_trace / vmm_initialize_workers / vmm_set_serial_driver_workers diagnostics and worker policy

Stream handles: stream.cuda_stream, with musa_stream fallback.

Safety checks on remap submit: unique slots, reserved VA, no in-flight remap, mapped slot must already have a pending release, host tensor pinned / offset 0 / large enough. Adopted mappings are never unmapped by context cleanup.

Debug: MEGATRON_VMM_REMAP_DEBUG=1 prints per-call remap/release stages (cuMemUnmap, cuMemSetAccess, remap wait, H2D launch wait).


Files

transformer_engine/pytorch/csrc/extensions/vmm_activation.cpp
transformer_engine/pytorch/csrc/extensions.h
transformer_engine/pytorch/csrc/extensions/pybind.cpp
transformer_engine/pytorch/vmm_activation.py
transformer_engine/common/include/transformer_engine/musify.h   # VMM aliases
tests/pytorch/test_vmm_activation.py

Checklist:

  • I have read and followed the contributing guidelines
  • The functionality is complete
  • I have commented my code, particularly in hard-to-understand areas
  • I have made corresponding changes to the documentation
  • My changes generate no new warnings
  • I have added tests that prove my fix is effective or that my feature works
  • New and existing unit tests pass locally with my changes

Suggested coverage:

  • Address stability: data_ptr() unchanged across unmap/remap and deferred H2D.
  • release_hooks_after + remap_and_copy_after repeated reload (test_repeated_async_reload_releases_context_off_host_callback): no host-callback deadlock, values match, context error==0.
  • Deferred path: remap_only_slot_after does not submit H2D; launch_remap_slot_h2d does; compute stream waits on the slot event.
  • Negative: using the tensor view after unmap and before remap must fault, not silently alias a caching-allocator reuse of the captured VA.
  • cuMemGetInfo (not PyTorch memory_allocated) when claiming physical pages returned.

@github-actions github-actions Bot added the community-contribution PRs from external contributor outside the core maintainers, representing community-driven work. label Sep 18, 2026
@greptile-apps

greptile-apps Bot commented Sep 18, 2026

Copy link
Copy Markdown
Contributor

RetriggerConfidence Score: 2/5

The PR is not safe to merge because serial-driver mode can deadlock, asynchronous close can release memory during an active H2D, and the stated MUSA API path is rejected.

Findings

  1. P1 Serial remap can deadlock
  2. P1 Close races active H2D
  3. P1 MUSA API path rejected
  4. P2 Duplicate slots release twice
  5. P2 Deferred H2D launches repeatedly

Summary

This PR adds fixed-address CUDA VMM activation slots, asynchronous release/remap workers, deferred H2D synchronization APIs, Python bindings, and an async reload test.

  • Reserves stable virtual addresses while replacing physical backing.
  • Moves VMM release and remap operations onto resident workers.
  • Adds whole-context and per-slot stream synchronization.
  • Adds tracing, worker-policy controls, and driver memory reporting.

Diagram

sequenceDiagram
  participant Caller
  participant CopyStream
  participant ReleaseWorker
  participant RemapWorker
  participant Slot

  Caller->>CopyStream: D2H activation copy
  Caller->>CopyStream: enqueue release callback
  CopyStream->>ReleaseWorker: release request
  ReleaseWorker->>Slot: unmap and release physical backing
  Caller->>RemapWorker: submit fixed-VA remap
  RemapWorker->>Slot: create, map, and set access
  alt Immediate H2D
    RemapWorker->>CopyStream: enqueue H2D and completion event
  else Deferred H2D
    Caller->>CopyStream: launch slot H2D
  end
  Caller->>CopyStream: enqueue consumer-stream wait
Loading

Reviews (1) · Last reviewed commit: "[pre-commit.ci] auto fixes from pre-comm..."

Comment on lines +875 to +878
bool ReleaseWorker::process_remap_batch(const std::shared_ptr<RemapHookContext> &context,
bool requeued) {
RemapWorker::process(context, /*serial_requeue=*/requeued);
return context->done.load(std::memory_order_acquire) != 0;

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P1 Serial remap can deadlock

In serial-driver mode, the first remap attempt uses requeued=false. If its release callback has not yet queued the matching release, the remap waits for that release instead of returning. Because the same worker must process the release, it can never make progress, permanently blocking this remap and subsequent waits. Every serial-worker remap attempt must remain requeueable while its release dependency is pending.

Comment on lines +1136 to +1149
void close() {
drain_remap();
drain_pending();
if (!reserved_) return;
if (mapped_) {
check_cuda(cuMemUnmap(address_, bytes_), "cuMemUnmap(close)");
mapped_ = false;
check_cuda(cuMemRelease(handle_), "cuMemRelease(close)");
handle_ = 0;
}
check_cuda(cuMemAddressFree(address_, bytes_), "cuMemAddressFree");
reserved_ = false;
address_ = 0;
}

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P1 Close races active H2D

close() waits only until an asynchronous H2D has been submitted, not until the GPU copy has completed, and then immediately unmaps and releases its destination. If a caller closes an allocation after remap_and_copy_after returns but before the copy stream completes, the GPU can continue writing to unmapped memory, causing a device fault or a corrupted transfer. The close and destructor paths must retain the mapping until the relevant completion event finishes.

Comment on lines +43 to +44
if self.device.type != "cuda":
raise ValueError(f"CUDA VMM allocation requires a CUDA device, got {self.device}")

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P1 MUSA API path rejected

The advertised MUSA path is unreachable through this Python API: construction rejects every device type except cuda, and _raw_cuda_stream checks only cuda_stream despite the stated musa_stream fallback. On a MUSA runtime without a complete torch.cuda alias, callers cannot construct or schedule this feature even if the native source is translated successfully.

Comment on lines +1297 to +1312
std::shared_ptr<ReleaseHookContext> release_hooks_after(
std::vector<std::shared_ptr<VMMActivationSlot>> slots, uintptr_t raw_stream) {
RECORD_USER_SCOPE("vmm::release_hooks_after");
// Ensure the worker thread exists before any host func can fire.
ReleaseWorker::instance();
auto context = std::make_shared<ReleaseHookContext>();
context->tls_state = std::make_shared<at::ThreadLocalState>();
for (auto &slot : slots) {
if (slot->pending_remap_) {
throw std::runtime_error("VMM batch release hook got a slot with an in-flight remap");
}
if (slot->pending_) {
throw std::runtime_error("VMM batch release hook got a slot with an in-flight release");
}
context->requests.push_back(slot->make_release_request());
if (context->requests.back().handle == 0) {

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P2 Duplicate slots release twice

The batch release helper accepts the same slot more than once, creating duplicate requests for one virtual address and allocation handle. The worker then attempts cuMemUnmap and cuMemRelease twice, records an asynchronous error on the shared context, and leaves the slot bookkeeping stale after the first release succeeds. Reject duplicate slots before installing the release context, as the batch remap helpers already do.

Comment on lines +1614 to +1620
void remap_slot_launch_h2d(const std::shared_ptr<RemapHookContext> &context, size_t slot_index,
uintptr_t raw_stream) {
RECORD_USER_SCOPE("vmm::remap_slot_launch_h2d");
if (!context->defer_copy) {
throw std::runtime_error("VMM slot H2D launch requires a deferred-copy context");
}
if (slot_index >= context->requests.size()) {

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P2 Deferred H2D launches repeatedly

A deferred-copy context allows launch_slot_h2d to be called repeatedly for the same request because copy_submitted is never checked. Each call queues another copy and event record, but lifetime retention is installed only when h2d_launched first equals the request count. After adoption and reference release, the first callback can therefore destroy the event and pinned host storage before a later duplicate copy completes. Reject a launch after that request has already been submitted.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

community-contribution PRs from external contributor outside the core maintainers, representing community-driven work.

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants