MORI CCO: Next-Generation Collective Communication Object for GPUs

Sep 29, 2026

Abstract glowing 3D computer chip

Collective communication is usually something your GPU kernels wait for, not something they do. Launch a compute kernel, wait, launch a communication op, wait, repeat — and the GPU idles at every transition. MORI CCO removes that boundary: it is a device-side collective communication object whose primitives are issued from inside your kernel, so compute and communication overlap at tile granularity within a single launch.

In this blog you will learn what MORI CCO is, how its symmetric memory model makes peer addressing nearly free, and how to call its primitives from HIP C++, FlyDSL, and Triton. We walk through runnable examples for each backend, show how torch.SymmetricMemory interoperates with CCO in both directions, and demonstrate a zero-copy Mixture-of-Experts (MoE) dispatch → expert → combine pipeline that allocates no intermediate buffers at all. If you build distributed inference or training kernels on AMD Instinct™ GPUs, this is the communication layer to try next.

What is MORI CCO?

MORI (Modular RDMA Interface) is a bottom-up, modular, and composable framework for building high-performance communication applications with a strong focus on RDMA + GPU integration. Inspired by the role of MLIR in compiler infrastructure, MORI provides reusable and extensible building blocks that make it easier for developers to adopt advanced techniques such as IBGDA (InfiniBand GPUDirect Async) and GDS (GPUDirect Storage).

CCO (Compute Communicate Object) is MORI’s next-generation collective communication object for GPUs. Unlike traditional host-driven collective libraries, CCO natively supports symmetric memory — every rank sees a coherent, globally addressable memory space without explicit host-side buffer registration. The symmetric heap design provides predictable memory layouts across all GPUs, enabling efficient pointer translation with minimal overhead since each rank’s memory has an identical structure at corresponding offsets. Built on this foundation, CCO lets GPU kernels drive communication directly: a single kernel (or a mega-kernel fusing multiple layers) can issue all communication primitives while still running compute, achieving true compute-communication overlap within one kernel launch.

MORI-CCO architecture: a flat symmetric address space sliced into per-rank slots, with symmetric windows mapped to each GPU's HBM and the LSA, SDMA, TDM, and IBGDA transports arming from one window registration
MORI-CCO architecture: a flat symmetric address space sliced into per-rank slots, with symmetric windows mapped to each GPU's HBM and the LSA, SDMA, TDM, and IBGDA transports arming from one window registration

Because every window is carved at the same slotOffset in every slot, a peer address is pure arithmetic:

		peer_va = winBase + peerLsaRank * perRankSize + offset

CCO hands out the address; the kernel does the load or store. One ccoWindowRegister call arms all transports at once.

Key Features

In-kernel compute-communication fusion

Traditional collective libraries work like this: launch a compute kernel, wait, launch a communication operation, wait, repeat. The GPU idles during transitions. A common workaround is to overlap compute and communication via two separate HIP streams, but this approach suffers from unpredictable resource contention - the GPU block scheduler cannot precisely partition compute unit (CU) resources between the two streams, so compute and communication kernels end up competing for the same CUs rather than cooperating.

CCO offers a new and more efficient solution: compute and communication execute within the same kernel, spanning both intra-node and inter-node GPU-driven communication in a single launch across multiple GPUs. A single kernel issues communication primitives alongside compute instructions, so compute and communication can be precisely coordinated within the same CUs — overlapping at tile granularity rather than at the coarse kernel level. When one stalls (for example, waiting on a network transfer), the other makes progress on the next tile, keeping the GPU fully utilized. This fine-grained in-kernel fusion fundamentally eliminates the scheduling inefficiency of the two-stream model, enabling higher throughput and more predictable performance.

Multi-language backend support

CCO’s communication primitives are accessible from multiple programming backends, giving you the freedom to choose the language that best fits your workflow:

  • HIP C++ — full low-level control for performance-critical paths
  • FlyDSL — high-level Python DSL for rapid prototyping and concise kernel authoring
  • Triton — integration with the broader PyTorch ecosystem and Triton-based kernel libraries

