howtf did a four-byte buffer crash the checkpoint worker?

1. The crash was in the wrong process

The first crash was in checkpointing. That was the problem.

We had just changed a network transport: the receive side of a collective-communications library. The change registered four bytes of host memory and used them as scratch for a GPUDirect RDMA visibility fence. Focused tests passed. Collectives produced the right answers. Performance looked normal. The new object was so small it barely felt like an allocation.

Then production training jobs began dying in code that had nothing to do with networking. Sometimes a checkpoint worker crashed walking metadata. Sometimes a data-loading child faulted dereferencing an object that was demonstrably valid in the parent. Stack traces pointed into serialization, input processing, allocator internals. The main training process ran on fine, which made the failures look stochastic. Rebuilding the same source could make a crash disappear, or move it somewhere else.

The useful clue was not the faulting instruction. It was the process boundary. Every victim was a child created with fork(). The same pointer, the same virtual address: readable in the parent, a segfault in the child.

That’s the howtf. Memory that survived in the parent and did not exist in the child, delivered by a four-byte change on the other side of the codebase. The feature was measured in bytes. Its side effect was measured in pages.

Scope note. This reconstructs a production incident from first-hand experience, with products, hardware generations, and deployment details deliberately blurred. Every mechanism claim is checked against public source: NVIDIA’s NCCL, rdma-core, and the Linux kernel, linked throughout. The public NCCL commit used below as the reference design also already contains the defense this post ends with. Hold that thought for §8.

2. Two kinds of done

The receive path used GPUDirect RDMA: the NIC (an RNIC, an RDMA-capable NIC) writes incoming data straight into GPU memory over PCIe, skipping the bounce through host RAM.

It is tempting to collapse every notion of “done” into one completion bit. On this path that is unsafe. NVIDIA’s GPUDirect documentation says it directly: even after a third-party device has issued its PCIe writes, a concurrently running GPU kernel can observe stale or partially written data. The DMA write has to be made visible to the scope that will consume it, which is why modern CUDA exposes cuFlushGPUDirectRDMAWrites() as an explicit operation. Transport completion and GPU visibility are two different claims.

The historical NCCL InfiniBand transport, the public skeleton of this design (and a returning character on this site: it last lost its device list to the linker), answers both. For transport completion: the receiver posts a receive work request that is really a notification credit, then advertises the GPU destination address, rkey, and size to the sender through a FIFO. The sender posts IBV_WR_RDMA_WRITE_WITH_IMM: the write lands the payload one-sided in GPU memory, and the immediate consumes the receiver’s credit, generating a completion whose imm_data carries the message size.

For GPU visibility: after that completion, the receiver posts a signaled IBV_WR_RDMA_READ on a small QP connected back to itself. The read’s remote side is the GPU receive buffer. Its local destination is a tiny registered host object. The value read does not matter. An RDMA read is a non-posted operation that cannot complete until its response comes back, and on this platform that response was the ordering point: the plugin API describes this callback as the flush that makes data received into CUDA memory visible to the GPU, and a later NCCL commit states the principle in one line: a CPU read of GPU BAR memory drains prior PCIe posted writes. (That is a platform-specific fence design, not a universal theorem about RDMA reads. The linked source shows its exact mechanics; don’t generalize it past them.)

The design carries an almost comic detail. The registered host object is an int, because the scratch word was declared as one. The read itself transfers a single byte:

struct ncclIbGpuFlush {
  int enabled;
  int hostMem;              // the four-byte scratch word
  struct ibv_mr* hostMr;
  struct ibv_sge sge;
  struct ibv_qp* qp;        // loopback QP
};

// ncclIbAccept():
NCCLCHECK(wrap_ibv_reg_mr(&rComm->gpuFlush.hostMr, rComm->verbs.pd,
                          &rComm->gpuFlush.hostMem, sizeof(int),
                          IBV_ACCESS_LOCAL_WRITE));
rComm->gpuFlush.sge.addr   = (uint64_t)&rComm->gpuFlush.hostMem;
rComm->gpuFlush.sge.length = 1;        // one byte transferred

// ncclIbFlush():
wr.wr.rdma.remote_addr = (uint64_t)data;   // the GPU receive buffer
wr.opcode              = IBV_WR_RDMA_READ;
wr.send_flags          = IBV_SEND_SIGNALED;

