[MSCCLPP-EP] Add C++ latency expert-parallel runtime by Binyang2014 · Pull Request #895 · microsoft/mscclpp

Binyang Li (Binyang2014) · 2026-09-03T00:49:20Z

Summary

add the public C++ MoERuntime API and latency dispatch/combine request types
add the current low-latency expert-major and rank-major CUDA dispatch/combine kernels
support BF16 and FP8 dispatch plus rank-local-reduce and direct-send combine modes
build the implementation as a standalone mscclpp_ep shared library with no Python dependency
add 8-GPU SM90+ C++ initialization, correctness, and CUDA-graph performance tests

Scope

empty ThroughputDispatchRequest and ThroughputCombineRequest types reserve names for the follow-up implementation
throughput request fields, prepare, notify, host runtime, and CUDA kernels are intentionally excluded
all Python bindings, wrappers, benchmarks, and tests are intentionally excluded for follow-up PRs

C++ test coverage

runtime initialization and unsupported throughput-mode validation
expert-major BF16 and FP8 with rank-local-reduce and direct-send combine
rank-major BF16 with rank-local-reduce and direct-send combine
CUDA-graph D+C performance for expert-major and rank-major, 32 tokens/rank, hidden 7168, 8 experts, top-8
tests skip unless launched with exactly eight local GPUs at SM90 or newer

Validation

./tools/lint.sh
C++-only mscclpp_ep builds for MSCCLPP_GPU_ARCHS=90 and native
Python-enabled configure builds both mscclpp_ep and mscclpp_py
external C++ consumer compiles and links against the installed public API
install contains libmscclpp_ep.so and the public EP headers
mpirun -np 8 ./build-ep-tests/bin/mp_unit_tests --filter=MoERuntimeTest: 4/4 passed
graph D+C: expert-major 51.08 us, rank-major 40.95 us

Feature/EP regression comparison

identical 8-GPU C++ harness, hidden 7168, 32 tokens/rank, BF16, rank-local-reduce, 50 D+C pairs/graph, 100 replays
split branch averages: expert-major 50.47 us, rank-major 40.75 us
feature/ep averages: expert-major 50.58 us, rank-major 41.21 us
six dispatch/combine correctness combinations pass on both branches

Add the public MoERuntime request API, latency dispatch/combine host runtime, and expert-major and rank-major CUDA kernels as a standalone mscclpp_ep library that does not depend on Python bindings. Preserve the unified throughput-facing API surface while rejecting throughput mode until its follow-up implementation. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

azure-pipelines · 2026-09-03T00:49:29Z

Azure Pipelines: There may be pipelines that require an authorized user to comment /azp run to run.

Callers can derive the topology distinction from isAvailable(), numRanks(), and numNvlRanks(); keep the initial C++ latency surface focused on the capability check. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Remove throughput-only prepare and notify methods and strip throughput request fields while retaining empty request types as placeholders for the follow-up implementation. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Replace the duplicate EP CUDA_CHECK macro with MSCCLPP_CUDATHROW and remove the unused throughput-only extended launch helper from the latency implementation. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Remove unused device helpers, align helper naming and inline annotations with MSCCL++ conventions, cap latency routing at top-8, and keep rank-major combine on its TMA path without the unreachable fallback state. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Cover latency runtime initialization, unsupported throughput requests, expert-major and rank-major dispatch/combine correctness across BF16 and FP8, and CUDA-graph D+C performance for 32 tokens per rank. Remove dead and redundant runtime state found during the kernel review. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Use 100 graph replays to reduce variance in the expert-major and rank-major 32-token D+C measurements. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Rename the mode-owned host context types to LatencyRuntimeContext and ThroughputRuntimeContext, and clarify the combine input buffer documentation. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

azure-pipelines · 2026-09-03T16:55:44Z

Azure Pipelines: There may be pipelines that require an authorized user to comment /azp run to run.

Rename the leader-only destination variable from completionRank to dstRank so the rank-major token-store path is easier to follow. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Binyang Li (Binyang2014) · 2026-09-03T17:38:49Z

/azp run mscclpp-ut

azure-pipelines · 2026-09-03T17:39:05Z

