nccl_ep: generation-tag LOW_LATENCY P2P signals to prevent stale-signal consumption

[Oseltamivir]

Fixes issue #2303: the LOW_LATENCY receive-timeout failure reported in #2303 (LL dispatch/combine receives time out on GB200/GB300 NVL72 while HT works).

Root cause

The LL count/flag protocol double-buffers signal slots and polls for a nonzero value, but signal values carry no generation: a signal left over from the previous use of the same parity buffer (two calls earlier) is bit-identical to a fresh one whenever the workload repeats — per-pair token counts are workload-determined, and finish flags are the constant 1. If any rank falls one full parity cycle behind (any scheduling hiccup suffices), its peers consume the stale signals, complete their calls without it, and lap it arbitrarily. The pipeline desynchronizes and wedges with the familiar systematic Warning: NCCL EP timeout for dispatch/combine receive → trap() → error 719.

We root-caused this with a device-side event ring (globaltimer-stamped records of DISP/COMB start, clean-done, count/flag sent/got, dumped at first timeout). At the wedge the four ranks were executing different call sequences — two ranks in dispatch #12, one in #11, one still in combine #9 — while receivers had recorded CNT_GOT "from" the lagging rank for dispatch calls it never executed: stale counts from two calls earlier, value-indistinguishable at a fixed workload.

This explains the platform asymmetry: on GB200/GB300 the failure is reproducible within hundreds of iterations at any token count (fast iterations make the 2-call lap easy), while on x86 single-node platforms the same defect surfaces only rarely as top-load combine corruption. Any µs-scale perturbation of the kernels (e.g. printf instrumentation) masks it — which made it look like a memory-ordering issue. It is not one: the CUDA-level primitives check out (we validated st.release.sys / ld.acquire.sys posted-store visibility across cuMem FABRIC and POSIX-FD peer mappings on GB200 in isolation), and neither system-scope atomics for the signal stores nor __threadfence_system() after the buffer cleans changes the outcome.

Fix

Tag every P2P signal value with a per-op generation and poll for the exact generation, so stale signals are inert:

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

The generation is stamped per launch (stream-ordered cudaMemcpyToSymbolAsync from pinned staging; a SEND-phase launch advances it, a RECV-only launch reuses it). The GIN path is untouched — its accumulating signals have different semantics. This also removes the protocol's sensitivity to clean-vs-arrival ordering on the recycled parity buffer.

Notes for review:

• Generation state is process-global in this change; moving it into per-group workspace (or plumbing it as a kernel argument) is a straightforward follow-up we're happy to do if you prefer that shape.
• numTokensSent + 1 fits 16 bits up to 65534 dispatch tokens per rank; an EP_HOST_ASSERT guard could be added if desired.

Validation

GB200 NVL72 (aarch64, CUDA 13.0 driver, nvcc 13.3, world within one NVLink domain):

• Before: LL wedges within hundreds of iterations (both world=4 single tray and world=8 across two trays), 100% reproducible.
• After: full decode ladder T=1..256, 2048 samples per point, dispatch+combine correctness checks passing; world=4 and world=8 validated.
• HT behavior unchanged.

One pre-existing rarity remains out of scope: an occasional combine-correctness flake at exactly max_dispatch_tokens_per_rank (top rung), which we also observe on x86 with the unmodified wheel.

[kwen2501]

Thanks for the PR. NCCL-EP is moved to the nccl-extensions repo. Would you like to migrate the PR over? We can also help making a copy.

[Oseltamivir]

Migrated to NVIDIA/nccl-extensions#6 — thanks @kwen2501.

The port is not a straight copy: it is rebased onto ll_ep.cuh / ll_ep_adapter.cu, and the generation now travels as a kernel argument (threaded like signalsBase) rather than a __device__ symbol, since the LL kernels over there are JIT-compiled and a symbol in the AOT module is not the one the kernel would read.

The root-cause description has also been corrected. This PR attributed the failure to a rank lapping a parity cycle and said LL worked on x86; both were wrong. The real trigger is two LL handles on one group aliasing the same parity signal slots — handle->ll.buffer_idx is per-handle while its layout offsets point into the per-group rdma_buffer — and it reproduces on x86 too. Full details in #2303 and on the new PR.

Closing in favour of the migrated one.