ibv_reg_mr() takes a byte address and a byte length. There is no four-byte minimum, no page rounding in the interface. Four bytes is what was asked for.

That registration was the trigger.

3. What ibv_reg_mr does when fork safety is armed

At the verbs API boundary, the registration line means: let the RNIC use these four bytes as a local write destination. On a fork-safe libibverbs, it also means something the signature never hints at: change the fork-inheritance policy of the entire virtual-memory page containing them.

Fork safety was armed in this stack, and that is not exotic. At the commit above, NCCL’s IB init calls ibv_fork_init() unconditionally. For everything else, rdma-core honors RDMAV_FORK_SAFE / IBV_FORK_SAFE from the environment, and real fleets set it: AWS’s EFA network plugin for NCCL exports RDMAV_FORK_SAFE=1 on your behalf and logs that it did. Training stacks fork constantly (data loaders, checkpoint writers), so fork safety was the responsible setting. What it actually does is the surprise.

ibv_fork_init() sets up an interval tree of registered ranges. From then on, every ordinary (non-ODP) registration passes through ibv_dontfork_range() before the provider ever sees it, and that function rounds the byte interval out to page boundaries and applies madvise(MADV_DONTFORK):

start = (uintptr_t) base & ~(range_page_size - 1);
end   = ((uintptr_t) (base + size + range_page_size - 1) &
         ~(range_page_size - 1)) - 1;
/* ... interval-tree bookkeeping, then, on the 0 -> 1 transition ... */
madvise(start, end - start + 1, MADV_DONTFORK);

Deregistration reference-counts the overlap and applies MADV_DOFORK when the last registration touching a range goes away. The accounting is careful. The granularity is the problem. On a 4 KiB-page machine:

requested MR length:        4 bytes
fork-policy side effect: 4096 bytes

And MADV_DONTFORK is blunt. Per madvise(2): “Do not make the pages in this range available to the child after a fork(2).” Not copy-on-write, not zero-filled. Per fork(2), the mapping is simply not inherited. The parent keeps the page. The child has a hole at that address, and the first dereference into the hole is a segfault.

So the full expansion of one innocent-looking line, on the fork-safe path, is:

Register four bytes for the RNIC, and quietly withhold the surrounding page from every child this process will ever fork.

The scratch word did not live on a private page. It was a small field inside a heap object, and the rest of its page belonged to whatever else the allocator had placed there. When that happened to include state a checkpoint or data-loading child would later walk, the child inherited an address space with a hole where its data should have been. The checkpoint object was never corrupted. Its page was withheld.

4. Why the page had to leave the child

Why would a library ever consider withholding your memory from your children a safety feature? Because for pinned DMA memory, the alternative was corruption. This section is the why; it is also the part of the mental model that survives past RDMA.

CPU copy-on-write works because the CPU asks permission. After fork(), parent and child share physical page P, and both mappings are write-protected. When either side writes, the CPU takes a page fault before the store lands, the kernel copies P to a fresh page Q, repoints the writer, and retries. The essential ingredient is the synchronous trap before the destructive write: at copy time, the old contents still exist.

Memory registration builds a second translation path that never asks. A classic (non-ODP) MR is designed to remove page faults from the fast path: at registration time the kernel resolves and pins the pages long-term (pin_user_pages_fast() with FOLL_LONGTERM, plus FOLL_WRITE for writable MRs), builds a scatter-gather list, and maps it into the device’s DMA address space. From then on the RNIC resolves RDMA addresses through its own MR translation to those pinned pages. It does not walk the process page tables, it cannot see a copy-on-write bit, and its DMA write cannot take a CPU page fault. If the process forks and a COW-shared P is still the RNIC’s target, the device writes P directly, and whichever process the kernel had decided should get “the old contents” gets the new bytes instead. Even a notification after the DMA would be too late: the fork-time snapshot is already destroyed.

This is an old problem, and MADV_DONTFORK is its old solution. It has been in Linux since 2.6.16, and madvise(2) still carries the original rationale: preventing copy-on-write from changing the physical location of a page, because “such page relocations cause problems for hardware that DMAs into the page.” The historical fork-safety policy is two sentences: do not try to make DMA participate in COW. Prevent the child from sharing the DMA page at all.

