# NVIDIA Hopper Architecture In-Depth > [!abstract] 概要 > NVIDIA の公式技術ブログによる、Hopper アーキテクチャと H100 Tensor Core GPU の機能解説である。 > 第4世代 Tensor Core、FP8、Transformer Engine、DPX、thread block cluster、分散共有メモリ、Tensor Memory Accelerator、HBM3、NVLink Network、MIG、PCIe Gen 5を、AI・HPCワークロード向けの性能向上という観点から説明する。 > 性能値の多くは2022年時点の出荷前推定値であり、実測値や後続製品の仕様と同一視してはならない。 ## 記事情報 - 媒体: NVIDIA Technical Blog - 著者: Michael Andersch - 公開日: 2022-03-22 - 対象: NVIDIA Hopper、GH100、H100 SXM5、H100 PCIe ## H100 の製品構成 H100 は A100 の後継にあたる第9世代データセンター GPU であり、AI・HPC・データ分析を主な対象とする。 フル実装の GH100 は8 GPC、72 TPC、144 SM、576 Tensor Core、6つの HBM スタック、60MB L2 キャッシュを持つ。 製品版 H100 SXM5 は132 SM、528 Tensor Core、80GB HBM3、50MB L2 キャッシュを持ち、PCIe 版は114 SM、456 Tensor Core、80GB HBM2e、50MB L2 キャッシュを持つ。 ![[_attachments/nvidia-hopper-architecture-in-depth/fig01-full-gh100.png]] （GH100 の構造。GPC、TPC、SM、L2 キャッシュ、HBM メモリコントローラ、NVLink の配置を示す。） 記事は H100 SXM5 の性能を A100 と比較し、FP8 Tensor Core で2,000 TFLOPS、構造化スパース性を使う場合で4,000 TFLOPS、FP16 Tensor Core で1,000 TFLOPS（スパース時2,000 TFLOPS）とする。 これらは記事自身が出荷前の暫定推定値と明記しているため、製品仕様や実運用ベンチマークとは区別する必要がある。 ## Tensor Core、FP8、Transformer Engine 第4世代 Tensor Core は、A100 と同じデータ型では SM あたりの MMA 性能を2倍にし、FP8 では前世代の16ビット浮動小数点形式に対して4倍のレートを実現すると説明される。 FP8 には、指数部4ビット・仮数部3ビットの E4M3 と、指数部5ビット・仮数部2ビットの E5M2 がある。 E4M3 は精度を、E5M2 は広いダイナミックレンジを優先する。 FP8 は FP16/BF16 と比べてデータ格納量を半分にし、理論上のスループットを2倍にする。 ![[_attachments/nvidia-hopper-architecture-in-depth/fig02-fp8-precision.jpg]] （FP8 の E4M3/E5M2 と、FP32 または FP16 への累積を含む Tensor Core のデータ形式。） [[Transformer Engine]] は、各 Transformer 層の出力統計と次層の情報を使い、FP8 と16ビット形式のどちらを使うか、どのスケーリング係数で変換するかを動的に決めるソフトウェアと専用ハードウェアの組み合わせである。 値を FP8 の表現可能な範囲へスケーリングし、必要に応じて再キャストすることで、精度を維持しながら低精度計算を利用する。 記事の比較条件では、A100 に対して AI 学習最大9倍、LLM 推論最大30倍の高速化を主張する。 ![[_attachments/nvidia-hopper-architecture-in-depth/fig05-transformer-engine.png]] （Transformer Engine が層ごとの統計、精度形式、スケーリング係数を使って計算形式を適応的に変換する流れ。） この主張は大規模 Transformer と NVIDIA の比較設定を前提とする。 後続の H800 マイクロベンチマークでは、形式変換や未対応演算、短い decode のメモリ律速によって FP8 の利点が小さくなる条件も報告されている（[[@2024__arXiv__Benchmarking and Dissecting the Nvidia Hopper GPU Architecture]]）。 ## DPX [[DPX]] は、動的計画法の内側ループで頻出する融合オペランドを高速化する命令群である。 記事は Smith-Waterman によるゲノム処理と、Floyd-Warshall による動的なロボット経路探索を例に挙げ、A100 に対して最大7倍の高速化を主張する。 この値はアルゴリズム全体ではなく、記事が想定する DP 処理の性能比較として読むべきである。 ## CUDA 実行階層と非同期性 Hopper は CUDA の実行階層に thread block cluster を追加する。 cluster は同時にスケジュールされることが保証された thread block の集合であり、GPC 内の複数 SM にまたがって協調実行する。 これにより、従来の thread block 単位より大きな局所性をプログラムから指定できる。 ![[_attachments/nvidia-hopper-architecture-in-depth/fig03-tensor-memory-accelerator.jpg]] （Tensor Memory Accelerator がテンソルの次元・ブロック座標・ストライドを記述子から扱う概念図。） [[Tensor Memory Accelerator]]（TMA）は、コピー記述子で指定した1〜5次元のテンソル領域を、少数の CUDA スレッドから非同期に global memory と shared memory の間で転送する。 各要素のアドレス生成、ストライド、境界計算をスレッド群からハードウェアへ移すことで、転送と計算の重ね合わせを容易にする。 非同期トランザクションバリアは、スレッドの到着数だけでなく転送バイト数も数え、データ生成と消費の同期を支える。 ## 分散共有メモリ thread block cluster 内では、スレッドが別 SM の shared memory を load、store、atomic 操作で直接参照できる。 [[分散共有メモリ]] はデータ交換を global memory 経由から SM-to-SM 通信へ近づけ、記事は global memory 方式に対して約7倍のデータ交換性能を主張する。 この機構は cluster、専用 SM-to-SM ネットワーク、非同期バリアを組み合わせたものであり、単なる共有メモリ容量の拡大ではない。 ## HBM3、L2 キャッシュ、MIG H100 SXM5 は80GBの HBM3 を搭載し、3TB/s超のメモリ帯域を提供する。 PCIe 版は80GBの HBM2e と2TB/s超の帯域を持つ。 50MBの L2 キャッシュは A100 の40MBから増加し、モデルやデータセットの再利用部分をキャッシュして HBM へのアクセスを減らす。 ![[_attachments/nvidia-hopper-architecture-in-depth/fig04-hbm3-memory.png]] （世代更新に伴う GPU DRAM 帯域の増加を示す図。H100 で HBM3 と2倍級の帯域を強調する。） 第2世代 [[GPU多重化(MPS・MIG)]] は、GPUを最大7つの隔離されたインスタンスへ分割する。 記事の比較では、A100 に対してインスタンスあたりの計算容量が約3倍、メモリ帯域が約2倍になる。 各インスタンスには専用の NVDEC・NVJPG と性能モニタがあり、MIG レベルの TEE による Confidential Computing も導入される。 ## NVLink、NVSwitch、PCIe H100 は第4世代 NVLink を18本搭載し、GPU I/O と shared memory アクセスに合計900GB/sの帯域を提供する。 第3世代 NVLink の A100 は12本・600GB/sであり、記事は H100 の通信帯域を A100 比1.5倍と説明する。 ![[_attachments/nvidia-hopper-architecture-in-depth/fig06-nvlink-switch-system.jpg]] （A100 の InfiniBand 接続型 SuperPOD と、NVLink Switch System で接続する H100 の256 GPU 構成の比較。） NVLink Network は、複数ノードの最大256 GPUを接続するスケールアップ用の拡張である。 通常の NVLink が共有アドレス空間を使うのに対し、NVLink Network はネットワークアドレス空間とアドレス変換ハードウェアを導入し、GPU 間のアドレス空間を分離する。 第3世代 NVSwitch と NVLink Switch System を組み合わせると、最大32ノード・256 GPUを2:1テーパードの fat tree で接続し、57.6TB/sの全対全帯域と FP8 スパースで1 exaFLOPの計算性能を狙える。 NVSwitch は multicast と [[NVIDIA SHARP]] による網内リダクションを用いて、`all_gather`、`reduce_scatter`、broadcast atomic などの集団通信を支援する。 記事は小さいブロックサイズの集団通信で、A100 上の NCCL に対して最大2倍のスループット向上を主張する。 H100 の PCIe Gen 5 x16 は合計128GB/s（片方向64GB/s）であり、A100 の PCIe Gen 4 の合計64GB/sから倍増する。 PCIe atomic 操作と SR-IOV にも対応し、CPU・GPU間同期や単一 GPU の仮想化を支援する。 ## NVIDIA Grace Hopper Superchip [[NVIDIA GH200]] は Grace CPU と Hopper GPU をチップ間インターコネクトで結合する構成として記事に登場する。 記事は CPU-GPU 間に900GB/sの帯域を持つインターコネクトを使い、PCIe Gen 5の約7倍、テラバイト規模データの処理で最大10倍の性能を目標にすると説明する。 この説明は後年の GH200 製品構成に接続する基礎情報だが、記事の性能値は発表時点の NVIDIA の推定値である。 ## 横断的知見 - **Hopper の性能向上は単一の演算器ではなく、精度・データ移動・局所性・通信の連鎖で構成される**: FP8 と Transformer Engine は演算量とデータ量を削減し、TMA と非同期バリアはデータ移動を計算へ重ね、cluster と分散共有メモリは SM 間の局所性を拡張する。後続の H800 測定でも、`wgmma`、TMA、分散共有メモリの選択が行列形状やクラスタサイズに依存することが示される。(Source: [[@2024__arXiv__Benchmarking and Dissecting the Nvidia Hopper GPU Architecture]]) - **理論ピークと実効性能の間には、ワークロードと実装条件による差がある**: 公式記事は FP8・Transformer Engine・NVLink の最大値を示す一方、後続研究は形式変換、短い decode のメモリ律速、小行列の共有メモリ待ち、電力上限などが実効性能を制約することを測定した。(Source: [[@2024__arXiv__Benchmarking and Dissecting the Nvidia Hopper GPU Architecture]]) - **GPU アーキテクチャのスケールアップはアドレス空間と通信プロトコルの設計問題でもある**: NVLink Network は帯域を増やすだけでなく、ネットワークアドレス空間と明示的なエンドポイント接続を導入する。後続の NVLink 世代でも、トポロジ、集団通信アルゴリズム、キャッシュ・チャネル挙動を分離して分析する必要がある。(Source: [[@2020__TPDS__Evaluating Modern GPU Interconnect - PCIe, NVLink, NV-SLI, NVSwitch and GPUDirect]], [[@2026__SIGCOMM__FabricPerf - Measuring NIC-less Scale-Up Network through GPU Communication Kernel Profiling]]) ## 未解決の問い - H100 の FP8 と Transformer Engine の最大性能値は、長文コンテキスト・小バッチ・decode 主体の LLM サービングでどこまで再現するか。 - TMA、thread block cluster、分散共有メモリを組み合わせたカーネルの性能モデルを、行列形状・クラスタ形状・メモリ再利用から自動生成できるか。 - NVLink Switch System の大規模集団通信で、NVSwitch の網内リダクションと NCCL の通信アルゴリズムをどう協調させるべきか。 - MIG の隔離性、TEE、性能モニタをマルチテナント推論の SLO と障害診断へどう接続するか。 ## 関連 - エンティティ: [[NVIDIA H100]] / [[NVIDIA Hopper]] / [[NVIDIA]] / [[NVIDIA GH200]] / [[Transformer Engine]] / [[Tensor Memory Accelerator]] - concept: [[テンソルコア]] / [[混合精度訓練]] / [[GPU最適化]] / [[DPX]] / [[分散共有メモリ]] / [[Thread Block Cluster]] / [[NVLink]] / [[PCIe性能]] / [[集合通信]] / [[メモリ階層とキャッシュ]] - 比較 source: [[@2024__arXiv__Benchmarking and Dissecting the Nvidia Hopper GPU Architecture]] ## 出典 - 原文: [[.raw/articles/nvidia-hopper-architecture-in-depth-2026-08-29.md]] - NVIDIA Technical Blog: https://developer.nvidia.com/blog/nvidia-hopper-architecture-in-depth/ - NVIDIA H100 Tensor Core GPU Architecture whitepaper: https://resources.nvidia.com/en-us-tensor-core