Skip to content

Sync with upstream deepseek-ai/DeepEP - #1

Open
pazqo wants to merge 52 commits into
neuralmagic:mainfrom
pazqo:sync-upstream-main
Open

Sync with upstream deepseek-ai/DeepEP#1
pazqo wants to merge 52 commits into
neuralmagic:mainfrom
pazqo:sync-upstream-main

Conversation

@pazqo

@pazqo pazqo commented Aug 26, 2026

Copy link
Copy Markdown

Summary

  • Merges latest changes from deepseek-ai/DeepEP main branch into the neuralmagic fork
  • Resolves pyproject.toml conflict by keeping neuralmagic's [project]/[build-system] config and adopting upstream's lint configuration

Test plan

  • Verify build still works with the existing pyproject.toml settings
  • Confirm no regressions from upstream changes

🤖 Generated with Claude Code

fzyzcjy and others added 30 commits September 1, 2025 17:06
@fzyzcjy
Update internode_ll.cu (deepseek-ai#376)
c78b9ed
@fzyzcjy
Speed up dispatch send by refining loop unrolling (deepseek-ai#385)
6a5d323 
* nits

* hack unrolled warp copy

* Revert "nits"

This reverts commit 3e1b28d.
@tianchongchong
Update intranode.cu (deepseek-ai#381)
c18eabd 
Each thread is responsible for one target rank
@LyricZhao
Fix bench warmup counts
174c209
@sphish
Fix tma mbarrier (deepseek-ai#399)
0970c4c 
* Fix mbarrier

* Remove redundant store
@LyricZhao
Update version number
ee5bd17
@sphish
Fix combine tma mbarrier
24324cc
@LyricZhao
Add self.group
130122a
@LyricZhao
Update version number
3e2c5d8
@sphish
Add pressure test mode for internode test (deepseek-ai#400)
ef70b83 
* Suppress kineto output

* Add pressure test mode

* Add `x_pure_rand_e4m3` test

* Add more results into hash value
@sphish
Fix bar.sync
1d3963d
@LyricZhao
Add some comments
e9d053c
@polarstormx
support topk10 in low latency kernel (deepseek-ai#403)
9af0e0d
@LyricZhao
Canonicalize TMA usages (deepseek-ai#410)
2012e31 
* Remove redundant TMA flushes

* Less barrier initialization overhead

* Simplify `elect_one_sync`

* Use `elect_one_sync` instead of lanes

* Minor fix

* Polish testing prints

* Refactor for internode kernels

* Better performance
@sphish
Fix hidden_size % 128 != 0 in intranode kernels (deepseek-ai#413)
abba6ad 
* Fix hidden_size % 128 != 0

* Add `align_down()` function

* Use the full warp to wait TMA store

* Support arbitrary hidden sizes in fp8 cast

* lint
@sphish
Support EP48/96 for internode kernels and update some config. (deepse…
c939644 
…ek-ai#421)
@yuantailing @yifeizhang-c
Make dtype of topk_idx configurable (deepseek-ai#422)
da6ca24 
Co-authored-by: Yifei Zhang <219273404+yifeizhang-c@users.noreply.github.com>
@jammycc @Gaojiaqi @CavalryXD
[Feature] Add Elasticity Support to DeepEP for Fault-Tolerant EP Infe…
c475ec6 
…rence (deepseek-ai#370)

* support EP elastic shrink

* Add test script for EP shrink failover

* add coll buffer and allgather api

* remove allgather api, merge barrier and clean kernels, rename elastic to shrink

* fix bug

* refine code, add shrink test script

* refine code

* refine code

* fix bug

* add/remove blank lines

---------

Co-authored-by: Jiaqi Gao <jiaqi.g@alibaba-inc.com>
Co-authored-by: Yichi Xu <xuyichi.xyc@alibaba-inc.com>
@sphish
Check QP_DEPTH for low latency kernels. (deepseek-ai#430)
d981409 
* Check `QP_DEPTH` for low latency kernels.

* Revert the default `QP_DEPTH`
@sphish
Support EP24 for internode kernels (deepseek-ai#432)
e02e4d2 
* Support EP24 for internode kernels.

* Skip check for round_scale test
@sphish
Add CPP and Python linter. (deepseek-ai#435)
79bbb54 
* Add format scripts.

* Formated

* Fix ruff check.
@sphish
Add format workflow (deepseek-ai#452)
12a0cc6 
* Add format workflow

* Remove debug code
@sphish
Update README
de3a15d
@sphish
Make format.sh runnable on macOS
3f4cf8d
@sphish
In format.sh, replace sed with perl.
987f8f0
@LyricZhao
Fix OOB (deepseek-ai#454)
4623c67 
* Fix OOB

* Add more assertions

* More checks on channels
@varun-sundar-rabindranath
support hidden-dim 3072 (deepseek-ai#458)
73b6ea4 
Signed-off-by: Varun Sundar Rabindranath <vsundarr@redhat.com>
Co-authored-by: Varun Sundar Rabindranath <vsundarr@redhat.com>
@LyricZhao
Update README (deepseek-ai#463)
2be0d4f
@sphish
Compatible with older versions of Python and Torch (deepseek-ai#469)
bfded34 
* Compatible with older versions of python and torch

* Format
@fzyzcjy
Support cuMem API in cross process shared memory management (deepseek…
9249c25 
…-ai#217)

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* more

* add flag

* add test

* fix

* more

* apply
yifeizhang-c and others added 22 commits November 5, 2025 15:33
@yifeizhang-c
Enable CUDA Graph for internode dispatch (deepseek-ai#438)
92fe2de
@dongwang4096
Update README.me for Hybrid-EP Experimental Branche (deepseek-ai#502)
1300811
@wangfakang
Update README for an experimental optimization branch (deepseek-ai#512)
b57e5e2 
Signed-off-by: wangfakang <fakangwang@gmail.com>
@MaoZiming
Add UCCL-EP to Community Forks section in README (deepseek-ai#529)
0a08009
@salmanmkc
Upgrade GitHub Actions for Node 24 compatibility (deepseek-ai#528)
3fcf25e 
Signed-off-by: Salman Muin Kayser Chishti <13schishti@gmail.com>
@yurekami Contributor @claude
Fix typo: NVHSMEM -> NVSHMEM in installation docs (deepseek-ai#544)
309785a 
🤖 Generated with [Claude Code](https://claude.com/claude-code)

Co-authored-by: Contributor <contributor@example.com>
Co-authored-by: Claude Opus 4.5 <noreply@anthropic.com>
@seth-howell
Nvshmem 3 5 runtime fixes (deepseek-ai#564)
29d31c0 
* DeepEP/csrc: Force CUDA RDC when NVSHMEM is used.

When DeepEP is against a source build of NVSHMEM starting
with NVSHMEM 3.3.20, multiple instances of nvshmemi_ibgda_device_state_d
are detected. One from NVSHMEM, another from the kernels.

Signed-off-by: Seth Howell <sethh@nvidia.com>

* DeepEP/csrc: Update RC selection for NVSHMEM 3.5.

The locations of QPs within the device state struct changed
with the addition of QP-specific APIs.

This change updates DeepEP to work with the new layout.

The ibgda_get_rc function is also templated so that main remains
backwards compatible with previous versions of NVSHMEM.

Tested with both internode and low_latency kernels on 4 H100 nodes.

Signed-off-by: Seth Howell <sethh@nvidia.com>

* Check Format

Signed-off-by: Seth Howell <sethh@nvidia.com>

---------

Signed-off-by: Seth Howell <sethh@nvidia.com>
@isytwu
Update README for mori-EP branch and mori community fork (deepseek-ai…
567632d 
…#578)
@LyricZhao @sphish
@alpha-baby fujianhao.fjh
[Public release 26/04] Introducing EPv2: faster EP, and Engram/PP/CP …
b306af0 
…supports (deepseek-ai#605)

* Merge with private repo

* Fix compilation

* Add readme

* Update docs

* Fix submodule

* Fix cu128 compilation

* enhance: add env:EP_NIC_NAME to config nic name (deepseek-ai#610)

* enhance: add env:EP_NIC_NAME to config nic name

* enhance: add env:EP_NIC_NAME to config nic name

* enhance: add env:EP_NIC_NAME to config nic name

* enhance: add env:EP_NIC_NAME to config nic name

* enhance: add env:EP_NIC_NAME to config nic name

---------

Co-authored-by: fujianhao.fjh <fujianhao.fjh@alipay.com>

---------

Co-authored-by: Shangyan Zhou <sy.zhou@deepseek.com>
Co-authored-by: AlphaBaby <fujianhao1997@qq.com>
Co-authored-by: fujianhao.fjh <fujianhao.fjh@alipay.com>
@sphish
Merge some updates (deepseek-ai#637)
a422b5a 
What's New
Bug Fixes

Fix a rare numerical error in multi-node hybrid combine
Features

Add CPU buffer support for Engram
Remove alignment requirement for all-gather
Improvements

Reuse torch NCCL comm by default to reduce GPU memory usage
Improve combine/combine-epilogue performance for certain cases
@MARD1NO
Fix internode dispatch args (deepseek-ai#641)
a1b3af3 
* fix missing num_worst_tokens arg

* fix
@hirschsn
Match "libnccl" instead of "nccl" in the library check (deepseek-ai#640)
bb29c6b 
This commit changes the library matching, so that it avoids accidentally
matching "/dev/shm/nccl-XXXXXX" shared memory mappings that might be present in
/proc/self/map.
@DanielDanyang
fix V2 initialization for pure single-node intra EP (deepseek-ai#630)
5616959
@MayDomine
fix(setup): resolve real NVSHMEM/NCCL SO names for pip-wheel installs (…
723716f 
…deepseek-ai#627)

When NVSHMEM and NCCL are installed via the official NVIDIA pip wheels
(`nvidia-nvshmem-cu12`, `nvidia-nccl-cu12`), the wheel only ships the
SONAME-suffixed library (e.g. `libnvshmem_host.so.3`, `libnccl.so.2`)
and omits the unversioned `libnvshmem_host.so` / `libnccl.so` symlink
that a system tarball install would create.

The build then fails at link time with:
  /usr/bin/ld: cannot find -l:libnvshmem_host.so: No such file or
directory
because `-l:NAME` is exact-name matching and refuses to consider
`libnvshmem_host.so.3` when only `libnvshmem_host.so` is requested.

`get_nvshmem_host_lib_name()` already existed for this exact reason,
but was never wired into `extra_link_args` — the link string still
hard-coded `-l:libnvshmem_host.so`. NCCL had the same hard-coded
`-l:libnccl.so` and the same wheel-vs-tarball mismatch, just no helper
to back it up.

This change:
  * generalises the helper into `_find_versioned_so(base_dir, prefix)`,
    preferring the unversioned symlink (tarball install) and falling
    back to the SONAME file (pip wheel install);
  * adds `get_nccl_lib_name` mirroring the NVSHMEM helper;
  * uses both helpers when assembling `extra_link_args` so the
    `-l:NAME` flag carries the real on-disk filename.

Validated by building a wheel inside a container that has only the
pip-wheel installs of `nvidia-nvshmem-cu12==3.6.5` and
`nvidia-nccl-cu13==2.30.4`. After the patch the link succeeds and
`ldd deep_ep/_C*.so` resolves to the wheel-shipped SONAME files.

No behaviour change for tarball installs that already provide the
unversioned symlink — the helper picks it up first.
@RayWang96
Add fence.proxy.async.shared::cta between mbarrier wait and TMA load. (…
d4f41e4 
…deepseek-ai#642)

* Add fence.proxy.async.shared::cta between mbarrier wait and TMA load.

* Move fence before "mbarrier_arrive"
@Eliasthunderdog @ZhenghangRen
add Experimental branch intro (deepseek-ai#664)
af9a040 
Co-authored-by: ZhenghangRen <zhenghangr@nvidia.com>
@sphish @LyricZhao
@tianr22 @interestingLSY @skip2004
Merge some updates (deepseek-ai#674)
099d5f2 
Co-authored-by: Chenggang Zhao <chenggangz@deepseek.com>
Co-authored-by: Rui Tian <tianr22@deepseek.com>
Co-authored-by: Shengyu Liu <shengyuliu@deepseek.com>
Co-authored-by: Yi Qian <qianyi@deepseek.com>
@LyricZhao
Fix SM90 compatibility (deepseek-ai#178)
60d4403
@LyricZhao @kgioioso
NCCL Device API compat: use runtime version for ncclDevCommCreate (de…
dd758ca 
…epseek-ai#688)

Store the runtime-sized device communicator behind jit::NoRefPtr so JIT launches pass its backing storage directly without host-side dereference.

Co-authored-by: Katie Gioioso <kgioioso@nvidia.com>
@EvadLee @sphish
fix: add system-scope release before the GIN barrier when scaleup spa…
11d1cab 
…ns NVLink and RDMA (deepseek-ai#715)

* fix: Add fence.acq_rel.sys before the GIN barrier for mixed-fabric scaleup

* fix: fence mixed-fabric writes in GIN barrier

---------

Co-authored-by: Shangyan Zhou <sy.zhou@deepseek.com>
@sphish
Add comments
01dc3aa
@pazqo @claude
Merge remote-tracking branch 'upstream/main'
b65fc3f 
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
@Gregory-Pereira Gregory-Pereira mentioned this pull request Aug 26, 2026
Merged
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.