{"id":994,"date":"2026-05-29T16:43:32","date_gmt":"2026-05-29T08:43:32","guid":{"rendered":"https:\/\/connectword.dpdns.org\/?p=994"},"modified":"2026-05-29T16:43:32","modified_gmt":"2026-05-29T08:43:32","slug":"meet-mkernel-a-multi-gpu-multi-node-fused-kernel-library-for-gpu-driven-communication","status":"publish","type":"post","link":"https:\/\/connectword.dpdns.org\/?p=994","title":{"rendered":"Meet mKernel: A Multi-GPU, Multi-Node Fused Kernel Library for GPU-Driven Communication"},"content":{"rendered":"

GPU communication overhead is a measurable bottleneck in production AI workloads. According to data cited by the mKernel project, communication can consume 43.6% of the forward pass and 32% of end-to-end training time<\/strong>. Across popular Mixture-of-Experts (MoE) models, inter-device communication can account for up to 47% of total execution time<\/strong>. Researchers from UC Berkeley\u2019s UCCL project have released mKernel, a library of persistent CUDA kernels that fuse intra-node NVLink communication, inter-node RDMA, and compute into a single kernel.<\/p>\n

The Problem: Host-Driven Communication<\/strong><\/h2>\n

The standard model for multi-GPU communication is host-driven<\/strong>: the CPU runs the control path and calls into a library like NCCL or NVSHMEM. The library issues the collective operation \u2014 an AllReduce, an AllGather, etc. \u2014 across GPUs. Compute and communication run on separate CUDA streams and overlap at kernel boundaries.<\/p>\n

The research team identifies two problems with this approach<\/strong>:<\/p>\n

(1) CPUs are not scaling with GPU compute. A GB300 NVL72 rack integrates 72 Blackwell Ultra GPUs and 36 Grace CPUs, delivering 720 PFLOP\/s FP8\/FP6, 1.44 EFLOP\/s FP4 Tensor Core performance, and 130 TB\/s of all-to-all intra-rack NVLink bandwidth. At those speeds, microsecond-scale host orchestration overhead \u2014 a cudaLaunchKernel<\/code> call, a CPU-side \u201call writes done\u201d check, an inter-stream event \u2014 shows up directly as pipeline bubbles<\/strong>.<\/p>\n

(2) Host-driven systems overlap compute and communication at coarse kernel boundaries. Finer-grained overlap at the tile or chunk level is not possible from the host side.<\/p>\n

The alternative is GPU-driven communication<\/strong>: the GPU itself triggers transfers, with communication fused into the same kernel as the compute. Most existing fused kernel libraries operate within a single node, or a single GPU. mKernel targets the multi-node case.<\/p>\n

What mKernel Does<\/strong><\/h2>\n

mKernel is a library of persistent CUDA kernels<\/strong>. Each kernel fuses intra-node NVLink communication, inter-node RDMA, and dense compute into a single kernel.<\/p>\n

Multi-GPU + multi-node, in one kernel<\/strong>: Both intra-node NVLink and inter-node RDMA live inside the same persistent kernel.<\/p>\n

Fine-grained intra-kernel overlap<\/strong>: Compute and communication overlap at tile\/chunk granularity, covering both intra-node and inter-node GPU communication.<\/p>\n

Persistent kernel with SM specialization<\/strong>: CTAs self-assign roles: compute<\/code>, intra-comm<\/code>, inter-send<\/code>, inter-reduce<\/code>. The number of SMs dedicated to each role is tunable per shape.<\/p>\n

GPU-driven networking built on libibverbs<\/code><\/strong>: mKernel uses GPU-initiated RDMA writes without depending on NCCL or NVSHMEM. The communication backend is written from scratch to maximize performance and support heterogeneous networking devices.<\/p>\n

The Five Fused Kernels<\/strong><\/h2>\n
\n\n\n\n\n\n\n\n\n\n
Kernel<\/th>\nWhat it fuses<\/th>\nDescription<\/th>\n<\/tr>\n<\/thead>\n
AllGather + GEMM<\/strong><\/td>\nAllGather \u2192 GEMM<\/td>\nEach rank holds a shard of A<\/code>. While ranks gather peers\u2019 shards over NVLink\/RDMA, the local GEMM consumes tiles as soon as they arrive.<\/td>\n<\/tr>\n
GEMM + AllReduce<\/strong><\/td>\nGEMM \u2192 AllReduce<\/td>\nComputes C = A @ B<\/code> and reduces partial outputs across all ranks in one launch. Output tiles are pushed into the reduction tree the instant they\u2019re produced.<\/td>\n<\/tr>\n
MoE Dispatch + GEMM<\/strong><\/td>\nAll-to-All dispatch \u2192 grouped GEMM<\/td>\nRoutes MoE tokens to their expert ranks (intra-node NVLink + inter-node all-to-all) and runs the per-expert grouped GEMM in the same kernel. Tokens are processed as soon as they land \u2014 no staging buffer round-trip.<\/td>\n<\/tr>\n
Ring Attention<\/strong><\/td>\nRing KV exchange \u2192 FlashAttention<\/td>\nSequence-parallel attention across ranks. Each step rotates a KV chunk around the ring while the local FlashAttention consumes the previously-received chunk. Compute and the ring send\/recv run concurrently inside a single persistent kernel.<\/td>\n<\/tr>\n
GEMM + ReduceScatter<\/strong><\/td>\nGEMM \u2192 ReduceScatter<\/td>\nComputes C = A @ B<\/code> and reduce-scatters the output. Each output tile is reduced and forwarded to its owning rank as soon as it is produced.<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n<\/figure>\n

\n

Evaluation Setup<\/strong><\/h2>\n<\/p>

The research team evaluated mKernel on two 2-node \u00d7 8-H200 clusters that differ only in their inter-node fabric:<\/p>\n

\n\n\n\n\n\n\n
Testbed<\/th>\nNodes \u00d7 GPUs<\/th>\nIntra-node<\/th>\nInter-node transport<\/th>\nNIC<\/th>\n<\/tr>\n<\/thead>\n
AWS EFA<\/strong><\/td>\n2 \u00d7 8 H200<\/td>\nNVLink<\/td>\nAWS EFA \/ SRD<\/td>\n16 \u00d7 200 Gb\/s EFA per node<\/td>\n<\/tr>\n
ConnectX-7<\/strong><\/td>\n2 \u00d7 8 H200<\/td>\nNVLink<\/td>\nInfiniBand<\/td>\n8 \u00d7 400 Gb\/s NVIDIA ConnectX-7 per node<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n<\/figure>\n

mKernel was benchmarked against NCCL, Triton-distributed, Flux, Mercury, MagiAttention, Transformer-Engine, and ring-flash-attention. The team notes that further benchmarking at larger scale is still in progress.<\/p>\n

Backends and Requirements<\/strong><\/h2>\n

mKernel supports two networking backends:<\/p>\n

\n\n\n\n\n\n\n
Backend<\/th>\nMacro<\/th>\nTransport<\/th>\nWhere it runs<\/th>\n<\/tr>\n<\/thead>\n
CX7<\/strong><\/td>\n-DINTERNODE_BACKEND_IBVERBS<\/code><\/td>\nlibibverbs RC<\/td>\nConnectX-7 \/ InfiniBand \/ RoCE<\/td>\n<\/tr>\n
EFA<\/strong><\/td>\n-DINTERNODE_BACKEND_EFA<\/code><\/td>\nlibibverbs + efadv (SRD)<\/td>\nAWS p5\/p5e (H200, EFA)<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n<\/figure>\n

Both backends share the same host-side API and the same on-GPU kernel. Only the proxy\/session implementation differs (session.h<\/code> for CX7, session_efa.h<\/code> for EFA). Requirements: NVIDIA Hopper GPUs (default build targets sm_90a<\/code>), CUDA 12.9, Python with PyTorch. The CX7 backend requires libibverbs development headers and libraries. The EFA backend requires AWS EFA installation with libfabric, libibverbs, efadv, and EFA headers under EFA_HOME=\/opt\/amazon\/efa<\/code> by default.<\/p>\n

Marktechpost\u2019s Visual Explainer<\/strong><\/h2>\n
\n
\n
UCCL<\/div>\n
mKernel<\/span> \u2014 Multi-GPU, Multi-Node Fused Kernels Guide<\/div>\n<\/div>\n
\n
<\/div>\n<\/div>\n
\n
\n

<\/p>\n

\n
01 \/ 07 \u2014 Overview<\/div>\n

What is mKernel<\/span>?<\/h2>\n

mKernel is an open-source library of persistent CUDA kernels from UC Berkeley\u2019s UCCL project. It fuses intra-node NVLink communication, inter-node RDMA, and dense compute into a single kernel.<\/p>\n

Most existing fused kernel libraries operate within a single node or a single GPU. mKernel is designed from the start to span node boundaries.<\/p>\n

\n
\n
43.6%<\/div>\n
of forward pass consumed by communication in production<\/div>\n<\/div>\n
\n
47%<\/div>\n
of total execution time in popular MoE models<\/div>\n<\/div>\n
\n
32%<\/div>\n
of end-to-end training time consumed by communication<\/div>\n<\/div>\n<\/div>\n<\/div>\n

<\/p>\n

\n
02 \/ 07 \u2014 The Problem<\/div>\n

Why Host-Driven<\/span> Communication Falls Short<\/h2>\n

The standard model is host-driven: the CPU calls NCCL or NVSHMEM, which issues collective operations across GPUs. The UCCL team identifies two problems.<\/p>\n

\n
\n
\"\u26a1\"<\/div>\n
CPUs are not scaling with GPUs.<\/strong> A GB300 NVL72 rack delivers 720 PFLOP\/s FP8\/FP6 and 1.44 EFLOP\/s FP4. At those speeds, microsecond-scale overhead from cudaLaunchKernel<\/code>, CPU-side sync checks, and inter-stream events shows up directly as pipeline bubbles.<\/div>\n<\/div>\n
\n
\"\ud83d\udd32\"<\/div>\n
Overlap is too coarse.<\/strong> Host-driven systems overlap compute and communication only at kernel boundaries. Finer-grained overlap at the tile or chunk level is not possible from the host side.<\/div>\n<\/div>\n
\n
\"\ud83d\udd00\"<\/div>\n
The answer: GPU-driven communication.<\/strong> The GPU itself triggers fine-grained transfers, fused into the same kernel as the compute.<\/div>\n<\/div>\n<\/div>\n<\/div>\n

<\/p>\n

\n
03 \/ 07 \u2014 Design<\/div>\n

Four Core Design Properties<\/span><\/h2>\n
\n
\n
\ud83d\udda7<\/div>\n
Multi-GPU + multi-node, in one kernel.<\/strong> Intra-node NVLink and inter-node RDMA both live inside the same persistent kernel.<\/div>\n<\/div>\n
\n
\"\ud83d\udd2c\"<\/div>\n
Fine-grained intra-kernel overlap.<\/strong> Compute and communication overlap at tile\/chunk granularity, covering both intra-node and inter-node communication.<\/div>\n<\/div>\n
\n
\"\u2699\"<\/div>\n
Persistent kernel with SM specialization.<\/strong> CTAs self-assign roles: compute<\/code>, intra-comm<\/code>, inter-send<\/code>, inter-reduce<\/code>. SM split is tunable per shape.<\/div>\n<\/div>\n
\n
\"\ud83d\udce1\"<\/div>\n
GPU-driven networking via libibverbs<\/code>.<\/strong> Uses GPU-initiated RDMA writes. No NCCL or NVSHMEM dependency. Communication backend is written from scratch.<\/div>\n<\/div>\n<\/div>\n<\/div>\n

<\/p>\n

\n
04 \/ 07 \u2014 Kernels<\/div>\n

The Five Fused Kernels<\/span><\/h2>\n
\n
\n
AllGather + GEMM<\/div>\n
AllGather \u2014> GEMM<\/div>\n
Each rank holds a shard of A<\/code>. The local GEMM consumes tiles over NVLink\/RDMA as they arrive \u2014 matmul starts before the collective finishes.<\/div>\n<\/div>\n
\n
GEMM + AllReduce<\/div>\n
GEMM \u2014> AllReduce<\/div>\n
Computes C = A @ B<\/code> and reduces partial outputs across all ranks in one launch. Output tiles enter the reduction tree the instant they are produced.<\/div>\n<\/div>\n
\n
MoE Dispatch + GEMM<\/div>\n
All-to-All dispatch \u2014> grouped GEMM<\/div>\n
Routes MoE tokens to expert ranks via NVLink + inter-node all-to-all, then runs per-expert grouped GEMM in the same kernel. No staging buffer round-trip.<\/div>\n<\/div>\n
\n
Ring Attention<\/div>\n
Ring KV exchange \u2014> FlashAttention<\/div>\n
Sequence-parallel attention across ranks. Each step rotates a KV chunk around the ring while the local FlashAttention consumes the previously-received chunk.<\/div>\n<\/div>\n
\n
GEMM + ReduceScatter<\/div>\n
GEMM \u2014> ReduceScatter<\/div>\n
Computes C = A @ B<\/code> and reduce-scatters the output. Each tile is reduced and forwarded to its owning rank as soon as it is produced.<\/div>\n<\/div>\n<\/div>\n<\/div>\n

<\/p>\n

\n
05 \/ 07 \u2014 Evaluation<\/div>\n

Evaluation Setup<\/span><\/h2>\n

Tested on two 2-node \u00d7 8-H200 clusters differing only in inter-node fabric.<\/p>\n\n\n\n\n\n\n
Testbed<\/th>\nNodes \u00d7 GPUs<\/th>\nInter-node<\/th>\nNIC<\/th>\n<\/tr>\n<\/thead>\n
AWS EFA<\/strong><\/td>\n2 \u00d7 8 H200<\/td>\nAWS EFA \/ SRD<\/td>\n16 \u00d7 200 Gb\/s EFA per node<\/td>\n<\/tr>\n
ConnectX-7<\/strong><\/td>\n2 \u00d7 8 H200<\/td>\nInfiniBand<\/td>\n8 \u00d7 400 Gb\/s CX7 per node<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

Both testbeds use NVLink intra-node. Benchmarked against: NCCL, Triton-distributed, Flux, Mercury, MagiAttention, Transformer-Engine, and ring-flash-attention. Larger-scale benchmarking is still in progress.<\/p>\n<\/div>\n

<\/p>\n

\n
06 \/ 07 \u2014 Backends & Requirements<\/div>\n

Backends & Requirements<\/span><\/h2>\n\n\n\n\n\n\n
Backend<\/th>\nTransport<\/th>\nWhere it runs<\/th>\n<\/tr>\n<\/thead>\n
CX7<\/strong><\/td>\nlibibverbs RC<\/td>\nConnectX-7 \/ InfiniBand \/ RoCE<\/td>\n<\/tr>\n
EFA<\/strong><\/td>\nlibibverbs + efadv (SRD)<\/td>\nAWS p5\/p5e (H200, EFA)<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n
\n
\n
\"\ud83d\udccb\"<\/div>\n
Requirements:<\/strong> NVIDIA Hopper GPUs (default sm_90a<\/code>), CUDA 12.9, Python with PyTorch. CX7 needs libibverbs headers. EFA needs libfabric, libibverbs, efadv under EFA_HOME=\/opt\/amazon\/efa<\/code>.<\/div>\n<\/div>\n
\n
\"\ud83d\udcdd\"<\/div>\n
License & Attribution:<\/strong> MIT licensed. MMA\/compute code adapted from ThunderKittens (HazyResearch).<\/div>\n<\/div>\n<\/div>\n<\/div>\n

<\/p>\n

\n
07 \/ 07 \u2014 Roadmap & Key Takeaways<\/div>\n

Roadmap & Key Takeaways<\/span><\/h2>\n
\n
\n
\u2705<\/div>\n
Fused GPU-driven multi-node kernels (AG+GEMM, GEMM+AR, MoE Dispatch+GEMM, Ring Attention, GEMM+RS)<\/div>\n<\/div>\n
\n
\u2705<\/div>\n
ConnectX-7 and AWS EFA backends<\/div>\n<\/div>\n
\n
\ud83d\udea7<\/div>\n
Full heterogeneous accelerator\/NIC support with topology-aware discovery, placement, routing<\/div>\n<\/div>\n
\n
\ud83d\udea7<\/div>\n
Inter-node megakernels: collapsing several fused steps into a single megakernel spanning a transformer layer<\/div>\n<\/div>\n
\n
\ud83d\udea7<\/div>\n
Blackwell GPU support<\/div>\n<\/div><\/div>\n
\n
\n
<\/div>\n
Fuses NVLink, inter-node RDMA, and compute into a single persistent CUDA kernel<\/div>\n<\/div>\n
\n
<\/div>\n
Five kernels: AllGather+GEMM, GEMM+AllReduce, MoE Dispatch+GEMM, Ring Attention, GEMM+ReduceScatter<\/div>\n<\/div>\n
\n
<\/div>\n
GPU-initiated RDMA via libibverbs<\/code> \u2014 no NCCL or NVSHMEM dependency<\/div>\n<\/div>\n
\n
<\/div>\n
Requires Hopper GPUs (sm_90a<\/code>) and ConnectX-7 or AWS EFA networking<\/div>\n<\/div><\/div>\n<\/div>\n<\/div>\n

\n <\/p><\/div>\n

<\/p>\n

\n