6
3

Delete article

Deleted articles cannot be recovered.

Draft of this article would be also deleted.

Are you sure you want to delete this article?

Fermiで時が止まっていた人のためのNVIDIA GPUアーキテクチャ進化史(2010→2026)

6
Last updated at Posted at 2026-07-17

はじめに

自分はG80からGT200、そしてFermiのあたりまでは、NVIDIA GPUの内部アーキテクチャをかなり熱心に追いかけていました。ユニファイドシェーダ、CUDAの登場、Fermiでの本格的なキャッシュ階層。ホワイトペーパーを読み込んでブロック図を眺めるのが楽しい時期でした。

ところがKepler以降、GPUの中身をきちんと追わなくなってしまいました。気がつけばGPUは「グラフィックスの主役」から「AIの主役」に変わっていて、FP4だのNVL72だのTMAだの、目新しい単語だらけになっています。

同じような方、いらっしゃるのではないでしょうか?

本記事は、Fermi(2010)あたりで知識が止まっているエンジニアを読者に想定して、Kepler以降の約15年の進化を整理する試みです。切り口として、何が追加されたかだけでなく、何が捨てられたかにも注目します。アーキテクチャの進化史は、追加の歴史であると同時に切り捨ての歴史でもあり、捨てたものにこそ設計思想の転換が表れると考えるためです。

なお、以下の点をあらかじめお断りしておきます。

  • 数値は原則としてNVIDIA公式ホワイトペーパーおよび公開されている解析論文に基づきますが、一部に実測報告値・報道値を含みます。
  • Rubin世代(2026)以降の情報は、公式発表とリーク・報道が混在しており確度にばらつきがあります。本文中でなるべく区別します。
  • 筆者はGPUベンダーの中の人ではないので、解釈には誤りが含まれるかもしれません。お気づきの点はコメントで指摘いただけると助かります。

長い記事になりましたので、目次から興味のある世代へ飛ぶ読み方も想定しています。

1. 全体マップ:15年を3つの転換点で読む

先に結論めいた見取り図を示します。筆者の理解では、Fermi以降の進化は次の3つの転換点で読むと整理しやすいです。

  1. Kepler(2012):賢いハードウェアをやめた転換点。ホットクロックを廃止し、命令の依存チェックをハードウェアからコンパイラへ移しました。
  2. Volta(2017):GPUがAIエンジンになった転換点。Tensor Coreの導入と、SIMT実行モデルそのものの拡張(独立スレッドスケジューリング)が行われました。
  3. Blackwell(2024):1枚のダイをやめた転換点。レティクル限界に到達し、2ダイ+超広帯域ダイ間接続という構成に移行しました。

もうひとつ、系統の分岐も押さえておくと以降が読みやすくなります。Pascal世代(2016)から、データセンタ向けの大型チップ(GP100、GV100、GA100、GH100、B200という系譜)と、グラフィックス向けチップ(GP104、TU102、GA102、AD102、GB202という系譜)は、同じアーキテクチャ名を冠しながら中身がかなり異なるものになっていきます。本記事では両方を追いますが、AI/HPC文脈では前者が主役です。

2. 出発点の確認:Fermi(2010)を覚えていますか?

まずFermiの姿を簡単に振り返ります。

  • SM構成:1つのSMに32個のCUDAコア、16個のLD/STユニット、4個のSFU。warpスケジューラは2個で、2つのwarpから同時に命令を発行できました。
  • レジスタファイル:32,768エントリ×32ビット/SM。
  • メモリ階層:64KBのオンチップメモリをL1キャッシュとshared memoryに48/16または16/48KBで分割して使う可変構成。その下に768KBの統合L2。GPUとして初めて本格的な読み書き可能キャッシュ階層を持った世代でした。
  • ECC:レジスタファイル、L1、L2、DRAMまでECCで保護(Tesla製品)。HPC市場への本気度の表れです。
  • ホットクロック:G80以来の伝統で、シェーダ(CUDAコア)はコアクロックの2倍で動作していました。GTX 480ではコア700MHz、シェーダ1401MHzです。
  • スケジューリング:マルチポートのレジスタスコアボードと依存チェッカによるハードウェアの動的な依存解決。どのwarpのどの命令が発行可能かをハードウェアが実行時に判定していました。
  • FP64:アーキテクチャとしてはFP32の1/2レートを実装(Tesla C2050等)。GeForce製品では1/8レートに制限されていました。

そしてFermiといえば、その難産ぶりも記憶に残っているのではないでしょうか。GF100は予定から半年ほど遅れ、フル512コアのうち1SMを無効化した480コア構成でGTX 480として出荷されました。TDP 250Wという当時としては異例の消費電力、高温、そして騒音。海外レビューでは、単体GPUカードとして測定史上もっとも騒がしかったという趣旨の評価が残っています。TSMC 40nmプロセスの立ち上がり問題(ビア不良など)と巨大ダイ(529mm²、約30億トランジスタ)の組み合わせは、歩留まりの面でも相当苦しかったと当時報じられていました。

ただし歴史は皮肉なもので、このFermi世代のGTX 580を2枚使ってAlexNetが学習された(2012年)ことが、後の深層学習ブームの実質的な起点になったとNVIDIA自身が振り返っています。汎用計算に振り切ったFermiの設計は、商業的には苦しみましたが、種は蒔かれていたわけです。

さて、ここからが本題です。Fermiで完成したかに見えた「賢く汎用的なGPU」は、次の世代で早くも大きく舵を切ります。

3. Kepler(2012):最初の大転換 —— 賢いハードウェアをやめる

Kepler(GTX 680のGK104、後に大型のGK110)は、Fermiの延長線を期待して見ると面食らう世代です。表面的にはSMあたりのCUDAコア数が32個から192個へ6倍になったのが目立ちますが、本質は2つの廃止にあります。

3.1 ホットクロックの廃止

G80から5年続いたシェーダ2倍クロックが、Keplerで廃止されました。GTX 680は単一クロックドメインの約1GHzで動作します。

GK110のホワイトペーパーには、この判断の理由がかなり率直に書かれています。趣旨をまとめると次のようになります。高いシェーダクロックは、少ないユニット数で目標スループットを達成するための面積最適化だった。しかし高クロック動作のためのクロッキングロジック自体が電力を消費する。Keplerでは電力効率を優先し、面積コストを払ってでも多数のユニットを低いクロックで動かす方を選んだ、と。

面積を節約するために周波数で稼ぐか、電力を節約するために並列度で稼ぐか。Fermiまでは前者、Kepler以降は後者です。トランジスタは増やせるが電力は増やせない、というダークシリコン時代の制約に対する最初の明確な回答だったと言えそうです。同クロックで比較すると、Keplerの96コアはFermi(GF104系)の48コアに相当する、という整理も海外の解析記事でなされています。コア数6倍という数字の額面どおりに演算能力が6倍になったわけではない点は、当時をご存じない方には注意が必要です。

3.2 ハードウェア依存チェックの廃止とcontrol words

もうひとつの廃止は、検証・設計者的にはこちらの方が面白いかもしれません。Fermiが持っていたマルチポートのレジスタスコアボードと依存チェッカが、Keplerで取り除かれました。

