Skip to content

nccl_ep: generation-tag LL P2P signals so sibling handles cannot alias them - #6

Closed
Oseltamivir wants to merge 1 commit into
NVIDIA:mainfrom
Oseltamivir:fix/ll-generation-tagged-signals
Closed

Oseltamivir wants to merge 1 commit into
NVIDIA:mainfrom
Oseltamivir:fix/ll-generation-tagged-signals

Conversation

@Oseltamivir

Copy link
Copy Markdown

Migrated from NVIDIA/nccl#2306 at @kwen2501's suggestion, now that NCCL-EP lives here. Fixes the LOW_LATENCY receive-timeout failure reported in NVIDIA/nccl#2303. Rebased onto this repo's layout (low_latency.cull_ep.cuh + ll_ep_adapter.cu) and reshaped along the way — see "What changed versus the old PR".

Root cause

The LL count/flag slots are addressed by offsets computed in handle->ll.layout, which point into group->rdma_buffer — per-group memory. The parity that selects between the two slots, handle->ll.buffer_idx, is per-handle state, flipped on each dispatch and combine (and it also picks signal_base).

So two handles created on one group with the same config compute identical offsets and alias the same parity-0/parity-1 count and flag slots, while each advances its own parity independently. The double-buffer invariant breaks across handles: one handle's "next buffer, safe to clean" is another handle's "current buffer, in flight". Cleans wipe live signals, and polls accept a sibling handle's leftover value — which is indistinguishable from the awaited one, because the values are workload-determined and finish flags are the constant 1. Counts and flags share a single offset (dispatch_rdma_recv_count_buffer_offset == combine_rdma_recv_flag_buffer_offset), so both are affected.

The visible symptom is the systematic NCCL EP timeout for dispatch/combine receive warnings followed by trap()cudaErrorLaunchFailure (719).

Evidence

Stock nccl4py[cu13]==0.3.1, GB200, 4 ranks on one tray, identical workload; the only variable is how many LL handles the caller creates on the group:

LL handles on the group result
1 rc=0, zero receive timeouts, combine correctness passing
2 64 dispatch + 6 combine receive timeouts → 719

ep_bench creates exactly one handle and so never hits this: it stayed green through eight configurations up to 20 000 iterations, including with its per-iteration MPI_Barrier removed, with a 500 ms stall injected on one rank, and with dispatch-only and combine-only timed loops. Rank stagger and loop shape are ruled out. Our own benchmark creates one handle per token count and interleaves them, so it fails almost immediately — and it fails the same way on x86 (H100, B200, B300), so this is not GB200/GB300- or MNNVL-specific.

The fix

Stamp every LL launch with a monotonically increasing generation, carry it into the signal value, and accept a signal only when its generation matches:

count: (gen << 16) | (numTokensSent + 1)     accept iff value >> 16 == gen
flag:  (gen << 16) | 1                        accept iff value >> 16 == gen

A sibling handle's or an older call's value then carries a different generation and is inert. The generation is never 0, so a freshly cleaned (all-zero) slot can never be mistaken for a signal. Ranks execute the same LL call sequence, so a per-process counter stays consistent across them; a SEND-phase launch advances it, a RECV-only launch reuses it. The GIN path is untouched — its accumulating signals have different semantics.

What changed versus the old PR

  • Rebased onto nccl_ep/device/ll_ep.cuh and ll_ep_adapter.cu.
  • The generation is now a kernel argument (signalGen, threaded exactly like the existing signalsBase) instead of a __device__ symbol set with cudaMemcpyToSymbol. That was necessary here — the LL kernels are JIT-compiled, so a symbol in the AOT module would not be the one the kernel reads — and it is also the cleaner shape reviewers asked about on the old PR.
  • The root-cause description is corrected. The old PR attributed the failure to a rank lapping a parity cycle, and claimed LL worked on x86. Both were wrong; details in [Issue]: OW_LATENCY dispatch/combine receives time out on GB200/GB300 NVL72 nccl#2303.

Validation

  • Runtime behaviour was validated on the pre-move code (contrib/nccl_ep at the commit the shipped wheel is built from), where the identical change turns a 100 %-reproducible wedge into full clean runs: decode ladders T=1…256, 2048 samples per point, dispatch and combine correctness passing, at world=4 (single tray, three runs) and world=8 (two trays).
  • On this tree the ported change builds cleanly AOT (library and ep_bench).
  • I could not exercise the JIT path end to end here: main's LL kernels reference GIN device APIs (ncclGin, ncclGin_SignalAdd, ncclGinSignal_t) that are not in NCCL 2.30.7, the newest published nvidia-nccl-cu13 wheel, so the runtime JIT compile fails. Unmodified main fails identically at the same sites in the same container, so this is an environment/version gap rather than something the patch introduces — but it does mean the JIT argument-list change here is compile-checked only through the AOT build, and a run on a tree with a matching NCCL would be worth doing before merge. Happy to re-run it if you can point me at the right NCCL revision.

