🚀
CPU/GPU SIMD・Warp/Wave 実行サイクル早見表（2025 年版）
TL;DR
- NVIDIA Maxwell／Pascal／GA102以降 … 1 warp (=32 threads) を 1 cycle で処理
- NVIDIA GA100／Turing／Volta … 1 warp を 2 cycle（16 threads ×2）
- AMD RDNA 3 からは 1CUあたりFP32 × 128 に倍増
- 表を見れば，各世代のレジスタ／共有メモリ量から大まかな Occupancy も掴める
はじめに
「このGPU、1SMで何本のwarpを同時に捌けるんだっけ？」
SIMD／GPGPU の性能チューニングでは、演算器の並列度と命令の完了サイクルを把握しておくと、レイテンシ隠蔽やリソース見積もりが楽になります。
いちいち資料を漁らなくて済むように、GPU(CPU)の主要アーキテクチャを1枚の表にまとめ直しました（2025/07 時点）。もともと2018年くらいに自分のlivedoorブログで公開していた内容をzennにもってきました。
GPU ― Warp/Wave 完了クロック & リソース
| アーキテクチャ | FP32 Core / SM・CU | Max Warp/ Wave | Warp/Wave 完了クロック† | レジスタ容量 / SM・CU | 共有メモリ(LDS) / SM・CU | ホワイトペーパー |
|---|---|---|---|---|---|---|
| NVIDIA | ||||||
| Tesla (G80) | 8 | 24 | 4 | ? | ? | ? |
| Fermi (GF110) | 32 | 48 | 2 | 128 KB | 64 KB(48 KB) | whitepaper |
| Kepler (GK110) | 192 | 64 | 2 | 256 KB | 64 KB(48 KB) | whitepaper |
| Maxwell (GM204) GTX 980 |
128 | 64 | 1 | 256 KB | 64–96 KB | whitepaper |
| Pascal (GP100) Tesla P100 |
64 | 64 | 1 | 256 KB | 64 KB | whitepaper |
| Pascal (GP104) GTX 1080 |
128 | 64 | 1 | 256 KB | 64 KB | whitepaper |
| Volta (GV100) Tesla V100 |
64 | 64 | 2 | 256 KB | 128 KB(96 KB) | whitepaper |
| Turing (TU102) RTX2080Ti |
64 | 32 | 2 | 256 KB | 96 KB(64 KB) | whitepaper |
| Ampere (GA100) NVIDIA A100 |
64 | 64 | 2 | 256 KB | 192 KB(164 KB) | whitepaper |
| Ampere (GA102) RTX 3090 |
128 | 48 | 1 | 256 KB | 128 KB(100 KB) | whitepaper |
| Ada Lovelace (AD102) RTX 4090 |
128 | 48 | 1 | 256 KB | 128 KB(100 KB) | whitepaper |
| Hopper (H100) NVIDIA H100 |
128 | 64 | 1 | 256 KB | 256 KB(228 KB) | whitepaper |
| Blackwell (GB100) NVIDIA B100 |
128 | 64 | 1 | 256 KB | 256 KB(228 KB) | なし |
| Blackwell (GB202) RTX 5090 |
128 | 48 | 1 | 256 KB | 128 KB | whitepaper |
| AMD | ||||||
| GCN (Tahiti) HD 7970 |
64 | 40 | 4 (Wavefront) | 256 KB | 64 KB | whitepaper |
| RDNA 1 (Navi 10) RX 5700 XT |
64 | 40 | 1 (Wave32) | 256 KB | 64 KB | whitepaper |
| RDNA 2 (Navi 21) RX 6900 XT |
64 | 32 | 1 (Wave32) | 256 KB | 64 KB | なし |
| RDNA 3 (Navi 31) RX 7900 XT |
128 | 32 | 1 (Wave32) | 384 KB | 64 KB | whitepaper |
| RDNA 4 (Navi 48) RX 9070 XT |
128 | 32 | 1 (Wave32) | 384 KB | 64 KB | 記事 |
† Warp/Wave 完了クロック
- NVIDIA: 1 warp = どの世代でも32 threads
- AMD: Wavefront = 64 threads
- AMD: Wave32 = 32 threads
- 「Max Warp/Wave」は 100 % Occupancy 時に 1 SM/CU が同時保持できる上限
- 共有メモリはon-chipの量を記載。()内はconfigureできる最大量。RDNAではWGP全体128KBを半割
- 表にはないがInt32:FP32の比がTuring1:1→Ampere1:2→Ada1:2→Blackwell1:1となっている
- AMD: GCN は 16-lane SIMD が 4 基あり、64 thread の wavefront を1つの SIMD が 4 クロックで実行するため、1 wavefront 全体では 4 cycles を要する
- AMD GCNでは100% ALUに必要なthread数が2 CU で 2464 = 512 threads
- AMD RDNA1では WGP(2CU) で 4*32 = 128 threads ということが読み取れる
- 低消費電力版Navi 33はレジスタ容量 256 KB / CU
CPU ― 1 コア 1 clock あたりの FMA/ADD スループット
こっちは更新するきが起きないのでlivedoorのやつからコピペ
| マイクロアーキテクチャ | 1 コア 1 clock あたりの積和スループット |
|---|---|
| Intel Core (Core 2 Duo 等) | 128 ADD + 128 MUL |
| Nehalem | 128 ADD + 128 MUL |
| Sandy Bridge | 256 ADD + 256 MUL |
| Ivy Bridge | 256 ADD + 256 MUL |
| Haswell | 256 FMA × 2 |
| Broadwell | 256 FMA × 2 |
| Skylake (Client) | 256 FMA × 2 |
| Kaby Lake | 256 FMA × 2 |
| Coffee Lake | 256 FMA × 2 |
| Skylake‑X (HEDT) | 512 FMA × 2 |
| Knights Landing (Xeon Phi) | 512 FMA × 2 |
| Zen 1 | 128 ADD × 2 + 128 MUL × 2 [1] |
| Zen 2 | 256 ADD × 1 + 256 FMA × 2 |
memo
GPU チューニングでは「命令レイテンシ隠蔽 メモリアクセスレイテンシ隠蔽 のために warp/wave を大量に並列発行する」のが基本です。しかしOccupancyは必ずしも100%を目指す必要はなく、これが低くともcp.async等のプリフェッチやソフトウェアパイプライニングなどでレイテンシ隠蔽をすることができます。スレッドグループ数を決めたあと本表の 「レジスタ / 共有メモリ量」 をみるとOccupancyが推定しやすくなります。
参考文献
- 4Gamer.net, “RDNA3 世代の CU は SIMD32 が 4 基に倍増” (2022‑12‑12)
- NVIDIA Developer Blog, “Volta/Turing Warp Scheduler and Execution Model”, 2019
- Agner Fog, Instruction tables & microarchitecture (v. 2025‑07)
- [Dream RDNA 3 アーキテクチャのさらなる詳細] (https://www.coelacanth-dream.com/posts/2022/11/16/amd-rdna_3-some-details/)
-
Zen 1 では MUL 演算器が FMA も実行可能。その場合 ADD パイプが停止。 ↩︎
Discussion