なぜ取り除けたのか。鍵は、演算パイプラインのレイテンシが決定論的(固定)であることです。FMAが常に固定サイクルで結果を返すなら、依存関係の解決に必要な情報はコンパイル時にすべて分かっています。そこでKepler以降は、コンパイラ(ptxas)が命令列の中に制御情報 —— いわゆるcontrol words —— を埋め込み、ハードウェア側は「このwarpはあと何サイクル待て」という指示に従って発行可否をマスクするだけの単純な回路になりました。GK110ホワイトペーパーでも、複雑で電力コストの高い複数のブロックを、事前に抽出されたレイテンシ情報を使う単純なハードウェアブロックに置き換えた、という趣旨の説明がされています。

ただし全部が静的になったわけではありません。メモリアクセスのようにレイテンシが可変の命令には引き続きスコアボードが使われますし、どのwarpを選ぶかというwarp間のスケジューリングは当然ハードウェアの仕事のままです。静的化されたのは、固定レイテンシ命令の依存解決という、もっとも高頻度でもっとも電力を食う部分です。

この制御語の埋め込み方は世代ごとに変わっていて、マイクロベンチマークによる解析論文(いわゆるDissecting系の論文)によれば、Keplerでは7命令ごとに1つの制御語、Maxwell/Pascalでは3命令ごとに1つ、Volta以降は128ビットの命令語1つに命令と制御情報が統合された形式になっているとされています。ISAのマニュアルには載らないこうした情報がマイクロベンチマークで暴かれていく過程は、読み物としても面白いのでおすすめです(参考文献参照)。

世代 制御情報の形式
Kepler 7命令ごとに64ビットの制御語を1個挿入
Maxwell / Pascal 3命令ごとに64ビットの制御語を1個挿入
Volta以降 128ビット命令語に命令と制御情報を統合(命令ごとに保持)

検証エンジニアの視点で一言添えると、これは複雑性をシリコンからツールチェーンへ移した取引です。CPUの世界でVLIWが夢見て果たせなかったことを、GPUは部分的に実現したとも言えます。成立した理由は明快で、GPUはレイテンシ隠蔽を投機やアウトオブオーダーではなく大量のwarp(スレッドレベル並列性)で賄えるからです。またこの取引の帰結として、SASS(実マシン命令)は世代間の互換性を持たず、PTXという中間表現が互換性を吸収する層として重要になっていきます。

3.3 SMXという巨大なSM、その功罪

Keplerでは、SMは SMX と呼ばれる192コア構成になりました。warpスケジューラは4個、それぞれが2命令を同時発行(dual-issue)できるため、理論上は毎サイクル8命令の発行が可能です。

ただ、この構成には無理もありました。192コアを4つのスケジューラで埋め切るには、warp内の命令レベル並列性(ILP)をdual-issueで引き出し続ける必要があり、コンパイラと実コードの相性に性能が左右されやすい構造だったのです。SMXの稼働率の低さは次のMaxwellで手当てされることになります。

商業的には、Keplerは大成功でした。GK104(294mm²)は、より大きなダイのRadeon HD 7970と互角以上に戦い、Fermi時代の「大きい・熱い・遅れる」という評判を一気に返上しています。

3.4 GK110:Hyper-QとDynamic Parallelism

コンピュート向け大型チップGK110(71億トランジスタ、551mm²、15SMX)では、FP64ユニットがSMXあたり64個搭載され、FP32の1/3レートを実現しました(GK104は1/24)。K20/K20X、そして初代GTX TITANに使われた世代です。

GK110で導入された機能のうち、後世に残った思想が2つあります。

  • Hyper-Q:ホストからGPUへの投入キューを32本のハードウェアキューに拡張し、複数プロセス(MPIランクなど)からの仕事を並行して受け付けられるようにした機構。GPUを複数の主体で共有するという考え方は、後のMPS、さらにはAmpereのMIGへとつながっていきます。
  • Dynamic Parallelism:GPU側のカーネルから子カーネルを起動できる機能。適応的な処理の再帰的分割などを想定したものでした。機能自体はその後の全世代に広がりましたが、当初のAPI設計は後年見直され、CUDA 12で旧APIは廃止・新設計に置き換えられています。

このほか、動作クロックを動的に引き上げるGPU Boostが入ったのもKepler世代です。

3.5 Keplerが捨てたもの・まとめ

  • ホットクロック(電力のため)
  • ハードウェアの依存チェック(電力と面積のため、コンパイラへ移管)
  • 結果として、Fermi的な「ハードウェアが実行時に賢く振る舞う」路線そのもの

以降のNVIDIA GPUは、この「単純なハードウェアを大量に並べ、賢さはコンパイラとスケジューラ階層に置く」という骨格の上に積み上がっていきます。

4. Maxwell(2014):分割統治と電力効率の極致

Maxwellは地味に見えて、マイクロアーキテクチャ的にはもっとも教科書的に面白い世代のひとつです。プロセスは28nmのままKeplerから据え置きなのに、電力あたり性能を約2倍にした、とNVIDIAは主張しました。プロセスの助けなしにマイクロアーキテクチャだけで達成した数字です。

4.1 SMM:4分割パーティション構造

MaxwellのSM(SMMと呼ばれます)は128コア構成ですが、単なる縮小版ではありません。128コアを32コア×4のパーティション(processing block)に分割し、各パーティションが専用のwarpスケジューラ、ディスパッチユニット、レジスタファイルを持つ構造になりました。

Keplerの192コアSMXでは、どのスケジューラの命令もSMX内の共有実行リソースへ届ける必要があり、大きなクロスバーと複雑な調停が必要でした。Maxwellではスケジューラと実行ユニットの対応を固定し、パーティション内で配線を閉じることで、クロスバーを削減し、電力と配線遅延を抑えています。NVIDIAはGM107のホワイトペーパーで、SMMは128コアでSMX(192コア)の約90%の性能を出す、という趣旨の説明をしています。コア数を減らして稼働率と効率を上げる、という割り切りです。

image.png

この「32コアのパーティション×4」という骨格は、細部を変えながら現在のBlackwellまで続いており、Maxwellが確立した構造がその後10年以上の標準になったと言ってよさそうです。

4.2 モバイル起源の設計

Maxwellの第1弾GM107は、ハイエンドではなく低消費電力帯(GTX 750 Ti、60W)から登場しました。ホワイトペーパー自体が電力制約環境向けの設計であることを掲げており、モバイル・小型PCを起点に設計してデスクトップへ展開するという、それまでと逆向きの製品展開でした。同世代のGPUコアはTegra X1にも採用され、これが後の組み込み・車載系への布石にもなっています。

4.3 L1/shared分離:振り子の始まり

メモリ階層では、FermiとKeplerで統合されていたL1キャッシュとshared memoryが分離されました。L1はテクスチャキャッシュと統合され、shared memoryはSMMあたり64KB(GM204では96KB)の専用領域になります。分割比を悩む必要がなくなり、shared memoryの実効容量が増えてoccupancyを確保しやすくなった、というのが狙いと説明されています。

