Der GPU-Kommunikationsaufwand ist ein messbarer Engpass bei KI-Workloads in der Produktion. Laut den vom mKernel-Projekt zitierten Daten kann die Kommunikation verbrauchen 43,6 % der Vorwärtspässe und 32 % der gesamten Trainingszeit. Bei gängigen Combination-of-Consultants-Modellen (MoE) kann die Kommunikation zwischen Geräten ausschlaggebend sein bis zu 47 % der gesamten Ausführungszeit. Forscher des UCCL-Projekts der UC Berkeley haben mKernel veröffentlicht, eine Bibliothek persistenter CUDA-Kernel, die knoteninterne NVLink-Kommunikation, RDMA zwischen Knoten und Rechenleistung in einem einzigen Kernel vereinen.

Das Downside: Host-gesteuerte Kommunikation

Das Standardmodell für die Multi-GPU-Kommunikation ist hostgesteuert: Die CPU führt den Steuerpfad aus und ruft eine Bibliothek wie NCCL oder NVSHMEM auf. Die Bibliothek gibt die kollektive Operation – ein AllReduce, ein AllGather usw. – über GPUs hinweg aus. Berechnung und Kommunikation laufen auf separaten CUDA-Streams und überlappen sich an Kernelgrenzen.

Das Forschungsteam identifiziert zwei Probleme bei diesem Ansatz:

(1) CPUs skalieren nicht mit GPU-Rechenleistung. Ein GB300 NVL72-Rack integriert 72 Blackwell Extremely-GPUs und 36 Grace-CPUs und liefert 720 PFLOP/s FP8/FP6, 1,44 EFLOP/s FP4 Tensor Core-Leistung und 130 TB/s gesamte Intra-Rack-NVLink-Bandbreite. Bei diesen Geschwindigkeiten entsteht ein Host-Orchestrierungsaufwand im Mikrosekundenbereich – a cudaLaunchKernel Aufruf, eine CPU-seitige „Alle Schreibvorgänge erledigt“-Prüfung, ein Inter-Stream-Ereignis – wird direkt angezeigt als Pipeline-Blasen.

(2) Hostgesteuerte Systeme überlappen Rechenleistung und Kommunikation an groben Kernelgrenzen. Eine feinkörnigere Überlappung auf Kachel- oder Blockebene ist von der Hostseite aus nicht möglich.

Die Different ist GPU-gesteuerte Kommunikation: Die GPU selbst löst Übertragungen aus, wobei die Kommunikation in denselben Kernel wie die Rechenleistung integriert wird. Die meisten vorhandenen Fused-Kernel-Bibliotheken arbeiten innerhalb eines einzelnen Knotens oder einer einzelnen GPU. mKernel zielt auf den Fall mit mehreren Knoten ab.

Was mKernel macht

mKernel ist eine Bibliothek von persistente CUDA-Kernel. Jeder Kernel vereint knoteninterne NVLink-Kommunikation, RDMA zwischen Knoten und dichte Datenverarbeitung in einem einzigen Kernel.

Multi-GPU + Multi-Knoten in einem Kernel: Sowohl knoteninterner NVLink als auch knoteninterner RDMA befinden sich im selben persistenten Kernel.

Feinkörnige Intra-Kernel-Überlappung: Rechen- und Kommunikationsüberschneidungen bei der Kachel-/Chunk-Granularität, die sowohl die GPU-Kommunikation zwischen Knoten als auch zwischen Knoten abdecken.

Persistenter Kernel mit SM-Spezialisierung: CTAs weisen sich selbst Rollen zu: compute, intra-comm, inter-send, inter-reduce. Die Anzahl der SMs, die jeder Rolle zugewiesen sind, kann je nach Type angepasst werden.

GPU-gesteuertes Netzwerk basiert auf libibverbs: mKernel verwendet GPU-initiierte RDMA-Schreibvorgänge, ohne von NCCL oder NVSHMEM abhängig zu sein. Das Kommunikations-Backend wurde von Grund auf neu geschrieben, um die Leistung zu maximieren und heterogene Netzwerkgeräte zu unterstützen.

Die fünf verschmolzenen Kerne

