Make the elastic GIN signal reads backend-agnostic (fix latent PROXY layout mismatch) - #706
Make the elastic GIN signal reads backend-agnostic (fix latent PROXY layout mismatch)#706KeitaW wants to merge 1 commit into
Conversation
…DAKI cast The GIN barrier and the PP send/recv wait both read a signal by casting the opaque GIN handle to `ncclGinGdakiGPUContext*`, indexing it by `contextId`, and loading `signals_table.buffer`: deep_ep/include/deep_ep/common/comm.cuh:166 deep_ep/include/deep_ep/impls/pp_send_recv.cuh:24 `ncclGin::_ginHandle` is an opaque, backend-owned pointer. NCCL reaches it only through `ncclGinCall<ncclGinApi_GetSignalPtr>`, whose per-backend specializations cast it to different structs: nccl_device/gin/gdaki/gin_gdaki.h ncclGinGdakiGPUContext[contextId] nccl_device/gin/proxy/gin_proxy.h ncclGinProxyGpuCtx_t[contextId] So the cast hardcodes one backend's layout. The two structs are independent declarations of different size, so `+ contextId` strides by the wrong amount over a PROXY array. Measured with sizeof/offsetof on the 2.30.7 headers (both structs are standard-layout): sizeof: GDAKI 88 PROXY 72 signal-array pointer offset: GDAKI 40 PROXY 40 At contextId == 0 the offsets coincide, so the cast reads the right pointer by luck rather than by contract. At contextId == 1 it loads 8 bytes from absolute byte 88+40 = 128, while the correct PROXY field is at 72+40 = 112; byte 128 is offset 56 of PROXY element 1, i.e. `ncclGinProxyGpuCtx_t::lastIssuedGet`. The wait would poll the per-peer get-index array as if it were the signal array. Today every call site takes QP 0, so the bug is latent rather than active (comm.cuh:156 constructs `ncclGin(nccl_dev_comm, 0, ...)` under the comment "Use QP 0 to do barrier" at comm.cuh:153). It is one refactor from live: the flush loop directly above already iterates `ncclGin(comm, i, ...)` for `i < ginContextCount` (comm.cuh:145-146), so moving the barrier off QP 0 would make the wait spin on an address that is never signalled until `timeout_while` fires `ptx::trap()`. Because CUDA 719 is sticky and gets attributed to whichever launch runs next, the resulting report would name the barrier as the faulting kernel while the real fault is the cast -- the same misattribution pattern as deepseek-ai#704. There is a second, independent consequence: `ncclGinGdakiGPUContext` is only declared when `NCCL_GIN_GDAKI_ENABLE` is set, which pulls in the DOCA GPUNetIO headers (nccl_device/gin/gin_device_api.h). On an NCCL install built without GDAKI, these two headers do not compile at all -- the type is incomplete. Fix: call the public `gin.readSignal(signal_idx, 64, memory_order_acquire)`, which dispatches on the negotiated backend. This keeps DeepEP's own `timeout_while` loop, so it is not blocked on NCCL adding a timeout to `waitSignal` (the TODO at comm.cuh:160); only the read is delegated. Equivalent by construction, not by inspection: - GDAKI's `GetSignalPtr` returns `{signals_table.buffer + id, offset = 0}` and `readSignal` computes `(raw - offset) & mask`; with `bits = 64` the mask is all ones, so the returned value is identical to the previous load. - Under PROXY the subtracted offset is a zero-initialised array that is only ever written by `resetSignal`, which DeepEP does not call. - `pp_send_recv.cuh` compares against a target that may be negative (`send_count - num_max_inflight_tensors + 1`), so the result is cast back to `int64_t` to keep that comparison signed. Verified, no behaviour claims beyond what was measured: - nvcc 12.8, sm_90, real NCCL device headers (2.30.7), DeepEP's own JIT flags, both call-site templates force-instantiated so the bodies are codegen'd: * PROXY-only (GDAKI disabled): before = does not compile ("expression must be a pointer to a complete object type" at both sites); after = compiles. * GDAKI enabled: before and after both compile. - Differential PTX on the GDAKI configuration, where both versions build, so the only variable is this change: the signal poll still emits exactly 4 system- scope 64-bit acquire loads, none at a weaker scope, and no new relaxed load. The one opcode difference is that `ld.acquire.sys.L1::no_allocate.global.u64` becomes `ld.acquire.sys.b64`, i.e. the L1-no-allocate cache hint on the spin poll is dropped. That hint is a performance property, and this change makes no performance claim -- it has not been measured on a multi-node GIN job. This is a correctness/portability change only. Signed-off-by: Keita Watanabe <keitaw09@gmail.com>
| // casting the opaque GIN handle to one backend's context struct: the | ||
| // handle layout differs per backend, so a direct cast is only correct | ||
| // for GDAKI and only at context 0. | ||
| const auto signal = gin.readSignal(signal_idx, 64, cuda::memory_order_acquire); |
There was a problem hiding this comment.
🔵 suggestion: Non-blocking: the previous inline-asm poll used ld.acquire.sys.L1::no_allocate.global.u64; cuda::atomic_ref::load (via readSignal) will emit ld.acquire.sys.b64 without the L1::no_allocate hint. Scope/ordering/width/count are preserved, but the cache hint on the busy-wait spin poll is lost. As acknowledged in the PR this is a performance property with no measurement made; if the hint is considered worth keeping on the poll, consider restoring it without reintroducing the cast.
🤖 v3
🤖 ds-review-bot Code Reviewv6The changes correctly replace backend-specific signal access with NCCL's dispatched API while preserving signed comparison semantics where required. v4v3The change makes the elastic GIN signal reads backend-agnostic by replacing the direct cast of the opaque
Both changes keep DeepEP's own The change is correct and I recommend it be merged. The one non-blocking observation (dropped L1::no_allocate cache hint on the spin poll) is captured as a suggestion comment below. Files reviewed: 2 |
|
We run the PROXY backend on EFA (p5en/H200, IBGDA unavailable), so this PR's fix path is the one we exercise — here is a measured PROXY data point, since the description notes the direct cast is only correct for GDAKI at context 0 and no PROXY measurement existed. Verdict: correctness-clean and performance-neutral at EP16 on proxy-Gin/EFA. Setup: 2× p5en.48xlarge (H200, 16 EFA NICs/node), EP16 = 2 nodes × 8 GPUs,
Pooled across the interleaved launches: dispatch +2.1%, combine +1.0% — inside launch-to-launch noise on this fabric (the two identical-binary BEFORE launches differ by 9.5% of median between themselves, which is why we interleaved and pooled rather than trusting a single launch per side). All 4 launches passed every gate: 8/8 ranks DONE per node, 0 timeouts, 0 tracebacks, 0 So on at least one non-GDAKI backend the fix costs nothing measurable at EP16 while closing the wrong-layout read. Happy to re-run at other shapes or EP32 when capacity allows. |
3 participants
Problem
The
elastic/GIN signal waits read a signal by castingncclGin::_ginHandletoncclGinGdakiGPUContext*, indexing that bycontextId, and loadingsignals_table.buffer:deep_ep/include/deep_ep/common/comm.cuh:166-167— the GIN device barrierdeep_ep/include/deep_ep/impls/pp_send_recv.cuh:24-26— the PP send/recv wait_ginHandleis an opaque, backend-owned pointer. NCCL only ever reaches it viancclGinCall<ncclGinApi_GetSignalPtr>, and the per-backend specializations cast it to different structs:_ginHandlepoints atncclGinGdakiGPUContext[contextId]nccl_device/gin/gdaki/gin_gdaki.hncclGinProxyGpuCtx_t[contextId]nccl_device/gin/proxy/gin_proxy.hSo the cast hardcodes one backend's layout. The two structs are independent declarations with different sizes, so
+ contextIdstrides by the wrong amount over a PROXY array. Measured withsizeof/offsetofon the 2.30.7 headers (both are standard-layout, so the offsets are the whole story):ncclGinGdakiGPUContextncclGinProxyGpuCtx_tsizeofsignals_table.buffer)signals)At
contextId == 0the two offsets coincide, so the cast reads the right pointer — by luck, not by contract. AtcontextId == 1the cast loads 8 bytes from absolute byte88 + 40 = 128, while the correct PROXY field sits at72 + 40 = 112. Byte 128 is offset 56 of PROXY element 1, which isncclGinProxyGpuCtx_t::lastIssuedGet— so the wait would poll NCCL's per-peer get-index array as if it were the signal array.There is a second, independent consequence.
ncclGinGdakiGPUContextis only declared whenNCCL_GIN_GDAKI_ENABLEis set, which pulls in the DOCA GPUNetIO headers (nccl_device/gin/gin_device_api.h). On an NCCL install built without GDAKI, these two headers do not compile at all — the type is incomplete, so the+ contextIdpointer arithmetic is ill-formed. Since the JIT passes only-I ${nccl_root}/include(csrc/jit/compiler.hpp),elastic/currently carries an implicit build dependency on a DOCA header tree.Why it is latent today, and how it would surface
Every call site currently takes QP 0 —
impls/barrier.cuh:23, andcomm.cuh:156(const ncclGin gin(nccl_dev_comm, 0, ...), under the comment "Use QP 0 to do barrier" atcomm.cuh:153) — which is exactly the coincidentally-correct case.It is one refactor from live. The flush loop immediately above the barrier already iterates all contexts:
with
num_qpsup tonccl_dev_comm.ginContextCount. Whenever the barrier follows the flush loop off QP 0, the wait would poll an address that is never signalled, spin totimeout_while, and hitptx::trap().That failure would be reported badly. CUDA 719 is sticky: once raised, every subsequent launch on the context returns it, so the error lands on whichever library makes the next launch-API call. The report would name the barrier while the actual fault is the cast — the same misattribution pattern described in #704.
Change
Call the public, backend-dispatched accessor instead:
Two files, one call site each, no signature or control-flow changes.
readSignalis declared on the publicncclGininterface (nccl_device/gin.h:344) and routes throughncclGinCall<ncclGinApi_GetSignalPtr>, so it picks the accessor for the backend NCCL actually negotiated. Notably this keeps DeepEP's owntimeout_whileloop — only the read is delegated. That is whywaitSignal()is not the right substitute here, and why this does not need to wait for the NCCL-side timeout support that theTODO(NCCL)atcomm.cuh:160refers to.Equivalent by construction, not by inspection:
GetSignalPtrreturns{signals_table.buffer + id, offset = 0}, andreadSignalcomputes(raw - offset) & mask. Withoffset == 0andbits == 64the mask is all ones, so the value is identical to what the removed load produced.signalOffsets, which is allocated zeroed and is written only byncclGinApi_ResetSignal. DeepEP never callsresetSignal, so the offset stays 0 and the comparison is unchanged.pp_send_recv.cuhcompares against a target that may legitimately be negative (send_count - num_max_inflight_tensors + 1).readSignalreturnsuint64_t, so the result is cast back toint64_tto keep that comparison signed — an unsigned compare there would change behaviour.Verification
Toolchain: nvcc 12.8 (V12.8.93),
-arch=sm_90, real NCCL 2.30.7 device headers staged in the install layout, using DeepEP's own JIT flags (-std=c++20 -O3 --expt-relaxed-constexpr --expt-extended-lambda). Both affected kernel templates are force-instantiated in the test TU so the bodies are actually type-checked and codegen'd — an uninstantiated template proves nothing.Compile matrix (the PROXY-only column is the negative control):
NCCL_GIN_GDAKI_ENABLE=0(PROXY only)expression must be a pointer to a complete object type, at both call sitesNCCL_GIN_GDAKI_ENABLE=1(DOCA headers present)Differential PTX, run on the GDAKI configuration — the only one where both versions build, so the change is the sole variable:
Scope, ordering, width and count of the poll are preserved. The one opcode difference:
ld.acquire.sys.L1::no_allocate.global.u64becomesld.acquire.sys.b64— i.e. the L1-no-allocate cache hint on the spin poll is dropped, because the hint came from the hand-written inline asm andcuda::atomic_ref::loaddoes not emit it.That hint is a performance property, and this PR makes no performance claim. It has not been measured on a multi-node GIN job. If the hint is considered worth keeping on the poll, it can be restored without going back to the cast, and I am happy to do that in this PR.
Also checked, with positive controls on every "absent" assertion: no remaining
_ginHandleuse or GDAKI-struct cast anywhere indeep_ep/; thetrap()sites, thetimeout_whilewrappers, and the timeoutprintftext are all unchanged.Scope
Correctness and portability only — no behaviour change on the GDAKI path, no performance claim. Independent of #704: that one touches
csrc/kernels/legacy/internode_ll.cu, this one touches only the twoelastic/headers. Both branches are single commits on the same base, the changed file sets are disjoint, andgit merge-treereports a conflict-free merge in both orders with both changes present in the resulting tree.