Memory Registration, All the Way Down, Part 3. Part 1 builds the RNIC translation path. Part 2 follows the production failure to a host page in MIGRATE_CMA.

By the end of Part 2, one fact still sounds contradictory:

CUDA had already pinned the host page.
RDMA then tried to pin the same page.
The second pin failed because Linux needed to move it.

The two pinning calls asked for different things. CUDA held the current page in place. RDMA requested a long-lived DMA mapping, which required Linux to check whether the page could safely stay in its current location.

The page was in CMA, where ordinary allocations must remain movable. Linux needed to move it before accepting the long-term pin, but the earlier CUDA pin prevented that. We’ll follow those checks from the page’s physical placement through to ENOMEM.

Scope note. The source walk uses Linux 6.x and public NVIDIA open-kernel-module releases. Helper names have changed across kernel versions, but the invariant is stable: long-term DMA pins cannot strand ordinary pages in memory that must remain migratable. The production kernel had internal patches, so reconstructed traces are illustrative rather than original logs.

1. A virtual buffer is not a physical neighborhood

A process allocates a 16 KiB buffer:

virtual address range


0x7f00_0000  +-------------------+
             | page 0            |
0x7f00_1000  +-------------------+
             | page 1            |
0x7f00_2000  +-------------------+
             | page 2            |
0x7f00_3000  +-------------------+
             | page 3            |
             +-------------------+

The virtual pages are contiguous. Their physical backing does not have to be:

virtual page 0 -> physical frame 91
virtual page 1 -> physical frame 8002
virtual page 2 -> physical frame 92
virtual page 3 -> physical frame 410

On a typical x86-64 Linux system, the base page size is 4 KiB. Linux identifies a physical base page by a page frame number, or PFN, and usually represents it with a struct page. A folio is a newer Linux abstraction for one or more physically contiguous base pages managed as one unit.

The CPU’s page table connects virtual pages to PFNs. A device using a registered MR has a separate translation, built at registration time.

CPU translation                      device translation


process VA                           MR IOVA + key
    |                                    |
CPU page table                           | MKey / MTT
    v                                    v
physical page P  <------------------- DMA address for P

If Linux replaces physical page P with page Q, it can update the CPU page table. It cannot silently update every third-party device translation unless the device subsystem participates.

That is why a device pin is more consequential than keeping a page in RAM. It stabilizes an address relationship outside the CPU MMU.

2. Zones and pageblocks are different layers of policy

Linux groups physical memory at several scales.

Zones

A zone is a large allocator domain based on addressing or migration constraints: ZONE_DMA, ZONE_DMA32, ZONE_NORMAL, ZONE_MOVABLE, and others depending on architecture.

A zone answers questions such as:

Can this device address the memory?
Can unmovable kernel allocations live here?
Is this region intended primarily for movable pages?

Pageblocks

Inside a zone, Linux divides memory into pageblocks. On common 4 KiB-page x86 systems, a pageblock is usually 2 MiB, though the exact size is architecture/configuration dependent.

fig. 1 · a zone is divided into pageblocks, and pageblocks carry policy
ZONE_NORMAL pageblock UNMOVABLE pageblock MOVABLE pageblock MOVABLE pageblock MIGRATE_CMA pageblock RECLAIMABLE usually 2 MiB each on 4 KiB-page x86 · migratetype = get_pageblock_migratetype(page) "the page was migrate CMA" means: its PFN fell inside a pageblock with this policy.

movable allocations expected CMA reserve: occupants must stay evictable

Each pageblock has a migratetype describing the kind of allocation it should serve:

MIGRATE_UNMOVABLE
MIGRATE_RECLAIMABLE
MIGRATE_MOVABLE
MIGRATE_HIGHATOMIC
MIGRATE_CMA
MIGRATE_ISOLATE

The migratetype describes the allocator’s policy for a pageblock.

A page from a MIGRATE_MOVABLE block is expected to be movable. A page from a MIGRATE_CMA block occupies physical space reserved for the Contiguous Memory Allocator and must remain removable when CMA needs the range back.

The phrase from the incident—“the page was migrate CMA”—means more precisely:

The page’s PFN fell inside a pageblock whose migratetype was MIGRATE_CMA.

That distinction matters when instrumenting the kernel. The relevant question is often get_pageblock_migratetype(page), not a single PageCma flag on the object.

3. Why physically contiguous memory is hard to allocate late