Kernel Was es verschmilzt Beschreibung
AllGather + GEMM AllGather → GEMM Jeder Rang enthält eine Scherbe A. Während die Ränge die Shards ihrer Kollegen über NVLink/RDMA sammeln, verbraucht der lokale GEMM Kacheln, sobald sie ankommen.
GEMM + AllReduce GEMM → AllReduce Berechnet C = A @ B und reduziert Teilleistungen über alle Ränge hinweg in einem Begin. Ausgabekacheln werden in dem Second, in dem sie erzeugt werden, in den Reduktionsbaum verschoben.
MoE-Versand + GEMM All-to-All-Versand → gruppiertes GEMM Leitet MoE-Tokens an ihre Expertenränge weiter (NVLink innerhalb des Knotens + All-to-All zwischen den Knoten) und führt das je Experte gruppierte GEMM im selben Kernel aus. Token werden verarbeitet, sobald sie landen – kein Staging-Puffer-Roundtrip.
Rufen Sie Aufmerksamkeit an Ring KV-Austausch → FlashAchtung Sequenzparallele Aufmerksamkeit über alle Ränge hinweg. Bei jedem Schritt rotiert ein KV-Block im Ring, während die lokale FlashAttention den zuvor empfangenen Block verbraucht. Compute und Ring Ship/Recv werden gleichzeitig in einem einzigen persistenten Kernel ausgeführt.
GEMM + ReduceScatter GEMM → Scatter reduzieren Berechnet C = A @ B und reduziert die Ausgabe. Jedes Ausgabeplättchen wird reduziert und an seinen eigenen Rang weitergeleitet, sobald es produziert wird.

Evaluierungs-Setup

Das Forschungsteam evaluierte mKernel auf zwei 2-Knoten × 8-H200-Clustern, die sich nur in ihrer Inter-Knoten-Struktur unterscheiden:

Prüfstand Knoten × GPUs Intra-Knoten Transport zwischen Knoten NIC
AWS EFA 2 × 8 H200 NVLink AWS EFA / SRD 16 × 200 Gbit/s EFA professional Knoten
ConnectX-7 2 × 8 H200 NVLink InfiniBand 8 × 400 Gbit/s NVIDIA ConnectX-7 professional Knoten

mKernel wurde mit NCCL, Triton-distributed, Flux, Mercury, MagiAttention, Transformer-Engine und Ring-Flash-Consideration verglichen. Das Crew stellt fest, dass ein weiteres Benchmarking in größerem Maßstab noch im Gange ist.

Backends und Anforderungen

mKernel unterstützt zwei Netzwerk-Backends:

Backend Makro Transport Wo es läuft
CX7 -DINTERNODE_BACKEND_IBVERBS libibverbs RC ConnectX-7 / InfiniBand / RoCE
EFA -DINTERNODE_BACKEND_EFA libibverbs + efadv (SRD) AWS p5/p5e (H200, EFA)

Beide Backends nutzen dieselbe hostseitige API und denselben On-GPU-Kernel. Nur die Proxy-/Sitzungsimplementierung unterscheidet sich (session.h für CX7, session_efa.h für EFA). Anforderungen: NVIDIA Hopper-GPUs (Commonplace-Construct-Ziele). sm_90a), CUDA 12.9, Python mit PyTorch. Das CX7-Backend erfordert libibverbs-Entwicklungsheader und -Bibliotheken. Das EFA-Backend erfordert eine AWS EFA-Set up mit libfabric, libibverbs, efadv und EFA-Headern darunter EFA_HOME=/decide/amazon/efa standardmäßig.

Der visuelle Erklärer von Marktechpost

01 / 07 — Übersicht

Was ist mKernel?

mKernel ist eine Open-Supply-Bibliothek persistenter CUDA-Kernel aus dem UCCL-Projekt der UC Berkeley. Es vereint knoteninterne NVLink-Kommunikation, knotenübergreifendes RDMA und dichte Rechenleistung in einem einzigen Kernel.

Die meisten vorhandenen Fused-Kernel-Bibliotheken arbeiten innerhalb eines einzelnen Knotens oder einer einzelnen GPU. mKernel ist von Anfang an darauf ausgelegt, Knotengrenzen zu überbrücken.

43,6 %

