> [!abstract] 概要 > GPU は AI・深層学習を含む現代の汎用ワークロードの計算需要に対応するため進化し続けている。 > 既存研究はさまざまな GPU 世代のマイクロアーキテクチャ指標を解明してきたが、最新の Hopper GPU が備える FP8 対応 Tensor Core、DPX、分散共有メモリの性能と動作特性には不明点が残っている。 > 本研究は、NVIDIA GPU の新しい命令セットアーキテクチャ(ISA)と CUDA API を用いて、Hopper GPU を対象としたベンチマーク研究を行う。 > 具体的には、Hopper・Ada・Ampere の3世代についてレイテンシとスループットを比較し、Hopper の DPX 命令、分散共有メモリ、FP8 Tensor Core を詳しく評価する。 > 提示したマイクロベンチマークは、Hopper が導入した GPU AI 機能ユニットとプログラミング機能の理解を深め、GPU プログラムの最適化と性能モデル化を支援する。 > 著者らの知る限り、Hopper 固有の Tensor Core の性能とプログラミング命令セットを解明する初めての試みである。 ## 論文情報 - タイトル: Benchmarking and Dissecting the Nvidia Hopper GPU Architecture - 著者: [[Weile Luo]]、[[Ruibo Fan]]、[[Zeyu Li]]、[[Dayou Du]]、[[Qiang Wang]]、[[Xiaowen Chu]] - 所属: [[Hong Kong University of Science and Technology, Guangzhou]]、[[Harbin Institute of Technology (Shenzhen)]] - 媒体: arXiv cs.AR - 発表日: 2024-02-21 - 識別子: arXiv:2402.13499v1 ## 概要 本論文は、Ampere の A100 PCIe、Ada Lovelace の RTX4090、Hopper の H800 PCIe を同じ測定枠組みで比較し、メモリ階層・Tensor Core・Transformer Engine・Hopper 固有の CUDA 機能を命令レベルからアプリケーションレベルまで評価する。 中心的な結論は、Hopper では従来の `mma` より `wgmma` を使うこと、計算規模に応じて低精度形式を選ぶこと、DPX・非同期データ移動・分散共有メモリをワークロードに合わせて使い分けることである。 ## 問題設定 GPU 世代更新で追加された機能は、公式仕様だけでは命令レベルのレイテンシ、実効スループット、データ移動との関係を判断しにくい。 特に Hopper では、FP8 Tensor Core、warp-group 単位の非同期 `wgmma`、DPX、Tensor Memory Accelerator(TMA)、SM 間の分散共有メモリが追加された。 そこで、異なる世代の基本性能を比較しつつ、新機能がどの実行条件で有効になるかを、再現可能なマイクロベンチマークで調べる。 ## 提案手法 ### 測定対象と階層 - メモリ階層: L1 キャッシュ、共有メモリ、L2 キャッシュ、グローバルメモリのレイテンシと帯域幅。 - Tensor Core: PTX の `mma`・`mma.sp`・`wgmma`・`wgmma.sp` を SASS に逆アセンブルし、完了レイテンシとスループットを測定。 - Transformer Engine: `te.Linear`、`te.TransformerLayer`、Llama の生成を、FP32・FP16・BF16・FP8 で比較。 - Hopper 機能: DPX 命令、Ampere の非同期コピーと Hopper の TMA、分散共有メモリ。 ### メモリ測定 ポインタ追跡型のマイクロベンチマークを使い、L1 には `ca`、L2 には `cg` 修飾子を付けて対象階層を温める。 グローバルメモリでは L2 より大きい領域を使い、TLB を温めた後に4スレッドが各8バイトを読む。 帯域幅測定では `float4` 相当のベクトル化アクセスを使い、アクセス量を経過時間で割る。 ### Tensor Core 測定 同期 `mma` は1 warp、非同期 `wgmma` は4 warpからなる warp group で発行し、CUDA カーネル内で各命令を1024回実行する。 レイテンシは命令発行から後続命令が結果を利用できるまでの完了レイテンシとし、スループットは `Total_OPS / Duration` で計算する。 GPU 周波数が命令ごとに変わりうるため、総クロック数ではなく実時間を分母に用いる。 ### Transformer Engine と LLM 生成 `te.Linear` では入力と重みを FP8 化し、最大絶対値をスケールとして `inp_fp8 = inp_fp16 / scale` を計算してから行列積を行い、出力を再スケールする。 この形式変換・量子化が小さな行列では大きなオーバーヘッドになる。 `te.TransformerLayer` では LayerNorm と MLP を融合し、層間の FP8 データ移動を可能にする一方、Softmax・GeLU は FP8 化されず、DotProductAttention は FlashAttention を使う。 Llama の評価では `nn.Linear` と RMSNorm を Transformer Engine の実装へ置き換え、ShareGPT の会話長から合成要求を作る。入力長と生成長は最大128、バッチサイズは8とする。 ### Hopper 固有機能 DPX は動的計画法で頻出する最小値・最大値演算を対象とし、単一スレッドでレイテンシ、ブロックで SM あたりのスループットを測る。 非同期データ移動では、同期コピーの `SyncShare` と二段パイプラインの `AsyncPipe` を比較する。 分散共有メモリでは、`cluster.map_shared_rank(SMEM, DST_BLOCK_RANK)` で別ブロックの共有メモリを参照し、レイテンシ、Ring-Based Copy(RBC)の帯域、分散共有メモリ版ヒストグラムを評価する。 ## 新規性 先行する GPU マイクロアーキテクチャ研究が Volta・Turing・Ampere の Tensor Core を主に扱ってきたのに対し、本論文は Hopper の `wgmma` と SASS 対応を命令レベルで調べる。 さらに、Tensor Core の比較を Transformer Engine と Llama 生成へ接続し、DPX・TMA・分散共有メモリを含む Hopper 固有の CUDA 機能を、Ampere・Ada との横断比較としてまとめている。 ## 実験設定 ベンチマーク対象は A100 PCIe(40GB HBM2e、108 SM、1555 GB/s)、RTX4090(24GB GDDR6X、128 SM、1008 GB/s)、H800 PCIe(80GB HBM2e、114 SM、2039 GB/s)である。 RTX4090 ではドライバ530.30.02・CUDA 12.1、A100 と H800 ではドライバ535.104.05・CUDA 12.2を使った。 **表III: 評価 GPU の構成** ![[_attachments/arxiv-2402.13499/table03-device-properties.png]] メモリと Tensor Core の評価には命令レベルのマイクロベンチマークを使い、AI 評価には Transformer Engine と Llama を使う。DPX・非同期コピー・分散共有メモリには CUDA サンプルを基にした実装を使う。 ## 実験結果 ### Hopper の構造と Tensor Core 命令 Hopper の構造は、複数 SM を含む Graphics Processing Cluster、SM 間ネットワーク、L2 キャッシュ/HBM グローバルメモリ、SM 内のレジスタ・Tensor Memory Accelerator・L1/共有メモリで構成される。 **図1: Hopper アーキテクチャ** ![[_attachments/arxiv-2402.13499/fig01-hopper-architecture.png]] `mma` は1 warpが同期実行し、`wgmma` は4 warpの warp group が非同期実行する。`wgmma` は行列を共有メモリから直接読み出せるため、`mma` のように全行列をレジスタへ置く必要がない。 **図2: `mma` と `wgmma` の命令形式** ![[_attachments/arxiv-2402.13499/fig02-mma-wgmma-instructions.png]] **表I: 世代別 Tensor Core の精度・プログラミング方式・実行モード** ![[_attachments/arxiv-2402.13499/table01-tensor-core-properties.png]] Hopper の `mma` は FP8 を扱わず、INT4 は Tensor Core 命令ではなく一連の `IMAD` として CUDA コア上で実行される。一方、`wgmma` は FP8 の E5M2/E4M3 と非同期実行を扱い、SASS では GMMA 系命令になる。 **表VI: Hopper の PTX 命令と SASS 命令の対応** ![[_attachments/arxiv-2402.13499/table06-sass-instructions.png]] H800 の `mma` は理論ピークの平均62.9%にとどまり、A100 は95%超だった。疎 `mma` の dense に対する平均高速化は H800 で1.42倍であり、RTX4090 では最大約2倍だった。 H800 の dense `wgmma` はゼロ初期化で理論ピークの95%超に達した。`N >= 64` ではピークに近いスループットを得たが、`N < 64` では共有メモリ待ちを計算で隠しにくくなった。 **表VII: dense/sparse `mma` のレイテンシとスループット** ![[_attachments/arxiv-2402.13499/table07-mma-results.png]] **表VIII: dense `wgmma` の測定結果** ![[_attachments/arxiv-2402.13499/table08-dense-wgmma-results.png]] **表IX: sparse `wgmma` の測定結果** ![[_attachments/arxiv-2402.13499/table09-sparse-wgmma-results.png]] **表X: `wgmma` の N 値による変化** ![[_attachments/arxiv-2402.13499/table10-wgmma-n-values.png]] ランダム値の行列では、H800 が350Wの電力上限に近づいて周波数を下げるため、ゼロ初期化より Tensor Core 性能が低下した。 `mma` の最大形状で測った H800 のエネルギー効率は、dense で A100 の1.60倍・RTX4090 の1.69倍、sparse でそれぞれ1.33倍・1.39倍だった。 **表XI: `mma` の消費電力とエネルギー効率** ![[_attachments/arxiv-2402.13499/table11-energy-efficiency.png]] ### メモリ階層 3 GPU の L1 と共有メモリの最大スループットは近かった。H800 の L2 キャッシュスループットは RTX4090 の2.6倍、A100 の2.2倍であり、測定値は理論性能の RTX4090 で92%、A100 で90%、H800 で91%に達した。 平均すると L2 レイテンシは L1 の6.5倍、グローバルメモリは L2 の1.9倍だった。 **表IV: メモリ階層のレイテンシ** ![[_attachments/arxiv-2402.13499/table04-memory-latency.png]] **表V: メモリ階層のスループット** ![[_attachments/arxiv-2402.13499/table05-memory-throughput.png]] ### Transformer Engine と LLM `te.Linear` は小さい行列では形式変換の割合が大きいが、行列サイズ `N=8192` で FP8 の利点が現れ、`N=16384` では H800 と RTX4090 の FP8 が FP16 のほぼ2倍のスループットになった。 **図3: FP8 `te.Linear` の実行時間内訳** ![[_attachments/arxiv-2402.13499/fig03-te-linear-time-breakdown.png]] **図4: `te.Linear` のスループット** ![[_attachments/arxiv-2402.13499/fig04-te-linear-throughput.png]] **表II: `te.TransformerLayer` の設定** ![[_attachments/arxiv-2402.13499/table02-transformer-layer-settings.png]] `te.TransformerLayer` では、FP16 は FP32 のほぼ2倍の速度だった。FP8 は hidden size が4096を超えると FP16 を上回るが、全演算とデータ移動が FP8 化されないため2倍にはならなかった。 **図5: `te.TransformerLayer` のレイテンシ** ![[_attachments/arxiv-2402.13499/fig05-te-transformer-layer-latency.png]] 短い入出力長の decode-only LLM 生成はメモリ律速であり、FP8 Tensor Core の計算性能が直接スループットへ反映されにくかった。4090 の Llama-3B では FP32/BF16/FP8 が414.08/425.19/429.31 token/s、H800 の Llama-2-13B では357.57/399.38/356.11 token/sだった。 **表XII: GPU・モデル・精度別の推論スループット** ![[_attachments/arxiv-2402.13499/table12-inference-throughput.png]] ### DPX、非同期データ移動、分散共有メモリ DPX は RTX4090 と A100 ではソフトウェアエミュレーションであり、H800 は一部の `relu`・16ビット演算で大きく高速化した。16ビット演算では最大13倍の高速化を確認したが、単純な `__viaddmax_s32` などでは世代間差が小さかった。 H800 では起動ブロック数が SM 数の整数倍を越えた直後に DPX スループットが落ち、整数倍付近で最大に戻った。この挙動から、DPX 加速ユニットは SM レベルにあると推測される。 **図6: DPX のレイテンシ** ![[_attachments/arxiv-2402.13499/fig06-dpx-latency.png]] **図7: DPX のスループット** ![[_attachments/arxiv-2402.13499/fig07-dpx-throughput.png]] 非同期コピーは小さいブロックで有効だった。8×8ブロックの `AsyncPipe` は `SyncShare` に対して H800 で39.5%、A100 で19.6%改善した。一方、H800 の32×32ブロックでは1.8%悪化し、大きなブロックのワープ並列性が同期コピーの待ち時間を隠した。 **表XIII: H800 の非同期コピー** ![[_attachments/arxiv-2402.13499/table13-async-copy-h800.png]] **表XIV: A100 の非同期コピー** ![[_attachments/arxiv-2402.13499/table14-async-copy-a100.png]] 分散共有メモリの SM-to-SM レイテンシは180クロックで、L2 キャッシュより32%短かった。RBC のピーク帯域はクラスタサイズ2で約3.27 TB/s、クラスタサイズ4で2.65 TB/sだった。クラスタ内ブロックが増えると帯域競合が強まり、帯域は低下した。 **図8: SM-to-SM 通信のスループット** ![[_attachments/arxiv-2402.13499/fig08-sm-to-sm-throughput.png]] ヒストグラムではブロックサイズ128の最適クラスタサイズが4、ブロックサイズ512では2だった。ビン数が増えると共有メモリ消費でアクティブブロック数が制限されるが、クラスタでビンを分割すると並列性を回復できる場合がある。 **図9: 分散共有メモリを使うヒストグラムの性能** ![[_attachments/arxiv-2402.13499/fig09-dsm-histogram.png]] ## 考察 Hopper の Tensor Core は、旧来の同期 `mma` の性能が必ずしも理論ピークに届かないため、Hopper 専用の非同期 `wgmma` と十分大きな `N` を選ぶことが重要である。 FP8 は行列規模が大きいと有効だが、形式変換・量子化・未対応演算・メモリ転送が残るため、decode-only LLM の短い要求では必ずしも FP16 を上回らない。 非同期実行と分散共有メモリの効果は、機能の有無だけで決まらず、ブロックサイズ、クラスタサイズ、命令レベル並列性、データ初期化、電力上限などの実行条件に依存する。 ## 強み / 弱点・課題 ### 強み - PTX 命令を SASS と実測結果に結び付け、Hopper 固有の `wgmma`・FP8・DPX・分散共有メモリを同じ論文で評価した点。 - 命令レベル、ライブラリレベル、Llama 生成のアプリケーションレベルを横断して、低精度化の効果が計算律速とメモリ律速で変わることを示した点。 - エネルギー効率と電力上限、クラスタサイズやブロックサイズによる性能変化まで測定した点。 ### 弱点・課題 - 実験 GPU は各アーキテクチャ1機種であり、H800 PCIe の結果を Hopper 全体へ一般化できない。 - 本文では完全な実験結果と再現コードが査読後に公開されると述べており、提示された結果だけではワークロードのばらつきや測定反復の詳細を検証しにくい。 - LLM 生成の入力長・出力長が最大128と短く、長文・大規模バッチ・十分に融合された Transformer Engine で FP8 がどこまで有効かは未解決である。 ## 出典 - PDF 原本: [[.raw/papers/arxiv-2402.13499.pdf]] - 抽出テキスト: [[.raw/papers/arxiv-2402.13499.txt]] - arXiv HTML: https://arxiv.org/html/2402.13499 - 著者らが参照する CUDA PTX 文書: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html - Transformer Engine: https://github.com/NVIDIA/TransformerEngine