[IlyasMoutawwakil]

Related issue

None. This supports the upcoming deepep expert-parallel strategy in transformers, discussed internally.

What does this PR do?

Adds DeepEP 2.5.0 as the Hub kernel kernels-community/deep-ep, preserving its expert-parallel dispatch/combine API. Only the small host extension is prebuilt; DeepJIT compiles communication kernels at runtime.

Review guide

The port is now reproduced by kernel-port, laid out like the other ported kernels (e.g. einops): deep-ep/port/ holds the inputs (recipe and overlays) and deep-ep/src/ the output, generated from them and never edited; the build flake sits at the kernel root and builds ./src. Start with the recipe and handwritten overlays: CI verifies that they reproduce src/ byte-for-byte, including provenance. Imported files still need attention when assessing upstream behavior, but the local porting changes are collected in these review surfaces:

1. Port recipe: review the source pins, retained files, guarded upstream edits, import/test rewrites, and generated build configuration. Each upstream replacement checks its expected occurrence count. Pins are DeepEP def865146e95aa9ea223e277c0a8a4c6b81e5ac7 and its DeepJIT dependency cba802df081985141eeee36803685775c183a047. NCCL is not vendored (see below).
2. Native bindings and ownership: torch_binding.cpp, deep_ep_holders.hpp, and compat/ replace pybind11 registration with Torch ops/custom classes and small c10::IValue shims. Focus on signatures, object/tensor lifetimes, deferred hooks, stream behavior, and generated operator namespaces. The recipe guards the corresponding small edits to upstream C++.
3. Python adapters and runtime setup: _C.py preserves the upstream Python API; _runtime.py defers dependency checks/JIT initialization until buffer construction; _nccl_symbols.py exposes PyTorch's already-loaded NCCL symbols before loading the extension. The extension neither links nor ships a separate NCCL library. Runtime JIT uses headers matching the runtime NCCL installation.
4. Tests and reproduction checks: upstream test adaptations are in the recipe; local CI tests are overlays. Tests load through kernels.get_kernel. The shared checker now supplies pinned vendor checkouts to the runner; its regression tests cover vendor checkouts, checkout reuse, and stale/extra generated files.

The remaining overlays hold the Hub card.