des Vorwärtsdurchlaufs, der durch die Kommunikation in der Produktion verbraucht wird

47 %

der gesamten Ausführungszeit in gängigen MoE-Modellen

32 %

der durch Kommunikation verbrauchten Finish-to-Finish-Schulungszeit

02 / 07 – Das Downside

Warum Hostgesteuert Die Kommunikation greift zu kurz

Das Standardmodell ist hostgesteuert: Die CPU ruft NCCL oder NVSHMEM auf, wodurch kollektive Operationen über GPUs hinweg ausgeführt werden. Das UCCL-Crew identifiziert zwei Probleme.

CPUs skalieren nicht mit GPUs. Ein GB300 NVL72-Rack liefert 720 PFLOP/s FP8/FP6 und 1,44 EFLOP/s FP4. Bei diesen Geschwindigkeiten liegt der Overhead im Mikrosekundenbereich cudaLaunchKernelCPU-seitige Synchronisierungsprüfungen und Inter-Stream-Ereignisse werden direkt als Pipeline-Blasen angezeigt.

🔲

Die Überlappung ist zu grob. Hostgesteuerte Systeme überlappen Rechenleistung und Kommunikation nur an Kernelgrenzen. Eine feinkörnigere Überlappung auf Kachel- oder Blockebene ist von der Hostseite aus nicht möglich.

🔀

Die Antwort: GPU-gesteuerte Kommunikation. Die GPU selbst löst feinkörnige Übertragungen aus, die in denselben Kernel wie die Rechenleistung integriert sind.

03 / 07 — Design

Vier-Kern-Design Eigenschaften

🖧

Multi-GPU + Multi-Knoten in einem Kernel. Intra-Node-NVLink und Inter-Node-RDMA befinden sich beide im selben persistenten Kernel.

🔬

Feinkörnige Intra-Kernel-Überlappung. Rechen- und Kommunikationsüberschneidungen bei Kachel-/Chunk-Granularität, die sowohl die knoteninterne als auch die knoteninterne Kommunikation abdecken.

⚙️

Persistenter Kernel mit SM-Spezialisierung. CTAs weisen sich selbst Rollen zu: compute, intra-comm, inter-send, inter-reduce. Der SM-Cut up ist je nach Type einstellbar.

📡

GPU-gesteuerte Vernetzung über libibverbs. Verwendet GPU-initiierte RDMA-Schreibvorgänge. Keine NCCL- oder NVSHMEM-Abhängigkeit. Das Kommunikations-Backend wurde von Grund auf neu geschrieben.

04 / 07 – Kernel

Die Fünf Fusionierte Kernel

AllGather + GEMM

AllGather —> GEMM

Jeder Rang enthält eine Scherbe A. Das lokale GEMM verbraucht Kacheln über NVLink/RDMA, sobald sie ankommen – Matmul beginnt, bevor das Kollektiv endet.

GEMM + AllReduce

GEMM —> AllReduce

Berechnet C = A @ B und reduziert Teilleistungen über alle Ränge hinweg in einem Begin. Ausgabekacheln werden in dem Second, in dem sie erzeugt werden, in den Reduktionsbaum aufgenommen.

MoE-Versand + GEMM

All-to-All-Versand —> gruppiertes GEMM

Leitet MoE-Token über NVLink + knotenübergreifendes All-to-All an die Expertenränge weiter und führt dann professional Experte gruppiertes GEMM im selben Kernel aus. Kein Staging-Puffer-Roundtrip.

Rufen Sie Aufmerksamkeit an

Ring KV-Austausch —> FlashAttention

Sequenzparallele Aufmerksamkeit über alle Ränge hinweg. Bei jedem Schritt rotiert ein KV-Block im Ring, während die lokale FlashAttention den zuvor empfangenen Block verbraucht.

GEMM + ReduceScatter

GEMM —> ReduceScatter

Berechnet C = A @ B und reduziert die Ausgabe. Jedes Plättchen wird reduziert und an seinen eigenen Rang weitergeleitet, sobald es produziert wird.

05 / 07 — Auswertung

Auswertung Aufstellen

Getestet auf zwei 2-Knoten × 8-H200-Clustern, die sich nur in der Inter-Node-Cloth unterscheiden.