Linux’s HugeTLB subsystem manages explicitly reserved huge pages. “TLB” refers to the CPU’s Translation Lookaside Buffer, a cache of recent virtual-to-physical translations. Mapping a large region with a huge page reduces page-table entries and TLB pressure compared with mapping the same bytes as thousands of 4 KiB pages. HugeTLB is distinct from Transparent Huge Pages: applications reserve or request its pages explicitly.

A 1 GiB HugeTLB page requires 1 GiB of contiguous physical address space. With 4 KiB base pages, that is 262,144 adjacent frames.

At boot, finding such a range is easy. After the machine has run for hours, physical memory looks more like this:

physical frames


[free][anon][free][page cache][kernel][free][anon][pinned][free]...

The machine may have many gigabytes free in total and still have no free 1 GiB run.

This is external fragmentation: sufficient capacity, insufficient contiguity.

Linux can compact memory by moving movable pages together and freeing larger extents. But some pages cannot move:

kernel allocations with embedded physical addresses
long-term DMA pins
some device mappings
hardware-reserved pages

One immovable 4 KiB page can spoil an otherwise free 1 GiB candidate.

That is why gigantic pages are commonly reserved at boot or backed by a dedicated CMA area.

4. What hugetlb_cma=6G creates

The Contiguous Memory Allocator (CMA) reserves physical ranges that can be evacuated and handed to users needing large contiguous allocations. The boot parameter from the incident was approximately:

hugetlb_cma=6G

Linux documents this as a CMA area used to allocate gigantic HugeTLB pages. On x86-64, the relevant gigantic page size is normally 1 GiB. The six-gigabyte size is a multiple of that unit.

At boot, Linux carves out physical ranges and marks their pageblocks for CMA use. A global size is apportioned across online NUMA nodes unless the command line specifies node-specific sizes. On a two-node host, hugetlb_cma=6G normally means up to roughly three GiB per node, subject to alignment and successful reservation—not one universal six-gigabyte interval.

physical memory, two-node example


NUMA node 0                         NUMA node 1
+-----------------------------+     +-----------------------------+
| ordinary memory             |     | ordinary memory             |
+-----------------------------+     +-----------------------------+
| ~3 GiB MIGRATE_CMA extents  |     | ~3 GiB MIGRATE_CMA extents  |
+-----------------------------+     +-----------------------------+

The reservation can provide up to six 1 GiB huge pages across the nodes, but it does not necessarily create those pages at boot. The per-node split also affects whether a later allocation lands in CMA.

The kernel can later ask CMA for a 1 GiB contiguous extent. Until then, leaving six GiB idle would waste memory, so CMA allows ordinary movable pages to occupy it temporarily.

CMA reserve while no huge page is requested


+--------------------------------------------------+
| anon | page cache | free | anon | free | ...     |
| all temporary occupants must remain movable      |
+--------------------------------------------------+

When a huge-page allocation arrives, CMA isolates the target range, migrates temporary occupants elsewhere, and returns the now-contiguous physical memory.

fig. 2 · evacuating the reserve for one gigantic page
before: lent-out reserve A B free C free D ordinary memory after: the promise is kept one contiguous physical extent · 1 GiB migrate A, B, C, D to ordinary pages this same migration, attempted later on one pinned page, is the whole incident.

CMA reserve temporary movable occupants

CMA can reclaim the range only while its temporary occupants remain movable.

5. Why the kernel consumed CMA first

Ordinary movable allocations can use either regular free memory or CMA free memory. Unmovable allocations usually cannot use CMA, because they would defeat its purpose. Stock Linux already has a conditional CMA fallback: in v6.6, eligible movable allocations may draw from CMA when CMA accounts for more than half of the zone’s free pages. The fleet’s CMA-first policy made that choice earlier and more frequent; it was an amplifier, not a prerequisite for the mechanism.

Suppose a machine has:

ordinary free memory:  1 GiB
CMA free memory:       5 GiB

If movable anonymous pages consume ordinary memory first, an unmovable allocation may later fail with five GiB still free in CMA. The unmovable page cannot go there, and Linux has no general mechanism to move existing anonymous pages into CMA to free ordinary space.

The CMA-first policy solves that local problem:

movable allocation arrives
    |
    +-- take CMA page first
    |
    +-- preserve ordinary memory for less flexible users

The public patch matching the incident says exactly that: movable pages can be migrated out of CMA later, so consume CMA first and preserve the memory that unmovable allocations need. The public proposal is evidence for the policy and rationale, not proof that its exact diff was the internal fleet change.