That solved the corruption problem. It also created our crash, because the unit it operates on is the page, and the page contained more than the four bytes anyone asked about. (Faulting device translations exist too, under on-demand paging, IOMMU page faults, and shared virtual addressing. Different memory model. The classic pinned MR is the one this transport used, and its whole point is the pre-resolved fast path.)

5. Why it looked random

Once the page layout is fixed, the mechanism is fully deterministic. The page layout is not fixed.

• Co-tenancy is a layout lottery. A field’s offset inside its struct is stable per build, but the allocation’s base address modulo page size, and its page neighbors, depend on allocator history: size classes, arena state, allocation order, thread timing, feature flags, ASLR. A rebuild reshuffles the lottery, which is why the crash moved between builds.
• The page keeps accepting tenants. The registered page’s free space stays in the allocator’s inventory. Objects malloc’d after the registration can move in next to the scratch word, and they vanish from children too.
• The child must touch the hole. A networking test that never forks cannot see it. A child that forks and immediately execs never touches inherited heap. A checkpoint worker does exactly the dangerous thing: it keeps executing in the inherited address space and walks parent-built state.

So “one binary failed and another did not” was directionally right, and the precise condition was: this build, times this allocator history, times this page placement, times whether the child dereferences the collateral. Every factor except the last one is invisible in source code.

6. Debugging backward from the process boundary

The investigation got short once we stopped asking “which object is corrupt” and started asking “what happened to this virtual page.”

SIGSEGV only in fork children
  -> same pointer valid in the parent          (kills use-after-free theories)
  -> diff /proc/<pid>/maps, parent vs child    (the page is absent, not remapped)
  -> round fault address down to page start
  -> grep /proc/<parent>/smaps: VmFlags has dc (VM_DONTCOPY: "do not copy on fork")
  -> inventory every ibv_reg_mr touching that page

The dc VmFlag is the kernel telling you, in its own handwriting, that someone madvised this range MADV_DONTFORK. From there the registration inventory is finite, and a four-byte MR sitting inside a shared heap page is hard to miss once you are actually looking for it.

One disambiguation worth writing down: a DONTFORK hole is absent, not zeroed. If a child observes zero-filled memory, that is a different mechanism (MADV_WIPEONFORK supplies zero pages on fork) or a tooling artifact from reading around the hole. The child of a DONTFORK page does not read zeros. It faults.

The closing experiment was the causal one: move the RDMA-owned scratch onto its own dedicated, page-rounded allocation, change nothing in checkpointing, and the crash disappears. Reintroduce the shared-page registration and it comes back. That experiment matters because it flips exactly the variable the hypothesis names, page co-tenancy, and nothing else.

7. Reproducing it, no RNIC required

The whole failure lives in one madvise call, so it reproduces in a hundred-odd lines of C on any Linux box, no RDMA hardware involved. The demo’s fake_reg_mr() does literally what fork-safe libibverbs does: round the byte range to page boundaries, madvise(MADV_DONTFORK). A 4-byte “registration” shares a page with state a forked child needs (demo/rdma-fork-dontfork):

page size: 4096 bytes
layout: flush_word at 0x29031000, neighbor_state at 0x29031004 (same page)
reg_mr: asked for 4 bytes at 0x29031000
reg_mr: marked DONTFORK [0x29031000, 0x29032000) = 4096 bytes
smaps:  VmFlags: rd wr mr mw me dc ac
parent: neighbor_state = "epoch=41 step=118000"
child:  dereferencing the same pointer...
parent: child killed by signal 11 (Segmentation fault)
parent: neighbor_state still "epoch=41 step=118000"

There is the whole story in nine lines: four bytes requested, 4096 marked, dc in VmFlags, the parent reading happily, the child dead on the same pointer. And the fix, same registration, dedicated page:

reg_mr: asked for 4 bytes at 0x721a000
reg_mr: marked DONTFORK [0x721a000, 0x721b000) = 4096 bytes
parent: neighbor_state = "epoch=41 step=118000"
child:  dereferencing the same pointer...
parent: child exited 0

8. The fixes: own the page, or let the kernel copy early