HIP C++ example

This kernel stores directly into peer memory through the symmetric window’s flat virtual address (VA)  - no send, no recv, no staging buffer:

		#include <hip/hip_runtime.h>
#include "mori/cco/cco.hpp"
using namespace mori::cco;

extern "C" __global__ void lsa_put_kernel(
    ccoDevComm* devComm,
    ccoWindowDevice* win,
    uint64_t src_off,
    uint64_t dst_off)
{
    if (threadIdx.x != 0 || blockIdx.x != 0) return;

    // LSA flat-VA: peer VA = winBase + peerLsaRank * (stride4G << 32) + offset
    uint64_t stride  = (uint64_t)win->stride4G << 32;
    uint64_t my_rank = devComm->lsaRank;

    for (int peer = 0; peer < devComm->lsaSize; peer++) {
        if (peer == my_rank) continue;
        uint64_t* dst = (uint64_t*)(
            win->winBase + (uint64_t)peer * stride + dst_off
        );
        uint64_t* src = (uint64_t*)(
            win->winBase + my_rank * stride + src_off
        );
        // LSA Direct Store
        *dst = *src;
    }
    __threadfence_system();
}

FlyDSL example

The same idea expressed as an SDMA all-gather in FlyDSL’s Python DSL:

		@flyc.kernel(known_block_size=[THREADS, 1, 1])
def scatter_kernel(dev_comm: Int64, win: Int64):
    sdma = cco.DevComm(dev_comm).sdma()
    my_rank = 0
    dst_off = GATHER_OFF + my_rank * NBYTES

    if fx.thread_idx.x == 0:
        for p in range(WS):
            if p != my_rank:
                sdma.put(
                    p, win, dst_off,
                    win, IN_OFF,
                    NBYTES, 0,
                    coop=cco.CoopScope.THREAD,
                )
        for p in range(WS):
            if p != my_rank:
                sdma.quiet(p, coop=cco.CoopScope.THREAD)

Triton example

A cross-node RDMA put with a signal notification to the peer, issued from a Triton kernel:

		@triton.jit
def gda_put_kernel(
    dev_comm,
    window,
    peer: tl.constexpr,
    dst_off: tl.constexpr,
    src_off: tl.constexpr,
    nbytes: tl.constexpr,
):
    cco.Gda.put(
        dev_comm, peer,
        window, dst_off,
        window, src_off,
        nbytes,
        signal_op=cco.SignalOp.INC,
        signal_id=0,
        signal_val=1,
        coop=cco.CoopScope.BLOCK,
        thread_mode=cco.ThreadMode.INDEPENDENT,
    )
    cco.Gda.flush_peer(dev_comm, peer, coop=cco.CoopScope.BLOCK)

@triton.jit
def gda_wait_kernel(dev_comm):
    cco.Gda.wait_signal(dev_comm, 0, 1, coop=cco.CoopScope.BLOCK)

# sender
gda_put_kernel[(1,)](
    dc.ptr, win.handle,
    peer=1, dst_off=0, src_off=0, nbytes=NBYTES,
    extern_libs=cco.get_extern_libs(), num_warps=1,
)
# receiver
gda_wait_kernel[(1,)](
    dc.ptr,
    extern_libs=cco.get_extern_libs(), num_warps=1,
)

MORI can be registered as a native torch.SymmetricMemory backend allocator, and conversely, symmetric memory allocated by PyTorch can be imported directly as CCO symmetric memory. You can flexibly choose to build on MORI’s native symmetric heap or on PyTorch’s symmetric memory abstraction depending on your integration needs.

The multi-backend design means CCO meets you where you already are — FlyDSL, HIP C++, and Triton can all adopt CCO without leaving their existing ecosystem. Combined with native torch.SymmetricMemory interoperability, CCO integrates into the broader AI/ML software stack with minimal friction, lowering the barrier for the community to build on top of it.

torch.SymmetricMemory × MORI

Registering MORI as the PyTorch symmetric memory backend takes one set_backend call. From there, a Triton kernel can push straight into peer windows:

		import torch