またL2キャッシュが大幅に増量されました(GM107で2MB)。128ビットという狭いメモリバスをオンチップキャッシュで補う設計で、後のAda Lovelaceの巨大L2を先取りする発想がすでに見えます。

なお、このL1/shared構成は後のVoltaで再び統合へ戻ります。統合→分離→再統合という振り子については6.4節で改めて触れます。

4.4 FP64の切り捨てとHPCの空白期

Maxwellにおける最大の切り捨てはFP64です。レートはFP32の1/32まで削減され、しかもKeplerのGK110に相当する大型FP64チップがMaxwell世代には存在しません。フラッグシップのGM200(TITAN X)ですら1/32のままでした。

このためHPC向けには、Kepler GK210を2ダイ載せたTesla K80が次のPascalまでの橋渡しを務めることになります。グラフィックスと(当時まだ小さかった)単精度系ワークロードに面積を全振りし、倍精度は世代ごと飛ばすという判断は、今から振り返ると製品系統分岐の予兆でした。

余談ですが、Maxwell世代ではラスタライズにタイルベースの処理が密かに導入されていたことが、2016年に外部の実測(David Kanter氏のテスト)で明らかになった、という経緯があります。ホワイトペーパーに書かれない変更が実測で暴かれた例として、マイクロアーキテクチャ好きには趣深い話です。

5. Pascal(2016):プロセス2段跳びと系統分岐の確定

Pascalは28nmから16nm FinFETへ、実質2世代分のプロセスジャンプを果たした世代です。そしてこの世代から、同じアーキテクチャ名の下でデータセンタ向けとグラフィックス向けが明確に別物になります。

5.1 GP100とGP104:同名別物

  • GP100(Tesla P100):610mm²、153億トランジスタ。SMは64 FP32コア+32 FP64コアという構成で、パーティションは2分割。FP64はFP32の1/2レートに復活しました。SMあたりのコア数を半分にした分、コアあたりのレジスタファイル容量は倍増しており、occupancyを確保しやすい構成です。
  • GP104(GTX 1080):Maxwell系の128コアSMを継承。FP64は1/32のままです。

両者を並べると、同名別物ぶりがよく分かります。

項目 GP100 GP104
SMあたりFP32コア 64 128
SMあたりFP64コア 32(1/2レート) 4(1/32レート)
SM内パーティション 2 4
レジスタファイル/SM 256KB(コアあたり2倍) 256KB
FP16 FP32の2倍速 1/64に制限
INT8内積(DP4A) なし あり
メモリ HBM2(732GB/s) GDDR5X(320GB/s)
NVLink 4リンク 160GB/s なし
主用途 学習・HPC グラフィックス・推論

以後、x00番台(GV100、GA100、GH100…)とグラフィックス系(TU102、GA102、AD102…)は、SM構成・メモリ・搭載ユニットのすべてが分岐していきます。

5.2 HBM2とNVLinkの登場

P100は、HBM2を搭載した世界初のGPUです(ホワイトペーパーに明記されています)。4スタック、4096ビット幅で16GB、732GB/s。TSMCのCoWoSによる2.5D実装で、ロジックダイとDRAMスタックをシリコンインターポーザ上に同居させる構成は、現在のAIアクセラレータの標準形の原点です。

同時に、GPU間専用インターコネクトNVLink 1.0が導入されました。4リンクで双方向合計160GB/s。PCIe(当時Gen3)の帯域がマルチGPU学習のボトルネックになるという課題への回答で、この独自インターコネクト路線が後のNVSwitch、さらにはラック全体を1つのGPUにするNVL72へと育っていきます。

5.3 FP16の導入と、コンシューマでのFP16切り捨て

GP100では、FP16を2要素ペア(half2)で処理してFP32の2倍のスループットを出す仕組みが入りました。深層学習の学習用途を意識した最初の精度拡張です。

一方でGP102/GP104(コンシューマ系)のFP16レートは、なんと1/64に制限されました。代わりに搭載されたのがDP4A/DP2A命令、すなわちINT8の4要素内積をフルレートで実行する命令です。学習はFP16でデータセンタGPU、推論はINT8で安価なGPU、という役割分担をハードウェアの仕様で誘導する構図が、この時点で出来上がっています。コンシューマFP16の制限は、FP64に続く精度面の切り捨ての第2弾と言えます。

5.4 Unified Memoryの実質化

CUDA 6以来のUnified Memoryは、Pascalでページマイグレーションエンジンとページフォルト機構を得て実用段階に入りました。49ビットの仮想アドレス空間でCPU/GPUのメモリを覆い、アクセス時にページ単位でオンデマンド移動、物理容量を超えるoversubscriptionも可能になっています。命令粒度のプリエンプションが入ったのもこの世代です。

6. Volta(2017):GPUがAIエンジンになった日

GV100は815mm²、211億トランジスタ。TSMC 12nm FFNで、レティクル限界に迫る巨大ダイ路線の始まりでもあります。しかしVoltaの本質はダイサイズではなく、2つのアーキテクチャ上の大事件にあります。

6.1 Tensor Core:行列演算の固定機能化

Voltaは各SMに8個、チップ全体で640個のTensor Coreを搭載しました。1個のTensor Coreは4x4x4の行列FMA(FP16入力、FP32累算)を実行し、1クロックあたり64回のFMAに相当します。V100全体でFP16行列演算約125 TFLOPSという、当時としては桁外れの数字でした。

image.png

アーキテクチャ的に重要なのは、GEMM(行列積)をSIMTのレーン演算から引き剥がし、専用の固定機能データパスに落としたという判断です。G80以来、GPUは固定機能を減らして汎用化する方向に進化してきました。Voltaはその流れを逆転させ、深層学習という特定ワークロードのために専用回路を積む方向へ舵を切りました。汎用化から再専用化へ。ダークシリコン時代には、電力効率のよい専用回路にトランジスタを使うのが合理的だという判断であり、以後のTensor Core増強の歴史はこの延長線上にあります。

6.2 独立スレッドスケジューリング:SIMTモデル自体の変更

もうひとつの大事件は、実行モデルの変更です。G80からPascalまで、warp内の32スレッドはプログラムカウンタ(PC)とコールスタックを1組だけ共有していました。分岐が発散(diverge)すると、片方のパスをマスクして順に実行し、再収束点まで他方は完全に止まる。これが古典的SIMTです。

Voltaでは、スレッドごとに独立したPCとコールスタックを持たせる独立スレッドスケジューリングが導入されました。実行自体は同一PCのスレッド群をまとめてSIMT実行する点は変わりませんが、発散した2つのパスを交互に進める(interleaveする)ことが可能になり、warp内の細粒度な同期 —— たとえばwarp内スレッド間でのロックのような、従来はデッドロックし得たパターン —— がstarvation-freeに書けるようになりました。

これはプログラミングモデル上の互換性問題も生みました。従来の暗黙のwarp同期(warp-synchronousプログラミング)に依存したコードは、Volta以降では正しく動く保証がなくなり、CUDA 9で__syncwarp()や*_sync系のプリミティブが導入されています。ハードウェアの都合が言語仕様に波及した例として、当時それなりの騒ぎになったと記憶しています。

