deepseek-ai/DeepEP

Fork 0

mirror of https://github.com/deepseek-ai/DeepEP.git synced 2026-11-05 00:47:32 +00:00

[PR #706] Make the elastic GIN signal reads backend-agnostic (fix latent PROXY layout mismatch) #721

New issue

Open

opened 2026-10-27 21:05:37 +00:00 by bsakaguchi · 0 comments

bsakaguchi commented

2026-10-27 21:05:37 +00:00

Owner

📋 Pull Request Information

Original PR: https://github.com/deepseek-ai/DeepEP/pull/706
Author: @KeitaW
Created: 7/28/2026
Status: 🔄 Open

Base: main ← Head: pr/gin-barrier-backend-agnostic-read

📝 Commits (1)

1c90922 elastic: read GIN signals through the backend-dispatched API, not a GDAKI cast

📊 Changes

2 files changed (+10 additions, -7 deletions)

View changed files

📝 deep_ep/include/deep_ep/common/comm.cuh (+5 -3)
📝 deep_ep/include/deep_ep/impls/pp_send_recv.cuh (+5 -4)

📄 Description

Problem

The elastic/ GIN signal waits read a signal by casting ncclGin::_ginHandle to ncclGinGdakiGPUContext*, indexing that by contextId, and loading signals_table.buffer:

deep_ep/include/deep_ep/common/comm.cuh:166-167 — the GIN device barrier
deep_ep/include/deep_ep/impls/pp_send_recv.cuh:24-26 — the PP send/recv wait

_ginHandle is an opaque, backend-owned pointer. NCCL only ever reaches it via ncclGinCall<ncclGinApi_GetSignalPtr>, and the per-backend specializations cast it to different structs:

backend	what `_ginHandle` points at	specialization
GDAKI	`ncclGinGdakiGPUContext[contextId]`	`nccl_device/gin/gdaki/gin_gdaki.h`
PROXY	`ncclGinProxyGpuCtx_t[contextId]`	`nccl_device/gin/proxy/gin_proxy.h`

So the cast hardcodes one backend's layout. The two structs are independent declarations with different sizes, so + contextId strides by the wrong amount over a PROXY array. Measured with sizeof/offsetof on the 2.30.7 headers (both are standard-layout, so the offsets are the whole story):

	GDAKI `ncclGinGdakiGPUContext`	PROXY `ncclGinProxyGpuCtx_t`
`sizeof`	88	72
byte offset of the signal-array pointer	40 (`signals_table.buffer`)	40 (`signals`)

At contextId == 0 the two offsets coincide, so the cast reads the right pointer — by luck, not by contract. At contextId == 1 the cast loads 8 bytes from absolute byte 88 + 40 = 128, while the correct PROXY field sits at 72 + 40 = 112. Byte 128 is offset 56 of PROXY element 1, which is ncclGinProxyGpuCtx_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. 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, so the + contextId pointer 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, and comm.cuh:156 (const ncclGin gin(nccl_dev_comm, 0, ...), under the comment "Use QP 0 to do barrier" at comm.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:

comm.cuh:145-146
    for (int i = global_warp_idx; i < num_qps; i += kNumSMs * kNumWarps)
        ncclGin(nccl_dev_comm, i, NCCL_GIN_RESOURCE_SHARING_CTA).flush(ncclCoopWarp());

with num_qps up to nccl_dev_comm.ginContextCount. Whenever the barrier follows the flush loop off QP 0, the wait would poll an address that is never signalled, spin to timeout_while, and hit ptx::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:

-            const auto gdaki = static_cast<struct ncclGinGdakiGPUContext*>(gin._ginHandle) + gin.contextId;
-            const auto signal_ptr = reinterpret_cast<uint64_t*>(__ldg(reinterpret_cast<uint64_t*>(&gdaki->signals_table.buffer))) + signal_idx;
             timeout_while<kNumTimeoutCycles>([=](const bool& is_last_check) {
-                const auto signal = ptx::ld_acquire_sys<uint64_t>(signal_ptr);
+                const auto signal = gin.readSignal(signal_idx, 64, cuda::memory_order_acquire);

Two files, one call site each, no signature or control-flow changes.

readSignal is declared on the public ncclGin interface (nccl_device/gin.h:344) and routes through ncclGinCall<ncclGinApi_GetSignalPtr>, so it picks the accessor for the backend NCCL actually negotiated. Notably this keeps DeepEP's own timeout_while loop — only the read is delegated. That is why waitSignal() is not the right substitute here, and why this does not need to wait for the NCCL-side timeout support that the TODO(NCCL) at comm.cuh:160 refers to.

Equivalent by construction, not by inspection:

GDAKI: its GetSignalPtr returns {signals_table.buffer + id, offset = 0}, and readSignal computes (raw - offset) & mask. With offset == 0 and bits == 64 the mask is all ones, so the value is identical to what the removed load produced.
PROXY: the subtracted offset comes from signalOffsets, which is allocated zeroed and is written only by ncclGinApi_ResetSignal. DeepEP never calls resetSignal, so the offset stays 0 and the comparison is unchanged.
pp_send_recv.cuh compares against a target that may legitimately be negative (send_count - num_max_inflight_tensors + 1). readSignal returns uint64_t, so the result is cast back to int64_t to 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 config	before	after
`NCCL_GIN_GDAKI_ENABLE=0` (PROXY only)	fails — `expression must be a pointer to a complete object type`, at both call sites	compiles (70,480 B object)
`NCCL_GIN_GDAKI_ENABLE=1` (DOCA headers present)	compiles (90,048 B)	compiles (91,024 B)

Differential PTX, run on the GDAKI configuration — the only one where both versions build, so the change is the sole variable:

	before	after
64-bit system-scope acquire loads on the signal poll	4	4
64-bit acquire loads at weaker scope	0	0
new relaxed-load opcode kinds introduced	—	0

Scope, ordering, width and count of the poll are preserved. The one opcode difference: 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, because the hint came from the hand-written inline asm and cuda::atomic_ref::load does 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 _ginHandle use or GDAKI-struct cast anywhere in deep_ep/; the trap() sites, the timeout_while wrappers, and the timeout printf text 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 two elastic/ headers. Both branches are single commits on the same base, the changed file sets are disjoint, and git merge-tree reports a conflict-free merge in both orders with both changes present in the resulting tree.

_{🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.}

## 📋 Pull Request Information **Original PR:** https://github.com/deepseek-ai/DeepEP/pull/706 **Author:** [@KeitaW](https://github.com/KeitaW) **Created:** 7/28/2026 **Status:** 🔄 Open **Base:** `main` ← **Head:** `pr/gin-barrier-backend-agnostic-read` --- ### 📝 Commits (1) - [`1c90922`](https://github.com/deepseek-ai/DeepEP/commit/1c90922f03a841223d0a6e81147315df78c57366) elastic: read GIN signals through the backend-dispatched API, not a GDAKI cast ### 📊 Changes **2 files changed** (+10 additions, -7 deletions) <details> <summary>View changed files</summary> 📝 `deep_ep/include/deep_ep/common/comm.cuh` (+5 -3) 📝 `deep_ep/include/deep_ep/impls/pp_send_recv.cuh` (+5 -4) </details> ### 📄 Description ### Problem The `elastic/` GIN signal waits read a signal by casting `ncclGin::_ginHandle` to `ncclGinGdakiGPUContext*`, indexing that by `contextId`, and loading `signals_table.buffer`: - `deep_ep/include/deep_ep/common/comm.cuh:166-167` — the GIN device barrier - `deep_ep/include/deep_ep/impls/pp_send_recv.cuh:24-26` — the PP send/recv wait `_ginHandle` is an opaque, backend-owned pointer. NCCL only ever reaches it via `ncclGinCall<ncclGinApi_GetSignalPtr>`, and the per-backend specializations cast it to *different* structs: | backend | what `_ginHandle` points at | specialization | |---|---|---| | GDAKI | `ncclGinGdakiGPUContext[contextId]` | `nccl_device/gin/gdaki/gin_gdaki.h` | | PROXY | `ncclGinProxyGpuCtx_t[contextId]` | `nccl_device/gin/proxy/gin_proxy.h` | So the cast hardcodes one backend's layout. The two structs are independent declarations with different sizes, so `+ contextId` strides by the wrong amount over a PROXY array. Measured with `sizeof`/`offsetof` on the 2.30.7 headers (both are standard-layout, so the offsets are the whole story): | | GDAKI `ncclGinGdakiGPUContext` | PROXY `ncclGinProxyGpuCtx_t` | |---|---|---| | `sizeof` | 88 | 72 | | byte offset of the signal-array pointer | 40 (`signals_table.buffer`) | 40 (`signals`) | At `contextId == 0` the two offsets coincide, so the cast reads the right pointer — by luck, not by contract. At `contextId == 1` the cast loads 8 bytes from absolute byte `88 + 40 = 128`, while the correct PROXY field sits at `72 + 40 = 112`. Byte 128 is offset 56 of PROXY element 1, which is `ncclGinProxyGpuCtx_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. `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, so the `+ contextId` pointer 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`, and `comm.cuh:156` (`const ncclGin gin(nccl_dev_comm, 0, ...)`, under the comment *"Use QP 0 to do barrier"* at `comm.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: ``` comm.cuh:145-146 for (int i = global_warp_idx; i < num_qps; i += kNumSMs * kNumWarps) ncclGin(nccl_dev_comm, i, NCCL_GIN_RESOURCE_SHARING_CTA).flush(ncclCoopWarp()); ``` with `num_qps` up to `nccl_dev_comm.ginContextCount`. Whenever the barrier follows the flush loop off QP 0, the wait would poll an address that is never signalled, spin to `timeout_while`, and hit `ptx::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: ```diff - const auto gdaki = static_cast<struct ncclGinGdakiGPUContext*>(gin._ginHandle) + gin.contextId; - const auto signal_ptr = reinterpret_cast<uint64_t*>(__ldg(reinterpret_cast<uint64_t*>(&gdaki->signals_table.buffer))) + signal_idx; timeout_while<kNumTimeoutCycles>([=](const bool& is_last_check) { - const auto signal = ptx::ld_acquire_sys<uint64_t>(signal_ptr); + const auto signal = gin.readSignal(signal_idx, 64, cuda::memory_order_acquire); ``` Two files, one call site each, no signature or control-flow changes. `readSignal` is declared on the public `ncclGin` interface (`nccl_device/gin.h:344`) and routes through `ncclGinCall<ncclGinApi_GetSignalPtr>`, so it picks the accessor for the backend NCCL actually negotiated. Notably this **keeps DeepEP's own `timeout_while` loop** — only the read is delegated. That is why `waitSignal()` is not the right substitute here, and why this does not need to wait for the NCCL-side timeout support that the `TODO(NCCL)` at `comm.cuh:160` refers to. Equivalent by construction, not by inspection: - **GDAKI:** its `GetSignalPtr` returns `{signals_table.buffer + id, offset = 0}`, and `readSignal` computes `(raw - offset) & mask`. With `offset == 0` and `bits == 64` the mask is all ones, so the value is identical to what the removed load produced. - **PROXY:** the subtracted offset comes from `signalOffsets`, which is allocated zeroed and is written *only* by `ncclGinApi_ResetSignal`. DeepEP never calls `resetSignal`, so the offset stays 0 and the comparison is unchanged. - **`pp_send_recv.cuh`** compares against a target that may legitimately be negative (`send_count - num_max_inflight_tensors + 1`). `readSignal` returns `uint64_t`, so the result is cast back to `int64_t` to 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 config | before | after | |---|---|---| | `NCCL_GIN_GDAKI_ENABLE=0` (PROXY only) | **fails** — `expression must be a pointer to a complete object type`, at both call sites | **compiles** (70,480 B object) | | `NCCL_GIN_GDAKI_ENABLE=1` (DOCA headers present) | compiles (90,048 B) | compiles (91,024 B) | **Differential PTX**, run on the GDAKI configuration — the only one where *both* versions build, so the change is the sole variable: | | before | after | |---|---|---| | 64-bit system-scope acquire loads on the signal poll | 4 | 4 | | 64-bit acquire loads at weaker scope | 0 | 0 | | new relaxed-load opcode kinds introduced | — | 0 | Scope, ordering, width and count of the poll are preserved. The one opcode difference: `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, because the hint came from the hand-written inline asm and `cuda::atomic_ref::load` does 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 `_ginHandle` use or GDAKI-struct cast anywhere in `deep_ep/`; the `trap()` sites, the `timeout_while` wrappers, and the timeout `printf` text 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 two `elastic/` headers. Both branches are single commits on the same base, the changed file sets are disjoint, and `git merge-tree` reports a conflict-free merge in both orders with both changes present in the resulting tree. --- <sub>🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.</sub>

bsakaguchi added the

pull-request

label

2026-10-27 21:05:37 +00:00

No labels

pull-request

No milestone

No project

No assignees

1 participant

Notifications

Due date

The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference

deepseek-ai/DeepEP#721

No description provided.

Rows
Columns