import torch.distributed as dist
import torch.distributed._symmetric_memory as symm_mem
import triton
import triton.language as tl
from mori.allocator import handle_type  # importing registers the "MORI" backend

# -- Setup ---------------------------------------------------------
dist.init_process_group("gloo")
torch.cuda.set_device(local_rank)
device = torch.device("cuda", local_rank)

symm_mem.set_backend("MORI")
symm_mem.enable_symm_mem_for_group(dist.group.WORLD.group_name)

# -- Allocate symmetric memory -------------------------------------
# recv is peer-writable; send is ordinary local memory.
recv = symm_mem.empty(world_size * chunk_elems, dtype=torch.int32, device=device)
send = torch.empty_like(recv)

# Rendezvous exchanges physical handles and returns peer pointers.
hdl = symm_mem.rendezvous(recv, dist.group.WORLD.group_name)

# -- Launch a Triton kernel that pushes directly into peer windows --
@triton.jit
def push_kernel(send_ptr, peer_ptrs, chunk_elems, rank_id, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    peer = pid  # one program per peer
    # Load the peer's base pointer from the device-side pointer array.
    dst = tl.load(peer_ptrs.to(tl.pointer_type(tl.uint64)) + peer)
    dst = dst.to(tl.pointer_type(tl.int32)) + rank_id * chunk_elems
    src = send_ptr + peer * chunk_elems
    offs = tl.arange(0, BLOCK)
    tl.store(dst + offs, tl.load(src + offs, mask=offs < chunk_elems), mask=offs < chunk_elems)

push_kernel[(world_size,)](send, hdl.buffer_ptrs_dev, chunk_elems, rank_id, BLOCK=1024)

Conversely, a tensor already allocated by torch.SymmetricMemory can be imported into CCO’s flat LSA space with zero copy via register_external_window(). CCO retains the tensor’s physical handle and maps it into the flat virtual-address space, so the buffer is immediately usable with all CCO device APIs (LSA, SDMA, GDA) — no data movement, no extra allocation:

		import torch
import torch.distributed._symmetric_memory as symm_mem
from mori.cco import Communicator

# Allocate via torch's symmetric memory (any backend).
t = symm_mem.empty(1024, dtype=torch.bfloat16, device=device)

# Import into CCO - zero copy, collective across all ranks.
with Communicator.init(nranks, rank, uid) as comm:
    win = comm.register_external_window(t.data_ptr(), t.nbytes)
    #  win.local_ptr  -> flat-VA alias for local + peer access
    #  win.handle     -> CCO window handle for device API (LSA / SDMA / GDA)

    # From here, use win.handle in any CCO device kernel:
    # e.g. cco.Window(win.handle).lsa_ptr(peer, offset) for direct peer reads.

Native collective zero-copy support

Traditional collective libraries rely on a send/recv + staging buffer model: data is first copied into a temporary buffer, transferred to the peer, and then copied out to the destination — introducing extra memory traffic and latency. CCO’s one-sided communication drops this cooperation entirely. Symmetric memory is the recv buffer — the sender writes directly into the receiver’s working memory, and operations like dispatch return tensor views into symmetric memory rather than allocating new tensors. Downstream modules (for example, MoE experts) consume the data in place, eliminating all extra local data movement on the receiver side and reducing HBM bandwidth pressure.

Zero-copy MoE dispatch → expert → combine

The following example shows a complete MoE expert-parallel pipeline using the EPv2 dispatch/combine op. Every intermediate tensor is a view into symmetric memory — no extra allocation, no extra copy:

		from mori.ops.dispatch_combine_v2 import EpDispatchCombineOp

op = EpDispatchCombineOp(config, comm)

# -- Step 1: Dispatch (all-to-all scatter) -------------------------
# Each rank sends its tokens to the ranks that own the selected experts.
# The dispatch kernel writes directly into each peer's symmetric receive
# buffer via one-sided P2P - no staging copy on either side.
recv_x, recv_w, recv_s, recv_i, total_recv, routing = op.dispatch(
    tokens, weights, scales, indices, return_routing=True
)
# recv_x is NOT a new tensor - it is a view into the symmetric arena's
# "disp_out" region.  The dispatching peers wrote here directly via LSA.

torch.cuda.synchronize()
dist.barrier()
total = int(total_recv.cpu().item())

# -- Step 2: Expert compute (in-place) -----------------------------
# Read dispatched tokens directly from symmetric memory.
expert_input = recv_x[:total]

# Point the expert's output buffer at the symmetric combine staging area.
# combine() will detect the matching data_ptr and skip its internal copy.
expert_output = op.combine_in_view()[:total]

# Run the expert GEMM: reads from disp_out, writes into out_tok - both
# in symmetric memory, zero local data movement.
expert_gemm(expert_input, expert_weights, out=expert_output)

# -- Step 3: Combine (all-to-all gather) ---------------------------
# Each rank reads its tokens back from the expert-owning peers via P2P.
# Because expert_output already lives in the combine staging buffer,
# combine() elides the d2d copy - data_ptr match detected automatically.
out, out_weights = op.combine(expert_output, weights, routing=routing)

In a traditional library this pipeline would allocate four intermediate buffers and perform two extra device-to-device copies (dispatch staging → recv, expert output → combine staging). With CCO, the dispatch kernel lands tokens directly into disp_out, the expert writes results directly into out_tok, and combine gathers from peers’ out_tok via P2P — zero extra HBM traffic end to end.

Note: The zero-copy path depends on combine() detecting a matching data_ptr. If you allocate the expert output yourself instead of taking the view returned by op.combine_in_view(), the op falls back to an internal device-to-device copy and you silently lose the saving. When in doubt, compare the data_ptr() values.

Multiple communication transports

CCO supports three transports — LSA (load/store over XGMI), GDA (device-initiated RDMA via IBGDA), and SDMA (DMA-engine offload) — usable together in a single kernel. Pick the right transport for each data-movement task:

Transport

Scope

Mechanism

Use it for

LSA / P2P

Intra-node

GPU load/store over XGMI

Low-latency signaling and small transfers

SDMA

Intra-node

DMA copy engines

Background bulk copies that keep CUs free

GDA

Cross-node

GPU-initiated NIC RDMA (IBGDA)

Bulk transfers with no host involvement

CCO is also NIC-agnostic: the same code runs unchanged on Mellanox ConnectX-7, Broadcom Thor2, and AMD Pollara, with the build system auto-selecting the right backend — write once, run on any supported NIC.

Examples

Mega-MoE: mega-kernel co-design with FlyDSL

FlyDSL’s high-level Python DSL and CCO’s device-side communication primitives are co-designed to enable mega-kernel authoring: FlyDSL handles kernel fusion, scheduling, and code generation while CCO provides the in-kernel communication backbone.

This co-design has been validated on MI455X, where FlyDSL has successfully prototyped a Mega-MoE kernel that fuses dispatch, expert GEMM, and combine into a single mega-kernel — eliminating all inter-kernel synchronization across the MoE pipeline. You write concise FlyDSL Python; the toolchain lowers it to a fused HIP kernel with embedded CCO calls, making mega-kernel development accessible without hand-writing thousands of lines of HIP C++.

GEMM + all-to-all fusion

CCO’s multi-transport design enables fine-grained fusion of GEMM with all-to-all (A2A) communication in both directions:

  • GEMM → A2A: the GEMM epilogue submits SDMA PUTs as output tiles become ready, or writes directly into peer LSA windows, overlapping communication with the remaining compute tiles.
  • A2A → GEMM: SDMA transfers prefill the next epoch while GEMM consumes the current one, or remote K-shards are consumed as they arrive via ready-counter polling.

On 8-GPU MI300X, fused LSA achieves up to 36.5% speedup over split pipelines, and fused SDMA achieves up to 34.8% speedup at large M — demonstrating the real-world benefit of CCO’s in-kernel compute-communication fusion across multiple transports. 

Summary

MORI CCO is a multi-backend, multi-transport communication object designed for true compute-communication fusion on GPUs. The same primitives are callable from FlyDSL, HIP C++, and Triton, sharing one symmetric memory model. Because CCO is a device-side API, compute and communication execute inside the same kernel launch — no host round-trip between phases. A single kernel can mix regular load/store for low-latency signaling with SDMA or GDA for bulk background transfers, keeping CUs free for compute while data moves in parallel. Along the way you saw how symmetric windows reduce peer addressing to a single offset computation, how PyTorch symmetric memory interoperates with CCO in both directions at zero copy, and how an MoE dispatch → expert → combine pipeline runs end to end without a single intermediate allocation.

This post is a first turn, not a finish line. The next round of work runs along two lines:

  • Broader compute-communication overlap — extending the mega-kernel co-design approach to more collective patterns and model architectures across multi-node multi-GPU topologies, including but not limited to Mega-MoE, GEMM + AllReduce, and beyond.
  • GPU-initiated storage and GPU-initiated SDMA for KV cache — leveraging GPUDirect Storage and GPU-initiated SDMA to accelerate KV cache transfer in long-context inference workloads.

Ready to try it? Start from the HIP C++, FlyDSL, or Triton snippet closest to your stack, swap your staging-buffer collective for a symmetric window, and measure what your kernels do when they stop waiting.

Footnotes

Disclaimers

Third-party content is licensed to you directly by the third party that owns the content and is not licensed to you by AMD. ALL LINKED THIRD-PARTY CONTENT IS PROVIDED “AS IS” WITHOUT A WARRANTY OF ANY KIND. USE OF SUCH THIRD-PARTY CONTENT IS DONE AT YOUR SOLE DISCRETION AND UNDER NO CIRCUMSTANCES WILL AMD BE LIABLE TO YOU FOR ANY THIRD-PARTY CONTENT. YOU ASSUME ALL RISK AND ARE SOLELY RESPONSIBLE FOR ANY DAMAGES THAT MAY ARISE FROM YOUR USE OF THIRD-PARTY CONTENT.

The information presented in this document is for informational purposes only and may contain technical inaccuracies, omissions, and typographical errors. The information contained herein is subject to change and may be rendered inaccurate for many reasons, including but not limited to product and roadmap changes, component and motherboard version changes, new model and/or product releases, product differences between differing manufacturers, software changes, BIOS flashes, firmware upgrades, or the like. Any computer system has risks of security vulnerabilities that cannot be completely prevented or mitigated. AMD assumes no obligation to update or otherwise correct or revise this information. However, AMD reserves the right to revise this information and to make changes from time to time to the content hereof without obligation of AMD to notify any person of such revisions or changes.

THIS INFORMATION IS PROVIDED “AS IS.” AMD MAKES NO REPRESENTATIONS OR WARRANTIES WITH RESPECT TO THE CONTENTS HEREOF AND ASSUMES NO RESPONSIBILITY FOR ANY INACCURACIES, ERRORS, OR OMISSIONS THAT MAY APPEAR IN THIS INFORMATION. AMD SPECIFICALLY DISCLAIMS ANY IMPLIED WARRANTIES OF NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR ANY PARTICULAR PURPOSE. IN NO EVENT WILL AMD BE LIABLE TO ANY PERSON FOR ANY RELIANCE, DIRECT, INDIRECT, SPECIAL, OR OTHER CONSEQUENTIAL DAMAGES ARISING FROM THE USE OF ANY INFORMATION CONTAINED HEREIN, EVEN IF AMD IS EXPRESSLY ADVISED OF THE POSSIBILITY OF SUCH DAMAGES.

AMD, the AMD Arrow logo, AMD Instinct, ROCm, and combinations thereof are trademarks of Advanced Micro Devices, Inc. Other product names used in this publication are for identification purposes only and may be trademarks of their respective companies. All other trademarks and product names referenced in this publication, including MORI, FlyDSL, PyTorch, Triton, InfiniBand, Mellanox ConnectX-7, and Broadcom Thor2, are the property of their respective owners.

© 2026 Advanced Micro Devices, Inc. All rights reserved.

Share:

Related Blogs

View All Blogs