DeepGEMM - 効率的なFP8行列乗算ライブラリ

細粒度スケーリングを備えた明確で効率的なFP8行列乗算ライブラリ、NVIDIA Hopper アーキテクチャ向けに最適化

DeepGEMM Matrix Multiplication Visualization

DeepGEMM アーキテクチャ設計

DeepGEMMは、NVIDIA Hopper テンソルコア向けに最適化された慎重に設計されたアーキテクチャを特徴とし、効率的なFP8行列乗算を可能にします。

DeepGEMM Architecture Diagram

完全JIT設計

DeepGEMMはインストール時のコンパイルが不要な完全なJust-In-Time(JIT)設計を採用しています。すべてのカーネルは軽量JITモジュールを使用して実行時にコンパイルされ、GEMM形状、ブロックサイズ、パイプラインステージをコンパイル時定数として扱い、レジスタを節約しコンパイラの最適化を可能にします。

Hopper TMA機能

Hopperアーキテクチャのテンソルメモリアクセラレータ(TMA)を完全に活用し、TMAロード、ストア、ブロードキャスト、ディスクリプタプリフェッチを含む高速で非同期のデータ移動を実現し、パフォーマンスを大幅に向上させます。

コア最適化技術

DeepGEMMは、従来のGEMMライブラリを超える複数の高度な最適化技術を採用し、優れたパフォーマンスを実現します。

DeepGEMM Optimization Techniques

永続的スレッド特殊化

CUTLASSデザインに従い、データ移動、テンソルコアMMA命令、CUDAコアプロモーションをオーバーラップさせ、計算効率を向上させます。

非整列ブロックサイズ

SMをより効率的に活用するために非整列ブロックサイズ(112など)をサポートします。例えば、M=256、N=7168の場合、BLOCK_M=128、BLOCK_N=112を使用すると、整列ブロックサイズの112個ではなく128個のSMが動作できます。

FFMA SASS インターリービング

コンパイルされたバイナリ内のFFMA命令を修正し、yieldビットとreuseビットを反転させることで、MMA命令とプロモーションFFMA命令のオーバーラップの機会を増やし、場合によっては10%以上のパフォーマンス向上を実現します。

統合最適化ブロックスケジューラ

非グループ化およびグループ化カーネルすべてに対応する1つのスケジューラで、ラスタライゼーションを採用してL2キャッシュの再利用を強化し、全体的なパフォーマンスを向上させます。

コード解説

DeepGEMMのコアカーネル関数はわずか約300行のコードで、クリーンな設計により、Hopper FP8行列乗算と最適化技術について学びやすくなっています。

// Simplified version of the core GEMM kernel
template <
    int BLOCK_M,
    int BLOCK_N,
    int BLOCK_K,
    int CLUSTER_M,
    int CLUSTER_N,
    int STAGES,
    int WARPGROUP_SIZE,
    bool ENABLE_INTERLEAVE
>
__global__ void gemm_fp8_fp8_bf16_nt_kernel(
    const void* __restrict__ lhs,
    const void* __restrict__ lhs_scale,
    const void* __restrict__ rhs,
    void* __restrict__ output,
    const void* __restrict__ rhs_scale,
    int m, int n, int k,
    int lhs_stride, int rhs_stride, int output_stride) {
    
    // Using TMA to load LHS, LHS scale factors, and RHS
    // Using CUDA cores for two-level accumulation (promotion)
    // Using TMA to store output
    
    // Persistent thread specialization implementation
    // Unaligned block size support
    // FFMA SASS interleaving optimization
}

コード構造の特徴

  • 異なるブロックサイズと最適化構成をサポートするテンプレート設計
  • 効率的なデータ移動のためのHopperアーキテクチャのTMA機能の完全活用
  • FP8テンソルコアの不正確な累積問題を解決する二段階累積
  • 計算とデータ移動のオーバーラップを最適化する永続的スレッド特殊化の実装
DeepGEMM Code Flow Diagram

インターフェース設計

  • 通常の密行列GEMM: deep_gemm.gemm_fp8_fp8_bf16_nt関数を使用
  • グループ化GEMM(連続レイアウト): m_grouped_gemm_fp8_fp8_bf16_nt_contiguous関数を使用、エキスパートが同じ形状を共有するMoEモデルに適しています
  • グループ化GEMM(マスクレイアウト): m_grouped_gemm_fp8_fp8_bf16_nt_masked関数を使用、CUDAグラフが有効な推論デコードフェーズで、CPUが各エキスパートが受け取るトークン数を知らない場合に使用

パフォーマンス

DeepGEMMのパフォーマンスは、様々な行列形状において専門家によって調整されたライブラリと同等かそれ以上で、場合によっては最大2.7倍の高速化を実現します。

DeepGEMM Performance Comparison

