[PR #677] Fix the CPU-assisted (GDRCopy) IBGDA path in the legacy kernels #697

New issue

Open

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

bsakaguchi commented

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

Owner

📋 Pull Request Information

Original PR: https://github.com/deepseek-ai/DeepEP/pull/677
Author: @tommy85
Created: 7/8/2026
Status: 🔄 Open

Base: main ← Head: fix-cpu-assisted-ibgda

📝 Commits (2)

91848ed Fix the CPU-assisted (GDRCopy) IBGDA path in the legacy kernels
5c4f2b8 Read ready_head in the quiet drain so it survives the async submit change

📊 Changes

1 file changed (+24 additions, -6 deletions)

View changed files

📝 csrc/kernels/legacy/ibgda_device.cuh (+24 -6)

📄 Description

docs/nvshmem.md (§ "Install GDRCopy and load the gdrdrv kernel module") documents the CPU-assisted IBGDA mode as the supported alternative "when modifying the driver regkeys is not an option". In practice that mode (NVSHMEM_IBGDA_NIC_HANDLER=cpu) has never worked with DeepEP: the async post-send branch in ibgda_submit_requests (csrc/kernels/legacy/ibgda_device.cuh) has two independent defects, either of which is fatal.

Bug 1 — submissions are invisible to the CPU proxy (first dispatch hangs).
The async branch CASes on qp->tx_wq.prod_idx. NVSHMEM's host code wires that pointer to a device-side management counter (ibgda_setup_gpu_state: rc_h[i].tx_wq.prod_idx = base_mvars_d_addr + prod_idx_offset), but the proxy thread polls a separate producer-slot array (nvshmemt_ibgda_progress reads qp_shared_object.prod_idx_mobject). The value DeepEP publishes is therefore never observed, no doorbell is ever rung, and the very first dispatch spins forever waiting for data (surfacing as the NCCL watchdog SIGABRT after 600 s).

Bug 2 — teardown deadlock on shared QPs.
The async branch also skips the ready_head CAS queue. NVSHMEM's own submit path orders every submission through mvars->tx_wq.ready_head, and DeepEP shares the same QP management memory with NVSHMEM. After a DeepEP workload, resv_head is far ahead while ready_head is still 0; the first NVSHMEM-side submission on a shared QP — the nvshmem_barrier_all() inside Buffer::destroy() — reserves a slot and then spins forever on atomicCAS(ready_head==0, expect resv_head). That spin has no timeout, so the process hangs in cuStreamSynchronize with the GPU at 100%.

Fix. Always order submissions through ready_head (matching NVSHMEM's protocol), and in the async mode publish the new producer index with a system-scope atomicMax on qp->tx_wq.bf — in the CPU-assisted modes NVSHMEM wires bf to the proxy-polled producer slot (see ibgda_get_device_qp), so this mirrors NVSHMEM's own ibgda_proxy_post_send exactly and needs no NVSHMEM-side change. System scope keeps the publish visible to the CPU observer also for NVSHMEM_IBGDA_NIC_HANDLER=cpu_host_memory, where the slot lives in host memory.

Validation. On our rig (NVSHMEM 3.4.5, 2 nodes, 4 RC QPs per rank, NVSHMEM_IBGDA_NIC_HANDLER=cpu with GDRCopy) tests/test_low_latency.py previously hung at the first dispatch; with only the pointer rewired it then deadlocked in Buffer::destroy(). With this patch the full run — benchmark plus teardown (barrier → free → finalize) — completes and exits 0, with benchmark numbers unchanged. The defect is pure protocol logic, independent of the NIC; on ConnectX it should reproduce by simply exporting NVSHMEM_IBGDA_NIC_HANDLER=cpu (NVSHMEM built with GDRCopy) and running tests/test_low_latency.py.

The GPU-direct doorbell path (NVSHMEM_IBGDA_NIC_HANDLER=gpu, the default) is unaffected: its submit branch already ordered through ready_head and does not take the async path.

_{🔄 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/677 **Author:** [@tommy85](https://github.com/tommy85) **Created:** 7/8/2026 **Status:** 🔄 Open **Base:** `main` ← **Head:** `fix-cpu-assisted-ibgda` --- ### 📝 Commits (2) - [`91848ed`](https://github.com/deepseek-ai/DeepEP/commit/91848ed26dcda97aae54bddbf3663257e678a388) Fix the CPU-assisted (GDRCopy) IBGDA path in the legacy kernels - [`5c4f2b8`](https://github.com/deepseek-ai/DeepEP/commit/5c4f2b8f5adcdbe57008c400658b119775b480c9) Read ready_head in the quiet drain so it survives the async submit change ### 📊 Changes **1 file changed** (+24 additions, -6 deletions) <details> <summary>View changed files</summary> 📝 `csrc/kernels/legacy/ibgda_device.cuh` (+24 -6) </details> ### 📄 Description [docs/nvshmem.md](https://github.com/deepseek-ai/DeepEP/blob/main/docs/nvshmem.md) (§ "Install GDRCopy and load the gdrdrv kernel module") documents the CPU-assisted IBGDA mode as the supported alternative "when modifying the driver regkeys is not an option". In practice that mode (`NVSHMEM_IBGDA_NIC_HANDLER=cpu`) has never worked with DeepEP: the async post-send branch in `ibgda_submit_requests` (csrc/kernels/legacy/ibgda_device.cuh) has two independent defects, either of which is fatal. **Bug 1 — submissions are invisible to the CPU proxy (first dispatch hangs).** The async branch CASes on `qp->tx_wq.prod_idx`. NVSHMEM's host code wires that pointer to a device-side management counter (`ibgda_setup_gpu_state`: `rc_h[i].tx_wq.prod_idx = base_mvars_d_addr + prod_idx_offset`), but the proxy thread polls a separate producer-slot array (`nvshmemt_ibgda_progress` reads `qp_shared_object.prod_idx_mobject`). The value DeepEP publishes is therefore never observed, no doorbell is ever rung, and the very first dispatch spins forever waiting for data (surfacing as the NCCL watchdog SIGABRT after 600 s). **Bug 2 — teardown deadlock on shared QPs.** The async branch also skips the `ready_head` CAS queue. NVSHMEM's own submit path orders every submission through `mvars->tx_wq.ready_head`, and DeepEP shares the same QP management memory with NVSHMEM. After a DeepEP workload, `resv_head` is far ahead while `ready_head` is still 0; the first NVSHMEM-side submission on a shared QP — the `nvshmem_barrier_all()` inside `Buffer::destroy()` — reserves a slot and then spins forever on `atomicCAS(ready_head==0, expect resv_head)`. That spin has no timeout, so the process hangs in `cuStreamSynchronize` with the GPU at 100%. **Fix.** Always order submissions through `ready_head` (matching NVSHMEM's protocol), and in the async mode publish the new producer index with a system-scope `atomicMax` on `qp->tx_wq.bf` — in the CPU-assisted modes NVSHMEM wires `bf` to the proxy-polled producer slot (see `ibgda_get_device_qp`), so this mirrors NVSHMEM's own `ibgda_proxy_post_send` exactly and needs no NVSHMEM-side change. System scope keeps the publish visible to the CPU observer also for `NVSHMEM_IBGDA_NIC_HANDLER=cpu_host_memory`, where the slot lives in host memory. **Validation.** On our rig (NVSHMEM 3.4.5, 2 nodes, 4 RC QPs per rank, `NVSHMEM_IBGDA_NIC_HANDLER=cpu` with GDRCopy) `tests/test_low_latency.py` previously hung at the first dispatch; with only the pointer rewired it then deadlocked in `Buffer::destroy()`. With this patch the full run — benchmark plus teardown (barrier → free → finalize) — completes and exits 0, with benchmark numbers unchanged. The defect is pure protocol logic, independent of the NIC; on ConnectX it should reproduce by simply exporting `NVSHMEM_IBGDA_NIC_HANDLER=cpu` (NVSHMEM built with GDRCopy) and running `tests/test_low_latency.py`. The GPU-direct doorbell path (`NVSHMEM_IBGDA_NIC_HANDLER=gpu`, the default) is unaffected: its submit branch already ordered through `ready_head` and does not take the async path. --- <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:22 +00:00

bsakaguchi referenced this issue

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

[PR #697] fix(test): remove dead num_worst_tokens test from internode file #714