The assumption underneath the optimization is:

anything allocated from CMA remains movable

A normal anonymous page satisfies that assumption—until a device pins it.

6. “Resident,” “locked,” and “pinned” are not synonyms

Linux has several ways to make memory less movable or less reclaimable.

Faulted-in / resident

A resident page currently has physical backing. It may still be reclaimed, swapped, migrated, or replaced later.

mlock()ed

mlock() asks Linux not to page the range out. It is a userspace residency policy. It is not the same device-DMA contract as GUP pinning, and mlocked pages can participate in some migration paths.

GUP reference

Historically, device drivers used get_user_pages()—GUP—to obtain references to userspace pages. A reference keeps the page object alive, but Linux struggled to distinguish short software access from device DMA pins.

FOLL_PIN

Modern drivers use pin_user_pages*(). Those wrappers set FOLL_PIN and account the page as DMA-pinned. For base pages, Linux uses a biased reference count; for large folios it may maintain a dedicated pin count.

This lets the VM recognize that the page is participating in a device mapping and requires special treatment.

FOLL_LONGTERM

FOLL_LONGTERM adds a lifetime and placement declaration:

This pin may persist long enough that normal VM operations cannot treat it as a brief interruption.

FOLL_LONGTERM describes the intended lifetime of a mapping, such as a classic RDMA MR. The API does not define a duration threshold in seconds.

FOLL_LONGTERM implies that Linux must reject or relocate pages whose location cannot safely be stranded for that lifetime.

The diagram compares these forms of residency and pinning:

fig. 3 · "pinned" is a ladder of contracts, not a Boolean
resident has physical backing right now mlock()ed not paged out · may still migrate GUP reference page object stays alive FOLL_PIN accounted device-DMA pin · page stays put FOLL_PIN + FOLL_LONGTERM + placement validated stronger contract only the top rung asks whether the page may stay put. CUDA's pin stopped one rung short.

DMA pin DMA pin with placement validation

A long-term pin also requires Linux to validate the page’s placement.

7. Why classic pinned RDMA MRs declare the long-term contract

A classic RDMA memory region can live for minutes, hours, or the lifetime of a process. The RNIC may issue DMA whenever a work request or remote packet references its key.

Linux RDMA core therefore begins ordinary userspace MR registration with:

FOLL_LONGTERM

and adds FOLL_WRITE if the device can write the pages.

The public ib_umem_get() path is conceptually:

check locked-memory allowance
allocate ib_umem
pin_user_pages_fast(FOLL_LONGTERM | maybe FOLL_WRITE)
build SG table
DMA-map SG table for RNIC

The long-term flag tells GUP that simply finding a present page is insufficient. The page’s kind and location must support a long-lived DMA pin.

CMA pages do not, in place.

8. Why a long-term pin cannot stay inside CMA

CMA needs to reclaim contiguous physical ranges by moving their occupants. A long-term DMA pin requires a stable physical destination for the device. Those requirements conflict when they apply to the same page:

CMA promise:
    "I can move this page when I need the physical range."


long-term DMA promise:
    "This page must stay at this physical DMA destination."

They cannot both hold indefinitely.

Linux resolves the conflict by moving the page before accepting the long-term pin.

fig. 4 · the repair path: move first, then pin
MIGRATE_CMA pageblock page P movable occupant ordinary memory page Q long-term pinned 1. allocate Q · 2. copy P -> Q 3. CPU mappings now point at Q 4. free P back to the reserve 5. pin Q long-term CMA keeps its movable contract the RNIC gets a stable DMA target both promises hold, because they end up on different pages.

movable page inside CMA relocated, long-term-pinned page

This is why a direct RDMA registration of an otherwise unpinned CMA page can succeed. The kernel is allowed to relocate it as part of registration.

The incident required an earlier pin.

9. The first pin: CUDA mapped host memory

The production evidence established a CUDA/NVIDIA pin on CPU-backed memory. The exact userspace operation is no longer recoverable: stock same-process NCCL can allocate mapped host memory directly, while file-backed NCCL paths can register an existing mapping.

The public NVIDIA R525.105.17 and R535.104.05 open-module source gives us an inspectable analogue for the latter case. Its host-page locking function builds GUP flags from writability and then calls NV_PIN_USER_PAGES(). Treat this as the source-backed lifetime model, not a recovered production stack.