G80以来の「1 warp = 1 PC」という前提の廃止。これも大きな切り捨てのひとつに数えてよいでしょう。

6.3 INT32データパスの分離とレイテンシ短縮

VoltaのSMは64 FP32コア+64 INT32コア+32 FP64コア+8 Tensor Coreという構成で、FP32とINT32のデータパスが分離され、同時実行できるようになりました。ループ内のアドレス計算(整数)と本体の浮動小数点演算をオーバーラップできるため、実効スループットが上がります。地味ですが実コードには効く改良です。

また、FMAの依存命令発行レイテンシがPascalの6サイクルから4サイクルに短縮されました。レイテンシ隠蔽に必要なwarp数が減るということで、occupancyの制約が緩む方向の改善です。

6.4 L1/shared再統合:振り子が戻る

Maxwellで分離されたL1とshared memoryは、Voltaで128KBの統合構成に戻りました。shared memoryとして最大96KBまで切り出せる可変carveout方式です。L1にshared memory並みの帯域・レイテンシ特性を持たせられるようになったため、統合して柔軟に使う方が得になった、と説明されています。

統合(柔軟性)→分離(単純さと容量)→再統合(回路技術の進歩で柔軟性のコストが下がった)。同じ仕様が行き来しているようでいて、その都度合理的な理由がある。仕様の振り子として、個人的に好きな事例です。

このほかVolta世代では、NVLinkが2.0(双方向300GB/s)へ、HBM2は900GB/sへ強化され、DGX-2ではGPU間を全対全で結ぶスイッチチップNVSwitchが初登場しています。

7. Turing(2018):コンシューマへの展開とRT Core

TuringはVoltaの設計をベースにグラフィックス向けへ再構成した世代です。海外レビューでは、TuringはVoltaの改造版であり、そのVoltaはPascalの大改造版である、という系譜整理がされていました。コンシューマ製品(GeForce RTX 20シリーズ)に初めてTensor Coreが載った世代でもあります。

7.1 RT Core:レイトレーシングの固定機能化

Turing最大の目玉はRT Coreです。レイトレーシングの計算量の大半を占めるBVH(Bounding Volume Hierarchy)のトラバーサルと、レイと三角形の交差判定を固定機能化しました。SMあたり1個搭載され、一次レイの処理でPascal比約10倍とされています。

Tensor Coreに続き、ここでも汎用シェーダから特定処理を引き剥がして専用回路化する流れが続いています。同時に、従来のTFLOPSでは性能を表現できなくなったため、NVIDIAはRTX-OPSという合成指標を持ち出しました(この指標自体はあまり定着しませんでしたが)。

7.2 Tensor Core第2世代:INT8/INT4

TuringのTensor CoreはFP16に加えINT8とINT4に対応しました。推論の量子化を意識した拡張です。ゲーム側ではこれを使うDLSS 1.0が始まりましたが、本格的に化けるのはもう少し先の話です。

なおこのINT4、後の世代でひっそり消える運命にあります(10.4節)。

7.3 uniform datapath:ホワイトペーパーに載らない追加

マイクロベンチマークによる解析論文(Dissecting Turing)で明らかにされた興味深い変更として、uniformレジスタファイルとuniform命令の追加があります。warp内で全スレッドが同じ値を持つ計算(ループカウンタやベースアドレスなど)を、32レーンで重複計算する代わりにスカラ的なデータパスで1回だけ処理する仕組みで、発想としてはAMD GCNのスカラユニットに近いものです。公式資料にはほぼ登場せず、SASSレベルの解析で存在が確認されたという点も含めて、味わい深い変更です。

7.4 SM構成とメモリ

TuringのSMは64 FP32+64 INT32(Volta譲りの並行実行)+8 Tensor Core+1 RT Coreで、L1/sharedは96KB統合。FP64はSMあたり2ユニットだけの1/32レートで、互換性のためだけの搭載と割り切られています。メモリはGDDR6(14Gbps)へ移行しました。

またRTX 2080以上にはNVLink(2リンク、双方向100GB/s)が搭載されました。用途は主にSLIでしたが、そのSLI自体がこの頃には下火で、コンシューマNVLinkの命脈は次の世代で早くも尽きることになります。

8. Ampere(2020):TF32という発明と、GPUを切る技術

8.1 二系統・二ファウンダリ

Ampereでは系統分岐がさらに進み、ついに製造委託先まで分かれました。

  • GA100(A100):TSMC 7nm、826mm²、542億トランジスタ。データセンタ専用。
  • GA102(RTX 3090等):Samsung 8nm(8N)、628mm²、283億トランジスタ。

GA100にはRTコアが搭載されず、表示出力もありません。データセンタ品からグラフィックスを切り捨てる流れは、この世代で本格化しました(次のHopperで極まります)。

8.2 TF32:互換性のための新フォーマット

A100の第3世代Tensor Coreは対応精度を大きく広げました。中でもTF32は発明と呼びたくなる仕様です。

TF32は符号1+指数8+仮数10の19ビットフォーマットで、指数部はFP32と同じ(=表現レンジが同じ)、仮数部はFP16と同じ(=精度はFP16並み)。既存のFP32コードを一切書き換えずに、入力を内部でTF32へ丸めてTensor Coreに流し、FP32で累算する。これにより「コード変更なしでV100のFP32比10倍前後」という売り文句を成立させました。数値フォーマットをソフトウェア互換性のための道具として設計した例で、ビット割りの妙を味わえる仕様です(ビットレイアウトは付録参照)。

このほか第3世代Tensor CoreはBF16、そしてFP64にも対応しました。FP64の行列演算をTensor Coreで実行することで、HPC向けにも19.5 TFLOPS(A100)を確保しています。さらに2:4の構造化スパース性(4要素中2要素が非ゼロという制約付き疎行列)のハードウェアサポートが入り、条件を満たせば実効スループット2倍、というカタログ表記が始まったのもこの世代です。以後のスペック表では、dense値とsparse値のどちらを読んでいるのか常に注意が必要になります。

8.3 MIG:1つのGPUを7つに切る

A100ではMIG(Multi-Instance GPU)が導入されました。1つのGPUを最大7個のインスタンスに分割し、それぞれにSM、L2スライス、メモリ帯域を専有割り当てして、性能干渉なく複数テナントに貸し出せる機構です。Kepler Hyper-Qに始まった「GPUを複数の主体で共有する」路線の到達点であり、クラウド時代のGPUに必須の機能になりました。

またA100では、グローバルメモリからshared memoryへレジスタファイルを経由せず直接コピーする非同期コピー命令(cp.async、SASSではLDGSTS)が入りました。データ搬送を計算スレッドから切り離していくこの方向性は、次のHopperのTMAで完成形に近づきます。L2は40MBに増量され、residency制御や圧縮も入りました。

8.4 GA102と、最後のNVLink付きGeForce