Open question

This makes stale and foreign values inert, which fixes the failure — but it is arguably a mitigation. The structural fix is to stop two handles sharing parity state and slots: per-handle signal regions, or one group-level parity counter that every handle advances. Happy to reshape it that way instead.

And if interleaving multiple LL handles on a single group is simply not intended to be supported, then the better outcome is a loud host-side error on the second ncclEpInitHandle for a group — today that usage costs a ~97 s timeout and a trap(), which is very hard to attribute.

…s them

The LL count and finish-flag slots are addressed by offsets held in
handle->ll.layout, which point into group->rdma_buffer -- per-group memory.
The parity that selects between the two slots, handle->ll.buffer_idx, is
per-handle state, flipped on every dispatch and combine.

Two handles created on one group with the same config therefore compute
identical offsets and alias the same parity-0/parity-1 slots while each
advances its own parity independently. The double-buffer invariant breaks
across handles: one handle's "next buffer, safe to clean" is another
handle's "current buffer, in flight", so cleans wipe live signals and polls
accept a sibling handle's leftover value. Nothing distinguishes that value
from the awaited one -- counts are workload-determined and finish flags are
the constant 1 -- so the receive completes against stale data or never
completes at all, surfacing as systematic "NCCL EP timeout for
dispatch/combine receive" warnings, then trap() and
cudaErrorLaunchFailure. Counts and flags share one offset
(dispatch_rdma_recv_count_buffer_offset ==
combine_rdma_recv_flag_buffer_offset), so both paths are affected.

Measured on GB200 with the stock nccl4py[cu13]==0.3.1 wheel, 4 ranks, same
workload, varying only the number of handles the caller creates on the
group: one handle completes with zero receive timeouts and correctness
passing, two handles produce 64 dispatch plus 6 combine receive timeouts
and error 719. ep_bench uses a single handle and stays green through eight
configurations up to 20000 iterations, including with its per-iteration
MPI_Barrier removed, with a 500 ms stall injected on one rank, and with
dispatch-only and combine-only timed loops -- so rank stagger and loop
shape are not involved. The failure reproduces on x86 as well as on
GB200/GB300; it is not MNNVL-specific.

Stamp every LL launch with a monotonically increasing generation, carry it
in the signal value, and accept a signal only when the generation matches:

  count: (gen << 16) | (numTokensSent + 1)   accept iff value >> 16 == gen
  flag:  (gen << 16) | 1                      accept iff value >> 16 == gen

A sibling handle's or an older call's value then carries a different
generation and is inert. The generation is never 0, so a freshly cleaned
all-zero slot can never look like a signal. Ranks run the same LL call
sequence, so a per-process counter stays consistent across them; a
SEND-phase launch advances it and a RECV-only launch reuses it. It travels
to the device as a kernel argument threaded exactly like signalsBase, which
the JIT-compiled kernels require. The GIN path is untouched.

This makes stale and foreign values inert rather than preventing the
aliasing itself; per-handle signal regions, or a single group-level parity
counter that all handles advance, would address it structurally.
@nv-lschneider

Copy link
Copy Markdown
Contributor

Thank you, @Oseltamivir, for your contribution and for identifying this issue — it's a genuine bug in the double-buffering logic.

The Problem

The current implementation assumes that dispatch and combine are always called in a specific order for a given group. This holds regardless of how many handles are in use, but your use case with two handles exposes the issue clearly. I've also encountered it when capturing dispatch and combine separately into CUDA graphs.

Your Proposed Fix

Your design using host-controlled counters does resolve the problem for your use case, and we appreciate the thoughtful approach.

That said, I believe a GPU-driven, CUDA-graph-based solution would be more comprehensive and cover a broader set of scenarios. We're still iterating on the design, as it does come with a performance trade-off. It may end up being opt-in, with the default behavior only guaranteeing safety for certain operation orderings.

Next Steps

We'll keep you updated as the design evolves. Thanks again for the effort you've put into this - we hear you and listen.

@nv-lschneider

Copy link
Copy Markdown
Contributor

We have now commit: 33c183a
Which enables a device drive double buffering semantics.
It should solve the issue that this PR was designed to address.

@Oseltamivir could you check if that commit fixes the issue, you experienced?
If yes, I think we can close this PR.
Thank you.

@Oseltamivir

Copy link
Copy Markdown
Author

W, thanks 🫡

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

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants