論文の概要: GC-MoE: Overlapping Expert-Parallel Token Dispatch with FP8/FP4 Grouped GEMM via CUDA Green Contexts
- arxiv url: http://arxiv.org/abs/2611.04222v1
- Date: Tue, 03 Nov 2026 03:04:04 GMT
- ステータス: 翻訳完了
- システム内更新日: 2026-11-11 01:01:14.945455
- Title: GC-MoE: Overlapping Expert-Parallel Token Dispatch with FP8/FP4 Grouped GEMM via CUDA Green Contexts
- Title(参考訳): GC-MoE: CUDA Green ContextsによるFP8/FP4 Grouped GEMMによるエキスパート並列トークンの重複分散
- Abstract要約: GC-MoEは、Hopper上のDeepSeek-V4.1Flashデコードのための融合演算子である。
Mega MoEは54点満点の精度でMega MoEを破る。
GC-MoEは148機中20機しかB300に配備していない。
- 参考スコア(独自算出の注目度): 6.990897820951724
- License: http://creativecommons.org/licenses/by/4.0/
- Abstract: Under expert parallelism, mixture-of-experts (MoE) layers pay all-to-all token dispatch, grouped GEMMs, and a combine step in every layer. In decode each rank holds only 1-256 tokens: dispatch latency and weight traffic dominate while tensor-core utilization is inherently low. In fused operators such as DeepGEMM's Mega MoE, dispatch and computation overlap only by sharing every SM: dispatch warps compete with TMA/MMA and epilogue warps for registers, shared memory, and issue slots; every phase shares one CTA resource budget, synchronization spans the whole SM array, and every replay ends with a workspace-clean pass and a cross-rank barrier. We present GC-MoE, an open-source FP8/FP4 MoE inference operator for NVIDIA Hopper and Blackwell GPUs that uses CUDA Green Contexts to split the device into a small communication partition and a large compute partition. Each role gets its own kernel, resource budget, and synchronization domain, so dispatch and GEMM tune independently. A three-node CUDA Graph (four on Hopper with fused shared experts) launches the dispatch and grouped-GEMM kernels as concurrent roots; a token-ring handshake starts computation on the first filled ring blocks, and double-buffered receive counters remove the trailing cross-rank barrier between replays. The compute partition runs operand-swapped tcgen05 block-scaled MMA on Blackwell and register-A WGMMA with in-register FP4 decoding on Hopper, with fused SwiGLU, requantization, and optional shared experts. On 8 H20 GPUs GC-MoE accelerates DeepSeek-V4.1-Flash decode by 1.188x on average on both the routed and shared-expert suites (peaks 1.296x and 1.353x); on 8 B300 GPUs it accelerates GLM-5.2 and DeepSeek-V4.1-Flash by 1.106x (1.123x with fused shared experts, peak 1.209x), beating Mega MoE at all 54 measured points at matched accuracy. GC-MoE reserves only 20 of 148 SMs on B300 for dispatch.
- Abstract(参考訳): 専門家の並列性の下では、Mix-of-experts(MoE)レイヤはすべてのトークンディスパッチ、GEMMのグループ化、すべてのレイヤのコンバインステップを支払う。
デコードでは、各ランクは1-256トークンのみを保持する: ディスパッチレイテンシとウェイトトラフィックが支配的であり、テンソルコアの利用は本質的に低い。
DeepGEMMのMega MoEのような融合演算子では、ディスパッチワープはTMA/MMAと競合し、レジスタ、共有メモリ、発行スロットが競合する。
本稿では,NVIDIA Hopper と Blackwell GPU 用のオープンソース FP8/FP4 MoE 推論演算子 GC-MoE について述べる。
各ロールには独自のカーネル、リソース予算、同期ドメインがあり、ディスパッチとGEMMは独立してチューニングされる。
3ノードのCUDAグラフ(4つは融合された共有専門家を持つホッパー)はディスパッチとグループ化されたGEMMカーネルを同時ルートとして起動し、トークンリングハンドシェイクは最初のリングブロックの計算を開始し、ダブルバッファの受信カウンタはリプレイ間の後続のクロスランク障壁を取り除く。
計算パーティションは、オペランドスワップされたtcgen05ブロックスケールのMMAをBlackwell上で実行し、レジスター-A WGMMAは、ホッパー上でFP4デコードされ、融合されたSwiGLU、再量子化、オプションで共有された専門家によって構成される。
8 H20 GPUでは、GC-MoEはDeepSeek-V4.1-Flashデコードの平均1.188xで、ルーティングおよび共有された専門家スイート(1.296xと1.353x)で加速し、8 B300 GPUではGLM-5.2とDeepSeek-V4.1-Flashを1.106x(1.123xは融合された共有専門家、ピーク1.209x)で加速し、Mega MoEを54の精度で上回った。
GC-MoEは148機中20機しかB300に配備していない。
関連論文リスト
- Evolving Agent Teams [58.34711503048237]
既存のコーディングエージェントは、実世界のエキスパートが書いたカーネル上で局所最適化に収束する。
ローカルな最適化から逃れるマルチエージェントフレームワークであるEvolving Agent Teamsを紹介する。
1)各ローカルの最適度を診断し、エージェントのチームを再構成する中央スーパーバイザーと、(2)学習文書を構成する永続ファイルシステムとを備えている。
論文 参考訳(メタデータ) (2026-10-27T08:46:21Z) - FlashMask-3: Efficient and Expressive Mask-Aware Distributed Attention [4.549299243411132]
FlashMask-3は、FlashAttention-3のTMA/WGMMAパイプライン上に3ロールのワープを特化したマスク対応のHopperカーネルを導入した。
4K--128Kコンテキストで12のマスクにまたがる1つのH100では、FlashMask、FlexAttention、MagiAttentionよりも45-141%、34-71%、および2--44%のTFLOPを実現している。
論文 参考訳(メタデータ) (2026-10-13T07:55:51Z) - Cascadia: Resident 975B MoE Inference on Eleven AI PCs [0.0]
我々は、11台のIntel Core Ultra X7 358H AI PC上で、カスカディアのInklingの常駐実行を示す。
Inklingのルーティングルールを保存し、OpenVINOの融合したiGPUプリミティブのための圧縮グラフを構築するカスタム常駐MoEエンジンにコントリビュートする。
ストリーミングパイプラインは同時生成をコーディネートし、キャプチャ状態のドラフト評価は、展開された数値パスと一致させる。
論文 参考訳(メタデータ) (2026-10-05T18:29:24Z) - VC-Attention: Value Smoothing and Softmax Casting for Low-bit Attention [47.30761996669105]
トレーニング不要な低ビットアテンションフレームワークであるVC-Attentionを提案する。
V-Smoothは、オンラインの軽量クラスタリングによってバリュートークンを再注文するので、融合ハードウェアブロック内のトークンは、共に量子化される。
ログドメインのスコアを直接E4M3確率符号に1つの融合乗算加算でマッピングし、FP32指数関数とフォーマット変換を排除した。
論文 参考訳(メタデータ) (2026-09-14T16:17:35Z) - Unfolding the Leech Lattice: Fused Multi-Shell Decoding and VRAM Layouts for 2-Bit LLM Weights [8.550323073861415]
リーチ格子ベクトル量子化は、最強の報告された2ビット品質を独自の評価プロトコルで保持する。
本報告では, バッチ1におけるデコード相GEMVの供用コストを計測する。
論文 参考訳(メタデータ) (2026-09-02T14:26:06Z) - Hand-Written PTX Tensor-Core GEMM Kernels: A Multi-Precision Study on NVIDIA L4 [0.7265954868708899]
コアカーネルは非同期データ移動から構築された低レベルのPTXパイプラインに依存している。
ほとんどのアプリケーションコードは、WMMA C++ APIを通じて間接的にCoresにアクセスする。
WMMAを手書きのPTXに置き換えるのはいつ頃なのか?
論文 参考訳(メタデータ) (2026-08-10T18:13:00Z) - FusionML: Prefill, Not Decode - Mechanism and Boundaries of CPU+GPU Co-Execution on Unified-Memory Apple Silicon [0.0]
私たち自身を含む以前の試みは、失敗に終わり、あるいは正確で確立された勝利を生み出しました。
MLXの遅延グラフスケジューラは、クロスストリーム作業を行う CPU ストリーム操作が1つの評価グラフ内で非マテリアル化された GPU 結果を消費する場合、行分割行列は、物質化された入力で x1.38 を高速に実行する。
デコードでは利益が得られず、共有帯域幅によるバウンドは追加されない。
論文 参考訳(メタデータ) (2026-07-24T11:42:17Z) - RoPE-Aware Bit Allocation for KV-Cache Quantization [52.099459337231345]
Block-GTQはTurboQuant-MSE上に構築されたキーキャッシュ量子化のためのビットアロケータである。
これは、RoPEクエリキーのロジットを10モデル診断パネルに保存する。
128Kコンテキストでfp16 FlashAttention2より1.34倍高速で動作する。
論文 参考訳(メタデータ) (2026-06-23T00:17:48Z) - dMoE: dLLMs with Learnable Block Experts [71.572316901001]
単純なブロックレベルのMoEフレームワークであるdMoEを提案する。
また,dMoEは,推論中に一意に活性化される専門家の数を著しく減少させることを示した。
また、メモリ使用量を76.64%から79.84%に減らし、1.14$times$から1.66$times$エンドツーエンドのレイテンシ高速化を実現している。
論文 参考訳(メタデータ) (2026-05-29T06:03:50Z) - Breaking the Memory Barrier: Near Infinite Batch Size Scaling for Contrastive Loss [59.835032408496545]
本稿では, コントラスト損失計算を任意の小ブロックに分割するタイルベースの戦略を提案する。
分散システムの階層構造を活用するためのマルチレベルタイリング戦略も導入する。
SOTAメモリ効率のソリューションと比較すると、同等の速度を維持しながら、メモリの2桁の削減を実現している。
論文 参考訳(メタデータ) (2024-10-22T17:59:30Z) - K-Net: Towards Unified Image Segmentation [78.32096542571257]
K-Netと名付けられたこのフレームワークは、学習可能なカーネルのグループによってインスタンスとセマンティックカテゴリの両方を一貫して分割する。
K-Netは双方向マッチングでエンドツーエンドでトレーニングすることができ、そのトレーニングと推論は自然にNMSフリーで、ボックスフリーである。
論文 参考訳(メタデータ) (2021-06-28T17:18:21Z)
関連論文リストは本サイト内にある論文のタイトル・アブストラクトから自動的に作成しています。
指定された論文の情報です。
本サイトの運営者は本サイト(すべての情報・翻訳含む)の品質を保証せず、本サイト(すべての情報・翻訳含む)を使用して発生したあらゆる結果について一切の責任を負いません。