Azure Pipelines: Successfully started running 1 pipeline(s).

Mahdieh Ghazi (mahdiehghazim) · 2026-09-04T14:29:06Z

+  auto& context = *latencyContext_;
+  EP_HOST_ASSERT(context.available_);
+  EP_HOST_ASSERT(context.deviceContext_.devicePtr_ != nullptr);
+  EP_HOST_ASSERT(maxTokensPerRank > 0 && maxTokensPerRank <= context.maxTokensPerRank_);


Should we also validate numTokens <= maxTokensPerRank here, consistent with launchLatencyDispatch()?

Mahdieh Ghazi (mahdiehghazim) · 2026-09-04T14:41:14Z

+  EP_HOST_ASSERT(mode == context.combineMode_);
+
+  LatencyStorageLayout allocationLayout(context.symmetricBuffer_, context.maxTokensPerRank_, hidden, context.numRanks_,
+                                        numExperts, numTopk, context.outputLayout_, context.combineMode_);


Are hidden, numExperts, and numTopk intended to vary between requests? I noticed they're specified both when constructing the runtime and in each request, and the request values don't appear to be required to match the runtime values

Mahdieh Ghazi (mahdiehghazim) · 2026-09-04T14:53:49Z

+    EP_HOST_ASSERT(input == allocationLayout.combineRecvBuffer_);
+  }
+
+  const Workload workload{.epoch_ = context.epoch_,


Does combine need to validate that its workload parameters match the preceding dispatch? Since it reuses the dispatch routing metadata/epoch, I'm wondering what happens if numTokens or other workload parameters differ between the two calls.

Mahdieh Ghazi (mahdiehghazim) · 2026-09-04T15:00:15Z

+
+constexpr int NumRanks = 8;
+constexpr int NumExperts = 8;
+constexpr int NumTopk = 8;


Could we add a couple more routing configurations here. we can also do it in a separate PR.

A case with NumExperts > NumTopk (for example, more experts per rank / a larger global expert count), since the current NumExperts == NumTopk == 8 means every token selects every expert.

A nonuniform routing case where different experts receive different numbers of tokens, including potentially some experts receiving no tokens. This would help exercise the dispatch/combine paths under more realistic expert-load distributions.

Mahdieh Ghazi (mahdiehghazim) · 2026-09-04T15:04:13Z

+}
+
+LatencyRuntimeContext::~LatencyRuntimeContext() noexcept(false) {
+  MSCCLPP_CUDATHROW(cudaDeviceSynchronize());


Is the device-wide synchronization intentional here? cudaDeviceSynchronize() also waits for unrelated GPU work. Could we synchronize only the EP work before freeing the runtime resources?

Mahdieh Ghazi (mahdiehghazim) · 2026-09-04T15:45:55Z

+  EP_HOST_ASSERT(symmetricBufferBytes_ % BufferAlignmentBytes == 0);
+  EP_HOST_ASSERT(maxTokensPerRank > 0);
+  EP_HOST_ASSERT(numExperts > 0 && numExperts % numRanks_ == 0);
+  EP_HOST_ASSERT(numTopk > 0 && numTopk <= MaxNumTopk);


Should we validate hidden against the supported latency hidden sizes here? The runtime allocates its buffers during construction, but unsupported hidden sizes are only rejected later when dispatch selects the kernel.

Mahdieh Ghazi (mahdiehghazim) · 2026-09-04T15:55:12Z

+  /// Global expert count.
+  int numExperts;
+  /// Expert ID used for invalid rank-major entries.
+  int invalidTokenExpertId;


Could we document the expected value/range for invalidTokenExpertId?

Binyang Li (Binyang2014) marked this pull request as draft September 3, 2026 02:41

Binyang Li (Binyang2014) and others added 4 commits September 3, 2026 02:49

Remove redundant EP internode availability API

fd153f6

Callers can derive the topology distinction from isAvailable(), numRanks(), and numNvlRanks(); keep the initial C++ latency surface focused on the capability check. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Limit initial EP API to latency mode

2b3208b

Remove throughput-only prepare and notify methods and strip throughput request fields while retaining empty request types as placeholders for the follow-up implementation. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Reuse MSCCL++ CUDA error handling in EP

ff69416

Replace the duplicate EP CUDA_CHECK macro with MSCCLPP_CUDATHROW and remove the unused throughput-only extended launch helper from the latency implementation. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Binyang Li (Binyang2014) changed the title ~~Add C++ latency expert-parallel runtime~~ [MSCCLPP-EP] Add C++ latency expert-parallel runtime Sep 3, 2026

Binyang Li (Binyang2014) and others added 3 commits September 3, 2026 05:41

Increase EP graph benchmark iterations

72a9570

Use 100 graph replays to reduce variance in the expert-major and rank-major 32-token D+C measurements. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Clarify EP runtime context naming

a910c1e

Rename the mode-owned host context types to LatencyRuntimeContext and ThroughputRuntimeContext, and clarify the combine input buffer documentation. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Binyang Li (Binyang2014) marked this pull request as ready for review September 3, 2026 16:55

Binyang Li (Binyang2014) requested a review from a team September 3, 2026 16:56

Clarify rank-major store destination

49075b5

Rename the leader-only destination variable from completionRank to dstRank so the rank-major token-store path is easier to follow. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Mahdieh Ghazi (mahdiehghazim) reviewed Sep 4, 2026

View reviewed changes

Navigation Menu

Uh oh!

[MSCCLPP-EP] Add C++ latency expert-parallel runtime - #895

[MSCCLPP-EP] Add C++ latency expert-parallel runtime#895
Binyang Li (Binyang2014) wants to merge 9 commits into
mainfrom
binyli/ep-latency-cpp

Binyang Li (Binyang2014) commented Sep 3, 2026 •

edited

Loading

Uh oh!

azure-pipelines Bot commented Sep 3, 2026

Uh oh!

azure-pipelines Bot commented Sep 3, 2026

Uh oh!

Binyang Li (Binyang2014) commented Sep 3, 2026

Uh oh!

azure-pipelines Bot commented Sep 3, 2026

Uh oh!

Mahdieh Ghazi (mahdiehghazim) Sep 4, 2026

Uh oh!

Mahdieh Ghazi (mahdiehghazim) Sep 4, 2026

Uh oh!

Mahdieh Ghazi (mahdiehghazim) Sep 4, 2026

Uh oh!

Mahdieh Ghazi (mahdiehghazim) Sep 4, 2026

Uh oh!

Mahdieh Ghazi (mahdiehghazim) Sep 4, 2026

Uh oh!

Mahdieh Ghazi (mahdiehghazim) Sep 4, 2026

Uh oh!

Mahdieh Ghazi (mahdiehghazim) Sep 4, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Uh oh!

Conversation

Binyang Li (Binyang2014) commented Sep 3, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Scope

C++ test coverage

Validation

Feature/EP regression comparison

Uh oh!

azure-pipelines Bot commented Sep 3, 2026

Uh oh!

azure-pipelines Bot commented Sep 3, 2026

Uh oh!

Binyang Li (Binyang2014) commented Sep 3, 2026

Uh oh!

azure-pipelines Bot commented Sep 3, 2026

Uh oh!

Mahdieh Ghazi (mahdiehghazim) Sep 4, 2026

Choose a reason for hiding this comment

Uh oh!

Mahdieh Ghazi (mahdiehghazim) Sep 4, 2026

Choose a reason for hiding this comment

Uh oh!

Mahdieh Ghazi (mahdiehghazim) Sep 4, 2026

Choose a reason for hiding this comment

Uh oh!

Mahdieh Ghazi (mahdiehghazim) Sep 4, 2026

Choose a reason for hiding this comment

Uh oh!

Mahdieh Ghazi (mahdiehghazim) Sep 4, 2026

Choose a reason for hiding this comment

Uh oh!

Mahdieh Ghazi (mahdiehghazim) Sep 4, 2026

Choose a reason for hiding this comment

Uh oh!

Mahdieh Ghazi (mahdiehghazim) Sep 4, 2026

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Binyang Li (Binyang2014) commented Sep 3, 2026 •

edited

Loading