The application-level invariant. Memory that may be registered must not share a page with state a forked child may need. That takes two properties, not one: a page-aligned base and a page-rounded dedicated extent. Alignment alone still lets the allocator move another tenant into the tail of the page.

This is not my invariant. It is in NCCL’s tree, with the incident’s lesson written out as a comment (src/include/alloc.h):

// Allocate memory to be potentially ibv_reg_mr'd. This needs to be
// allocated on separate pages as those pages will be marked DONTFORK
// and if they are shared, that could cause a crash in a child process
static ncclResult_t ncclIbMalloc(void** ptr, size_t size) {
  size_t page_size = sysconf(_SC_PAGESIZE);
  ...
  int ret = posix_memalign(&p, page_size, size_aligned);

And here the scope note’s planted thought pays off: at the very commit this post has been quoting, the IB transport’s connection structs are already allocated through this helper. In fact they always were: the helper, comment and all, is present in the first public commit that added the IB transport (2.3.5-5, September 2018). No released public NCCL ever had the shared-page hazard. The comment is scar tissue from a lesson learned before open-sourcing, and the transport in this story re-derived that lesson independently. That is what makes this worth writing down as a class rather than a bug: any codebase that registers small heap objects on a fork-safe verbs stack re-derives it on schedule.

Two footnotes on the invariant. First, it degrades gracefully: on the old path the only memory a child loses is RDMA-owned state it should never touch anyway. Second, huge pages have their own trap: fork-safe rounding uses the base page size unless RDMAV_HUGEPAGES_SAFE is also set, which rdma-core’s own docs call required if the application uses huge pages at all.

The kernel fix. Linux eventually implemented the direct thing: copy early. Since Linux 5.9 for ordinary PTEs and 5.12 for hugetlb (both by Peter Xu), fork() detects pages that may be DMA-pinned and copies them for the child at fork time, before any future DMA can touch them. The parent and the RNIC keep the pinned page and the established translation. The child gets a correct snapshot. No fault in the DMA path was ever needed, only a copy moved earlier in time.

rdma-core exposes the boundary between the two worlds: ibv_is_fork_initialized() returns IBV_FORK_UNNEEDED when the kernel reports copy-on-fork support (over RDMA netlink, not a version sniff), and on such kernels ibv_fork_init() short-circuits to a no-op. On a heterogeneous fleet, though, the oldest kernel in the pool sets your rules. The allocator invariant costs almost nothing and is correct in both worlds. Keep it.

9. What generalizes

Completion must name a scope. “The operation completed” is not a systems statement. Completed at the sender, at the receiver’s CQ, in host memory, or for a GPU consumer are different claims, and most ordering bugs start when two components use the same word for different ones.

A byte-range API can carry page-range consequences. The interface accepted four bytes. The enforcement unit was a page. The same shape hides under cache-line false sharing, huge-page mappings, IOMMU granules, and filesystem blocks: the unit you asked in is not necessarily the unit the system acts in.

Process topology is part of the interface. The transport passed its tests because the data movement was correct. Production added fork(), inherited heaps, checkpoint children. Registered memory couples to all of it, which makes ibv_reg_mr() a lifecycle contract with the whole process tree, not a permission slip for one NIC.

Fix with invariants, not layouts. “Move this integer until the crash stops” survives until the next rebuild. “Registered memory owns its pages and contains nothing a child needs” survives allocator changes, rebuilds, and time. The fix that lasts is the one you can state without mentioning an address.

Epilogue

Every local decision in this incident was reasonable. The GPU needed a visibility fence. The fence needed a registered local buffer, and four bytes was honestly all it needed. Libibverbs needed to keep pinned DMA pages from tearing forked children, and page granularity was the only granularity fork() offers. The checkpoint worker expected inherited memory to be there. Every layer kept its contract.

The contracts did not compose, because four bytes and one page were treated as the same unit. That was the bug.

The hardware-free reproducer (the madvise model, the crash, and the fixed layout) lives at demo/rdma-fork-dontfork. Mechanism sources are pinned commits linked inline: NCCL’s historical IB transport, rdma-core’s fork tracking, the kernel’s umem pinning path, and Peter Xu’s copy-on-fork series.