The observed page placement also tells us which broad NVIDIA sysmem path was involved. Driver-owned NVIDIA pages use GFP_KERNEL in the public R535 allocator and are not movable, so they cannot originate from a MIGRATE_CMA pageblock. By contrast, the public RM OS-descriptor path imports a userspace virtual range through RmCreateOsDescriptor() and os_lock_user_pages(). Anonymous userspace faults use GFP_HIGHUSER_MOVABLE. A traced MIGRATE_CMA backing page is therefore evidence for movable userspace memory imported and pinned by NVIDIA, not for the driver’s own unmovable page allocator.

Simplified:

flags = writable ? FOLL_WRITE : 0;
pin_user_pages(address, count, flags, pages);

The public path does not add FOLL_LONGTERM in those releases.

Therefore, if the host page already belongs to a CMA pageblock, the first pin can follow this path:

NCCL host allocation
    -> Linux supplies a page from MIGRATE_CMA
    -> CUDA/NVIDIA pins it with FOLL_PIN
    -> no long-term placement check
    -> page remains physically inside CMA

The page is now immobile in practice while CUDA holds the pin, but it entered that state without passing the placement validation designed for long-lived DMA.

The first pin stabilized the current page. It did not first establish that the page lived in a location safe for a long-term pin.

10. The second pin: RDMA asks for placement correctness

Later, NCCL registers the same host range with the ConnectX RNIC:

ibv_reg_mr_iova2()
    -> mlx5_ib_reg_user_mr()
    -> ib_umem_get()
    -> pin_user_pages_fast(FOLL_LONGTERM)

Linux now examines the backing page and finds:

pageblock migratetype = MIGRATE_CMA

The long-term GUP path cannot leave it there. It collects the unpinnable/movable page, drops the temporary pin it took while inspecting the range, and attempts synchronous migration to an acceptable page.

A simplified control flow is:

try long-term pin
    |
    v
find page not valid for long-term placement
    |
    v
unpin temporary GUP references
    |
    v
migrate page to ordinary memory
    |
    +-- success -> retry pin on new page
    |
    +-- failure -> return error

On Linux 6.x, the source contains helpers with names such as check_and_migrate_movable_pages() or their folio-oriented successors. The migration reason is MR_LONGTERM_PIN.

The migration can repair the placement if CUDA has not already pinned the page. With the earlier pin still held, that repair fails.

11. Why the earlier pin prevents migration

To migrate page P to page Q, Linux must know that it controls every relevant mapping and reference to P.

The migration operation roughly does this:

1. isolate P from normal VM lists
2. allocate Q
3. freeze / validate references to P
4. copy contents P -> Q
5. replace CPU mappings
6. transfer metadata
7. free P

A device pin is an external promise that some DMA translation still points to P.

Linux cannot solve that by updating the CPU page tables:

CPU mapping after migration:     VA -> Q
existing device mapping:         DMA -> P

If it freed and reused P, the device could corrupt an unrelated allocation. If it kept P, CMA would not regain the contiguous range. Neither is acceptable.

FOLL_PIN accounting makes these hidden device references visible enough for migration to fail rather than silently corrupt data. The extra page references or pin count prevent the migration code from freezing the page in the expected state.

page P expected references: VM mappings + migration reference
page P actual references:   expected + CUDA DMA pin


actual != expected
    -> cannot safely move P

The earlier CUDA pin prevented the move that RDMA’s placement check required.

12. Why the error is ENOMEM

The deepest failure is “could not migrate this page into an acceptable location.” Linux reports that from the long-term GUP migration helper as -ENOMEM when migration does not complete.

That value propagates:

migrate_pages() does not migrate every page
    -> long-term GUP returns -ENOMEM
    -> ib_umem_get() returns ERR_PTR(-ENOMEM)
    -> mlx5_ib_reg_user_mr() returns error
    -> userspace provider returns NULL
    -> errno = ENOMEM
    -> NCCL prints "Cannot allocate memory"

Migration needs a suitable replacement page and a source page that can be moved. Free RAM alone does not make the source movable.

This is one of several meanings hidden behind MR ENOMEM:

memory capacity exhausted
memlock quota exceeded
kernel metadata allocation failed
MKey resource unavailable
long-term page placement could not be established  <-- this incident

A useful error message would have preserved the last one.

13. The ordering bug in one page

The entire mechanism fits in two timelines.

Safe ordering

The first device declares the long-term lifetime:

