[Not ready] MSCCLPP EP implementation - #852
Draft
Binyang Li (Binyang2014) wants to merge 134 commits into
Draft
Conversation
|
Azure Pipelines: There may be pipelines that require an authorized user to comment /azp run to run. |
Contributor
There was a problem hiding this comment.
Pull request overview
This PR introduces the MSCCL++ Expert-Parallel (EP) extension (LL + HT backends) with a Python-facing API, adds a unified benchmarking harness for multiple EP implementations (mscclpp, NCCL-EP, DeepEP, FlashInfer), and extends PortChannel/proxy plumbing to support remote 64-bit atomicAdd, alongside build/CI/docs updates to enable and validate the new functionality.
Changes:
- Add
src/ext/epEP extension (C++/CUDA + nanobind) and Python frontend (python/mscclpp/ep) for MoE dispatch/combine. - Add EP benchmark backends + helpers under
test/python/ep/, including an optional CUPTI kernel timer. - Add PortChannel
atomicAddtriggers and connection implementations (CUDA-IPC / IB / Ethernet), plus new unit tests; update build options, docs, and CI defaults.
Reviewed changes
Copilot reviewed 74 out of 74 changed files in this pull request and generated 4 comments.
Show a summary per file
| File | Description |
|---|---|
| test/python/ep/ep_bench_nccl.py | New NCCL-EP backend for unified EP benchmarking harness. |
| test/python/ep/ep_bench_mscclpp.py | New MSCCL++ EP backend benchmark driver with optional validation and CUDA graph options. |
| test/python/ep/ep_bench_flashinfer.py | New FlashInfer MoeAlltoAll backend adapter for benchmark comparisons. |
| test/python/ep/ep_bench_deepep.py | New DeepEP v2 ElasticBuffer backend adapter for benchmark comparisons. |
| test/python/ep/ep_bench_common.py | Shared MPI/torch bootstrap, input generation, dtype helpers, and validation utilities for EP benches. |
| test/python/ep/cupti_kernel_timer.cpp | Optional in-process CUPTI timer library for low-perturbation kernel timing. |
| test/python/ep/CMakeLists.txt | Standalone CMake build for EP bench binary + CUPTI timer. |
| test/mscclpp-test/CMakeLists.txt | Add tma_pipeline_perf test binary gated to CUDA and sm90+ arch lists. |
| test/mp_unit/port_channel_tests.cu | Add concurrent PortChannel atomicAdd test coverage across IPC/IB/Ethernet. |
| test/mp_unit/mp_unit_tests.hpp | Add testAtomicAdd declaration to PortChannelOneToOneTest. |
| src/ext/ep/runtime_base.hpp | New EP runtime base class defining common topology/availability surface. |
| src/ext/ep/README.md | New EP extension architecture/build/validation documentation. |
| src/ext/ep/moe_runtime.hpp | Declare createMoERuntime(...) factory for LL/HT EP backends. |
| src/ext/ep/moe_runtime.cc | Implement createMoERuntime(...) to construct LL or HT runtime. |
| src/ext/ep/low_latency/config.cuh | LL dispatch configuration, workspace views, kernel config caching utilities. |
| src/ext/ep/ll_runtime.hpp | LL runtime class declaration + dispatch/combine host entry points. |
| src/ext/ep/ll_runtime.cc | LL runtime implementation: symmetric buffer setup + dispatch/combine wrappers. |
| src/ext/ep/include/quantization.cuh | FP8 E4M3 quantization helpers for LL dispatch. |
| src/ext/ep/include/launch.cuh | Cooperative launch macros and rank-switch helpers for EP kernels. |
| src/ext/ep/include/exception.cuh | EP exception/assert/CUDA_CHECK helpers for torch-free kernel/runtime code. |
| src/ext/ep/include/constants.cuh | EP constants/timeouts/alignment + Torch restriction removal for kernels. |
| src/ext/ep/include/api.cuh | EP private kernel API: enums, workload/context structs, dispatch/combine declarations. |
| src/ext/ep/ht_runtime.hpp | HT runtime class declaration for direct-mapped pool + dispatch/combine protocol. |
| src/ext/ep/ht_runtime.cc | HT runtime implementation: pool mapping, notify/dispatch/combine plumbing. |
| src/ext/ep/high-throughput/runtime.cu | HT runtime helper kernels (e.g., barrier). |
| src/ext/ep/high-throughput/layout.cu | HT routing layout construction kernel (counts + token-in-rank). |
| src/ext/ep/high-throughput/dispatch.cu | HT notify/dispatch kernels with cooperative grid sync. |
| src/ext/ep/high-throughput/config.cuh | HT config constants and pool sizing helpers. |
| src/ext/ep/high-throughput/combine.cu | HT TMA pipeline combine kernel implementations. |
| src/ext/ep/config.hpp | Shared EP config helpers + LL symmetric buffer layout definitions. |
| src/ext/ep/CMakeLists.txt | Build mscclpp_ep_cpp nanobind extension; enforce sm90+ arch selection. |
| src/ext/ep/bindings.cpp | nanobind module exposing MoERuntime + LL/HT raw-pointer entry points. |
| src/ext/CMakeLists.txt | Wire EP extension into overall ext build when enabled. |
| src/core/unix_socket.cc | Add shutdown helper and destructor cleanup for UnixSocketServer lifecycle. |
| src/core/port_channel.cc | Add proxy handling for atomicAdd triggers; adjust stop/drain behavior. |
| src/core/include/unix_socket.hpp | Add UnixSocketServer destructor and shutdown declaration. |
| src/core/include/context.hpp | Add proxy atomic stream/context members and atomicAdd API to CudaIpcStream. |
| src/core/include/connection.hpp | Add atomicAdd virtual API to connection implementations. |
| src/core/include/communicator.hpp | Update comments related to last-recv-item lifecycle behavior. |
| src/core/context.cc | Implement CudaIpcStream destructor and document atomic stream sync caveats. |
| src/core/connection.cc | Implement Connection::atomicAdd and transport-specific atomicAdd paths. |
| src/core/communicator.cc | Refactor ordered recv/connect/semaphore futures without helper wrapper. |
| src/core/atomicadd_kernel.cu | Add device kernel + CudaIpcStream::atomicAdd implementation. |
| python/test/test_gpu_buffer_pool_nvls_zero.py | New NVLS zero-copy allreduce timing test using GpuBufferPool + CUDA graphs. |
| python/mscclpp/ep/utils.py | New EP Python utility helpers (pointer views, object broadcast/gather, etc.). |
| python/mscclpp/ep/types.py | New EP public Python dataclasses/types for dispatch/combine handles and config. |
| python/mscclpp/ep/communicator.py | New high-level Python MoECommunicator dispatch/combine API. |
| python/mscclpp/ep/_cpp.py | EP extension loader/shim exporting nanobind symbols. |
| python/mscclpp/ep/init.py | EP package public exports. |
| python/mscclpp/_core/comm.py | Fix torch-group UniqueId broadcast to handle variable-size payloads. |
| python/csrc/core_py.cpp | Release GIL around Bootstrap send/recv bindings for better concurrency. |
| pyproject.toml | Enable EP extension build in scikit-build CMake defines. |
| include/mscclpp/port_channel_device.hpp | Add PortChannel device-side atomicAdd trigger emission helpers. |
| include/mscclpp/gpu.hpp | Extend HIP CUDA-driver API type/function shims (CUcontext, cuCtx*). |
| include/mscclpp/gpu_data_types.hpp | Add bf16<->fp32 packed conversions and guard bf16 alias conflicts. |
| include/mscclpp/fifo_device.hpp | Reserve trigger type==0 for atomicAdd semantics. |
| include/mscclpp/core.hpp | Add public Connection::atomicAdd API declaration. |
| docs/quickstart.md | Document .[cuda12,ep] extra and include EP in “all extras” example. |
| CMakeLists.txt | Add MSCCLPP_BUILD_EXT_EP option to top-level build configuration. |
| .github/workflows/mscclpp-lang.yml | Adjust CI LD_LIBRARY_PATH to avoid CUDA compat lib mismatch. |
| .github/workflows/lint.yml | Move lint jobs to ubuntu-24.04 runners. |
| .github/agents/device-code-reviewer.agent.md | Add device-code-reviewer agent guidance document. |
| .devcontainer/devcontainer.json | Update devcontainer base image to cuda13.0 tag. |
Comment on lines
+61
to
+72
| // Note: proxyAtomicStream_ is NOT synced here. The atomicAdd kernels are fire-and-forget | ||
| // operations that complete asynchronously on the GPU. Syncing them here would deadlock | ||
| // because sync() is called from the proxy thread while the main thread may hold the | ||
| // device context via cudaStreamSynchronize() on the test kernel's stream. | ||
| // | ||
| // TODO(#796): As a side effect, `Connection::flush()` does not order/complete pending | ||
| // remote `atomicAdd` operations on the CUDA-IPC transport, so PortChannel flush no | ||
| // longer guarantees that a peer kernel sees the updated value. EP currently relies on | ||
| // higher-level signaling (PortChannel signal/wait, FIFO drain) for ordering, but a | ||
| // correct fix needs a deadlock-free way to drain `proxyAtomicStream_` here. Carried | ||
| // over from the DeepEP `chhwang/dev-atomic-add-cleanup` cherry-pick; revisit before | ||
| // this lands on `main`. |
Comment on lines
+21
to
+35
| #if !defined(MSCCLPP_DEVICE_HIP) | ||
| // On CUDA, the proxy thread cannot launch kernels or perform stream operations on the | ||
| // primary context without deadlocking with the main thread's cudaStreamSynchronize(). | ||
| // The CUDA runtime uses a per-context lock; the main thread holds it while waiting for | ||
| // the test kernel, and the proxy thread needs it to launch the atomicAdd kernel. | ||
| // A separate CUDA context avoids this contention. | ||
| // | ||
| // TODO(#796): `dst` is a CUDA-IPC mapping registered in the primary/runtime context, so | ||
| // launching this kernel from `proxyAtomicCtx_` is technically UB (device pointers are | ||
| // context-scoped). It works in practice on current drivers because the IPC handle aliases | ||
| // the same physical allocation, but a correct fix would either (a) avoid the separate | ||
| // context (e.g. break the deadlock differently) or (b) re-open the IPC mapping inside | ||
| // `proxyAtomicCtx_`. Carried over from the DeepEP `chhwang/dev-atomic-add-cleanup` | ||
| // cherry-pick; revisit before this lands on `main`. | ||
| if (!proxyAtomicCtx_) { |
Comment on lines
+761
to
+771
| if (isAtomicAdd && received && size == sizeof(int64_t)) { | ||
| // Atomic add: receive the value, read-modify-write on GPU memory | ||
| int64_t addValue; | ||
| recvSocket_->recvUntilEnd(&addValue, sizeof(int64_t), &closed); | ||
| received &= !closed; | ||
|
|
||
| if (received) | ||
| mscclpp::gpuMemcpy(ptr + (recvSize / sizeof(char)), recvBuffer_.data(), messageSize, cudaMemcpyHostToDevice); | ||
| recvSize += messageSize; | ||
| if (received) { | ||
| int64_t current; | ||
| mscclpp::gpuMemcpy(reinterpret_cast<char*>(¤t), ptr, sizeof(int64_t), cudaMemcpyDeviceToHost); | ||
| current += addValue; | ||
| mscclpp::gpuMemcpy(ptr, reinterpret_cast<char*>(¤t), sizeof(int64_t), cudaMemcpyHostToDevice); | ||
| } |
Comment on lines
+646
to
+650
| // Step 3: Block 0 signals remote that all adds are done, flushes, then waits for remote. | ||
| if (blockIdx.x == 0) { | ||
| portChan.signal(); | ||
| portChan.flush(); | ||
| portChan.wait(); |
![github-advanced-security[bot]](https://avatars.githubusercontent.com/in/57789?s=60&v=4){kind=small}
Binyang Li (Binyang2014)
marked this pull request as ready for review
August 11, 2026 23:02
|
Azure Pipelines: There may be pipelines that require an authorized user to comment /azp run to run. |
Binyang Li (Binyang2014)
changed the title
Binyang Li (Binyang2014)
changed the base branch from
main
to
feature/tma-bulk-memory-channel
August 11, 2026 23:02
Binyang Li (Binyang2014)
added a commit
that referenced
this pull request
Aug 13, 2026
## Summary Adds pointer-based 1-D bulk-copy primitives for NVIDIA sm_90+ in `mscclpp/bulk_device.hpp`: - `BulkBarrier` for reusable async-load completion - `bulkLoad` and `bulkStore` for global ↔ shared copies - `bulkReduceStore<T>` for copy-engine reduction into global or peer memory - commit, completion-wait, source-reuse-wait, and proxy-fence helpers - `isBulkSupported()` / `is_bulk_supported` capability queries The API operates on global pointers, including peer-mapped pointers. It does not add `MemoryChannel` methods: EP kernels use raw peer-pointer arrays for data movement and channels only for synchronization. The API covers all TMA instructions used by #852 and `qinghuazhou/unified_ep_bench_ht_python`. Unsupported device targets fail at compile time instead of silently doing nothing. `BulkBarrier` storage remains host-visible so launch code can size dynamic shared memory. The store API separates destination completion from source-tile reuse for pipelined kernels. `bulkReduceStore` currently supports `Add` for `float`, `__nv_bfloat16`, and `uint32_t`; unsupported combinations fail a `static_assert`. ## Tests - Single-GPU tests: load, multi-source gather, barrier reuse, pipelined store, and reduction - Multi-rank EP-shaped tests: - staged dispatch to peer memory - double-buffered multi-source combine - push combine with remote reduction - Doxygen/Sphinx coverage for the new API - Python binding build and import | Platform | Result | |---|---| | H200, sm_90 | `unit_tests` 40/40; `mp_unit_tests` 60/60; patterns 3/3 at 6 ranks | | GB200, sm_100 | `unit_tests` 40/40; non-IB `mp_unit_tests` 32/32; patterns 3/3 at 4 ranks | | MI300X, ROCm | Results match the merge-base: 23 passed / 9 skipped; non-IB mp tests 24 passed / 3 pre-existing failures | ## Limitation Concurrent reduction from several GPUs into the same peer address produced exact results across repeated H200 and GB200 runs, but cross-peer atomicity was not found in the PTX documentation. The header documents this as empirical rather than guaranteed.
Binyang Li (Binyang2014)
force-pushed
the
feature/ep
branch
from
August 13, 2026 22:05
c7a335a to
a61ee11
Compare
Binyang Li (Binyang2014)
marked this pull request as draft
August 13, 2026 22:09
Port DeepEP's high-throughput MoE dispatch/combine kernels onto MSCCL++
as an optional build target `mscclpp_ep_cpp`, gated by -DMSCCLPP_BUILD_EXT_EP
(OFF by default). Sources are lifted from DeepEP branch
`chhwang/dev-atomic-add-cleanup` and rebased onto upstream MSCCL++ APIs;
the NVSHMEM / IBGDA dependencies are replaced with `PortChannel` +
`MemoryChannel` + the new `Connection::atomicAdd` primitive.
Scope
-----
Intranode (NVLink-only):
* `Buffer` ctor/dtor: cudaMalloc nvl workspace, export IPC handle,
allocate FIFO + peer-pointer tables, start `ProxyService`.
* `sync()`: import peer IPC handles, upload peer pointer table,
build `MemoryDevice2DeviceSemaphore` + `MemoryChannel` per peer.
* `get_dispatch_layout`, `intranode_dispatch`, `intranode_combine`
ported verbatim (torch::Tensor ABI preserved).
Internode HT (NVLink + RDMA):
* `sync()` RDMA branch: cudaMalloc RDMA buffer + `bootstrap->barrier()`
(replacing NVSHMEM symmetric-heap allocation); register with
`all_transport`, exchange via `sendMemory`/`recvMemory`, build 12 IB
QPs/peer + 16 semaphores/peer + 16 port channels/peer.
* Full `internode.cu` port (notify_dispatch / dispatch / cached_notify
/ combine / get_dispatch_layout). The 4 raw `ChannelTrigger` atomic
sites are rewritten to call the new
`PortChannelDeviceHandle::atomicAdd(offset, value)` API; the single
`nvshmem_fence()` is replaced with `__threadfence_system()` (remote
visibility guaranteed by the subsequent port-channel barrier).
* `internode_dispatch` / `internode_combine` host code ported, with
the torch tensor marshalling and CPU spin-wait on mapped counters.
Low-latency (pure RDMA):
* Not ported. `low_latency_dispatch`, `low_latency_combine`,
`clean_low_latency_buffer`, `get_next_low_latency_combine_buffer`
throw `std::runtime_error`; the Python frontend refuses to
construct a Buffer with `low_latency_mode=True`.
Python layer
------------
* New pybind11 + libtorch Python extension `mscclpp_ep_cpp` (separate
from the nanobind `_mscclpp` because the EP ABI carries
`torch::Tensor` / `at::cuda::CUDAStream`).
* `mscclpp.ext.ep.Buffer` mirrors `deep_ep.Buffer`; exchanges device
IDs, IPC handles and the bootstrap UniqueId over the user's
`torch.distributed` process group before calling `sync()`.
* `mscclpp.ext` auto-imports `ep` if the extension is built.
Build
-----
* `src/ext/ep/CMakeLists.txt`: finds Python + Torch; warns and skips if
`CMAKE_PREFIX_PATH` doesn't point at `torch.utils.cmake_prefix_path`.
Falls back to Torch's bundled pybind11 if a standalone pybind11 is not
installed. Links `libtorch_python` explicitly (without it, `import
mscclpp_ep_cpp` fails with `undefined symbol: THPDtypeType`).
* Top-level `CMakeLists.txt` exposes the `MSCCLPP_BUILD_EXT_EP` option
(default OFF).
Tests
-----
* `test/python/ext/ep/test_ep_smoke.py`: skipped if the extension isn't
built. Covers Config round-trip, low-latency size hint, and the LL
construction guard. Multi-rank functional tests still to do on H100.
Notes
-----
* Builds against the preceding "atomic add" commit which adds
`Connection::atomicAdd` and `PortChannelDeviceHandle::atomicAdd` to
upstream MSCCL++.
* Intranode path verified end-to-end (build + import + smoke tests).
* Internode HT is code-complete but requires real IB hardware to
validate; see `src/ext/ep/README.md` for the detailed port plan and
remaining LL migration.
Port DeepEP's pure-RDMA low-latency (LL) MoE kernels from
csrc/kernels/internode_ll.cu (branch chhwang/dev-atomic-add-cleanup)
into the MSCCL++ EP extension. NVSHMEM / IBGDA device primitives are
replaced with MSCCL++ PortChannelDeviceHandle operations:
nvshmemx_barrier_all_block() -> port-channel signal+wait ring
nvshmemi_ibgda_put_nbi_warp(...) -> lane-0 PortChannel.put(...)
nvshmemi_ibgda_amo_nonfetch_add(...) -> lane-0 PortChannel.atomicAdd(...)
The atomicAdd path relies on the MSCCL++ Connection::atomicAdd /
PortChannelDeviceHandle::atomicAdd API cherry-picked from branch
chhwang/new-atomic-add; the LL dispatch path uses a signed delta
(-num_tokens_sent - 1) which the new int64_t signature supports.
Changes:
* New file src/ext/ep/kernels/internode_ll.cu (~530 lines) with the
three kernels clean_low_latency_buffer, dispatch<kUseFP8,...>,
combine<...> plus their launchers. rdma_buffer_ptr is threaded
through the launchers so the kernel can translate virtual addresses
into registered-memory offsets expected by MSCCL++.
* kernels/api.cuh: replace the single stub signature with full LL
launcher prototypes.
* buffer.cc: replace the four LL throw-stubs
(clean_low_latency_buffer, low_latency_dispatch,
low_latency_combine, get_next_low_latency_combine_buffer) with
torch-Tensor implementations ported from DeepEP/csrc/deep_ep.cpp.
* Drop src/ext/ep/internode_stub.cc and its CMake entry.
* python/mscclpp/ext/ep/buffer.py: remove the low_latency_mode=True
NotImplementedError guard; update docstring.
* test/python/ext/ep/test_ep_smoke.py: rename
test_low_latency_rejected -> test_low_latency_buffer_construct
to reflect that LL construction is now accepted.
* src/ext/ep/README.md: update status matrix, document the
NVSHMEM -> MSCCL++ translation table, and list the known
limitations.
This is a structural port: the kernels compile, link, and pass the
single-rank smoke tests, but end-to-end behaviour on multi-node H100
is not yet validated. Two known caveats:
1. Performance will NOT match IBGDA because MSCCL++ port channels
use a CPU proxy; this port is for functional parity, not latency.
2. Buffer::sync() in LL mode only connects peers that share the
same local GPU id (DeepEP convention), so the LL kernels assume
a one-GPU-per-node topology (num_ranks == num_rdma_ranks).
Multi-GPU-per-node LL layouts will need a follow-up in sync().
Tested:
cmake --build build -j --target mscclpp_ep_cpp # builds clean
pytest test/python/ext/ep/test_ep_smoke.py # 3 passed
Three issues blocked end-to-end intranode validation across multiple ranks. This commit fixes them and adds a 2/4/8-rank functional test. 1. Combine receiver: OOB __shared__ read In the combine receiver warp, the wait loop evaluated `channel_tail_idx[recv_lane_id] <= expected_head` before the `expected_head >= 0` guard. `channel_tail_idx` is a shared array of size `kNumRanks`, but the loop runs on all 32 lanes of a warp, so lanes with `recv_lane_id >= kNumRanks` indexed out of bounds. compute-sanitizer reported "Invalid __shared__ read of size 4 bytes" at combine<bf16,2,768>+0xdd0, surfaced asynchronously as cudaErrorIllegalAddress at the kernel launch site. Swap the operands so the rank-bounds check short-circuits the shared read. 2. Python bindings: UniqueId ABI `mscclpp::UniqueId` is a `std::array<uint8_t, N>` which pybind11 auto-converts to a Python `list`, silently overriding any `py::class_<UniqueId>` wrapper. Expose `create_unique_id` / `connect` as lambdas that produce/consume `py::bytes` and memcpy into a local `UniqueId`. Also coerce `bytes`->`bytearray` at the Python call site for `sync()` whose signature expects `pybind11::bytearray`. 3. Python frontend: communicator required for NVL-only sync `Buffer::sync()` uses `communicator->connect(ipc_config, ...)` on the pure-NVLink path, so the communicator must be initialized even when `num_rdma_ranks == 1` and `low_latency_mode == False`. Always broadcast the unique id and call `runtime.connect()` before `sync()`. Validation on a single H100x8 node via torchrun: - 2 ranks: dispatch 195 tokens, combine diff=0 - 4 ranks: dispatch 371 tokens, combine diff=0 - 8 ranks: dispatch 456 tokens, combine diff=0 Test harness added at test/python/ext/ep/test_intranode_multirank.py.
The `internode` kernels index device-side port channel handles as
`port_channel_handles[channel_id * num_ranks + peer_rank]`, where
`peer_rank` is a global rank in [0, num_ranks). `Buffer::sync` was
building that table by iterating `std::unordered_map<int, MemoryId>`
(and similarly for connections/semaphores), which yields hash order
rather than ascending rank order. Once the cross-node fan-out grew
beyond a single peer, a local rank's trigger for peer `r` landed on
the semaphore/memory pair of a different peer, so RDMA puts and
atomic tail updates went to the wrong destination and the forwarder
spun on a tail counter that never advanced.
Changes:
- Build `sema_ids` and `port_channel_handles` by iterating
`for (int r = 0; r < num_ranks; ++r)` and looking up the
connection / memory id for rank `r`, skipping ranks excluded by
low-latency mode (inserting a placeholder handle so the stride
stays `num_ranks`).
- Tag the RDMA-phase `sendMemory`/`recvMemory`/`connect` calls with
`kRdmaTag = 1` so they do not collide with NVL-phase tag-0
traffic between the same pair of ranks.
- Drop an unused `r` local in the NVL setup loop.
With this fix and a matched `libmscclpp.so` on both nodes, the
2-node x 8-GPU internode HT dispatch path completes successfully
(`[dispatch] OK`). Combine is still under investigation.
Also adds `test/python/ext/ep/test_internode_multirank.py`, a
torchrun-based 2-node functional test that exercises
`get_dispatch_layout` -> `internode_dispatch` -> `internode_combine`
and validates per-source-rank token values end-to-end.
Two issues prevented internode HT combine from completing on 2x8 H100: 1. Wrong prefix matrices passed to internode_combine. Combine runs in the reverse direction of dispatch, so it must consume the receiver-side matrices returned by dispatch (recv_rdma_channel_prefix_matrix, recv_rdma_rank_prefix_sum, recv_gbl_channel_prefix_matrix), not the sender-side rdma_channel_prefix_matrix / gbl_channel_prefix_matrix. This matches DeepEP's deep_ep/buffer.py::internode_combine handle unpacking. Without the fix the NVL forwarder's 'NVL check' timed out because token_start_idx/token_end_idx were computed against the wrong per-channel layout. 2. Cross-rank race between dispatch and combine. Even with the correct matrices, launching combine immediately after dispatch deadlocked the forwarder NVL check (tail stuck one short of expected_head) because peers still had in-flight dispatch proxy traffic while fast ranks had already started combine. A torch.cuda.synchronize() + dist.barrier() between the two calls makes the test pass deterministically on 16 ranks (combine diff == 0, max|expected| up to 60.0). The barrier in the test is a workaround; the real fix belongs in Buffer::internode_dispatch / Buffer::internode_combine so the dispatch->combine handoff fully fences outstanding proxy work across ranks. Marked with an XXX comment in the test.
Refresh status docs and comments now that internode HT dispatch and combine have been validated end-to-end on 2 nodes x 8 H100 GPUs via test/python/ext/ep/test_internode_multirank.py (all 16 ranks recover their per-rank token payloads with zero diff). - src/ext/ep/README.md: consolidate the previously duplicated README into a single document; mark intranode and internode HT dispatch and combine as validated in the status table; add a 'Running the tests' section with torchrun examples for both the intranode and the 2x8 internode setups; record the dispatch->combine torch.cuda.synchronize() + dist.barrier() requirement under Known limitations; mark Phase 2 DONE and keep Phase 3 (LL) as structural port, untested. - python/mscclpp/ext/ep/buffer.py: update the module docstring and the Buffer constructor docstring to say internode HT is validated and clarify that LL mode is untested on multi-node hardware. - src/ext/ep/buffer.cc: drop the stale 'NVSHMEM support not yet ported' and 'low-latency paths still stubbed' comments. mscclpp_ep does not use NVSHMEM at all (PortChannel/MemoryChannel replace it), and the LL paths are a structural port that is present but untested, not stubbed. Note validation on 2x H100x8 in the internode section header.
- Buffer::sync no longer drops non-same-GPU-id peers in low_latency_mode. DeepEP's original filter was safe because its LL path used NVSHMEM; this port drives LL via PortChannel so the kernel indexes port_channel_handles[local_expert*num_ranks + dst_rank] for every dst_rank. All peers now get a real memory/connection/semaphore/port channel entry. - Add test/python/ext/ep/test_low_latency_multirank.py (LL dispatch+combine functional round-trip, BF16 only). Works cross-node in DeepEP's 1-GPU-per-node topology. - Known limitation documented in src/ext/ep/README.md and the test docstring: intra-node 8-GPU LL currently hangs because every peer transfer routes through the CPU proxy over IB loopback between distinct HCAs on the same host, and (separately) CudaIpcConnection::atomicAdd is a 64-bit op which mis-aligns the 32-bit rdma_recv_count slots when used for same-node peers. Proper fix needs a mixed-transport LL variant (MemoryChannel for same-node, PortChannel for cross-node) or 64-bit counters.
Gated behind MSCCLPP_EP_BENCH=1 to keep correctness runs fast. Reports per-iter latency (max across ranks, CUDA-event timed) and aggregate effective bandwidth (sum across ranks, dispatch+combine payload bytes). Tunable via MSCCLPP_EP_BENCH_WARMUP / _ITERS / _TOKENS / _HIDDEN. Bench reuses the Buffer allocated for the correctness phase and self-skips if the requested hidden exceeds the per-peer NVL/RDMA budget.
Previously the optional benchmark measured full round-trip latency. Split it to time dispatch alone (N iters) and combine alone (N iters reusing one dispatch output), reporting per-phase latency (max across ranks) and aggregate effective bandwidth (sum across ranks). Applies to intranode HT, internode HT, and the (currently unreachable on intra-node 8-GPU) LL test. Internode HT keeps the sync+barrier guard between dispatch and combine but excludes it from either phase's timing.
…o int64 The low-latency dispatch/combine kernels signal recv counts via MSCCL++ PortChannel.atomicAdd, which lowers to IB IBV_WR_ATOMIC_FETCH_AND_ADD. That opcode requires the remote address to be 8-byte aligned, but LowLatencyLayout packed the per-expert signaling slots as int32. Odd slots landed at offset %8 == 4; the NIC silently dropped those atomics and the target rank spun forever in recv_hook (observed: even->odd direction works, odd->even does not, across all tested topologies including 2-rank intra-node, 8-rank intra-node, and 2-node 1-GPU-each). Widen dispatch_rdma_recv_count_buffer / combine_rdma_recv_flag_buffer to int64_t, update clean kernel + kernel signatures + next_clean pointers accordingly, and add int64_t overloads for st_na_release / ld_acquire_sys_global in utils.cuh. Also drop the bogus self CUDA-IPC connection in Buffer::sync() that was previously skewing the cross-rank buildAndAddSemaphore handshake order; the kernel's same-rank branch uses a direct warp copy and never touches the self port-channel slot (filled with a zero-initialized placeholder so the [local_expert*num_ranks + dst_rank] indexing still holds).
Dropping the self ipc_cfg connection caused cudaErrorInvalidResourceHandle on multi-node launches. Keep the self connection (needed by other code paths that assume every rank is in the connections map) but continue to skip the self slot in the semaphore + port-channel construction loops so the kernel's [local_expert*num_ranks + dst_rank] indexing hits only peer handles; the self slot is a zero-initialized placeholder since the kernel's same-rank branch uses a direct warp copy.
Extend direct-fabric HT launch and runtime support to 16 ranks. Decouple TMA contributor staging from global rank discovery and restore adaptive warp geometry. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Support E4M3 payloads with FP32 scales per 32 elements, return global token-major expert IDs with num_experts sentinels, and simplify QuantConfig. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 59760d2a-e65f-44b6-b2e6-9c271b834d7e
Use linear UE8M0 scale bytes per 32 elements, global token-major expert IDs, optimized Blackwell conversion targets, and updated validation. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 59760d2a-e65f-44b6-b2e6-9c271b834d7e
Make the invalid token expert sentinel configurable and preserve tiny nonzero MXFP8 blocks for UE8M0 code zero. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 59760d2a-e65f-44b6-b2e6-9c271b834d7e
#836) ## Summary Adds `run_ep_bench_python.py` — an in-process low-latency EP benchmark that drives **both** the mscclpp EP (`MoECommunicator`) and NVIDIA NCCL-EP Python APIs through one shared paired dispatch→combine loop, so the two backends are timed with byte-for-byte identical methodology (matching `mscclpp_ep_bench.cu` output). ## Changes - New `test/python/ep/run_ep_bench_python.py` (renamed from `ep_bench_unified.py`), MPI/mpi4py bootstrap, `--backend {mscclpp,nccl,both}`, optional in-process CUPTI kernel timing. - `run_ep_bench.py`: dropped the redundant Python `mscclpp` (`ep_bench_ll.py`) backend; deleted `ep_bench_ll.py` (superseded by the unified script). - Updated the mscclpp LL API usage for the current (IPC-capable) runtime (`num_rdma_bytes=0`; removed `get_low_latency_rdma_size_hint`). - Docstring: clarified mscclpp LL supports both CUDA-IPC (NVLink) and RDMA/IB, and added single-node and validated 2-node (shared MNNVL fabric) launch examples with placeholder IPs/iface. Both backends run at 1 and 2 nodes over the shared NVLink fabric. --------- Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> Co-authored-by: Binyang Li <binyli@microsoft.com>
) The kineto per-kernel attribution in _kineto_kernel_us matched the "dispatch"/"combine" substring against the full demangled kernel name. The rank-major combineKernel is templated on DispatchLayout (combineKernel<.., DispatchLayout::RANK_MAJOR>), so its name contains the substring "dispatch" and was wrongly summed into the dispatch bucket. This only surfaced under --cuda-graph, where the combine kernel also appears in the dispatch profiling pass, doubling the reported dispatch kernel time (e.g. 16->32 us at 1 node). Match on the function name with template arguments stripped (text before the first "<") so DispatchLayout / CombineMode template params no longer collide.
This PR adds rank-major output support to the MSCCL++ EP low-latency (LL) path, wiring it through the CUDA kernels/runtime, Python bindings, and the Python tests/benchmarks. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Clarify the EP backend phases and reduce duplicated low-level machinery: - use MSCCL++ memory-channel semaphores for HT device barriers - rename HT routing-count/exchange launchers and document each phase - size LL workspace exactly instead of reserving a fixed 32 MiB - replace blanket cooperative-launch macros with scoped launch configs - remove unused PTX helpers and the obsolete EP constants header - make EP default-on for supported CUDA/Python builds and detect concrete native GPU architectures Also remove invalid EP install-extra docs, a non-importable GPUBufferPool standalone test, and the orphan tma_pipeline_perf CMake target. Restore Unix socket lifecycle code to origin/main and fix torch subgroup bootstrap source handling. Validated with 32-GPU LL rank-major and 16-GPU HT benchmarks, plus 2/4-GPU HT and LL BF16/FP8/direct-send correctness runs. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Apply the repository clang-format 18 style to the cooperative combine launch macro so tools/lint.sh passes for both C++ and Python. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…t converted to larger type' Co-authored-by: Copilot Autofix powered by AI <62310815+github-advanced-security[bot]@users.noreply.github.com>
Expose the sm_90+ bulk copy engine (cp.async.bulk / cp.reduce.async.bulk) as a pointer-based device API in a new header, mscclpp/bulk_device.hpp. The primitives are channel-free by design. Bulk copies move bytes between global and shared memory and do not care whether the global side is local or peer mapped; channels answer where a peer's memory is and how to synchronize with it. Those concerns compose through a plain pointer. The expert-parallel kernels that motivate this feature gather from arrays of raw peer pointers and use channels only for signal/wait, so binding the primitives to a channel would not serve them. This also matches how SwitchChannel exposes multimem: a public, pointer-based primitive is the foundation. Surface: BulkBarrier load completion, with caller-held phase parity bulkLoad global -> shared, tracked by a barrier bulkStore shared -> global bulkReduceStore<T,Op> shared -> global, accumulating at the destination bulkStoreCommit close the current bulk group bulkStoreWait<N> stores have landed bulkStoreWaitSource<N> source tiles are reusable, stores may be in flight bulkFence order generic shared accesses against the async proxy isBulkSupported() host capability query, alongside isNvlsSupported() Notes on specific choices: - MSCCLPP_BULK_AVAILABLE gates the declarations, so unguarded use on an unsupported target is a compile error rather than a silent no-op. This follows the existing __CUDA_ARCH__ >= 900 call-site guards used for NVLS. - BulkBarrier storage is declared on every target, including host compilation, because host code must size the dynamic shared memory that holds barriers. Only the operations are gated. - expect, arrive and arriveAndExpect are separate so a multi-source gather can accumulate N loads against one barrier and wait once. Every expect for a batch must precede the arrival that completes the arrival count; arriveAndExpect carrying the batch total is documented as the recommended form. - init() includes the proxy fence that publishes the barrier to the async proxy. relaxedInit() omits it so an array of barriers can be set up under one fence, mirroring signal() and relaxedSignal(). - wait() advances the phase, so a barrier is initialized once and reused rather than reinitialized per batch. - bulkStoreWait and bulkStoreWaitSource are distinct because a double-buffered store pipeline needs to refill a tile without draining the store. Tests. test/unit/bulk_tests.cu covers the primitives on a single GPU. The reduction tests seed the destination so they distinguish accumulate from overwrite. test/mp_unit/bulk_pattern_tests.cu adds BulkPatternTest, three multi-rank kernels shaped after expert-parallel dispatch and combine: a staged push, a multi-source pull and reduce that is double buffered across chunks, and the same reduction expressed as a push using bulkReduceStore. Verified on H200 (sm_90): unit_tests 40/40, mp_unit_tests 60/60 at 2 ranks, BulkPatternTest 3/3 at 6 ranks. Guard behavior confirmed: guarded code builds at sm_80 and for multi-arch sm_80+sm_90, unguarded code fails to compile at sm_80, and an unsupported reduction type fails its static_assert.
<cuda_bf16.h> was included only where MSCCLPP_BULK_AVAILABLE is 1, so __nv_bfloat16 could not be named in a host translation unit and bulkReduceStore<__nv_bfloat16> could not be instantiated from one. The unit test only compiled because gpu_utils.hpp happened to pull the type in. Gate the include on MSCCLPP_DEVICE_CUDA instead, matching switch_channel_device.hpp. Plain host builds without the CUDA toolkit are unaffected. Found while validating on GB200 (sm_100), where the header is used without that incidental include. Verified on GB200 (sm_100, CUDA 13.0, aarch64): 11/11 standalone checks including peer-memory load, store and reduce, and 4-GPU concurrent accumulate into one buffer over 50 iterations. Guard behavior holds under CUDA 13: guarded code builds at sm_80 and multi-arch sm_80+sm_100a, unguarded code fails to compile at sm_80, plain g++ reports sizeof(BulkBarrier)=8. Re-verified on H200 (sm_90): unit_tests 40/40, mp_unit_tests 60/60 at 2 ranks.
Three cleanups found by running things that had not been run. Doxygen strips everything guarded by MSCCLPP_BULK_AVAILABLE, because the macro derives from __CUDA_ARCH__ and doxygen does not define it. Every directive added to cpp_api.rst therefore failed: WARNING: doxygenfunction: Cannot find function "mscclpp::bulkLoad" in doxygen xml output for project "mscclpp" from directory: ./doxygen/xml Add MSCCLPP_BULK_AVAILABLE=1 to PREDEFINED in the Doxyfile, next to the existing MSCCLPP_DEVICE_COMPILE and MSCCLPP_DEVICE_CUDA entries that exist for the same reason. Sphinx now emits no bulk warnings and every symbol renders. Bind isBulkSupported() as is_bulk_supported and export it, matching is_nvls_supported. Replace the unit test's homegrown compute-capability check with isBulkSupported(). The mp_unit tests already used it; having two ways to ask the same question is what the host query exists to avoid. Also document that the cross-device atomicity of bulkReduceStore() is established by measurement rather than by the PTX documentation, so callers do not take it as guaranteed. Verified on H200: docs build clean of bulk warnings and all ten symbols present in the generated HTML; unit_tests 40/40; mp_unit_tests 60/60 at 2 ranks; Python bindings build and mscclpp.is_bulk_supported() returns True.
Where MSCCLPP_BULK_AVAILABLE is 0, BulkBarrier has no operations, only the storage declared so host code can size shared memory holding barriers. Clang then warns on every ROCm translation unit that includes the header: include/mscclpp/bulk_device.hpp:139:23: warning: private field 'mbar_' is not used [-Wunused-private-field] Mark the member maybe_unused. No effect where the operations exist. Verified on MI300X (ROCm 7.2, gfx942): warning count for this field 0, and both suites match the merge-base exactly -- unit_tests 32 ran / 23 passed / 9 skipped, mp_unit_tests --filter=-Ib 29 ran / 24 passed / 3 failed. The three failures are CommunicatorTest.BasicWrite, .WriteWithDeviceSemaphores and .WriteWithHostSemaphores, which fail identically on the merge-base; this node has no IB device. Re-verified on H200: unit_tests 40/40, mp_unit_tests 60/60.
Mark kernel parameters [[maybe_unused]] at their declarations instead of adding #else blocks with (void)param casts when bulk copy is unavailable. This keeps the unsupported-target path declarative and removes 25 lines of warning-only code. Verified with nvcc at sm_80, the H200 bulk tests (6/6 unit, 3/3 multi-rank), and the MI300X ROCm build (no unused-parameter warnings; unit tests 23 passed / 9 skipped).
Co-authored-by: Binyang Li <binyli@microsoft.com>
## Summary Assign low-latency dispatch epochs once in `MoELowLatencyRuntime` and pass the epoch to dispatch and combine through `Workload`. Previously, each CUDA block derived the epoch from device workspace state while block 0 updated that state without a grid-wide ordering guarantee. Late blocks could therefore observe a different epoch and wait indefinitely on readiness flags. ## Changes - Add a host-owned dispatch epoch counter to `MoELowLatencyRuntime`. - Increment the epoch once before each dispatch. - Pass the same epoch to dispatch and its subsequent combine operation. - Use the workload epoch for rank-major synchronization. - Remove device-side epoch generation, publication, and workspace storage. This gives every block in a launch one immutable epoch value. ## Validation - Passed eager rank-major tests at capacities 1 and 2. - Passed eager expert-major tests at capacity 2. - Passed bit-exact 2-rank and 4-rank multi-operation CUDA Graph replays. - Passed all 10 SGLang dispatcher tests. - Completed six 32-rank serving rounds with 40/40 responses and no scheduler watchdog timeouts.
Unify latency and overlap under one runtime, organize kernels by dispatch and combine algorithms, and expose one opaque dispatch handle. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Keep conflicted files aligned with the latest main branch after rebasing feature/ep. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Binyang Li (Binyang2014)
force-pushed
the
feature/ep
branch
from
August 21, 2026 22:08
9bd94e1 to
8e688ee
Compare
…llow-ups (#858) ## Summary Follow-up changes on top of the merged unified Python EP benchmark (#836), addressing the outstanding review comments and refreshing the CUDA-graph methodology. Benchmarks all four EP backends — mscclpp EP, NVIDIA NCCL-EP, DeepEP V2, FlashInfer — through one shared in-process harness with identical dispatch→combine timing. ## Changes (by review comment) - **Move the kernel-name parse into the backend files** — each backend owns `parse_kineto_kernels(...)`; the harness holds zero per-library kernel knowledge (shared `sum_matching_kernel_us()` in ep_bench_common). - **remove sync here** — dropped the redundant warmup `stream.synchronize()`. - **Move cudaGraph logic to run_ep_bench_python / unified timing** — capture is owned by the harness (`_capture_paired_graph`); each `setup_*` returns a uniform `{dispatch, combine, teardown, barrier, graph}` dict and one shared loop times all four backends. Per-library specifics stay in the backend closures. - **run multiple iterations inside the cuda graph** — added `--iters-per-graph N` (sglang bench_moe_ep.py pattern): N dispatch→combine iterations per graph, reported per iteration. ## Other fixes - DeepEP CUDA-graph gate keyed on transport (EP_DISABLE_GIN), not node count — verified graph capture at 1/2/4 nodes on GB200 NVL72. - Force the paired kineto pass whenever the harness captures a single graph (fixes mscclpp "captured 0" under --cuda-graph). - Removed the C++ LL bench, shell driver, and CUPTI timer. - Clarified EP_KINETO_SEPARATE=0 is the required (not "legacy") CUDA-graph mode. ## Validation All four backends: correct kernel-only dispatch/combine in eager + CUDA-graph, genuine capture (no eager fallback), consistent per-iteration numbers at --iters-per-graph 1 and 10, on GB200 NVL72. --------- Co-authored-by: Changho Hwang <changhohwang@microsoft.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: seagater <7475084+seagater@users.noreply.github.com> Co-authored-by: Binyang Li <binyli@microsoft.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: Caio Rocha <caiorocha@microsoft.com> Co-authored-by: Copilot Autofix powered by AI <62310815+github-advanced-security[bot]@users.noreply.github.com> Co-authored-by: Caio Rocha <164253795+caiocbr@users.noreply.github.com> Copilot-Session: 15f71a84-4219-4ae9-a87e-e5fab4205de6 Copilot-Session: 59760d2a-e65f-44b6-b2e6-9c271b834d7e
…ified Python EP benchmark (#860) ## Summary Adds the mscclpp high-throughput (HT) EP backend to the unified in-process Python benchmark, on top of the review-comment follow-ups from #858. Benchmarks all five backends through one shared harness with identical dispatch→combine timing: mscclpp LL, **mscclpp HT (TOKEN_MAJOR / RANK_MAJOR)**, NCCL-EP, DeepEP V2, FlashInfer. ## Changes - **New `mscclpp-ht` backend** (`--backend mscclpp-ht`): drives `MoECommunicator` in `HIGH_THROUGHPUT` mode. Supports `--ep-layout token_major` (default) and `rank_major`; `--validate` checks rank-major reduces bit-exactly to the token-major reference. - Lives in `ep_bench_mscclpp.py` alongside the LL `setup_mscclpp` (both drive the same `MoECommunicator` API), sharing `parse_kineto_kernels`. - **CUDA-graph capture enabled** for HT: the harness captures the cached dispatch (`previous_handle=` → no host-side notify wait) + combine as one graph, reusing the harness-owned `_capture_paired_graph` and `--graph-group-size`. - Runtime resolved to the official rank-major HT support (#857) + multi-node NVL/NVLS (#855) from feature/ep; supersedes the branch's local rank-major prototype. ## Validation (GB200 NVL72) - HT cuda-graph captured at **1 and 2 nodes**, both TOKEN_MAJOR and RANK_MAJOR, d8704 e512 k8 t4096 — all single-graph, no eager fallback. Kernel Total (D+C): 1n ~934 µs, 2n ~1406 µs; rank-major ≈ token-major (no measurable overhead). - The four LL/other backends still run eager and cuda-graph correctly through the merged harness (numbers unchanged from #858). --------- Co-authored-by: Changho Hwang <changhohwang@microsoft.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: seagater <7475084+seagater@users.noreply.github.com> Co-authored-by: Binyang Li <binyli@microsoft.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 15f71a84-4219-4ae9-a87e-e5fab4205de6 Copilot-Session: 59760d2a-e65f-44b6-b2e6-9c271b834d7e
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
No description provided.