deepseek-ai/DeepEP

Fork 0

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

[PR #791] Make symmetric-window ordering configurable: opt into relaxed ordering to recover ~6x hybrid inter-node EP bandwidth #788

New issue

Open

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

bsakaguchi commented

2026-10-27 21:06:21 +00:00

Owner

📋 Pull Request Information

Original PR: https://github.com/deepseek-ai/DeepEP/pull/791
Author: @sssssakuraaaaa
Created: 10/23/2026
Status: 🔄 Open

Base: main ← Head: fix/relax-symmetric-window-ordering

📝 Commits (1)

863328c Make symmetric-window ordering configurable (EP_RELAXED_WIN_ORDERING)

📊 Changes

1 file changed (+14 additions, -1 deletions)

View changed files

📝 csrc/kernels/comm/context.cpp (+14 -1)

📄 Description

TL;DR

The main symmetric window is registered with NCCL_WIN_STRICT_ORDERING
(csrc/kernels/comm/context.cpp). On hybrid inter-node EP this costs ~6x
dispatch/combine bandwidth. The strict flag was added in #674 to fix the ordering
hazard in #661, but that hazard only bites when PCIe native atomics
(PCI_ATOMIC_MODE=4) are combined with relaxed ordering. This PR keeps strict as the
default and adds an opt-in (EP_RELAXED_WIN_ORDERING=1) to restore the fast
NCCL_WIN_DEFAULT path on platforms where relaxed ordering is safe.

Background: two optimizations that are each fine alone

DeepEP's inter-node path leans on two independent hardware optimizations:

PCIe Relaxed Ordering (RO) on the IB MR. The large token payload goes over RDMA
write; RO lets posted writes land out of order, saturates PCIe bandwidth and reduces
head-of-line blocking. This is why the window was originally registered with
NCCL_WIN_DEFAULT (RO on).
PCIe native atomics for the tail signal (PCI_ATOMIC_MODE=4). dispatch/combine
signal progress with a large number of per-channel×rank RDMA atomics. Mode 4 uses
the PCIe-native AtomicOp (hardware-offloaded, high throughput); mode 0 falls back to a
slower NIC RMW ext-atomic, so the README recommends mode 4.

The hazard (#661) and the fix (#674)

Used together, they create a data-ordering risk (see #661):

per PCIe spec, an RO posted write can be bypassed by a subsequent non-posted request;
an RO FAA (the tail-signal atomic) does not wait for the write-combining buffer to flush;
so the receiver's tail becomes visible before the payload data lands → the forward
warp polls the tail, reads the recv buffer, and gets stale / not-yet-landed data →
silent data corruption (critical in training).

#674 fixed this by switching the window to NCCL_WIN_STRICT_ORDERING, which maps to
NCCL_NET_MR_FLAG_FORCE_SO and disables IBV_ACCESS_RELAXED_ORDERING at MR registration.
That is correct, but heavy-handed: it serializes all RDMA writes to the window, not just
the signal path.

Cost of the blanket strict ordering

Measured on B300 + NCCL 2.32.3 + CUDA 13.0, EP16 (2 nodes × 8 GPU) hybrid,
num_sms=32, EP_OVERRIDE_RDMA_SL=162, token=4096:

NCCL_WIN_STRICT_ORDERING (current): dispatch SO 23 GB/s, 2745 us; combine SO ~28 GB/s
NCCL_WIN_DEFAULT: dispatch SO 130 GB/s, 469 us; combine SO 124 GB/s
v2.0.0 baseline (pre-#674, already DEFAULT): dispatch SO 135 GB/s, 452 us; combine SO 126 GB/s

Full token sweep with DEFAULT (1024→16384) gives dispatch SO 82/101/129/152/168 GB/s, on par
with v2.0.0; the strict build stays ~23 GB/s across the sweep. The regression is independent
of SM count — it is a memory-ordering/PCIe limit, not a compute one, so raising num_sms
does not help.

This PR

The #661 hazard is specific to the RO + PCIe-native-atomic (PCI_ATOMIC_MODE=4) combination.
When PCIe native atomics are not used (e.g. PCI_ATOMIC_MODE=0), relaxed ordering is both
faster and, in our testing, correct.

Rather than unconditionally revert #674 (which would re-expose PCI_ATOMIC_MODE=4 users to the
original corruption), this PR:

keeps NCCL_WIN_STRICT_ORDERING as the default (behavior unchanged);
lets operators opt into NCCL_WIN_DEFAULT via EP_RELAXED_WIN_ORDERING=1 on platforms
where RO is known to be safe (uses the existing get_env helper / EP_ env convention);
leaves the GIN-only RDMA-storage window (NCCL_WIN_GIN_ONLY) untouched.

const int main_win_flags = get_env<int>("EP_RELAXED_WIN_ORDERING", 0)
                         ? NCCL_WIN_DEFAULT : NCCL_WIN_STRICT_ORDERING;
NCCL_CHECK(ncclCommWindowRegister(comm, raw_window_ptr, symmetric_memory->num_sym_bytes,
                                  &window, main_win_flags));

Correctness check

With the relaxed window (EP_RELAXED_WIN_ORDERING=1, PCI_ATOMIC_MODE=0), rebuilt and ran the
full matrix without --skip-check: 144 cases (72 fp8 + 72 bf16) all pass, 0 assertion errors /
0 mismatches on B300 + NCCL 2.32.3.

Caveats

Do not set EP_RELAXED_WIN_ORDERING=1 together with PCI_ATOMIC_MODE=4 — that is exactly
the combination #661 flagged as unsafe.
We only validated relaxed ordering on B300 + NCCL 2.32.3 with PCI_ATOMIC_MODE=0; we cannot
guarantee it on every platform / root complex / BIOS, which is why the default stays strict and
this is strictly opt-in.
A more surgical long-term fix would keep RO on and add explicit fence/barrier logic on just the
signal path (so write-before-signal visibility holds even under RO), rather than serializing the
whole window. That is a larger change and out of scope here. Happy to follow up if maintainers
are interested.
The IBV_ACCESS_RELAXED_ORDERING linkage is from reading NCCL source; not independently
re-verified at the binary level.

Repro

HW: NVIDIA B300, 2 nodes × 8 GPU (EP16), hybrid mode (--allow-hybrid-mode 1)
SW: NCCL 2.32.3 (GIN backend), CUDA 13.0
Test: tests/ep/test_ep.py, experts=256, hidden=7168, topk=8, num_sms=32,
EP_OVERRIDE_RDMA_SL=162, token sweep

_{🔄 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/791 **Author:** [@sssssakuraaaaa](https://github.com/sssssakuraaaaa) **Created:** 10/23/2026 **Status:** 🔄 Open **Base:** `main` ← **Head:** `fix/relax-symmetric-window-ordering` --- ### 📝 Commits (1) - [`863328c`](https://github.com/deepseek-ai/DeepEP/commit/863328c45666a3bbd848d9b8690d94c8b9773f71) Make symmetric-window ordering configurable (EP_RELAXED_WIN_ORDERING) ### 📊 Changes **1 file changed** (+14 additions, -1 deletions) <details> <summary>View changed files</summary> 📝 `csrc/kernels/comm/context.cpp` (+14 -1) </details> ### 📄 Description ## TL;DR The main symmetric window is registered with `NCCL_WIN_STRICT_ORDERING` (`csrc/kernels/comm/context.cpp`). On hybrid inter-node EP this costs **~6x** dispatch/combine bandwidth. The strict flag was added in #674 to fix the ordering hazard in #661, but that hazard only bites when PCIe **native atomics** (`PCI_ATOMIC_MODE=4`) are combined with relaxed ordering. This PR keeps strict as the default and adds an opt-in (`EP_RELAXED_WIN_ORDERING=1`) to restore the fast `NCCL_WIN_DEFAULT` path on platforms where relaxed ordering is safe. ## Background: two optimizations that are each fine alone DeepEP's inter-node path leans on two independent hardware optimizations: 1. **PCIe Relaxed Ordering (RO) on the IB MR.** The large token payload goes over RDMA write; RO lets posted writes land out of order, saturates PCIe bandwidth and reduces head-of-line blocking. This is why the window was originally registered with `NCCL_WIN_DEFAULT` (RO on). 2. **PCIe native atomics for the tail signal (`PCI_ATOMIC_MODE=4`).** dispatch/combine signal progress with a large number of per-`channel`×`rank` RDMA atomics. Mode 4 uses the PCIe-native AtomicOp (hardware-offloaded, high throughput); mode 0 falls back to a slower NIC RMW ext-atomic, so the README recommends mode 4. ## The hazard (#661) and the fix (#674) Used **together**, they create a data-ordering risk (see #661): - per PCIe spec, an RO posted write can be bypassed by a subsequent non-posted request; - an RO FAA (the tail-signal atomic) does not wait for the write-combining buffer to flush; - so the receiver's **tail becomes visible before the payload data lands** → the forward warp polls the tail, reads the recv buffer, and gets **stale / not-yet-landed data** → silent data corruption (critical in training). #674 fixed this by switching the window to `NCCL_WIN_STRICT_ORDERING`, which maps to `NCCL_NET_MR_FLAG_FORCE_SO` and disables `IBV_ACCESS_RELAXED_ORDERING` at MR registration. That is correct, but heavy-handed: it serializes **all** RDMA writes to the window, not just the signal path. ## Cost of the blanket strict ordering Measured on **B300 + NCCL 2.32.3 + CUDA 13.0**, EP16 (2 nodes × 8 GPU) hybrid, `num_sms=32`, `EP_OVERRIDE_RDMA_SL=162`, token=4096: - `NCCL_WIN_STRICT_ORDERING` (current): dispatch SO **23 GB/s**, **2745 us**; combine SO ~28 GB/s - `NCCL_WIN_DEFAULT`: dispatch SO **130 GB/s**, **469 us**; combine SO **124 GB/s** - v2.0.0 baseline (pre-#674, already DEFAULT): dispatch SO 135 GB/s, 452 us; combine SO 126 GB/s Full token sweep with DEFAULT (1024→16384) gives dispatch SO 82/101/129/152/168 GB/s, on par with v2.0.0; the strict build stays ~23 GB/s across the sweep. The regression is independent of SM count — it is a memory-ordering/PCIe limit, not a compute one, so raising `num_sms` does not help. ## This PR The #661 hazard is specific to the RO + PCIe-native-atomic (`PCI_ATOMIC_MODE=4`) combination. When PCIe native atomics are **not** used (e.g. `PCI_ATOMIC_MODE=0`), relaxed ordering is both faster and, in our testing, correct. Rather than unconditionally revert #674 (which would re-expose `PCI_ATOMIC_MODE=4` users to the original corruption), this PR: - keeps `NCCL_WIN_STRICT_ORDERING` as the **default** (behavior unchanged); - lets operators opt into `NCCL_WIN_DEFAULT` via **`EP_RELAXED_WIN_ORDERING=1`** on platforms where RO is known to be safe (uses the existing `get_env` helper / `EP_` env convention); - leaves the GIN-only RDMA-storage window (`NCCL_WIN_GIN_ONLY`) untouched. ```cpp const int main_win_flags = get_env<int>("EP_RELAXED_WIN_ORDERING", 0) ? NCCL_WIN_DEFAULT : NCCL_WIN_STRICT_ORDERING; NCCL_CHECK(ncclCommWindowRegister(comm, raw_window_ptr, symmetric_memory->num_sym_bytes, &window, main_win_flags)); ``` ## Correctness check With the relaxed window (`EP_RELAXED_WIN_ORDERING=1`, `PCI_ATOMIC_MODE=0`), rebuilt and ran the full matrix without `--skip-check`: **144 cases (72 fp8 + 72 bf16) all pass, 0 assertion errors / 0 mismatches** on B300 + NCCL 2.32.3. ## Caveats - Do **not** set `EP_RELAXED_WIN_ORDERING=1` together with `PCI_ATOMIC_MODE=4` — that is exactly the combination #661 flagged as unsafe. - We only validated relaxed ordering on B300 + NCCL 2.32.3 with `PCI_ATOMIC_MODE=0`; we cannot guarantee it on every platform / root complex / BIOS, which is why the default stays strict and this is strictly opt-in. - A more surgical long-term fix would keep RO on and add explicit fence/barrier logic on just the signal path (so write-before-signal visibility holds even under RO), rather than serializing the whole window. That is a larger change and out of scope here. Happy to follow up if maintainers are interested. - The `IBV_ACCESS_RELAXED_ORDERING` linkage is from reading NCCL source; not independently re-verified at the binary level. ## Repro - HW: NVIDIA B300, 2 nodes × 8 GPU (EP16), hybrid mode (`--allow-hybrid-mode 1`) - SW: NCCL 2.32.3 (GIN backend), CUDA 13.0 - Test: `tests/ep/test_ep.py`, experts=256, hidden=7168, topk=8, num_sms=32, `EP_OVERRIDE_RDMA_SL=162`, token sweep --- <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:06:22 +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#788

No description provided.

Rows
Columns