fig. 5 · safe ordering: the strongest contract goes first
page P from CMA CUDA asks long-term- aware pin Linux migrates P -> Q still movable CUDA pins Q RDMA registers Q success the declaration arrives first the placement check runs while the page can still be moved.

Incident ordering

The first device pins without the placement declaration:

fig. 6 · incident ordering: the weak pin gets there first
page P from CMA CUDA pins P without FOLL_LONGTERM RDMA asks for long-term pin Linux must migrate P CUDA pin prevents migration -ENOMEM no placement check · P stays in CMA same actors as fig. 5, one missing flag at step two. everything after it is forced.

The failure depends on the order of the pinning calls. They do not have to run simultaneously: once the first pin holds this physical page, the later placement check can fail deterministically.

The apparent randomness comes from whether the allocation landed in CMA and when the two registrations occurred.

14. The NVIDIA source change that mirrors the mechanism

The public source provides a useful before-and-after.

R525 and R535

The host-page lock path uses pin_user_pages with writability flags, but no FOLL_LONGTERM.

R525/R535 public path:
    FOLL_PIN + maybe FOLL_WRITE

R555

The R555.42.02 public source adds FOLL_LONGTERM on x86 before calling the same pinning API.

R555 public path:
    FOLL_PIN + FOLL_LONGTERM + maybe FOLL_WRITE

The newer code also contains a workaround for kernels where a large long-term GUP request can hit an allocation limit while building VMA metadata, retrying in smaller chunks on ENOMEM.

The public commit history does not link this change to the Meta incident. It does show the placement check moving to the first pin:

before:
    pin first, discover incompatible placement later


after:
    declare long-term intent at the first NVIDIA host-page pin

If the page is in CMA, the first pin can now trigger migration while the page is still movable.

15. Why disabling CMA was the right production fix

Several fixes are theoretically possible.

Make the first pin long-term aware

This addresses the contract at its source. A modern driver path that uses FOLL_LONGTERM before establishing the CUDA host mapping should avoid stranding the page in CMA.

But upgrading a fleet driver is not always an immediate or isolated change, and it still leaves a memory policy on nodes that do not need it.

Force the allocation outside CMA

An allocator could request unmovable memory or use a dedicated pool. That is difficult from userspace and may have broader fragmentation costs. NCCL should not need to know the fleet’s HugeTLB reservation policy to allocate a small host buffer.

Share one registration/lifetime object

DMA-BUF and other exporter/importer models reduce duplicated, independent lifetime decisions. But stock NCCL’s host LL path was ordinary host memory, not an exported CUDA device allocation.

Remove the irrelevant reserve

The GPU fleet did not need the six-gigabyte HugeTLB CMA area. Disabling it eliminated the problematic placement and returned the memory to ordinary allocation policy.

hugetlb_cma=0

Without that reserve, registration no longer needed to move pages out of CMA:

no CMA page
    -> no need to migrate before long-term pin
    -> first and second pins can agree on the same ordinary page

The production result—no recurrence for weeks and months—confirmed that choice.

16. A hardware-independent reproducer

The core mechanism does not require an H100 or a ConnectX-7. It requires:

1. a userspace page backed by CMA
2. a first GUP pin without FOLL_LONGTERM
3. a second long-term GUP pin

A rigorous lab can use a VM, a small kernel module, and Soft-RoCE.

16.1 Create CMA memory

Boot with a CMA area. To mirror production closely, use hugetlb_cma on a sufficiently large VM; for a smaller lab, a generic cma= reserve can exercise the same pageblock property.

There are two ways to make the test page land there. On an unmodified stock kernel, consume enough ordinary free memory that the allocator’s conditional CMA fallback engages. Alternatively, apply a CMA-first policy to raise the hit rate and mirror production. The first route is the stronger proof: the two-pin mechanism does not depend on the fleet-specific policy; that policy only made the placement common.

16.2 Find a userspace page in MIGRATE_CMA

Fault anonymous pages until a probe reports:

get_pageblock_migratetype(page) == MIGRATE_CMA

A test module can expose a debug ioctl that resolves a userspace address and prints its PFN and pageblock migratetype. This avoids relying on restricted /proc/pagemap access.

16.3 Take the first pin

The test module calls:

pin_user_pages(..., no FOLL_LONGTERM)

and holds the returned page reference.

16.4 Attempt a real RDMA MR