コンシューマ側のGA102は、SMあたりのFP32を128に倍増しました。ただし内訳は64の専用FP32+64のFP32/INT32兼用データパスで、整数命令が混じるとFP32スループットは落ちます。CUDAコア数2倍というマーケティングの数字と実効性能の乖離が話題になった世代です。

そしてRTX 3090は、NVLinkコネクタを備えた最後のGeForceになりました。NVIDIAは2021年初頭以降の新規SLIプロファイル提供終了を公式に告知しており、マルチGPUレンダリングという文化はここで事実上幕を下ろしています。

9. Ada Lovelace(2022):巨大L2という判断

9.1 TSMC 4Nへの回帰と密度ジャンプ

AdaのAD102はTSMC 4N(5nm系カスタム)で製造され、608mm²に763億トランジスタ。Samsung 8nmのGA102から一気に密度が約2.7倍になりました。フルチップで144SM、18432 CUDAコアという規模です。

9.2 L2 96MB:バス幅ではなくSRAMで解く

Adaの設計判断として面白いのがL2キャッシュです。GA102の6MBから96MB(フルチップ)へ、実に16倍。メモリバス幅を広げる代わりに巨大なオンダイSRAMで帯域とレイテンシの問題を解く構成で、レイトレーシングのような不規則アクセスに特に効きます。AMDがRDNA 2で先行したInfinity Cacheと同じ方向の解であり、Maxwellの「狭バス+大L2」の拡大再生産とも読めます。

9.3 SER、第3世代RT Core、DLSS 3

  • SER(Shader Execution Reordering):レイトレーシングでは、レイごとにヒットするマテリアルが異なるためシェーダ実行が発散しがちです。SERはシェーダの起動をハードウェア支援で動的に並べ替え、似た処理をまとめて実行することで発散を減らします。シェーダ処理で最大2〜3倍、実ゲームでも数十%の改善とされています。
  • 第3世代RT Core:Opacity Micromap、Displaced Micro-Meshといった、ジオメトリ表現をRTハードウェア側に寄せる機構が追加されました。
  • Optical Flow Acceleratorの強化とDLSS 3:フレーム間の動きベクトルをハードウェアで求め、AIでフレームそのものを生成する(フレーム生成)機能が始まりました。レンダリングしていないフレームを表示するという、性能指標の定義自体を揺さぶる機能です。
  • 第4世代Tensor Core:HopperからFP8対応が持ち込まれています。

9.4 NVLink削除:I/OをAIに明け渡す

RTX 4090ではNVLinkが削除されました。発表直後の記者会見でJensen Huang氏は、NVLinkを外したのはそのI/Oを他の用途に使う必要があったからで、その面積にできる限りのAI処理を詰め込んだ、という趣旨の説明をしています。ダイエッジのSerDes面積とパッケージピンは有限で、何かを載せるには何かを外すしかない。切り捨ての理由がここまで率直に語られるのは珍しく、記録しておきたい発言です。

10. Hopper(2022):データ搬送の再設計

H100(GH100:814mm²、800億トランジスタ、TSMC 4N)は、Voltaに次ぐ大改造世代です。SM数はフル144(H100 SXMでは132有効)、SMあたりFP32は128、L1/sharedは256KBまで拡大しました。しかしHopperの本質は演算器の数ではなく、データをどう運ぶかの再設計にあります。

10.1 TMA:アドレス計算をスレッドから剥がす

TMA(Tensor Memory Accelerator)は、グローバルメモリとshared memoryの間で1次元〜5次元のテンソルブロックを非同期転送する専用エンジンです。従来はコピーのためのアドレス計算やループをスレッド群が実行していましたが、TMAではディスクリプタ(コピー記述子)を1つ渡せば、アドレス生成・境界処理・転送をハードウェアが済ませ、その間スレッドは計算に専念できます。Ampereのcp.asyncをさらに推し進め、DMAエンジンをSMの中に持ち込んだ形です。

10.2 thread block clusterとDSMEM:SMに近傍構造を導入

CUDAの実行階層(thread → block → grid)に、新しくclusterという階層が追加されました。同一クラスタ内のthread blockは同じGPC(GPU Processing Cluster)内のSM群に同時にスケジュールされることが保証され、SM同士が専用ネットワークで互いのshared memoryを直接read/write/atomicアクセスできます。これが分散shared memory(DSMEM)です。

外部の実測論文では、SM間のshared memoryアクセスはL2経由より3割ほど低レイテンシ(180サイクル前後)と報告されています。それまでフラットな多数として扱われてきたSMアレイに、物理的な近さを利用する階層構造が持ち込まれたわけで、NoC・フロアプラン設計の観点からも興味深い変化です。プログラミングモデルがダイ上の物理配置を意識し始めた、と言い換えてもよいかもしれません。

10.3 FP8とTransformer Engine

第4世代Tensor CoreはFP8に対応しました。E4M3(符号1+指数4+仮数3)とE5M2(符号1+指数5+仮数2)の2形式で、順伝播の重み・活性化にはレンジより精度のE4M3、勾配にはレンジ重視のE5M2、という使い分けが想定されています。

ただし8ビットで学習を回すのは容易ではなく、そこを支えるのがTransformer Engineです。層ごとに統計を取りながらFP8と16ビットを動的に選択し、スケーリング係数を自動調整するハードウェア+ソフトウェアの複合機構で、精度を保ったままFP8の速度を取り出すことを狙っています。数値フォーマットの選択がコンパイル時ですらなく実行時の動的制御になった、という点で象徴的な機能です。

10.4 追加されたもの、消えたもの

追加としてはほかに、動的計画法(Smith-Waterman、Floyd-Warshallなど)向けのDPX命令(最大7倍)、GPUとして初のConfidential Computing対応(NVLink上のトラフィック暗号化を含む)、非同期トランザクションバリアなどがあります。インターコネクトはNVLink 4.0(双方向900GB/s)、メモリはHBM3で3.35TB/s(SXM)、ホスト接続はPCIe Gen5になりました。Grace CPUとコヒーレントに接続するNVLink-C2C(900GB/s)を使ったGrace Hopperスーパーチップもこの世代です。

一方で消えたものも明確です。GH100にはRTコアがなく、表示出力もなく、グラフィックス動作が可能なTPCはチップ内にわずか2つだけ残されています(互換性のためと思われます)。グラフィックス向けにはL40系という別SKUを立て、フラッグシップからはグラフィックスをほぼ完全に切り捨てました。G80以来「GPUはグラフィックスのついでに計算もできる」だったものが、Hopperに至って「計算専用機が申し訳程度にグラフィックスも残している」まで反転したことになります。

もうひとつ、地味な切り捨てとしてINT4があります。Turing/AmpereのTensor Coreが対応していたINT4は、Hopperで対応から外れました。低ビット化の本命は整数ではなく浮動小数点(FP8、そして次のFP4)である、という方向付けがここでなされたと見ることができます。

11. Blackwell(2024–2025):レティクルの壁を越える

11.1 モノリシックの終わり:デュアルダイとNV-HBI

Hopperの814mm²という数字が示すとおり、NVIDIAのフラッグシップはすでに露光装置のレティクル限界(1ショットで露光できる約800mm²強)に張り付いていました。世代ごとの規模拡大をダイサイズで稼ぐ道は、物理的に尽きたわけです。