Prüfstand Knoten × GPUs Zwischenknoten NIC
AWS EFA 2 × 8 H200 AWS EFA / SRD 16 × 200 Gbit/s EFA professional Knoten
ConnectX-7 2 × 8 H200 InfiniBand 8 × 400 Gbit/s CX7 professional Knoten

Beide Testumgebungen verwenden NVLink innerhalb des Knotens. Im Vergleich: NCCL, Triton-distributed, Flux, Mercury, MagiAttention, Transformer-Engine und Ring-Flash-Consideration. Ein größeres Benchmarking ist noch im Gange.

06 / 07 – Backends & Anforderungen

Backends & Anforderungen

Backend Transport Wo es läuft
CX7 libibverbs RC ConnectX-7 / InfiniBand / RoCE
EFA libibverbs + efadv (SRD) AWS p5/p5e (H200, EFA)

📋

Anforderungen: NVIDIA Hopper-GPUs (Commonplace sm_90a), CUDA 12.9, Python mit PyTorch. CX7 benötigt libibverbs-Header. EFA benötigt libfabric, libibverbs, efadv unter EFA_HOME=/decide/amazon/efa.

📝

Lizenz & Namensnennung: MIT-lizenziert. MMA/Compute-Code angepasst von ThunderKittens (HazyResearch).

07 / 07 – Roadmap und wichtige Erkenntnisse

Roadmap & Wichtige Erkenntnisse

Fusionierte GPU-gesteuerte Multi-Node-Kernel (AG+GEMM, GEMM+AR, MoE Dispatch+GEMM, Ring Consideration, GEMM+RS)

ConnectX-7- und AWS EFA-Backends

🚧

Vollständige heterogene Beschleuniger-/NIC-Unterstützung mit topologiebewusster Erkennung, Platzierung und Routing

🚧

Interknoten-Megakernel: Zusammenführung mehrerer verschmolzener Schritte zu einem einzigen Megakernel, der sich über eine Transformatorschicht erstreckt

🚧

Blackwell-GPU-Unterstützung

Verschmelzt NVLink, knotenübergreifendes RDMA und Rechenleistung in einem einzigen persistenten CUDA-Kernel

Fünf Kernel: AllGather+GEMM, GEMM+AllReduce, MoE Dispatch+GEMM, Ring Consideration, GEMM+ReduceScatter

GPU-initiiertes RDMA über libibverbs – keine NCCL- oder NVSHMEM-Abhängigkeit

Erfordert Hopper-GPUs (sm_90a) und ConnectX-7 oder AWS EFA-Netzwerke

Wichtige Erkenntnisse

  • mKernel vereint knoteninternes NVLink, knoteninternes RDMA und Rechenleistung in einem einzigen persistenten CUDA-Kernel.
  • Der Kommunikationsaufwand macht bis zu 47 % der Ausführungszeit in MoE-Modellen professional zitierter Produktionsdaten aus.
  • Fünf Kernel sind enthalten: AllGather+GEMM, GEMM+AllReduce, MoE Dispatch+GEMM, Ring Consideration und GEMM+ReduceScatter.
  • GPU-initiiertes RDMA wird direkt über implementiert libibverbs – keine NCCL- oder NVSHMEM-Abhängigkeit.
  • Erfordert derzeit Hopper-GPUs (sm_90a) und ConnectX-7 oder AWS EFA-Netzwerk; Blackwell-Unterstützung steht auf der Roadmap.

Schauen Sie sich das an Repo Und Technische Particulars. Sie können uns auch gerne weiter folgen Twitter und vergessen Sie nicht, bei uns mitzumachen 150.000+ ML SubReddit und Abonnieren Unser E-newsletter. Warten! Bist du im Telegram? Jetzt können Sie uns auch per Telegram kontaktieren.

Möchten Sie mit uns zusammenarbeiten, um Ihr GitHub-Repo ODER Ihre Hugging Face Web page ODER Produktveröffentlichung ODER Ihr Webinar usw. zu bewerben? Vernetzen Sie sich mit uns


Von admin

Schreibe einen Kommentar

Deine E-Mail-Adresse wird nicht veröffentlicht. Erforderliche Felder sind mit * markiert