Create an rdma_rxe software RDMA device and call ibv_reg_mr() on the same range. rdma_rxe is Linux’s software implementation of the RDMA over Converged Ethernet (RoCE) verbs interface, so it exercises RDMA core’s memory-registration path without requiring a physical RNIC. The provider still enters RDMA core’s ordinary ib_umem_get() path and requests FOLL_LONGTERM, even though the final data movement is software-emulated.

Expected result:

ibv_reg_mr() -> NULL
errno        -> ENOMEM

16.5 Run the causal matrix

Case A
CMA page + no first pin
    -> RDMA migrates page
    -> registration succeeds


Case B
CMA page + first non-long-term pin
    -> migration blocked
    -> registration fails


Case C
CMA page + first FOLL_LONGTERM pin
    -> first pin migrates page
    -> later registration succeeds


Case D
ordinary page + first non-long-term pin
    -> no CMA migration required
    -> registration succeeds

That matrix is more valuable than reproducing one failure. It proves each edge of the hypothesis separately.

16.6 Trace the transition

Useful tracepoints or probes include:

pin_user_pages_fast
long-term GUP helper
get_pageblock_migratetype
migrate_pages
ib_umem_get
mlx5/rxe user-MR registration

A reconstructed successful repair path would look like:

long-term pin requested
  pageblock = MIGRATE_CMA
  migrate P -> Q
  retry GUP
  pageblock(Q) != MIGRATE_CMA
  success

The failing case differs by one fact:

migrate P -> Q
  source P has an existing device pin
  migration fails
  -ENOMEM

17. Better telemetry for memory registration

The incident took months partly because each layer discarded context.

A production-quality registration event should preserve:

userspace address and length
allocation class: host / CUDA / DMA-BUF
protocol and direction
MR API used
process and communicator identity
underlying errno
memlock usage and limit
pageblock migratetype for host pages
whether FOLL_LONGTERM was requested
DMA-BUF exporter name
BAR1 usage for GPU mappings
MKey/provider failure stage

Not every field belongs in every log line. A structured trace or error report can gather them conditionally on failure.

In this case, a registration failure report needed to preserve pageblock=MIGRATE_CMA and the failed migration, rather than stopping at “NCCL system error.”

18. What generalizes

Pinning is a contract, not a property

Asking “is the page pinned?” is incomplete. Ask:

Who pinned it?
Through which API?
With FOLL_PIN or only a reference?
Was FOLL_LONGTERM declared?
Which device translation depends on it?
Who owns invalidation?

Two callers can both say “pinned” and still disagree about placement and lifetime.

Physical placement is part of device correctness

Most application code treats physical memory as an implementation detail. Device DMA makes placement observable. ZONE_MOVABLE, CMA, device memory, DAX, and filesystem-backed pages all carry rules that a long-term pin must respect.

Fleet roles need different memory policy

HugeTLB reserves, IOMMU modes, NUMA balancing, transparent huge pages, and reclaim settings can affect DMA-heavy workloads. Sharing a kernel across fleet roles does not mean they all need the same boot policy.

19. The invariant that would have prevented the incident

A page must pass its strongest lifetime and placement contract before any subsystem makes it immovable.

For this incident:

strongest contract = long-term device DMA

Therefore one of the following must happen first:

- allocate from a location valid for long-term pins;
- request FOLL_LONGTERM on the first pin so Linux can relocate it;
- export one shared lifetime object used by all devices;
- or remove the movable reserve from a workload that does not need it.

What cannot safely happen is:

pin now under a weak placement contract
validate long-term placement later

20. A note on modern stacks

This series describes an R525/R535-era stack. It should not be read as a claim that an unchanged failure path exists on every current cluster.

Three later changes matter:

  • In public R555 source on x86, NVIDIA’s host-page import adds FOLL_LONGTERM. That moves the placement check to the first NVIDIA pin and closes the exact missing-contract window on that public path.
  • Starting with R560, NVIDIA makes the open kernel-module flavor the default and suggested installation on supported GPUs, making DMA-BUF the normal direction for modern GPUDirect deployments when the rest of the stack supports it.
  • NCCL 2.19 introduced explicit user-buffer registration for NVLS through ncclCommRegister() / ncclMemAlloc(), and later releases expanded registration to more paths. Modern NCCL can therefore create and reuse registrations differently from the internal-buffer flow described for 2.17.

Those changes rearrange or close this particular path; they do not make memory-lifetime contracts irrelevant. Older driver branches, vendor forks, and third-party device pins can still create the same general ordering error. The first debugging step should always be to identify the exact allocation, registration API, driver flavor, and GUP flags on the versions actually running.

Source map