Blackwell(B200)の答えは、レティクル限界サイズのダイを2枚並べ、NV-HBI(NVIDIA High-Bandwidth Interface)と呼ばれる双方向10TB/sのダイ間接続で結合し、ソフトウェアからは単一のGPUに見せるというものでした。合計2080億トランジスタ(1ダイ約1040億)、プロセスはTSMC 4NP。パッケージングにはCoWoS-L(RDLインターポーザ+局所シリコンブリッジ)が使われています。NVIDIAの担当者は、2つのダイは1つの巨大なGPUとして振る舞う、キャッシュコヒーレンシも透過だ、という趣旨の説明をしています。

image.png

AMDがMI200/MI300で先行したチップレット化に、NVIDIAは「ソフトウェアから見て完全に1個」という形で追随した、と整理できます。1枚のダイという30年来の前提の切り捨てです。

11.2 第5世代Tensor Core:FP4とNVFP4

演算精度はさらに下がり、第5世代Tensor CoreはFP6(E2M3/E3M2)とFP4(E2M1:符号1+指数2+仮数1)に対応しました。FP4は非ゼロの絶対値が6通りしかない、もはや表引きに近い世界です。

これを実用にするのがブロックスケーリングです。OCP標準のMXFP4は32要素ごとにE8M0(2の冪のみの8ビット指数)の共有スケールを持ちますが、NVIDIA独自のNVFP4は16要素ごとにFP8(E4M3)のスケール、さらにテンソル全体にFP32のスケールという2段構成を取ります。ブロックが小さく、スケールが2の冪に限定されないぶん量子化誤差を抑えやすい設計で、実効ビット数は1要素あたり約4.5ビットになります。第2世代Transformer Engineがこれらのマイクロスケーリング形式を前提に再設計されました。

命令の発行形態も世代ごとに別物になっています。Volta〜AmpereのMMA命令はwarp(32スレッド)同期で発行され、オペランドとアキュムレータはレジスタファイルに置かれていました。Hopperではwarpgroup(4 warp=128スレッド)単位の非同期命令wgmmaが導入され、オペランドをshared memoryから直接読めるようになります。そしてBlackwellのtcgen05系命令では、1スレッドが発行して完了を非同期に待つ形のMMAになり、オペランド・アキュムレータの置き場も、TMEM(Tensor Memory)と呼ばれるSMあたり256KBの新設専用メモリへ移ったと外部のマイクロベンチマーク解析で報告されています。レジスタファイル中心からTMEM中心へのデータパスの作り替えで、Tensor Core周りはもはやSMの中の別プロセッサの様相です。カーネルの書き方も、TMAと非同期バリア、warp specializationを組み合わせたproducer/consumer型のパイプライン構成が標準になっており、SIMTコア側のブロック図がHopper以降ほとんど変わらない一方で、この行列演算パスは毎世代作り直されている、というのがデータセンタGPUの実像です。

11.3 NVLink 5とNVL72:ラックが1つのGPUになる

NVLink 5.0はGPUあたり双方向1.8TB/s(リンク数18のまま、リンクあたり速度を224Gbps SerDesで倍増)。第4世代NVSwitchと組み合わせたGB200 NVL72では、72基のBlackwell GPUと36基のGrace CPUを1本のNVLinkドメインに収め、集約帯域130TB/s、液冷ラック1本がソフトウェアから単一の巨大GPUとして見える構成になっています。メモリはHBM3eで192GB・8TB/s(B200)。

さらにBlackwell Ultra(B300/GB300)では、12段スタックのHBM3eで288GB、電力は1チップ1400Wに達し、ホスト接続はPCIe Gen6を初採用しています。H100の700Wから2世代で倍、空冷はもはや選択肢にありません。

11.4 歩留まり問題:CTE不整合とマスク変更

Blackwellの立ち上がりは平坦ではありませんでした。2024年、CoWoS-Lパッケージにおける歩留まり問題が報じられます。アナリストの分析では、GPUダイ・シリコンブリッジ・RDLインターポーザ・基板の熱膨張係数(CTE)の不整合による反りが原因とされました。NVIDIAはGPUのマスク変更(トップメタル層とバンプの再設計、機能ロジックは不変)で対応し、Jensen Huang氏は2024年10月のイベントで、Blackwellには設計上の欠陥があった、機能はしたが歩留まりを下げた、100% NVIDIAの責任だ、と率直に認めています(TSMCとの不和報道はfake newsだ、とも)。当初計画のB100が事実上取り下げられ、構成変更版に置き換えられたという報道もありました。

Fermiの40nm立ち上げ苦戦から15年、今度はパッケージングが歩留まりの主戦場になったという対比は、先端実装の時代を象徴していると感じます。

11.5 RTX Blackwell(RTX 50):neural shadingへ

コンシューマ側のGB202(RTX 5090)は750mm²・922億トランジスタのモノリシックで、TSMC 4N。フル192SM(24576コア)のうちRTX 5090では170SMが有効です。メモリはPAM3信号方式のGDDR7になり、512ビット幅・28Gbps相当で約1.8TB/sに達します。L2はフル128MB(5090では96MB)。TDPは575Wです。

アーキテクチャ面では、SMがneural shading(シェーダ内から小さなニューラルネットを呼ぶCooperative Vectors系の処理)向けに最適化され、AI処理とグラフィックスのスケジューリングを司るAMP(AI Management Processor)と呼ばれる専用プロセッサがフロントエンドに追加された、とされています。DLSS 4ではモデルがTransformerベースになり、1フレームから最大3フレームを生成するマルチフレーム生成と、表示エンジン側でペーシングを制御するFlip Meteringが導入されました。RT Coreは第4世代で、三角形クラスタ単位の交差エンジン(Mega Geometry対応)に置き換わっています。

レンダリングとニューラル推論の境界は、この世代で本格的に溶け始めたように見えます。

11.6 FP64の停滞という静かな切り捨て

なお、B200のFP64(Tensor Core)性能は40〜45 TFLOPS程度とされ、H100の67 TFLOPSから見て後退しています。AI向け低精度に演算資源を寄せた結果がFP64にまで及んだ形で、伝統的HPCコミュニティでは議論になっているようです。Fermiが1/2レートのFP64でHPCに殴り込みをかけてから15年、優先順位は完全に入れ替わりました。

12. Rubin(2026):HBM4世代 —— 執筆時点の情報

