kvcache-ai/Mooncake

Fork 0

mirror of https://github.com/kvcache-ai/Mooncake.git synced 2026-11-05 01:17:55 +00:00

[GH-ISSUE #2630] [RFC]: NCCL device API in Mooncake TENT #689

New issue

Closed

opened 2026-10-29 21:18:11 +00:00 by bsakaguchi · 10 comments

bsakaguchi commented

2026-10-29 21:18:11 +00:00

Owner

Originally created by @xiaofanl-nvidia on GitHub (Jun 26, 2026).
Original GitHub issue: https://github.com/kvcache-ai/Mooncake/issues/2630

Changes proposed

Proposal and motivation

Add a TENT NCCL transport backend built on NCCL's device API.
This proposal is aligned with roadmap here: https://github.com/kvcache-ai/Mooncake/issues/1883 and RFC here: https://github.com/kvcache-ai/Mooncake/issues/2123 by offering a widely adopted solution for device-initiated communication, with easy-to-maintain abstractions and production quality validation.

Leverage NCCL's device API

NCCL's device API abstracts hardware topology and features into a clean, cohesive API suite. By adopting it, Mooncake gains access to capabilities such as NVLink SHARP (NVSwitch multicast), GPU-initiated networking (GIN) over multiple network backends (IB, RoCE, EFA, and more), CFT (Compute Fabric Transport), Counted Writes, CuteDSL/cuTiles integration, and others, all without the effort of rebuilding from scratch.

Fine-grained comm/compute fusion

Supporting device-initiated communication allows per-layer, fine-grained data transfer instead of relying on the CPU to initiate network traffic. Beyond the usual benefits — fewer CPU threads, fewer kernel launches, and fewer staging buffers — it also makes comm/compute fusion possible.

Build custom collectives

NCCL's own collectives are built on this same device-API surface, so integrating it gives Mooncake the same proven foundation for building high-performance custom collectives. Ready-to-use reference implementations live at https://github.com/NVIDIA/nccl/tree/master/contrib, with minimal examples at https://github.com/NVIDIA/nccl/tree/master/docs/examples/06_device_api.

Leverage DeepEPv2 / nccl-ep

DeepEPv2 and nccl-ep make heavy use of NCCL's device API, so integrating it lets Mooncake gain the same benefits and build its own EP solutions on the same foundation.

Leverage the NCCL ecosystem

Taking on NCCL as a dependency lets Mooncake easily use other capabilities from the ecosystem — high-performance host-side collectives, one-sided put/signal APIs, immediate support for new NVIDIA hardware, topology, and CUDA features, and other NCCL extensions such as nccl-ep, nccl-m2n, nccl-ubx, and more.

Implementation details

Cross-node transfers use one communicator per peer session with multiple GIN contexts for striping. WRITE uses the device-side ncclGin::put with signal/ack, and READ uses the device-side ncclGin::get.
For same-node P2P / LSA-connected peers, the backend bypasses GIN and uses NCCL's LSA peer-pointer APIs (ncclGetPeerDevicePointer) with CUDA D2D copies, giving an NVLink fast path for both READ and WRITE.

Design notes

Requires NCCL 2.30+ for the device-side GIN GET (which also requires GDAKI and is not supported on Ampere or directNIC platforms).
Uses one communicator per peer session with configurable GIN contexts.
Detects LSA membership from NCCL teams after communicator init.

POC API and integration example

https://github.com/akhillanger/Mooncake/pull/1 — NCCL backend for TENT

Note: this is a simple example and may not be the best design, since it uses a host API path to launch a kernel that then uses the device API. Ideally we should either build more fused ops on the device-API backend or extend TENT to invoke the device API directly. We're happy to work with the Mooncake developers to design such APIs to allow NCCL device API to be used by Mooncake applications seamlessly.

Before submitting a new issue...

Make sure you already searched for relevant issues and read the documentation

Originally created by @xiaofanl-nvidia on GitHub (Jun 26, 2026). Original GitHub issue: https://github.com/kvcache-ai/Mooncake/issues/2630 ### Changes proposed ## Proposal and motivation Add a TENT NCCL transport backend built on NCCL's device API. This proposal is aligned with roadmap here: https://github.com/kvcache-ai/Mooncake/issues/1883 and RFC here: https://github.com/kvcache-ai/Mooncake/issues/2123 by offering a widely adopted solution for device-initiated communication, with easy-to-maintain abstractions and production quality validation. ### Leverage NCCL's device API NCCL's device API abstracts hardware topology and features into a clean, cohesive API suite. By adopting it, Mooncake gains access to capabilities such as NVLink SHARP (NVSwitch multicast), GPU-initiated networking (GIN) over multiple network backends (IB, RoCE, EFA, and more), CFT (Compute Fabric Transport), Counted Writes, CuteDSL/cuTiles integration, and others, all without the effort of rebuilding from scratch. ### Fine-grained comm/compute fusion Supporting device-initiated communication allows per-layer, fine-grained data transfer instead of relying on the CPU to initiate network traffic. Beyond the usual benefits — fewer CPU threads, fewer kernel launches, and fewer staging buffers — it also makes comm/compute fusion possible. ### Build custom collectives NCCL's own collectives are built on this same device-API surface, so integrating it gives Mooncake the same proven foundation for building high-performance custom collectives. Ready-to-use reference implementations live at https://github.com/NVIDIA/nccl/tree/master/contrib, with minimal examples at https://github.com/NVIDIA/nccl/tree/master/docs/examples/06_device_api. ### Leverage DeepEPv2 / nccl-ep DeepEPv2 and nccl-ep make heavy use of NCCL's device API, so integrating it lets Mooncake gain the same benefits and build its own EP solutions on the same foundation. ### Leverage the NCCL ecosystem Taking on NCCL as a dependency lets Mooncake easily use other capabilities from the ecosystem — high-performance host-side collectives, one-sided put/signal APIs, immediate support for new NVIDIA hardware, topology, and CUDA features, and other NCCL extensions such as nccl-ep, nccl-m2n, nccl-ubx, and more. ## Implementation details Cross-node transfers use one communicator per peer session with multiple GIN contexts for striping. WRITE uses the device-side ncclGin::put with signal/ack, and READ uses the device-side ncclGin::get. For same-node P2P / LSA-connected peers, the backend bypasses GIN and uses NCCL's LSA peer-pointer APIs (ncclGetPeerDevicePointer) with CUDA D2D copies, giving an NVLink fast path for both READ and WRITE. ### Design notes - Requires NCCL 2.30+ for the device-side GIN GET (which also requires GDAKI and is not supported on Ampere or directNIC platforms). - Uses one communicator per peer session with configurable GIN contexts. - Detects LSA membership from NCCL teams after communicator init. ## POC API and integration example https://github.com/akhillanger/Mooncake/pull/1 — NCCL backend for TENT Note: this is a simple example and may not be the best design, since it uses a host API path to launch a kernel that then uses the device API. Ideally we should either build more fused ops on the device-API backend or extend TENT to invoke the device API directly. We're happy to work with the Mooncake developers to design such APIs to allow NCCL device API to be used by Mooncake applications seamlessly. ### Before submitting a new issue... - [x] Make sure you already searched for relevant issues and read the [documentation](https://kvcache-ai.github.io/Mooncake/)

bsakaguchi closed this issue

2026-10-29 21:18:11 +00:00

bsakaguchi commented

2026-10-29 21:18:12 +00:00

Author

Owner

@github-actions[bot] commented on GitHub (Jun 26, 2026):

Thanks for opening this issue, @xiaofanl-nvidia!

Field	Value
Issue	#2630
GitHub user ID	`204747896`
Reporter	@xiaofanl-nvidia

A maintainer will triage this when possible. To help us respond faster, please include:

Mooncake version or commit SHA
Environment (OS, CUDA/driver, RDMA stack if relevant)
Steps to reproduce and expected vs. actual behavior

Useful links: Documentation · Contributing guide

This message was posted automatically by the issue bot.

@github-actions[bot] commented on GitHub (Jun 26, 2026): Thanks for opening this issue, @xiaofanl-nvidia! | Field | Value | |-------|-------| | **Issue** | #2630 | | **GitHub user ID** | `204747896` | | **Reporter** | @xiaofanl-nvidia | A maintainer will triage this when possible. To help us respond faster, please include: - Mooncake version or commit SHA - Environment (OS, CUDA/driver, RDMA stack if relevant) - Steps to reproduce and expected vs. actual behavior Useful links: [Documentation](https://kvcache-ai.github.io/Mooncake/) · [Contributing guide](https://github.com/kvcache-ai/Mooncake/blob/main/CONTRIBUTING.md) > This message was posted automatically by the issue bot.

bsakaguchi commented

2026-10-29 21:18:12 +00:00

Author

Owner

@alogfans commented on GitHub (Jun 29, 2026):

Thanks for your proposal! I have added some comments in https://github.com/akhillanger/Mooncake/pull/1 for reference.

@alogfans commented on GitHub (Jun 29, 2026): Thanks for your proposal! I have added some comments in https://github.com/akhillanger/Mooncake/pull/1 for reference.

bsakaguchi commented

2026-10-29 21:18:13 +00:00

Author

Owner

@staryxchen commented on GitHub (Jun 29, 2026):

Thanks for the detailed proposal and POC @xiaofanl-nvidia! We've reviewed the implementation in https://github.com/akhillanger/Mooncake/pull/1 and want to share some architectural thoughts.

Analysis of the current POC approach

The POC validates that NCCL device API (GIN put/get, LSA peer pointers) can be integrated into TENT — that's valuable. However, we believe the current approach has a fundamental shape mismatch:

The existing TENT Transport interface is host-initiated — CPU submits tasks, polls for completion via getTransferStatus. The POC wraps device-initiated capabilities inside this host-initiated model, which reintroduces the overhead that device-initiated communication is meant to eliminate:

Per-transfer host ThreadPool dispatch
Per-WRITE cross-node RPC on the data path (postRemoteWaitSignal)
Lazy window initialization with dummy buffer allocation (GPU memory waste)
cudaEvent polling for completion (pulling device-side events back to host)

As the POC authors noted: "Ideally we should either build more fused ops on the device-API backend or extend TENT to invoke the device API directly." We agree this is the right direction.

Proposed direction: a separate device-initiated transport abstraction

Rather than fitting device-initiated communication into the existing Transport interface, we propose a new DeviceTransport category that respects the different execution model:

Design principles:

Separate control plane (host) from data plane (device) — connection setup, window registration can be slow/RPC-heavy; data path must be pure device-side with zero host intervention
Eager resource setup — windows registered at buffer-add time, not lazily on first transfer
Expose device-side primitives — user kernels call put/get/signal/wait directly, enabling true comm/compute fusion

Sketch:

// Host-side: setup only (can be slow, can RPC)
class DeviceTransport {
    Status install(config);
    Status connect(peer_info, DeviceChannelHandle& handle);
    Status register_window(void* buf, size_t len, window_t& win);
    Status deregister_window(window_t win);
    Status disconnect(DeviceChannelHandle handle);
    TentDeviceChannel get_device_channel(DeviceChannelHandle handle);
};

// Device-side: data path (inlined in user kernels)
struct TentDeviceChannel {
    __device__ void put(window_t dst, size_t dst_off, window_t src, size_t src_off, size_t len);
    __device__ void get(window_t remote, size_t remote_off, window_t local, size_t local_off, size_t len);
    __device__ void signal(int peer, int signal_id);
    __device__ void wait_signal(int signal_id, uint64_t expected);
    __device__ void flush();
};

Usage (what comm/compute fusion actually looks like):

__global__ void fused_kernel(..., TentDeviceChannel chan) {
    compute_layer(input, output);
    chan.put(remote_win, off, local_win, 0, size);
    chan.signal(peer, sig_id);
    compute_next_layer(...);       // overlaps with transfer
    chan.wait_signal(sig_id, val); // wait only when needed
}

Relationship to existing Transport:

TENT
 ├── Transport (host-initiated): RDMA, TCP, NVLink, ...
 └── DeviceTransport (device-initiated): NCCL-GIN, (future: CX-8, ...)

Both share the control plane for peer discovery (segment registry, metadata store), but data planes are completely independent.

Open design questions

We'd like to discuss these with the community before committing to a design:

Pure device-side vs hybrid — Some use cases are purely device-initiated (EP fusion kernels). Others are host-decided but device-executed (e.g., KV cache prefetch where the host decides what to transfer but the GPU does the actual move). Should DeviceTransport cover both, perhaps by providing framework-managed kernels for the hybrid case?
Abstraction timing — If NCCL GIN is the only backend in the short term, should we define the DeviceTransport interface now, or start with a thin NCCL-specific wrapper and extract the interface when a second backend (e.g., CX-8 device-side verbs) arrives?
Relationship to EP frameworks — DeepEPv2/nccl-ep already build scheduling logic on top of NCCL device API. Should TENT's DeviceTransport sit below them (providing primitives) or above (providing higher-level abstractions)?

We're happy to collaborate on the design. Would appreciate thoughts from @xiaofanl-nvidia and the broader community on these questions.

@staryxchen commented on GitHub (Jun 29, 2026): Thanks for the detailed proposal and POC @xiaofanl-nvidia! We've reviewed the implementation in https://github.com/akhillanger/Mooncake/pull/1 and want to share some architectural thoughts. ### Analysis of the current POC approach The POC validates that NCCL device API (GIN put/get, LSA peer pointers) can be integrated into TENT — that's valuable. However, we believe the current approach has a fundamental shape mismatch: **The existing TENT `Transport` interface is host-initiated** — CPU submits tasks, polls for completion via `getTransferStatus`. The POC wraps device-initiated capabilities inside this host-initiated model, which reintroduces the overhead that device-initiated communication is meant to eliminate: - Per-transfer host ThreadPool dispatch - Per-WRITE cross-node RPC on the data path (`postRemoteWaitSignal`) - Lazy window initialization with dummy buffer allocation (GPU memory waste) - cudaEvent polling for completion (pulling device-side events back to host) As the POC authors noted: *"Ideally we should either build more fused ops on the device-API backend or extend TENT to invoke the device API directly."* We agree this is the right direction. ### Proposed direction: a separate device-initiated transport abstraction Rather than fitting device-initiated communication into the existing `Transport` interface, we propose a new `DeviceTransport` category that respects the different execution model: **Design principles:** 1. **Separate control plane (host) from data plane (device)** — connection setup, window registration can be slow/RPC-heavy; data path must be pure device-side with zero host intervention 2. **Eager resource setup** — windows registered at buffer-add time, not lazily on first transfer 3. **Expose device-side primitives** — user kernels call `put/get/signal/wait` directly, enabling true comm/compute fusion **Sketch:** ```cpp // Host-side: setup only (can be slow, can RPC) class DeviceTransport { Status install(config); Status connect(peer_info, DeviceChannelHandle& handle); Status register_window(void* buf, size_t len, window_t& win); Status deregister_window(window_t win); Status disconnect(DeviceChannelHandle handle); TentDeviceChannel get_device_channel(DeviceChannelHandle handle); }; // Device-side: data path (inlined in user kernels) struct TentDeviceChannel { __device__ void put(window_t dst, size_t dst_off, window_t src, size_t src_off, size_t len); __device__ void get(window_t remote, size_t remote_off, window_t local, size_t local_off, size_t len); __device__ void signal(int peer, int signal_id); __device__ void wait_signal(int signal_id, uint64_t expected); __device__ void flush(); }; ``` **Usage (what comm/compute fusion actually looks like):** ```cuda __global__ void fused_kernel(..., TentDeviceChannel chan) { compute_layer(input, output); chan.put(remote_win, off, local_win, 0, size); chan.signal(peer, sig_id); compute_next_layer(...); // overlaps with transfer chan.wait_signal(sig_id, val); // wait only when needed } ``` **Relationship to existing Transport:** ``` TENT ├── Transport (host-initiated): RDMA, TCP, NVLink, ... └── DeviceTransport (device-initiated): NCCL-GIN, (future: CX-8, ...) ``` Both share the control plane for peer discovery (segment registry, metadata store), but data planes are completely independent. ### Open design questions We'd like to discuss these with the community before committing to a design: 1. **Pure device-side vs hybrid** — Some use cases are purely device-initiated (EP fusion kernels). Others are host-decided but device-executed (e.g., KV cache prefetch where the host decides *what* to transfer but the GPU does the actual move). Should `DeviceTransport` cover both, perhaps by providing framework-managed kernels for the hybrid case? 2. **Abstraction timing** — If NCCL GIN is the only backend in the short term, should we define the `DeviceTransport` interface now, or start with a thin NCCL-specific wrapper and extract the interface when a second backend (e.g., CX-8 device-side verbs) arrives? 3. **Relationship to EP frameworks** — DeepEPv2/nccl-ep already build scheduling logic on top of NCCL device API. Should TENT's `DeviceTransport` sit below them (providing primitives) or above (providing higher-level abstractions)? We're happy to collaborate on the design. Would appreciate thoughts from @xiaofanl-nvidia and the broader community on these questions.

bsakaguchi commented

2026-10-29 21:18:13 +00:00

Author

Owner

@xiaofanl-nvidia commented on GitHub (Jun 30, 2026):

Thanks for the detailed review. I'm on board with a separate DeviceTransport interface. The three principles you list map well onto how NCCL's device API is intended to be used: the host sets up the communicator and registers windows, and the device path does put/get/signal/wait with no host intervention (true for GIN's GDAKI mode; we have a CPU-proxy mode for compatibility reason).

Here are my quick thoughts on the 3 questions you posted above:

Pure vs hybrid: Both seems useful to provide from Mooncake -

The pure device-initiated case (e.g. EP) is the user kernel calling the device APIs directly.
The hybrid case (host decides what to move, GPU executes) could be a thin host-launched, optimized transfer kernel using the same device API surface.
I'd recommend Mooncake shipping a small set of such kernels for the hybrid path (e.g. a generic device copy / scatter-gather) so frameworks that just want fast data movement don't have to write CUDA, while leaving the pure path open for frameworks that want fully custom fusion ops.
Note that this is very similar to the strategy in NCCL itself - NCCL ships both high-performance collective ops and the lower level device APIs which we use to build the collectives.

Abstraction timing: put/get/signal/wait/flush is a fairly universal one-sided device-primitive set, so it should generalize to other backends without much interface churn, which makes defining it early low-risk.
BTW - NCCL device API's GIN layer already supports CX8. Could you share more info about "CX-8 device-side verbs"?
Relationship to EP frameworks: I'd imagine we can put DeviceTransport below the EP frameworks: it provides the primitives, and Mooncake-EP layer can build dispatch/combine scheduling on top. From what I can see in Mooncake-EP today, it uses ibgda similar to DeepEPv1. Once we have the NCCL GIN layer, we can work together to port all the elasticity features onto Mooncake-EPv2 and leverage NCCL GIN. Recently we added timeout to all the GIN APIs to prepare for this, and we also ported the ActiveMask idea to nccl-ep to prove the path.

One related topic to get your feedback: NCCL depends on the communicator abstraction, which requires collective bootstrapping, instead of pair-wise setup. That's deliberate - it lets the runtime exploit global topology (PXN, symmetric heap, NVLS), which is where the performance for collectives and EP comes from. The trade-off is heavier setup and dynamic membership needing explicit support. Do you foresee additional constructs/features we should consider adding to NCCL to make this fit well with Mooncake's requirements in this area?

In terms of concrete next steps, we will modify the PoC based on the new direction and share perf data once it's carefully validated.

@xiaofanl-nvidia commented on GitHub (Jun 30, 2026): Thanks for the detailed review. I'm on board with a separate DeviceTransport interface. The three principles you list map well onto how NCCL's device API is intended to be used: the host sets up the communicator and registers windows, and the device path does put/get/signal/wait with no host intervention (true for GIN's GDAKI mode; we have a CPU-proxy mode for compatibility reason). Here are my quick thoughts on the 3 questions you posted above: 1. Pure vs hybrid: Both seems useful to provide from Mooncake - - The pure device-initiated case (e.g. EP) is the user kernel calling the device APIs directly. - The hybrid case (host decides what to move, GPU executes) could be a thin host-launched, optimized transfer kernel using the same device API surface. I'd recommend Mooncake shipping a small set of such kernels for the hybrid path (e.g. a generic device copy / scatter-gather) so frameworks that just want fast data movement don't have to write CUDA, while leaving the pure path open for frameworks that want fully custom fusion ops. Note that this is very similar to the strategy in NCCL itself - NCCL ships both high-performance collective ops and the lower level device APIs which we use to build the collectives. 2. Abstraction timing: put/get/signal/wait/flush is a fairly universal one-sided device-primitive set, so it should generalize to other backends without much interface churn, which makes defining it early low-risk. BTW - NCCL device API's GIN layer already supports CX8. Could you share more info about "CX-8 device-side verbs"? 3. Relationship to EP frameworks: I'd imagine we can put DeviceTransport below the EP frameworks: it provides the primitives, and Mooncake-EP layer can build dispatch/combine scheduling on top. From what I can see in Mooncake-EP today, it uses ibgda similar to DeepEPv1. Once we have the NCCL GIN layer, we can work together to port all the elasticity features onto Mooncake-EPv2 and leverage NCCL GIN. Recently we added timeout to all the GIN APIs to prepare for this, and we also ported the ActiveMask idea to nccl-ep to prove the path. One related topic to get your feedback: NCCL depends on the communicator abstraction, which requires collective bootstrapping, instead of pair-wise setup. That's deliberate - it lets the runtime exploit global topology (PXN, symmetric heap, NVLS), which is where the performance for collectives and EP comes from. The trade-off is heavier setup and dynamic membership needing explicit support. Do you foresee additional constructs/features we should consider adding to NCCL to make this fit well with Mooncake's requirements in this area? In terms of concrete next steps, we will modify the PoC based on the new direction and share perf data once it's carefully validated.

bsakaguchi commented

2026-10-29 21:18:13 +00:00

Author

Owner

@staryxchen commented on GitHub (Jun 30, 2026):

Thanks for the thoughtful response! Some follow-up thoughts:

1. Hybrid path placement

Agreed that the hybrid scenario (host decides what to transfer, GPU executes) is a real need. However, I'd lean toward placing it under the existing host-initiated Transport as an NCCL host-API backend, rather than inside DeviceTransport.

Rationale: DeviceTransport's core value proposition is "zero host intervention on the data path." The hybrid scenario is still fundamentally host-initiated (host constructs descriptor → submit); it's only the execution layer that moves to GPU. Mixing both under one abstraction blurs DeviceTransport's contract — when a user sees DeviceTransport, they should immediately know "this is for calling directly inside my kernel."

Concretely:

DeviceTransport: purely device-initiated, exposes TentDeviceChannel for user kernels to call inline
Existing host-initiated Transport + NCCL backend: host-initiated, uses NCCL host API (ncclSend/Recv or host-launched GIN kernels) under the hood, preserves submit/poll semantics for the upper layer

Both can share the same NCCL communicator and window resources — only the entry point differs. This keeps the layering clean and aligns with NCCL's own strategy of offering both collective ops and lower-level device APIs.

2. Abstraction timing & CX-8 clarification

Agreed that put/get/signal/wait/flush as data-plane primitives are stable enough to define now.

On CX-8: you're right — NCCL GIN already abstracts the underlying hardware including CX-8. We mentioned "CX-8 device-side verbs" as a potential second backend, but since GIN already covers it, there's no clear need to bypass NCCL for bare device verbs. This further supports your point: NCCL GIN is the only backend in the near term, so the interface can be designed close to GIN's semantics without over-generalizing.

That said, the control plane (setup semantics) needs more discussion — which leads to the communicator topic below.

3. EP layering

Agreed that DeviceTransport should sit below the EP framework and provide primitives. One practical concern: EP performance optimizations are heavily coupled to NCCL's internal topology awareness (NVSwitch multicast, GIN connection placement, etc.). The DeviceTransport abstraction should be thin enough not to hide these backend-specific capabilities. We may need some form of backend-specific escape hatch so the EP layer can access topology info or NCCL-specific configuration when needed.

That said, we are not the maintainers of Mooncake-EP. @UNIDY2002 — would appreciate your thoughts on how DeviceTransport primitives would fit beneath the EP layer, and whether a thin abstraction vs direct NCCL access works for the EP use case.

4. Communicator model — the core tension

This is the most critical topic. Frankly, there's a fundamental tension between NCCL's communicator model and Mooncake's architecture:

Mooncake's core scenario is distributed KV cache storage, which is characterized by:

Nodes dynamically joining/leaving (elastic scaling, fault recovery)
Communication topology is pair-wise and changes at runtime (which prefill node sends KV cache to which decode node is decided by the scheduler on the fly)
Connections established/torn down on demand

NCCL's communicator model:

N nodes collectively initialize with fixed membership
Runtime exploits global topology for routing optimization
Dynamic membership changes are expensive (requires rebuild or comm split)

The POC uses a 2-rank communicator per peer session — essentially using collective machinery to simulate pair-wise connections. This works but isn't elegant: one communicator per pair means O(N²) communicators, each going through full collective bootstrap.

What would be ideal for Mooncake:

Lightweight communicator creation/teardown — ncclCommInitRank currently involves global barriers and topology discovery, which is heavy for 2-rank pair-wise sessions. Is there a lighter-weight path?
Dynamic membership — nodes joining/leaving without rebuilding the entire communicator. Does NCCL have a roadmap for this (similar to MPI dynamic process management)?
Lazy connection establishment — GIN connections established only when two nodes actually need to communicate, rather than at communicator init time for the full mesh.

Mooncake likely needs to distinguish two usage patterns:

KV cache transfers: pair-wise, dynamic, needs lightweight sessions
EP dispatch/combine: collective, relatively stable membership, suits N-rank communicator

These may require different setup paths but could share the same device-side data-plane primitives.

Next steps

Looking forward to the revised PoC and perf data. Before that, would it be useful to align on the communicator setup approach first? This determines the DeviceTransport control-plane design. If possible, we'd also welcome any insight into NCCL's plans around dynamic membership / lightweight communicator creation — it's critical to our architectural choices here.

@staryxchen commented on GitHub (Jun 30, 2026): Thanks for the thoughtful response! Some follow-up thoughts: ### 1. Hybrid path placement Agreed that the hybrid scenario (host decides what to transfer, GPU executes) is a real need. However, I'd lean toward placing it under the existing **host-initiated Transport** as an NCCL host-API backend, rather than inside DeviceTransport. Rationale: DeviceTransport's core value proposition is "zero host intervention on the data path." The hybrid scenario is still fundamentally host-initiated (host constructs descriptor → submit); it's only the execution layer that moves to GPU. Mixing both under one abstraction blurs DeviceTransport's contract — when a user sees DeviceTransport, they should immediately know "this is for calling directly inside my kernel." Concretely: - **DeviceTransport**: purely device-initiated, exposes `TentDeviceChannel` for user kernels to call inline - **Existing host-initiated Transport + NCCL backend**: host-initiated, uses NCCL host API (ncclSend/Recv or host-launched GIN kernels) under the hood, preserves submit/poll semantics for the upper layer Both can share the same NCCL communicator and window resources — only the entry point differs. This keeps the layering clean and aligns with NCCL's own strategy of offering both collective ops and lower-level device APIs. ### 2. Abstraction timing & CX-8 clarification Agreed that `put/get/signal/wait/flush` as data-plane primitives are stable enough to define now. On CX-8: you're right — NCCL GIN already abstracts the underlying hardware including CX-8. We mentioned "CX-8 device-side verbs" as a potential second backend, but since GIN already covers it, there's no clear need to bypass NCCL for bare device verbs. This further supports your point: NCCL GIN is the only backend in the near term, so the interface can be designed close to GIN's semantics without over-generalizing. That said, the control plane (setup semantics) needs more discussion — which leads to the communicator topic below. ### 3. EP layering Agreed that DeviceTransport should sit below the EP framework and provide primitives. One practical concern: EP performance optimizations are heavily coupled to NCCL's internal topology awareness (NVSwitch multicast, GIN connection placement, etc.). The DeviceTransport abstraction should be thin enough not to hide these backend-specific capabilities. We may need some form of backend-specific escape hatch so the EP layer can access topology info or NCCL-specific configuration when needed. That said, we are not the maintainers of Mooncake-EP. @UNIDY2002 — would appreciate your thoughts on how DeviceTransport primitives would fit beneath the EP layer, and whether a thin abstraction vs direct NCCL access works for the EP use case. ### 4. Communicator model — the core tension This is the most critical topic. Frankly, there's a fundamental tension between NCCL's communicator model and Mooncake's architecture: **Mooncake's core scenario is distributed KV cache storage**, which is characterized by: - Nodes dynamically joining/leaving (elastic scaling, fault recovery) - Communication topology is pair-wise and changes at runtime (which prefill node sends KV cache to which decode node is decided by the scheduler on the fly) - Connections established/torn down on demand **NCCL's communicator model**: - N nodes collectively initialize with fixed membership - Runtime exploits global topology for routing optimization - Dynamic membership changes are expensive (requires rebuild or comm split) The POC uses a 2-rank communicator per peer session — essentially using collective machinery to simulate pair-wise connections. This works but isn't elegant: one communicator per pair means O(N²) communicators, each going through full collective bootstrap. What would be ideal for Mooncake: 1. **Lightweight communicator creation/teardown** — `ncclCommInitRank` currently involves global barriers and topology discovery, which is heavy for 2-rank pair-wise sessions. Is there a lighter-weight path? 2. **Dynamic membership** — nodes joining/leaving without rebuilding the entire communicator. Does NCCL have a roadmap for this (similar to MPI dynamic process management)? 3. **Lazy connection establishment** — GIN connections established only when two nodes actually need to communicate, rather than at communicator init time for the full mesh. Mooncake likely needs to distinguish two usage patterns: - **KV cache transfers**: pair-wise, dynamic, needs lightweight sessions - **EP dispatch/combine**: collective, relatively stable membership, suits N-rank communicator These may require different setup paths but could share the same device-side data-plane primitives. ### Next steps Looking forward to the revised PoC and perf data. Before that, would it be useful to align on the communicator setup approach first? This determines the DeviceTransport control-plane design. If possible, we'd also welcome any insight into NCCL's plans around dynamic membership / lightweight communicator creation — it's critical to our architectural choices here.

bsakaguchi commented

2026-10-29 21:18:13 +00:00

Author

Owner

@UNIDY2002 commented on GitHub (Jun 30, 2026):

Thanks for bringing this up. I think this discussion aligns closely with the Device API / DeviceTransport direction that is already underway in Mooncake.

We have already landed the initial Device API support in TE (#2333) and switched Mooncake-EP to use that path (#2382). The goal there is exactly to expose device-side communication primitives to EP while keeping the control plane reusable.

My suggestion would be to position NCCL device API as one backend implementation under the same DeviceTransport / Device API framework, alongside the existing IBGDA / P2P / proxy-based backends. I think this gives Mooncake a clean multi-backend device communication layer and makes it easier for Mooncake-EP to benefit from different device-side transports through one common abstraction.

So I’m very supportive of integrating NCCL device API, and I’d be happy to help align it with the existing DeviceTransport work so we can build this out together.

@UNIDY2002 commented on GitHub (Jun 30, 2026): Thanks for bringing this up. I think this discussion aligns closely with the Device API / DeviceTransport direction that is already underway in Mooncake. We have already landed the initial Device API support in TE ([#2333](https://github.com/kvcache-ai/Mooncake/pull/2333)) and switched Mooncake-EP to use that path ([#2382](https://github.com/kvcache-ai/Mooncake/pull/2382)). The goal there is exactly to expose device-side communication primitives to EP while keeping the control plane reusable. My suggestion would be to position NCCL device API as one backend implementation under the same DeviceTransport / Device API framework, alongside the existing IBGDA / P2P / proxy-based backends. I think this gives Mooncake a clean multi-backend device communication layer and makes it easier for Mooncake-EP to benefit from different device-side transports through one common abstraction. So I’m very supportive of integrating NCCL device API, and I’d be happy to help align it with the existing DeviceTransport work so we can build this out together.

bsakaguchi commented

2026-10-29 21:18:13 +00:00

Author

Owner

@xiaofanl-nvidia commented on GitHub (Jul 4, 2026):

Just a quick update that we have started prototyping the integration to add NCCL device API to DeviceTransport under TE.
We should be able to share a preliminary (but working) PR by next week. @UNIDY2002 we'll appreciate your review & hands-on help if you are available. We want to lay the groundwork before asking for your time :)

Also, I saw the RFE to add more collectives in Mooncake: https://github.com/kvcache-ai/Mooncake/issues/2716
Given how expensive it is to optimize and validate these ops, we want to encourage Mooncake team to leverage NCCL device/host APIs for this work and review NCCL's implementations and examples, including interesting community contributed work that are quite novel under https://github.com/NVIDIA/nccl/tree/master/contrib We are happy to discuss more in a meeting if you have questions.

@staryxchen - to address your comments above - we agree that the communicator model and the pair-wise/dynamic connection model both have their strengths and weaknesses. Lately we have experimented with a few features that make communicators more capable of workloads that require dynamic membership. Some examples include:

New APIs such as ncclCommShrink/Grow that can quickly shrink and grow a communicator under certain fault tolerance scenarios.
ActiveRank mask idea in nccl-ep, similar to mooncake-ep.
RMA (remote memory access) APIs such as ncclPutSignal/ncclWaitSignal that enables one-sided 0-SM network transfers.

We will continue looking into ways that can make communicators more flexible and lightweight, while taking advantage of its innate benefits.
On the other hand, we are also looking from the other direction to see if the pair-wise/dynamic connection construct should be something we directly support in NCCL. The challenge is to ensure this construct works well with the rest of NCCL. E.g. ways to convert such connections into communicator and vise-versa; how memory registration/device API could be designed etc.

Here are some concrete suggestions based on where we are now:

We should consider integrating NCCL's device API with the communicator construct:
- We believe it is immediately helpful to Mooncake's effort to build high performance collectives, EP and various fine-grained comm/compute fusion use cases as described in this RFC.
- We can help with the integration and benchmarking to demonstrate the immediate benefits.
During this collaboration, we would love to hear your feedback and learn more about your use cases that could require more light-weight, dynamic constructs, as this can help motivate our future design. We can also collaborate further to add this new feature in both NCCL and Mooncake.

@xiaofanl-nvidia commented on GitHub (Jul 4, 2026): Just a quick update that we have started prototyping the integration to add NCCL device API to DeviceTransport under TE. We should be able to share a preliminary (but working) PR by next week. @UNIDY2002 we'll appreciate your review & hands-on help if you are available. We want to lay the groundwork before asking for your time :) Also, I saw the RFE to add more collectives in Mooncake: https://github.com/kvcache-ai/Mooncake/issues/2716 Given how expensive it is to optimize and validate these ops, we want to encourage Mooncake team to leverage NCCL device/host APIs for this work and review NCCL's implementations and examples, including interesting community contributed work that are quite novel under https://github.com/NVIDIA/nccl/tree/master/contrib We are happy to discuss more in a meeting if you have questions. @staryxchen - to address your comments above - we agree that the communicator model and the pair-wise/dynamic connection model both have their strengths and weaknesses. Lately we have experimented with a few features that make communicators more capable of workloads that require dynamic membership. Some examples include: - New APIs such as ncclCommShrink/Grow that can quickly shrink and grow a communicator under certain fault tolerance scenarios. - ActiveRank mask idea in nccl-ep, similar to mooncake-ep. - RMA (remote memory access) APIs such as ncclPutSignal/ncclWaitSignal that enables one-sided 0-SM network transfers. We will continue looking into ways that can make communicators more flexible and lightweight, while taking advantage of its innate benefits. On the other hand, we are also looking from the other direction to see if the pair-wise/dynamic connection construct should be something we directly support in NCCL. The challenge is to ensure this construct works well with the rest of NCCL. E.g. ways to convert such connections into communicator and vise-versa; how memory registration/device API could be designed etc. Here are some concrete suggestions based on where we are now: - We should consider integrating NCCL's device API with the communicator construct: - We believe it is immediately helpful to Mooncake's effort to build high performance collectives, EP and various fine-grained comm/compute fusion use cases as described in this RFC. - We can help with the integration and benchmarking to demonstrate the immediate benefits. - During this collaboration, we would love to hear your feedback and learn more about your use cases that could require more light-weight, dynamic constructs, as this can help motivate our future design. We can also collaborate further to add this new feature in both NCCL and Mooncake.

bsakaguchi commented

2026-10-29 21:18:14 +00:00

Author

Owner

@akhillanger commented on GitHub (Jul 7, 2026):

@UNIDY2002 here is NCCL backend for TE Device Transport. https://github.com/akhillanger/Mooncake/pull/3
Please take a look and see if this is in desired shape. After an initial review we can open an official PR.

@akhillanger commented on GitHub (Jul 7, 2026): @UNIDY2002 here is NCCL backend for TE Device Transport. https://github.com/akhillanger/Mooncake/pull/3 Please take a look and see if this is in desired shape. After an initial review we can open an official PR.

bsakaguchi commented

2026-10-29 21:18:14 +00:00

Author

Owner

@UNIDY2002 commented on GitHub (Jul 8, 2026):

@akhillanger Thanks for sharing this. Overall, I think this is a good direction and it aligns well with the DeviceTransport / Device API direction we have been building in Mooncake.

I’m glad to see NCCL device API integrated as a backend under the existing DeviceTransport framework. I think this is the right way to grow Mooncake toward a multi-backend device communication layer. Happy to help review and align it further!

@UNIDY2002 commented on GitHub (Jul 8, 2026): @akhillanger Thanks for sharing this. Overall, I think this is a good direction and it aligns well with the DeviceTransport / Device API direction we have been building in Mooncake. I’m glad to see NCCL device API integrated as a backend under the existing DeviceTransport framework. I think this is the right way to grow Mooncake toward a multi-backend device communication layer. Happy to help review and align it further!

bsakaguchi commented

2026-10-29 21:18:15 +00:00

Author

Owner

@xiaofanl-nvidia commented on GitHub (Jul 17, 2026):

Given that we are in PR stage with https://github.com/kvcache-ai/Mooncake/pull/2852, I'll close out this RFC.
Thank you for the insightful discussions. We'll reach out for future design discussions on new feature requests to NCCL.
For NCCL's future roadmap, please refer to announcements like this one: https://github.com/NVIDIA/nccl/issues/2272

@xiaofanl-nvidia commented on GitHub (Jul 17, 2026): Given that we are in PR stage with https://github.com/kvcache-ai/Mooncake/pull/2852, I'll close out this RFC. Thank you for the insightful discussions. We'll reach out for future design discussions on new feature requests to NCCL. For NCCL's future roadmap, please refer to announcements like this one: https://github.com/NVIDIA/nccl/issues/2272

bsakaguchi referenced this issue

2026-10-29 21:30:00 +00:00

[PR #689] [MERGED] add instructions for running on AMD GPU #1745