NCCL comes from kernel-builder (huggingface/kernels#951): every CUDA kernel builds against the NCCL of its Torch. DeepEP 2.5's host code uses GIN and EFA-GDA APIs from NCCL 2.31/2.32, while kernel-builder's Torch ships 2.30.7, so flake.nix passes an overlay that builds NCCL 2.32.3 with mkNccl (from the nvidia-nccl-cu13 wheels) and gives it to Torch. Those wheel headers are byte-identical to the 69 headers the recipe used to assemble from NCCL's Git sources, so the extension compiles the same code. The overlay can go once kernel-builder's Torch ships NCCL >= 2.32.

Regeneration instructions pin the experimental kernel-port runner and explain how to reproduce/check the tree. Future port changes should update the recipe/overlays and regenerate the output together. Normal builds use the checked-in generated files.

This PR retains version-specific Torch builds and the original Torch ops/custom-class bindings. There is no Torch stable-ABI conversion or DLPack adapter in the recipe. Adding the recipe leaves implementation sources byte-for-byte unchanged and preserves the existing build settings; the generated manifest's formatting/order changes. Stable ABI remains separate follow-up work requiring performance validation. DeepEP's runtime JIT also means the ahead-of-time GPU build optimization discussed in #999 does not apply here.

Runtime requirements

• Hopper/Blackwell GPUs with NVLink; RDMA for multi-node operation.
• NCCL >= 2.32.3 loaded by PyTorch, with matching headers available for runtime JIT. A PyTorch installation with an older NCCL needs an appropriate NCCL upgrade; see CARD.md.
• CUDA toolkit >= 13.1 (nvcc) and a C++20 host compiler at runtime. The prebuilt host extension has a CUDA 13.0 build minimum.

Requirements are checked at first buffer construction with actionable errors. Compiled GPU kernels are cached by DeepJIT.

Validation

• NCCL from kernel-builder (latest commit): the recipe edit and the matching generated files (build.toml, deleted headers, port-provenance.json recomputed with kernel-port's hashing) were made by hand, without nix; the 125 CI-script tests pass. The build and the reproduction check run in CI on this commit; the GPU results below predate it.
• Reproduction: the shared CI checker regenerates both DeepEP and the existing einops port byte-for-byte, including provenance. GitHub reproduction check passes; all 127 CI-script tests pass locally, and the CI-script test job passes on GitHub.
• Generated package: config validation, CMake build/install, all six CI smoke tests, and a two-GPU deferred dispatch/combine case pass after regeneration, using B200 GPUs, Torch 2.14 + CUDA 13.0, NCCL 2.32.3, and CUDA 13.2 for runtime JIT. The requested build-and-stage workflow's build also passed.
• Earlier port validation: all 36 two-GPU EP cases passed, covering FP8/BF16, expanded layout, deferred epilogues, cached dispatch, and previous events. Additional upstream tests passed against both upstream and this port: buffer allocation, weight prefetch, gradient reduction, barrier/gate, bucket all-gather, PP, and Engram. CI smoke tests also verified the expected buffer-construction error with NCCL 2.30.7.
• Limits: multi-node/RDMA has not been tested. The recipe-only change does not establish new performance results; it preserves the existing implementation and build settings.

Checklist

• This PR is linked to an issue that was discussed and approved
• I have tested these changes locally
• New/changed functionality has test coverage
• LLM disclosure:
  • I did not use an LLM to create this PR.
  • I used an LLM for assistance while creating this PR.
  • This PR was mostly or completely generated by an LLM.

[github-actions]

This PR changes source code for the following kernel(s): deep-ep.

If you have the kernel-bot build permissions, then trigger /kernel-bot build deep-ep yourself. Otherwise, a maintainer will review the PR and will take an appropriate action. Thanks for your patience.

[IlyasMoutawwakil]

/kernel-bot build-and-stage deep-ep

[github-actions]

Build request processed.

Command: /kernel-bot build-and-stage deep-ep
Mode: build and stage
Target branch: pr-1231
PR head SHA: 76ce3dec5300798352f4c1881ca7bb3381e82f45
Workflows: build.yaml, build-mac.yaml, build-windows.yaml

Dispatched (1):

• deep-ep (build.yaml): https://github.com/huggingface/kernels-community/actions/runs/38708572733

Hub uploads:

• Hub repo: https://huggingface.co/kernels/kernels-staging/deep-ep

[sayakpaul]

I think we can remove the figures?

[IlyasMoutawwakil]

is it common practice to remove the version ? i think it's useful since the hub versioning starts at 0/1

[sayakpaul]

Some thoughts:

• Possible to use Stable ABI here?
• Possible to follow the strategy of #999 to speed up the builds?
• Possible to use kernel-port tool to see how far we can get with it? Cc: @drbh

I will also let @danieldk comment here.

[IlyasMoutawwakil]

@sayakpaul

• A stable abi seems feasible, my agent got one built with torch 2.14 and working with 2.15, but we need to decide for a minimum version bcz the implementation depends on what that version offers in terms of api. it also comes with changes that diverge from upstream.
• doesn't seem to apply since DeepEP jit compiles at runtime.
• also feasible and very nice feature ! trying it out rn !

[IlyasMoutawwakil]

hmm so yeah i had this feeling that stable-abi-ing might cause perf regression due to the stable abi apis being meant for versatility and not necessarily dispatching to the maximum perf route. the initial prototypes my agent came up with pass all tests but don't hold perf parity. will investigate more but i'd be wary of any stable-abi rewrite we did in the past without comparing performance. @sayakpaul @vasqu

DeepEP stable ABI performance experiment

Decision: retain the original bindings on the PR. Neither stable ABI prototype is ready to ship under the requirement to preserve performance.

The functional minimum reached is Torch 2.10. That does not establish performance parity. The 2.10 prototype passes the stable ABI checker with Python limited ABI 3.9, using the host manylinux_2_39 check; the host toolchain still prevents establishing manylinux_2_28 compatibility. Its source passed eight CI tests and 36 two-GPU deferred EP cases on actual Torch 2.10, plus an external-storage/view ownership probe.

Measurements used Torch 2.14 + CUDA 13.0 and NCCL 2.32.3 on two B200 GPUs (physical GPUs 4 and 5), with CUDA 13.2 runtime JIT. Each paired run loads the original and one candidate in both rank processes, uses the same inputs, group, GPU placement and CPU affinity, warms both implementations, and alternates measurement order across ten repetitions of 100 calls. The stable builds use -O3. Inference mode is enabled. Measurements include current_stream_wait; CUDA event elapsed time therefore includes host launch gaps and rank coordination. Reported values are the maximum across ranks of each rank’s median per-call elapsed time.

Representative workload: 4096 tokens/rank, hidden size 7168, 256 experts, top-k 6. These are local two-GPU measurements, not a multi-node or production-scale characterization.

Prototype	Allocation on comm stream	Operation	Original us	Candidate us	Change
2.10	False	dispatch	305.5	396.6	+29.8%
2.10	False	cached_dispatch	275.5	338.5	+22.9%
2.10	False	combine	192.8	231.3	+19.9%
2.10	True	dispatch	307.3	397.9	+29.5%
2.10	True	cached_dispatch	279.5	338.7	+21.2%
2.10	True	combine	192.6	234.2	+21.6%
2.14	False	dispatch	302.4	460.5	+52.3%
2.14	False	cached_dispatch	268.2	390.3	+45.5%
2.14	False	combine	192.3	297.3	+54.6%
2.14	True	dispatch	304.6	461.0	+51.3%
2.14	True	cached_dispatch	269.7	403.6	+49.6%
2.14	True	combine	191.1	300.8	+57.4%

The 2.10 prototype regresses about 20–30% on this workload; the 2.14 prototype also regresses. Raising the minimum to 2.14 does not solve the issue in these implementations. The host-side paths show additional time, while separately profiled GPU epilogues remain close. Communication kernel durations can include rank waiting and must not be read as pure GPU compute time. These results do not prove stable ABI itself necessarily entails a regression; they reject these prototypes as production conversions.

[IlyasMoutawwakil]

/kernel-bot build-and-stage deep-ep

[github-actions]

Build request processed.

Command: /kernel-bot build-and-stage deep-ep
Mode: build and stage
Target branch: pr-1231
PR head SHA: 15af174b9b5efae518b5175a71dec3351ebedf8c
Workflows: build.yaml, build-mac.yaml, build-windows.yaml

Dispatched (1):

• deep-ep (build.yaml): https://github.com/huggingface/kernels-community/actions/runs/38884855792

Hub uploads:

• Hub repo: https://huggingface.co/kernels/kernels-staging/deep-ep

[IlyasMoutawwakil]

kernel-port used !

[drbh]

I need to take a deeper look at the overall kernel implementation but I'm not sure this is using the kernel-porting too in its intended way...

for example kernels that are ported with the porting tool should roughly have the shape

└── kernel (deep-ep)
    ├── port
    │   ├── overlay
    │   ├──  ...
    │   └── port.kdl
    └── src
        ├── build.toml
        ├──  ...
        └── torch-ext/deep-ep

where the port directory contains the inputs, and the src contains the output (completely generated from the inputs and not edited)

**the einops kernel is a good example of this

based on the structure of this kernel its doesnt seem like the output is generated from the porting tool.

since the kernel seems pretty complex and likely needs non trivial changes to be built with the kernel builder - is probably not the best candidate for the porting tool, so unless this can be restructured to align with the tooling i'd recommend that we remove the port from this pr

[IlyasMoutawwakil]

i see, thanks @drbh, i only pointed the agent at the tool without an example, the tool probably doesn't explain how the codebase should be structured and that's why it treated it as just another folder within the kernel codebase rather than outside of it. will restructure it and see if nccl can be used fom torch as @danieldk said on slack🙏

[IlyasMoutawwakil]

checked: deepep 2.5's host code uses GIN symbols that only exist from nccl 2.31/2.32 (custom-stride/resource-sharing connections, gin-only windows, VA signal ops, EFA-GDA transport), and kernel-builder's torch pins nccl 2.30.4, so we can't build against it yet. if kernel-builder bumps nccl to 2.32 we can drop the vendored headers (runtime already uses torch's nccl) @danieldk

[vasqu]

Any updates here? Wasnt there some plan to use an alternative to nccl in general, i vaguely remember something

[IlyasMoutawwakil]

there was a pr that simplifies nccl handling huggingface/kernels#951 using it

[IlyasMoutawwakil]

/kernel-bot build deep-ep

[IlyasMoutawwakil]

/kernel-bot build-and-stage deep-ep

[github-actions]

Build request processed.

Command: /kernel-bot build deep-ep
Mode: build only
Target branch: pr-1231
PR head SHA: bcdbab5a726aadabc87e57c608a9433799019d06
Workflows: build.yaml, build-mac.yaml, build-windows.yaml

Dispatched (1):

• deep-ep (build.yaml): https://github.com/huggingface/kernels-community/actions/runs/41622125068

[github-actions]

Build request processed.

Command: /kernel-bot build-and-stage deep-ep
Mode: build and stage
Target branch: pr-1231
PR head SHA: bcdbab5a726aadabc87e57c608a9433799019d06
Workflows: build.yaml, build-mac.yaml, build-windows.yaml

Dispatched (1):

• deep-ep (build.yaml): https://github.com/huggingface/kernels-community/actions/runs/41622158305

[IlyasMoutawwakil]

/kernel-bot build-and-stage deep-ep

[github-actions]

Build request processed.

Command: /kernel-bot build-and-stage deep-ep
Mode: build and stage
Target branch: pr-1231
PR head SHA: d3c78f8f48ac31c7eadfda7b8f9bcdf021fa0ccf
Workflows: build.yaml, build-mac.yaml, build-windows.yaml

Dispatched (1):

• deep-ep (build.yaml): https://github.com/huggingface/kernels-community/actions/runs/41622659150

Hub uploads:

• Hub repo: https://huggingface.co/kernels/kernels-staging/deep-ep

[drbh]

we can probably remove the readme since it should be generated from the the CARD.md

[IlyasMoutawwakil]

removed, the useful bits (regenerate/check commands, what's pruned) are in the deep-ep section of AGENTS.md now

[drbh]

nice to see the port being used!

long term we want to avoid explicit string replacement as much as possible, this is because the replace ops are a bit brittle and will require a bunch of changes if the upstream moves.

however, in the short term if you are open to maintaining these replace ops, I think they are acceptable and move us in the right direction.

**in general its already a benefit to use the port with replace over manually porting the files since all of the string changes are contained within this single file and are reproducible (instead of being spread over many files)

note: the porting tool is going to be updated (in this pr huggingface/kernels#946 and some upcoming ones) which should help make it easier to write ports without large replace ops. once those land, maybe we can revisit these ops and make them a bit more robust to upstream changes.

[IlyasMoutawwakil]

yeah happy to maintain them for now, will revisit once kernels#946 and the follow-ups land

[drbh]

we can probably prune this file in port.kdl since it comes from upstream but doesnt make sense to copy into the kernel here

[IlyasMoutawwakil]

pruned, see below

[drbh]

maybe we can prune upstream files the kernel doesn't use: lint/CI config (.clang-format, .editorconfig, .github/workflows/format.yml, format.sh, requirements-lint.txt), the stale .gitmodules (DeepJIT is inlined), unreferenced figures/ (1.2 MB), and the top-level deep_ep/*.py, which duplicate torch-ext/deep_ep (only deep_ep/include is used by build.toml).

Suggested change

	prune keep=".clang-format,.editorconfig,.gitignore,.gitmodules,.github/workflows/format.yml,third-party/tilelang_ops/**,LICENSE,README.md,format.sh,requirements-lint.txt,csrc/**,deep_ep/**,tests/**,figures/**"
	prune keep=".gitignore,LICENSE,README.md,csrc/**,deep_ep/include/**,tests/**,third-party/tilelang_ops/**"

[IlyasMoutawwakil]

applied both prunes and regenerated src/, only deletions, check_ports passes byte for byte

[drbh]

same set as above, plus only the DeepJIT headers are built, so CMakeLists.txt, csrc/python_api.cpp, deep_jit/ and tests/ under third-party/deep_jit can go. src/ needs regenerating after this (port/regenerate.py).

Suggested change

	prune keep=".clang-format,.editorconfig,.gitignore,.gitmodules,.github/workflows/format.yml,third-party/tilelang_ops/**,LICENSE,README.md,format.sh,requirements-lint.txt,csrc/**,deep_ep/**,tests/**,figures/**,torch-ext/**,third-party/deep_jit/**"
	prune keep=".gitignore,LICENSE,README.md,csrc/**,deep_ep/include/**,tests/**,third-party/tilelang_ops/**,torch-ext/**,third-party/deep_jit/include/**,third-party/deep_jit/README.md"

note: we need this second prune after the vendoring since it prunes the third party stuff

[drbh]

i think we can drop regenerate.py. it does the same thing as scripts/check_ports.py (run kernel-port on port.kdl and byte compare against src/ with the same ignore rules), and python3 scripts/check_ports.py deep-ep already checks just this kernel the same way ci does

the bigger issue is that it hard codes --vendor python=... --vendor deepjit=... and takes the checkouts as args, while check_ports.py reads source/vendor from port.kdl and clones the pinned commits itself. so if a sync adds or renames a vendor, or someone passes a checkout at the wrong commit, it can pass locally and fail in ci (or the other way around). einops/port doesnt have one either, so this would be the only per kernel script

the only extra thing it gives us is --output to write the generated tree for copying into src/, and that isnt deep-ep specific. would be nicer as a --write flag on check_ports.py that replaces <kernel>/src instead of comparing, then every port gets the same regenerate command

pushed --write in #1299, its one commit (b90e73c) on top of this branch so you can either cherry pick it or rebase onto it

to drop it:

• delete deep-ep/port/regenerate.py
• point deep-ep/port/README.md (~L39) at python3 scripts/check_ports.py --write deep-ep
• update the deep-ep section in AGENTS.md (~L662) which says to regenerate with regenerate.py

note: kernel-port drops the exec bit (format.sh comes out 644) and compare() only checks contents, so ci doesnt catch it. doesnt matter here once the prune suggestion lands since those files go away

[IlyasMoutawwakil]

dropped, merged your --write branch in and pointed AGENTS.md at check_ports.py --write deep-ep. thanks!