本記事執筆時点(2026年7月)での次世代、Rubin/Vera Rubinについても触れておきます。ここから先は公式発表(GTC 2025、CES 2026など)とリーク・報道が混在するため、数値の確度は一段落ちるものとしてお読みください。

  • Rubin GPU:コンピュートダイ2枚+I/Oダイ2枚のチップレット構成を4倍レティクルサイズのCoWoS-Lに実装、合計3360億トランジスタとされています。プロセスはTSMC N3系。I/O機能を別ダイに分離したのは、Blackwellのデュアルダイからさらに一歩進めた分割です。
  • メモリ:HBM4を288GB搭載し、帯域は22TB/s級と報じられています。HBM3e比でざっくり2倍超です。
  • 演算:NVFP4での推論性能50 PFLOPS(dense換算では35 PFLOPSとの情報)。第3世代Transformer Engineは、層ごとに動的に圧縮率・精度を調整する適応的な2段マイクロブロックスケーリングを持つとされています。
  • インターコネクト:NVLink 6で双方向3.6TB/s。ラック構成はVera Rubin NVL144で、この144はGPUパッケージ数ではなくコンピュートダイ数(72パッケージ×2ダイ)を数える新しい流儀です。数え方が変わった点はスペック比較時の罠になりそうです。
  • CPU:相方のVeraは、Neoverse系だったGraceと異なりOlympusと呼ばれる自社カスタムコア88基(176スレッド)構成とされています。
  • 派生品Rubin CPX:2025年9月に発表された、HBMの代わりにGDDR7を128GB搭載するprefill特化のGPUです。推論には、コンテキストを読み込むprefillフェーズが演算律速、トークンを生成するdecodeフェーズがメモリ帯域律速という性質の違いがあり、前者を安価なメモリ構成の専用チップへ分業させる(disaggregated serving)流れを、ソフトウェアだけでなくチップのラインナップ側で支える製品と言えます。NVFP4で30 PFLOPS、attentionの指数関数演算はGB300比3倍とされ、Rubin本体と混載するVera Rubin NVL144 CPXラック(全体でNVFP4 8 EFLOPS級)とともに2026年末の提供が予告されています。広帯域が過剰な場面ではHBMを捨てる、という割り切りは、本記事で追ってきた切り捨ての系譜の最新の一例かもしれません。
    なお、CPXの状況については以下の記事を参照ください。

  • スケジュール:2026年6月に量産開始、2026年後半からパートナー提供と報じられています。さらに先のRubin Ultra(2027、コンピュートダイ4枚、HBM4eで1TB、NVL576ラックは600kW級)はロードマップ上の計画であり、確度は低めです。

エコシステム面では、NVLinkをサードパーティのCPUやアクセラレータに開放するNVLink Fusionが2025年に発表されており、独自インターコネクトが業界標準の座を狙う段階に入ったようにも見えます。

13. 捨てられた技術・総まとめ

ここまでに登場した切り捨てを一覧にします。

技術・機構 導入 廃止・転換 理由(要旨) 行き先・代替
ホットクロック(シェーダ2倍クロック) G80(2006) Kepler(2012) クロッキングロジックの電力 多数コア×低クロック
ハードウェア依存チェック(固定レイテンシ命令) 〜Fermi Kepler(2012) 電力・面積。レイテンシが決定論的 コンパイラのcontrol words
L1/shared統合(第1期) Fermi Maxwellで分離 shared実効容量とoccupancy Voltaで再統合(振り子)
コンシューマの高レートFP64 Fermi(GeForceは1/8) Kepler以降縮小、Maxwellで1/32、現在1/64 ダイ面積をグラフィックス/AIへ データセンタ系チップに集約
コンシューマの高レートFP16 GP100(2倍速) GP102/104で1/64に制限 製品セグメント分離 INT8 DP4A、後にTensor Core
warp単一PC(古典SIMT) G80 Volta(2017) 細粒度同期、divergence柔軟化 per-thread PC+明示同期
INT4(Tensor Core) Turing(2018) Hopper(2022)で対応外 低ビット化の本命はFP系へ FP8、FP4/NVFP4
SLI(マルチGPUレンダリング) 2004(SLI名復活) 2021新規プロファイル終了 効果と開発コストが見合わず 単一巨大GPU+AIアップスケール
コンシューマNVLink Turing Ada(RTX 4090)で削除 I/O面積をAI処理に転用(公式発言) PCIeのみ
データセンタGPUのRTコア・表示出力 〜GV100相当 GA100/GH100で削除・最小化 AI/HPC特化 L40系などグラフィックスSKUへ分離
Dynamic Parallelism旧API GK110(2012) CUDA 12で旧API廃止 設計刷新 新Dynamic Parallelism API
Hyper-Q(単体機能として) GK110(2012) 発展的解消 より強い分離が必要に MPS、そしてMIGへ
モノリシック巨大ダイ 〜GH100(814mm²) Blackwell(2024) レティクル限界 デュアルダイ+NV-HBI、Rubinで4チップレット

こうして並べると、切り捨ての理由は驚くほど一貫しています。電力(ホットクロック、依存チェック)、面積の機会費用(FP64、NVLinkコネクタ)、そして物理限界(レティクル)。技術が劣っていたから捨てられたものはほとんどなく、より価値の高い用途にリソースを譲った結果である、というのが筆者の読み後感です。

14. 世代横断スペック表

主要ダイの数値をまとめます。トランジスタ数・面積の一部は実測報告値を含みます。

世代 代表ダイ プロセス トランジスタ 面積 SM数(フル) FP32/SM FP64比 Tensor Core NVLink(双方向/GPU)
Fermi GF100 2010 TSMC 40nm 約30億 529mm² 16 32 1/2(Tesla) なし なし
Kepler GK110 2012 TSMC 28nm 71億 551mm² 15 192 1/3 なし なし
Maxwell GM200 2015 TSMC 28nm 80億 601mm² 24 128 1/32 なし なし
Pascal GP100 2016 TSMC 16nm FF 153億 610mm² 60 64 1/2 なし 160GB/s
Volta GV100 2017 TSMC 12nm FFN 211億 815mm² 84 64 1/2 第1世代 300GB/s
Turing TU102 2018 TSMC 12nm FFN 186億 754mm² 72 64 1/32 第2世代 100GB/s(GeForce)
Ampere GA100 2020 TSMC 7nm 542億 826mm² 128 64 1/2 第3世代 600GB/s
Ampere GA102 2020 Samsung 8N 283億 628mm² 84 128 1/64 第3世代 112.5GB/s(3090のみ)
Ada AD102 2022 TSMC 4N 763億 608mm² 144 128 1/64 第4世代 なし
Hopper GH100 2022 TSMC 4N 800億 814mm² 144 128 1/2 第4世代 900GB/s
Blackwell B200 2024 TSMC 4NP 2080億(2ダイ) レティクル級×2 非公表(2ダイ) 128 縮小傾向 第5世代 1800GB/s
Blackwell GB202 2025 TSMC 4N 922億 750mm² 192 128 1/64 第5世代 なし
Rubin(参考) Rubin 2026 TSMC N3系 3360億(4チップレット、報道込み) 4倍レティクル実装 224(報道値) 非公表 非公表 第6世代相当 3600GB/s

SMの中身がどう変わってきたかも一覧にします。INT32の列は、専用データパスを持つか(Volta以降)、FP32とレーンを共用するかを示します。