パフォーマンスのハイライト

  • 通常のGEMM:特定の形状で最大2.7倍の高速化
  • 連続レイアウトのグループ化GEMM:最大1.2倍の高速化
  • マスクレイアウトのグループ化GEMM:同じく最大1.2倍の高速化
  • H800 SXM5でテストした計算パフォーマンスは最大1358 TFLOPS
  • メモリ帯域幅の高い利用率、最大2668 GB/s達成

DeepGEMMの主な特徴

卓越したパフォーマンス

DeepGEMMはNVIDIA Hopperアーキテクチャ上で専門家によって調整されたライブラリと比較して最大2.7倍の高速化を実現し、1350 TFLOPSを超える計算性能を達成します

高度な最適化

DeepGEMMは永続的スレッド特殊化、非整列ブロックサイズ、FFMA SASSインターリービングを採用し、最大の計算効率を実現します

柔軟な統合

DeepGEMMは明確で十分に文書化されたインターフェースを通じて、深層学習フレームワークや科学計算ライブラリとシームレスに統合します

細粒度スケーリング

DeepGEMMは最適化されたブロックスケジューリングとラスタライゼーション技術により、小さな行列から大規模計算まで効率的にスケールします

ネイティブFP8サポート

DeepGEMMはNVIDIA Hopperテンソルコア向けに最適化された一流のFP8精度サポートを提供します

高メモリ帯域幅

DeepGEMMは最大2668 GB/sのメモリ帯域幅利用率を達成し、ハードウェア能力を最大限に活用します

クリーンなコード設計

DeepGEMMは明確な設計パターンを持つ簡潔な300行のコアカーネルを特徴とし、理解と拡張が容易です

最適化されたグループGEMM

DeepGEMMは連続レイアウトとマスクレイアウトの両方のグループGEMM操作に優れており、混合エキスパート(MoE)モデルに最適です

よくある質問

DeepGEMMとは何か、何が特別なのですか?
DeepGEMMはNVIDIA Hopperアーキテクチャ向けに最適化された効率的なFP8行列乗算ライブラリです。特別な点は、クリーンな設計(コアコードはわずか約300行)、卓越したパフォーマンス(専門家調整ライブラリと比較して最大2.7倍の高速化)、永続的スレッド特殊化やFFMA SASSインターリービングなどの高度な最適化技術にあります。
DeepGEMMはどのハードウェア向けに最適化されていますか?
DeepGEMMは特にNVIDIA Hopperアーキテクチャ向けに最適化されており、HopperのテンソルメモリアクセラレータTMA機能とFP8行列乗算用のテンソルコアを最大限に活用しています。H800 SXM5 GPUでテストされ、最大1358 TFLOPSの計算性能を達成しています。
DeepGEMMはどのようにパフォーマンス向上を実現していますか?
DeepGEMMは複数の高度な技術を通じてパフォーマンスを実現しています:コンパイル時定数を持つ完全JIT設計、SMをより効率的に利用するための非整列ブロックサイズ、FFMAのSASSインターリービング、FP8精度問題を解決する二段階累積、L2キャッシュの再利用を強化するラスタライゼーションを採用した統合最適化ブロックスケジューラなどです。
DeepGEMMはどのようなGEMM操作をサポートしていますか?
DeepGEMMは3つの主要なGEMM操作をサポートしています:通常の密行列GEMM(deep_gemm.gemm_fp8_fp8_bf16_nt関数を使用)、連続レイアウトのグループ化GEMM、マスクレイアウトのグループ化GEMMです。後者2つは特に混合エキスパート(MoE)モデルに有用です。
DeepGEMMは混合エキスパート(MoE)モデルに適していますか?
はい、DeepGEMMはMoEモデルに非常に適しています。連続レイアウトとマスクレイアウトの両方のグループ化GEMM操作を提供し、MoEアーキテクチャの特定の計算パターン向けに最適化されています。これらの操作は他のライブラリと比較して最大1.2倍の高速化を実現できます。
DeepGEMMはどの精度をサポートしていますか?
DeepGEMMは主にFP8精度に焦点を当てており、これは効率的なAIモデル推論にますます重要になっています。FP8テンソルコアで発生する可能性のある不正確な累積問題を解決するために二段階累積アプローチを実装し、パフォーマンスと精度の両方を確保しています。
DeepGEMMのメモリ帯域幅利用率は他のライブラリと比較してどうですか?
DeepGEMMは最大2668 GB/sという優れたメモリ帯域幅利用率を達成しています。これはHopperのTMA機能を効率的に使用して、TMAロード、ストア、ブロードキャスト、ディスクリプタプリフェッチを含む高速で非同期のデータ移動を実現しているためです。
DeepGEMMの実装から自分のCUDAコードを最適化するために学ぶことはできますか?
もちろんです!DeepGEMMのコアカーネルはわずか約300行のコードでクリーンな設計を持ち、優れた学習リソースとなります。高度なCUDA最適化技術、Hopperアーキテクチャ機能の効率的な利用、他の高性能計算タスクにも適用できる巧妙な行列乗算アプローチを示しています。