> [!abstract] 概要 > 科学研究の急速な発展により、より大きな計算能力が必要となっており、その一部は GPU によって解決されている。 > 本論文は、十分に検討したマイクロベンチマークによって GPU の性能機能を調べ、現代の NVIDIA Blackwell アーキテクチャをマイクロアーキテクチャの観点から分析する。 > メモリ階層、SM 実行パイプライン、FP4 および FP6 精度をサポートする第5世代 Tensor Core を含む SM サブコアなどの主要サブシステムを明らかにする。 > NVIDIA GPU の主要機能を理解するため、レイテンシ、スループット、キャッシュ挙動、スケジューリングの詳細を調べ、Blackwell の設計における微妙なチューニング指標を明らかにする。 > 包括的な分析のため、Blackwell と前世代 Hopper を、それぞれ GeForce RTX 5080 と H100 PCIe を用いて比較する。 > 結果を評価・比較し、世代間の改善と性能退行の両方を示す。 > さらに、異なるワークロードにおける電力効率とエネルギー消費の役割を調べる。 > 本研究の知見は、アプリケーション開発者、コンパイラ作成者、性能エンジニアが Blackwell ベースのプラットフォーム向けにワークロードを最適化するための実用的な指針を提供し、GPU アーキテクチャ研究の蓄積に新たなデータを加える。 ## 論文情報 - タイトル: Dissecting the NVIDIA Blackwell Architecture with Microbenchmarks - 著者: Aaron Jarmusch、Nathan Graddon、Sunita Chandrasekaran - 所属: University of Delaware, Department of Computer Information Sciences - 媒体: arXiv, cs.DC - 初回公開: 2025-07-14。取得した PDF は v2(2025-07-21) - arXiv: [2507.10789](https://arxiv.org/abs/2507.10789) - 比較対象: Blackwell の GB203(GeForce RTX 5080) と Hopper の GH100(H100 PCIe) - コード: ブラインド査読のため現時点では非公開。査読プロセスと結果の後に公開予定と記載されている。 ## 概要 本論文は、PTX と CUDA で記述したマイクロベンチマークを用いて、Blackwell の SM 実行パイプライン、Tensor Core、メモリ階層を Hopper と比較する。 GB203 は FP4/FP6、統合 INT32/FP32 実行ユニット、低いワープ数での高 ILP という特徴を示すが、評価した FP8 D-GEMM では H100 が高い実効スループットを示した。 したがって Blackwell の理論的な低精度能力だけでは実ワークロードの性能を予測できず、命令形式、ワープ数、データ供給、カーネル選択、電力状態を合わせて評価する必要がある。 ## 問題設定 GPU の公開仕様は実行ユニット数や対応精度を示すが、依存命令のレイテンシ、命令の完了スループット、キャッシュ境界、ワープスケジューラの挙動、実ワークロードでの電力特性までは十分に説明しない。 そこで本論文は、Blackwell の GB203 と Hopper の GH100 について、compute-bound と memory-bound の双方を対象に、単一命令から Transformer 推論までを同じ比較の流れで調べる。 測定する中心指標は、依存命令の true latency、独立命令を重ねた completion latency、SM あたりのスループット、メモリレイテンシ、帯域幅、電力である。 ## アーキテクチャ比較 両 GPU は CUDA コアのプログラミングモデルを共有するが、設計目標は異なる。 GH100 は大規模 AI 訓練と科学計算向けで、HBM2e、大きなオンチップ資源、高い実行スループットを重視する。 GB203 はゲーム、レンダリング、小バッチ推論を含む電力制約下のコンシューマー向けで、キャッシュ容量と倍精度能力の一部をクロック周波数・効率と交換する。 | 実行ユニット | GH100 | GB203 | |---|---:|---| | アーキテクチャ | Hopper | Blackwell | | GPU 名 | H100 PCIe | GeForce RTX 5080 | | FP32 / SM | 128 | 統合 INT32/FP32 ユニット | | INT32 / SM | 64 | 統合 INT32/FP32 ユニット | | FP64 / SM | 64 | 2 | | Tensor Core | 第4世代 | 第5世代(FP4・FP6対応) | | Transformer Engine | 第1世代 | 第2世代 | (表I。GH100 と GB203 の実行ユニット) | メモリ単位 | GH100 | GB203 | |---|---|---| | L0 i-cache | 分離パーティション | 統合 | | レジスタファイル / SM | 256 KB | 256 KB | | L1 cache / SM | 256 KB(統合) | 128 KB(統合) | | Shared Memory / SM | 228 KB(統合) | 統合 | | L2 cache | 50 MB(2パーティション) | 65 MB(1パーティション) | | Global Memory | 80 GB(HBM2e) | 16 GB(GDDR7) | (表II。GH100 と GB203 のキャッシュ階層とパーティション) Blackwell ではワープスケジューリングが改善され、発散ワークロードのワープディスパッチレイテンシが短縮されると論じられる。 命令互換性については、Hopper の `wgmma` は Blackwell では利用できず、Blackwell では `tcgen05` がワープグループ計算の新しい選択肢になる。 ただし、本評価対象の GB203 では `tcgen05` がまだ実装されていないため、Tensor Core の比較には両 GPU の `mma` を用いる。 ## マイクロベンチマーク方法 レイテンシは、直列化した依存命令列で測る true latency と、独立命令の重なりを含む completion latency に分ける。 スループットは、実測実行時間と命令数から、SM あたりクロックサイクルごとに完了した命令数として求める。 PTX カーネルを CUDA ファイル中のインライン命令にせず個別ファイルとして実行し、コンパイル時最適化の影響を避ける。 さらに各 PTX カーネルから生成された SASS を確認し、命令が最適化で消えていないことを検証する。 クロックは読み取り専用特殊レジスタ `%clock64` で読み、命令の前後の差を計測する。 図1に相当する PTX の測定コードは、画像ではなくコードブロックとして転記する。 ```ptx .reg .u64 %start, %end; ld.param.u64 %time_ptr, [param_time]; mov.u64 %start, %clock64; mad.lo.s32 r1, r1, r2, r3; mov.u64 %end, %clock64; ``` このコードは `mad.lo.s32` の前後で `%clock64` を読み、差分から命令のクロック数を求める。 最初の実行で GB203 にのみ高いレイテンシが現れた場合は、キャッシュウォームアップの影響を考慮して表から除外している。 ## Compute パイプライン ### INT32・FP32 実行ユニット GB203 の統合ユニットは INT32 と FP32 の命令を動的に扱えるため、混合ワークロードの遊休サイクルを減らす可能性がある。 一方、同一クロックサイクルでは INT32 または FP32 のどちらか一方しか実行できず、混合命令列にはハザードが生じうる。 `fma`(FP32) と `mad`(INT32) を用い、純粋 INT32、純粋 FP32、2 種類の混合命令列を各 1024 回実行した。 | GPU | Pure INT32 | Pure FP32 | Mixed 1 | Mixed 2 | Pure FP64 | |---|---:|---:|---:|---:|---:| | GB203 | 4 / 16.97 | 4 / 7.97 | 15.96 / 14 | 26.28 / 18 | 63.57 / 11 | | GH100 | 4 / 16.69 | 4 / 7.86 | 31.62 / 16 | 43.54 / 20 | 8.04 / 13 | (表III。true latency / completion latency。原文見出しの「GH203」は本文に合わせて GB203 と解釈) 純粋 INT32/FP32 の true latency は両 GPU とも 4 サイクルで、純粋カーネルの completion latency は GH100 がわずかに低い。 対して混合命令列では GB203 の方が低く、統合 INT32/FP32 実行ユニットが混合計算に有利である可能性を示す。 **図2: 依存命令列の総サイクル** ![[_attachments/arxiv-2507.10789/fig02-dependent-total-cycles.png]] (図2。GB203 と GH100 の INT32、FP32、FP64 ワークロードについて、命令数に対する総サイクルを示す。8 命令付近で両アーキテクチャの総サイクルが急に変化する。) **図3: 依存命令列のスループット** ![[_attachments/arxiv-2507.10789/fig03-dependent-throughput.png]] (図3。命令数に対するスループット。短い依存列では ILP 不足で低く、命令数の増加に伴ってスケジューラがレイテンシを隠蔽しスループットが上がる。) ### FP64 実行ユニットとワープスケジューラ 両 GPU の FP64 ユニットは INT32/FP32 ユニットから物理的に分離されている。 GB203 は SM あたり 2 個、GH100 は 64 個であり、1024 個の依存 FP64 命令では GH100 のレイテンシが低い。 GB203 で依存命令を 2 個だけ実行するとレイテンシは 37.5 サイクルまで下がるため、著者らは 2 ユニットが主に命令形式のサポートを担い、計算を FP32 ユニットまたは Tensor Core で代替する設計の可能性を示す。 依存列の測定では、GH100 は短いレイテンシ制約列に強く、GB203 は規則的で高い ILP を持つカーネルに適するという解釈が示される。 ## 第5世代 Tensor Core ### 対応形式と命令 GB203 の第5世代 Tensor Core は FP4、FP6、FP8、INT8、FP16、BF16、TF32、FP64 を扱う。 GH100 の第4世代は FP8、INT8、FP16、BF16、TF32、FP64 を扱うが、FP4 と FP6 には対応しない。 GB203 は `mma`、`wmma`、`tcgen05`、GH100 は `mma`、`wmma`、`wgmma` を対応命令として持つ。ただし評価時点の GB203 では `tcgen05` は未実装である。 | | GB203(第5世代) | GH100(第4世代) | |---|---|---| | 対応データ形式 | FP4, FP6, FP8, INT8, FP16, BF16, TF32, FP64 | FP8, INT8, FP16, BF16, TF32, FP64 | | MMA 命令 | `mma`, `wmma`, `tcgen05` | `mma`, `wmma`, `wgmma` | (表IV。Tensor Core の対応形式と MMA 命令) ### MMA のタイルと SASS 対応 MMA は GEMM のための行列積累算であり、命令ごとに `M×N×K` のタイル形状を指定する。 例として `mma.sync.aligned.m16n8k32.f32.f16.f16.f32` は、16×32 と 32×8 の入力から 16×8 の出力タイルを計算し、FP16 入力を FP32 で累積・出力する。 Blackwell の FP4/FP6 MMA には CUDA 12.9 の `.kind::f8f6f4` 接尾辞が必要で、指定しない場合や GH100 で使った場合は PTX エラーになる。 | 形式 | データ形式 | PTX 命令の主要部 | |---|---|---| | `e2m1` | FP4 | `.m16n8k32.row.col.f32.e2m1.e2m1.f32` | | `e3m2` | FP6 | `.m16n8k32.row.col.f32.e3m2.e3m2.f32` | | `e2m3` | FP6 | `.m16n8k32.row.col.f32.e2m3.e2m3.f32` | | `e4m3` | FP8 | `.m16n8k32.row.col.f32.e4m3.e4m3.f32` | | `e5m2` | FP8 | `.m16n8k32.row.col.f32.e5m2.e5m2.f32` | (表V。評価したデータ形式と PTX 命令) GH100 では各 `mma.sync` が HMMA へ変換された。 Blackwell では FP8 と FP6 の入力が QMMA へ変換され、FP4 入力も意図された OMMA ではなく QMMA になった。 FP8 のブロック尺度に `ue8m0` を使った場合は OMMA が観測され、現在のソフトウェアでは FP4 に QMMA がフォールバックしている可能性が示される。 ### 精度、電力、ワープ数 低精度形式は表現ビット数を減らすことでメモリフットプリントとスループットを改善するが、動的範囲と数値表現力とのトレードオフがある。 | データ形式 | Blackwell | Hopper | |---|---:|---:| | FP4 e2m1 | 16.753 | n/a | | FP6 e2m3 | 39.383 | n/a | | FP6 e3m2 | 46.723 | n/a | | FP8 e4m3 | 46.661 | 55.823 | | FP8 e5m2 | 46.806 | 55.786 | (表VI。データ形式ごとの電力使用量(W) / 性能毎ワット) 表VI の数値では FP4 の電力が 16.753 W と最も低く、FP6 は 39 W 超、FP8 は 46 W 超である。 同じ FP8 形式では GH100 が約55 W、GB203 が約46 Wであり、Blackwell の低精度時の電力効率を示す。 **図4: 精度形式とワープ数に対する Tensor Core スループット** ![[_attachments/arxiv-2507.10789/fig04-tensor-throughput.png]] (図4。GB203 の FP4/FP6/FP8 と GH100 の FP8 のスループットを、ワープ数ごとに比較する。GB203 は評価した低精度形式で一貫して高い。) **図5: 精度形式とワープ数に対する Tensor Core レイテンシ** ![[_attachments/arxiv-2507.10789/fig05-tensor-latency.png]] (図5。GB203 は特に FP4/FP6 で低レイテンシを維持し、GH100 はワープ数の増加に伴って段階的にレイテンシが増加する。) 低精度 MMA では、GH100 の持続スループット条件が ILP=5、29 アクティブワープ、GB203 が ILP=6、25 アクティブワープだった。 GB203 はより少ないワープ数でもスレッド内の独立 `mma` 命令を活用でき、6 ILP・32 ワープで 11 TFLOP/s 超に達した。 ILP=1、ワープ数1での completion latency は GB203 が全形式で 1.21094 サイクル、GH100 が 1.65625 サイクルだった。 この結果から、Blackwell は高 ILP・低精度・規則的な制御フローに、Hopper はより大きな同時実行性と深いバッファリングに適するという設計上の差が論じられる。 ## メモリサブシステム メモリの評価は、ホストとデバイス間転送を除き、レジスタ、共有メモリ、L1/L2 cache、global memory を対象にする。 ランダムな直列ポインタ追跡でレイテンシを分離し、共有メモリと L1 ではアクセスストライドとワープ数、L2 ではワープレベル同時実行、global memory では持続転送帯域を測る。 **図6: メモリ階層のレイテンシ** ![[_attachments/arxiv-2507.10789/fig06-memory-hierarchy-latency.png]] (図6。データサイズの増加に対するレイテンシ。GB203 は約128 KB・約30 MB、GH100 は約256 KB・約60 MB付近で段階的な境界が現れ、L1、L2、global memory の領域に対応する。) L1 cache のヒットレイテンシは両 GPU とも約30〜40サイクルでほぼ同じである。 一方、統合 L1/shared memory の容量は GH100 が最大256 KB/SM、GB203 が128 KB/SMで、動的 shared memory の上限はそれぞれ約227 KB/SMと約99 KB/SM、静的な既定値は双方48 KB/SMである。 **図7: 共有メモリのワープスケーリング** ![[_attachments/arxiv-2507.10789/fig07-shared-memory-latency.png]] (図7。共有メモリのレイテンシをワープ数とストライド1/4で比較する。GB203 は1〜4ワープで低く、6〜32ワープでは GH100 が有利で、GB203 のストライド4では増加が急になる。) **図8: L1 cache のワープスケーリング** ![[_attachments/arxiv-2507.10789/fig08-l1-cache-latency.png]] (図8。L1 cache のレイテンシをワープ数とストライド1/4で比較する。GB203 はストライド1の2〜11ワープでわずかに低いが、ストライド4では増加が急になる。) 共有メモリはワープ数とストライドへの感度が高く、特に GB203 ではバンク競合と小さいメモリパーティションの飽和が影響する可能性がある。 L1 は共有メモリよりワープ数に対して平坦で、GH100 は高いワープ圧力と密な再利用に有利、GB203 は少数ワープの低レイテンシ経路に有利と解釈される。 ### L2 cache GH100 の L2 は GPC の一部を担当する2つの独立パーティション、GB203 は全 GPC が共有する単一構成である。 標準的な L2 ヒットのレイテンシは GB203 が約358サイクル、GH100 が約273サイクルだった。 GH100 は両パーティションが飽和すると31〜45 MBの範囲で約508サイクルまで増加するが、GB203 は65 MBの容量により基準レイテンシをより大きなフットプリントまで維持する。 **図9: L2 cache のワープスケーリング** ![[_attachments/arxiv-2507.10789/fig09-l2-latency-warp-scaling.png]] (図9。低いワープ数では GH100 が有利だが、高いワープ数では GB203 が追いつく。32ワープ時の1ワープあたりサイクルは GB203 約128.4k、GH100 約128.9kである。) 1〜4ワープでは GH100 が約43.5kサイクル、GB203 が約49kサイクルで、8〜16ワープでも GH100 が有利だった。 16〜32ワープでは GB203 が追いつき、32ワープでわずかに上回る。 したがって GH100 は低〜中程度の同時実行性とレイテンシ感度の高いワークロード、GB203 は全負荷時の帯域バウンドな処理に有利という整理になる。 ### Global memory ![[_attachments/arxiv-2507.10789/fig10-memory-bandwidth.png]] (図10。メモリ階層の帯域幅。GH100 は読み取り15.8 TB/s、書き込み2.2 TB/s、GB203 は読み取り8.2 TB/s、書き込み1.6 TB/sのピークを示す。) GH100 のピーク読み取り帯域は15.8 TB/s、GB203 は8.2 TB/s、書き込み帯域はそれぞれ2.2 TB/sと1.6 TB/sだった。 global memory に移るデータサイズは GB203 が約71 MB超、GH100 が約55 MB超で、レイテンシはそれぞれ約876.7サイクルと約658.7サイクルだった。 HBM2e を使う GH100 は、GDDR7 を使う GB203 より帯域とレイテンシで有利である。 ## マイクロベンチマークのケーススタディ ### Dense GEMM FP8 入力の D-GEMM を `cuBLASLt` と `__nv_fp8_e4m3` で実装し、$D=A^T \times B+C$ を計算した。 A と B は FP8、C は bfloat16、D は FP8で、32 MBのワークスペースを使い、各構成を100回実行して平均した。 行列サイズは1024、2048、4096、8192を組み合わせ、TFLOPS は $2\times M\times N\times K / runtime$ で求めた。 **図11: D-GEMM の実行時間** ![[_attachments/arxiv-2507.10789/fig11-gemm-runtime.png]] (図11。各 M×N×K 形状の実行時間。8192×8192×8192 の GB203 の4.710 msはグラフから省略されている。) | 行列サイズ | Hopper(TFLOP/s) | Blackwell(TFLOP/s) | |---|---:|---:| | 8192×8192×8192 | 0.887 | 0.233 | | 2048×2048×2048 | 0.554 | 0.191 | | 2048×2048×4096 | 0.674 | 0.192 | | 2048×4096×8192 | 0.759 | 0.217 | | 1024×1024×1024 | 0.239 | 0.134 | (表VII。GH100 と GB203 の D-GEMM スループット) Hopper は評価した全形状で低い実行時間を示し、8192×8192×8192 では0.887 TFLOP/sで、Blackwell の0.233 TFLOP/sのほぼ4倍だった。 Blackwell は大きな形状でレイテンシのスパイクが現れ、理論上 FP8 の計算率が高くても、現時点のカーネル選択、メモリ階層利用、スケジューリングが実効性能を制限する可能性がある。 **図12: D-GEMM の電力** ![[_attachments/arxiv-2507.10789/fig12-gemm-power.png]] (図12。M=N=K の行列サイズに対する平均電力。Hopper は概ね58〜60 Wで最大68 W、Blackwell は形状依存の変動が大きく最大114.4 Wに達する。) Hopper の電力プロファイルは比較的平坦で、最大形状でも68 Wだった。 Blackwell は平均80 W超の領域と114.4 Wのピークを示し、512×512×512だけは Hopper より大幅に低い電力だった。 低いスループットと高い電力が組み合わさるため、現行ソフトウェアとカーネル実装では D-GEMM の性能毎ワットは Hopper が有利である。 ### Transformer 推論 実ワークロードの例として、TensorRT を使った Transformer 推論を実装した。 小規模で FP8/FP4 量子化経路と互換性がある GPT-NeoX を用い、best、normal、FP16、FP8 の精度設定を各100回実行して平均した。 | 精度設定 | Hopper(W) | Blackwell(W) | |---|---:|---:| | FP32 | 60.24 | 58.82 | | FP16 | 57.64 | 47.78 | | FP8 | 57.69 | 45.14 | | Best | 60.15 | 61.03 | (表VIII。精度設定ごとの平均推論電力) Hopper は精度設定をまたいで57〜60 Wと安定した。 Blackwell は FP32 の58.82 Wから FP8 の45.14 Wへ電力が低下したが、TensorRT の Best 設定では61.03 Wへ増加した。 この結果は、Blackwell が精度を下げた推論で電力を抑えられる一方、設定された最適エンジンが常に低電力になるとは限らないことを示す。 ## 考察 - Blackwell の世代改善は一様ではない。FP4/FP6 の対応、統合 INT32/FP32 ユニット、少数ワープ・高 ILP の低精度 MMA では GB203 が有利だが、FP64、HBM による global memory、短い依存列、評価した大規模 FP8 D-GEMM では GH100 が有利な条件がある。 - GPU の仕様表にある精度対応や理論演算率は、実効性能を単独では決めない。PTX から SASS への変換、MMA のタイル形状、ワープ数、ILP、キャッシュ境界、カーネル選択、電力・温度状態が結果を変える。 - GB203 の単一 L2 は低い同時実行性では GH100 の分割 L2 より高レイテンシだが、極端なワープ圧力では容量と集約帯域が有利に働く。レイテンシと帯域のどちらを重視するかで最適な GPU が変わる。 - 本論文の数値は H100 PCIe と RTX 5080 の特定構成、現時点の CUDA/TensorRT ソフトウェア、特定のカーネル実装に依存する。Blackwell 全体やデータセンター向け GB200/B200 の性能へそのまま一般化できるとは書かれていない。 ## 強み / 弱点・課題 ### 強み - 公開仕様だけでは分からない Blackwell の FP64、FP4/FP6 MMA、SASS 対応、キャッシュ境界、ワープスケジューリングを、命令レベルと実ワークロードの両方から測定している。 - Hopper との比較により、世代改善だけでなく性能退行、理論性能と実効性能の乖離、電力の変動も同じ枠組みで示している。 - PTX の個別ファイル化、SASS 確認、各設定の反復平均により、コンパイル時最適化と測定ノイズを抑えようとしている。 ### 弱点・課題 - コードはブラインド査読中のため公開されておらず、追試に必要なベンチマーク実装と設定を利用できない。 - 比較は GB203/RTX 5080 と GH100/H100 PCIe の一組に限られ、同じアーキテクチャ内の SKU 差、冷却、電力上限、ドライバ・CUDA バージョンの感度は評価されていない。 - Transformer 推論ではレイテンシやスループットではなく主に電力を報告しており、FP4/FP8 の品質への影響も評価していない。 - 一部の説明は「可能性がある」「おそらく」というマイクロアーキテクチャ推論であり、NVIDIA の内部仕様による確定的な説明ではない。 ## 関連 - 概念: [[GPUマイクロベンチマーク]] / [[GPU最適化]] / [[性能測定]] / [[テンソルコア]] / [[メモリ階層とキャッシュ]] / [[Instruction-Level Parallelism (GPU)]] / [[量子化]] / [[LLM推論]] - エンティティ: [[NVIDIA Blackwell]] / [[RTX 5080]] / [[NVIDIA H100]] / [[NVIDIA Hopper]] / [[NVIDIA]] / [[CUDA]] / [[TensorRT]] / [[GPT-NeoX]] - 関連ソース: [[@2024__arXiv__Benchmarking and Dissecting the Nvidia Hopper GPU Architecture]] / [[@2025__arXiv__Efficient Fine-Grained GPU Performance Modeling for Distributed Deep Learning of LLM]] ## 出典 - [[.raw/papers/arxiv-2507.10789.pdf]] - [[.raw/papers/arxiv-2507.10789.txt]] - [arXiv abstract and metadata](https://arxiv.org/abs/2507.10789)