世代(代表ダイ) FP32/SM INT32/SM FP64/SM Tensor/SM RT/SM L1+shared/SM
Fermi(GF100) 32 FP32と共用 実効1/2レート - - 64KB(統合・可変)
Kepler(GK110) 192 共用 64 - - 64KB(統合・可変)
Maxwell(GM200) 128 共用 4 - - 分離(shared 96KB)
Pascal(GP100) 64 共用 32 - - 分離(shared 64KB)
Volta(GV100) 64 64(分離) 32 8(第1) - 128KB(統合)
Turing(TU102) 64 64(分離) 2 8(第2) 1 96KB
Ampere(GA100) 64 64(分離) 32 4(第3) - 192KB
Ampere(GA102) 128 64(FP32と兼用) 2 4(第3) 1 128KB
Hopper(GH100) 128 64(分離) 64 4(第4) - 256KB
Ada(AD102) 128 64(FP32と兼用) 2 4(第4) 1 128KB
Blackwell(GB202) 128 128(全レーンFP32/INT32兼用) 2 4(第5) 1 128KB

Tensor Coreは世代ごとに1ユニットの規模が拡大しているため、個数の減少(8→4)は能力の減少ではない点に注意してください。SMあたりの行列演算能力はむしろ世代ごとに増えています。

代表世代のSM内部を模式図で並べると次のようになります。

image.png

インターコネクトの世代もまとめます。

NVLink世代 初出 双方向帯域/GPU 備考
1.0 P100 2016 160GB/s 4リンク
2.0 V100 2017 300GB/s 6リンク。DGX-2で初代NVSwitch(2018)
3.0 A100 2020 600GB/s 12リンク
4.0 H100 2022 900GB/s 18リンク。第3世代NVSwitchはSHARP演算内蔵
5.0 B200 2024 1800GB/s 224Gbps SerDes。NVL72で72GPU単一ドメイン
6.0 Rubin 2026 3600GB/s(報道値) NVLink Fusionで外部にも開放へ

電力の推移も並べておきます。Maxwellで一度下がったあと、Volta以降は一方向に増え続けているのが分かります。

データセンタ系 TDP グラフィックス系 TDP
2010 - - GTX 480 250W
2012 K20X 235W GTX 680 195W
2014-15 - - GTX 980 165W
2016 P100(SXM) 300W GTX 1080 180W
2017-18 V100(SXM2) 300W RTX 2080 Ti 250W
2020 A100(SXM) 400W RTX 3090 350W
2022 H100(SXM) 700W RTX 4090 450W
2024-25 B200 / B300 1000W / 1400W RTX 5090 575W

メモリの系譜も一行で:GDDR5 → GDDR5X(Pascal)→ GDDR6(Turing)→ GDDR6X(Ampere/Ada)→ GDDR7(RTX 50、PAM3)。HBM系はHBM2(P100)→ HBM2e(A100)→ HBM3(H100)→ HBM3e(H200/B200/B300)→ HBM4(Rubin)。

15. 数値フォーマットのビットレイアウト

エンジニア的にいちばん血が騒ぐところなので、付録としてまとめます。SはSign、EはExponent、MはMantissaです。

image.png

  • TF32は、FP32の指数(レンジ)とFP16の仮数(精度)を組み合わせた互換性フォーマットです。
  • FP8の2形式は、順伝播(精度重視のE4M3)と逆伝播の勾配(レンジ重視のE5M2)での使い分けが想定されています。
  • FP4(E2M1)単体では実用に耐えないため、ブロックスケーリングと組で使われます。
    • MXFP4(OCP標準):32要素ブロック+E8M0(2の冪のみ)共有スケール。実効約4.25bit/要素。
    • NVFP4(NVIDIA独自):16要素ブロック+FP8(E4M3)スケール、さらにテンソル全体にFP32スケールの2段構成。実効約4.5bit/要素。ブロックが小さくスケールが2の冪に縛られないぶん、量子化誤差の面で有利とされています。

image.png

Tensor Core世代と対応精度のマトリクスも載せておきます。

精度 第1世代(Volta) 第2世代(Turing) 第3世代(Ampere) 第4世代(Hopper/Ada) 第5世代(Blackwell)
FP16 対応 対応 対応 対応 対応
INT8 - 対応 対応 対応 対応
INT4 - 対応 対応 廃止 -
TF32 - - 対応 対応 対応
BF16 - - 対応 対応 対応
FP64(TC) - - 対応(GA100) 対応(GH100) 対応(縮小傾向)
FP8 - - - 対応 対応
FP6/FP4/NVFP4 - - - - 対応

16. おわりに

長々と書いてきましたが、15年をあらためて俯瞰すると、次の3つの流れに集約されるように思います。

1つめは、賢さの引っ越しです。実行時のハードウェア(スコアボード)からコンパイラ(control words)へ、さらにHopper以降はTransformer Engineのような実行時の動的制御へ。賢さは消えたのではなく、もっとも電力効率のよい場所を探して移動し続けています。動的→静的→また動的、という往復も、L1/sharedの振り子と同じで、その時々の制約下での最適解の移動です。

2つめは、汎用と専用の振り子です。G80からFermiへの流れは固定機能を捨てて汎用化する歴史でしたが、Volta以降はTensor Core、RT Core、TMA、DPXと専用回路を積み増す歴史に反転しました。ムーアの法則の減速と電力制約の下では、使うトランジスタを選ぶしかない、ということなのでしょう。

3つめは、スケールの主戦場の移動です。周波数(ホットクロック)→ダイ面積(GV100〜GH100)→パッケージ(Blackwellのデュアルダイ)→ラック(NVL72)。1つの物理限界に達するたびに、一段外側の階層でスケールを稼ぐ。Rubinのコンピュート/I/Oダイ分離やNVLink Fusionを見ると、次の主戦場はさらに外側、データセンタ全体の設計に移りつつあるようです。

この進化がDVの難易度をどう変えてきたかも気になるところです。依存チェックのコンパイラ移管はハードウェア単体の状態空間を減らした一方で、ツールチェーンとの契約(コンパイラが正しい制御語を吐く前提)という検証境界の外側を生みました。クラスタ/DSMEM、デュアルダイのコヒーレンシ、NVLinkドメインとなると、もはやチップ検証というよりシステム検証で、UVMのスコープで扱える範囲を超えていきます。このあたりの話は、いずれ別記事で掘ってみたいテーマです。

Fermiで時が止まっていた筆者と同じ境遇の方に、この15年の見取り図として少しでもお役に立てば幸いです。誤りの指摘、補足、思い出話など、コメントでお待ちしています。

参考文献

一次資料(NVIDIA公式ホワイトペーパー・技術ブログ):

マイクロベンチマークによる解析論文(Dissecting系):

このほか、Maxwell各ホワイトペーパー、Hopper/Blackwellアーキテクチャホワイトペーパー、Hot Chips・GTC発表資料、Real World Tech、Chips and Cheese、AnandTechの各アーキテクチャ解説を参照しました。URLは執筆時点のものです。リンク切れの際は資料名で検索してください。

6
3
0

Register as a new user and use Qiita more conveniently

  1. You get articles that match your needs
  2. You can efficiently read back useful information
  3. You can use dark theme
What you can do with signing up
6
3

Delete article

Deleted articles cannot be recovered.

Draft of this article would be also deleted.

Are you sure you want to delete this article?