# テンソルコア ## 定義 テンソルコア(Tensor Core)は、NVIDIA GPU に搭載された専用の行列積累算(matrix multiply-accumulate、MMA)ユニット。Volta アーキテクチャ(V100、2017 年)で初めて実用化され、主に深層学習の行列演算を高速化する目的で設計された。 NVIDIA V100 の仕様(Haidar+ SC 2018 より): - V100 PCIe の FP16-TC 理論性能: 112 テラFLOP/秒(理論ピーク)、実用値 85 テラFLOP/秒 - 演算単位: 1 クロックサイクルで $D = A \times B + C$ を 4×4 行列として計算(64 回の FMA 混合精度演算) - **A と B は FP16(入力)**、**C と D は FP16 または FP32(出力・蓄積)** - **乗算は FP16 で実行し、蓄積は FP32 で行う**(FP16 均一演算より数値的に優れる) - SM(ストリーミングマルチプロセッサ)あたり 8 個、V100 全体では 640 個 テンソルコアを使う GEMM(`hgemm-TC`)の特性: - 正方行列 GEMM: FP64 の約 12 倍の速度 - ランク-k 更新(LU 分解で使われる tall-skinny GEMM): FP64 の約 6 倍(データ転送コストのため正方より遅い) - FP16 均一 GEMM よりも数値安定性が高い(蓄積が FP32 のため) ## テンソルコアと HPC の関係 深層学習向けに設計されたテンソルコアをHPC の数値線形代数に適用するには技術的課題がある(Haidar+ SC 2018): - 低精度許容の深層学習と異なり、HPC ソルバーは FP64 精度の解を要求する - FP16 の数値範囲(最大 65,504、単位丸め誤差 5×10^-4)はアンダーフロー・オーバーフローを起こしやすい - LU 分解のパネル分解や trsm は数値的に感度が高く、全 FP16 化は不安定になる これらの課題を回避するため、SC 2018 では「マルチ精度 LU」を提案した。GEMM のみテンソルコアを使い、数値感度の高い計算部分(パネル分解・trsm)は FP32 に留めるという協調設計(co-design)戦略である。 ## 横断的知見 - **Hopper の第4世代 Tensor Core は、従来の精度拡張に加えて FP8 と Transformer Engine によるデータ量削減を性能軸へ加えた**: H100 は FP8 の E4M3/E5M2、FP16/BF16、TF32、FP64、INT8 を扱い、公式記事は FP8 のデータ格納量半減と理論スループット倍増を説明する。後続の実測は、精度形式の選択だけでなく `wgmma` の発行方式、行列形状、形式変換コストが実効性能を決めることを示す。(Source: [[@2022__NVIDIA Developer Blog__NVIDIA Hopper Architecture In-Depth]], [[@2024__arXiv__Benchmarking and Dissecting the Nvidia Hopper GPU Architecture]]) - **蓄積精度が安定性を決める**: FP16-TC と FP16(均一演算)の最大の違いは GEMM の蓄積精度。FP32 蓄積により、クラスタ特異値・複素固有値を持つ行列でも反復精密化が収束する。数値安定性は「入力精度」だけでなく「蓄積精度」に大きく依存する。(Source: [[@2018__SC__Harnessing GPU Tensor Cores for Fast FP16 Arithmetic to Speed up Mixed-Precision Iterative Refinement Solvers]]) - **深層学習で生まれたハードウェアが HPC を変える**: テンソルコアは深層学習のニーズに応じて実装されたが、その性能差は HPC の LU 分解にも直接適用できる。演算密度が高い GEMM ルーティンを多用する密線形代数は、テンソルコアの恩恵を最大限に受けられる応用領域である。アーキテクチャ世代(Volta→Turing→Ampere→Hopper)とともにテンソルコアの精度サポート(FP16→TF32→BF16→FP8)が拡張されており、各世代での HPC 再評価が期待される。(Source: [[@2018__SC__Harnessing GPU Tensor Cores for Fast FP16 Arithmetic to Speed up Mixed-Precision Iterative Refinement Solvers]]) - **アテンション演算でのテンソルコア利用率は世代とともに向上するが 70-75% で頭打ちになる**: [[FlashAttention]] シリーズは Volta/Ampere の FP16 テンソルコア（FA1/FA2、利用率 25-73%）から Hopper の FP8 テンソルコア（FA3、FP16 で 75%・FP8 で約 1.2 PFLOP/秒）を経て Blackwell の BF16 テンソルコア（FA4、71%）へと進化した。テンソルコアスループットが 1 PFLOP/秒（H100）→ 2.25 PFLOP/秒（B200）と倍増する一方で、softmax に必要な指数関数ユニットのスループットは据え置き（16 ops/clock/SM）であり、非 MMA 演算が利用率の天井を決定する構造的制約となっている。(Source: [[@2024__arXiv__FlashAttention-3 - Fast and Accurate Attention with Asynchrony and Low-precision]], [[@2026__arXiv__FlashAttention-4 - Algorithm and Kernel Pipelining Co-Design for Asymmetric Hardware Scaling]]) - **Blackwell のテンソルメモリ（TMEM）は蓄積器のレジスタ圧力問題を解決した**: Hopper では MMA の蓄積器がレジスタに書き込まれ、タイルサイズの拡大を制約していた。Blackwell の TMEM（256 KB/SM）は MMA 出力を直接書き込む専用メモリであり、FA4 はこれを活用して 128×128 タイル（Hopper の 64×128 の 2 倍面積）で動作する。テンソルコアの「出力先」が世代ごとに変わる（レジスタ → TMEM）ことが、アルゴリズム設計を根本的に変える。(Source: [[@2026__arXiv__FlashAttention-4 - Algorithm and Kernel Pipelining Co-Design for Asymmetric Hardware Scaling]]) - **Rubin は「精度拡張」に加え「K 次元スループット倍増」という第 2 の軸でテンソルコアを進化させた**: これまでの世代交代は主に精度サポート拡張(FP16→TF32→BF16→FP8→NVFP4)とテンソルコアの出力先(レジスタ→TMEM)の変化として観察されてきたが、[[NVIDIA Rubin GPU]] は 1 クロックあたり処理できる K 次元データ量を倍にすることで GEMM の K ループ反復回数を半減させる(Blackwell 4 反復→Rubin 2 反復)という、精度とは独立した第 2 のスケーリング軸を導入した。FA4 が指摘した「テンソルコアスループットが伸びても指数関数ユニット(softmax)のスループットが据え置きだと利用率が頭打ちになる」という構造的制約に対し、Rubin は指数関数スループットも同時に引き上げる(FP32 で 2 倍・BF16/FP16 で 4 倍)ことで、テンソルコアと softmax のスループット比を保ったまま両輪で高速化する設計を取る。(Source: [[@2026__NVIDIA Developer Blog__Inside NVIDIA Rubin GPU Architecture Powering the Era of Agentic AI]], [[@2026__arXiv__FlashAttention-4 - Algorithm and Kernel Pipelining Co-Design for Asymmetric Hardware Scaling]]) - **テンソルコアの6世代進化(V100→Rubin)は「発行に必要なスレッド数の縮小」と「発行と実行の脱結合」という2つの独立した軸で整理できる**: [[AI Chip Architectures]]は本ページが精度・出力先(レジスタ→TMEM)・K次元スループットという軸で追ってきた進化に、命令発行(issue)側の変化という第4の軸を加える。Voltaの`mma.sync`は32スレッドのwarpが同期発行しwarp全体が完了までブロックするのに対し、Hopperの`wgmma.mma_async`は128スレッドのwarp-groupが非同期発行し(operand Bはシェアードメモリ記述子経由)、Blackwellの`tcgen05.mma`ではオペランドA/Bが共にシェアードメモリ記述子(またはTMEM直読)、アキュムレータDがTMEMに着地するため単一スレッドが発行するだけでよく即座に他の作業へ戻れる。CTA-pair変種では2つのSMにまたがるペアの一方の1スレッドが協調MMAを発行し256×256×16タイルを構成する。この非同期化の進行が、softmax・マスキング・次タイルのプリロードをmatmulの裏で回すFlashAttention系カーネル(warp specialization、producer/consumer分離をmbarrierで同期)を成立させる直接の技術的前提である。(Source: [[AI Chip Architectures]]) - **AMD Matrix Core(MFMA)はNVIDIA Tensor Coreと対照的に、6世代を通じて発行スコープをwavefrontに固定し続けている**: [[AI Chip Architectures]]は、CDNA 1(2020)からCDNA 4/MI355X(2025)まで全てのMFMA命令が「1つのwave64が単一の行列演算を発行し、4つのSIMDが協調してそれを駆動する」というwavefront スコープを維持していることを指摘する。NVIDIAがTensor Coreの発行スコープをwarp(32)→warp-group(128)→単一スレッドへ縮小してきたのとは正反対の経路であり、AMDが唯一導入した緩和策はCDNA 4の「LDSからのMFMA転置ロード」(TMAに小規模に相当するオペランド供給の効率化)に留まる。この差が、matmulの裏でsoftmax等をオーバーラップさせる能力(Overlap)の構造的な差として現れ、純粋なdense GEMM(訓練の内側ループ)では両者の差が小さい一方、Transformer attentionのような不規則なワークロードではAMDがソフトウェアで手作業のパイプライン構築を強いられる。(Source: [[AI Chip Architectures]]) - **Hopper の Tensor Core は命令形式の移行自体が性能条件になる**: `mma` は1 warpの同期命令で H800 では理論ピークの平均62.9%だったのに対し、4 warp の warp group が非同期実行する `wgmma` はゼロ初期化で理論ピークの95%超に達した。Hopper では精度形式の追加だけでなく、発行単位・オペランド供給・非同期性を含むプログラミングモデルの選択が実効性能を左右する。(Source: [[@2024__arXiv__Benchmarking and Dissecting the Nvidia Hopper GPU Architecture]]) - **疎行列化の利点は命令形式とデータ供給経路に依存する**: H800 の sparse `wgmma` では、共有メモリから2:4枝刈りを実行する SS モードのレイテンシが144クロック、枝刈り済みレジスタを直接読む RS モードが128クロックだった。疎性の理論的な演算削減だけではなく、メタデータ処理と共有メモリ帯域が性能を制約する。(Source: [[@2024__arXiv__Benchmarking and Dissecting the Nvidia Hopper GPU Architecture]]) - **Tensor Core の世代差は、精度形式の追加と命令発行モデルの差を分けて評価すべきである**: Blackwell の GB203 は FP4/FP6 を追加し、ILP=6・25アクティブワープで持続条件に達したのに対し、Hopper の GH100 は ILP=5・29アクティブワープだった。低精度形式の追加は、数値形式だけでなく、少ないワープで独立命令を重ねるスケジューリング条件も変える。(Source: [[@2024__arXiv__Benchmarking and Dissecting the Nvidia Hopper GPU Architecture]], [[@2025__arXiv__Dissecting the NVIDIA Blackwell Architecture with Microbenchmarks]]) - **PTX の命令名と実際の SASS 命令は一致しない場合があり、低精度 Tensor Core の性能説明には両方の観測が必要である**: Blackwell の FP8/FP6 `mma` は QMMA に、FP4 は評価条件によって QMMA または OMMA に変換された。FP4/FP6 対応を仕様表だけから判断すると、現在のソフトウェア経路と実行性能を取り違える可能性がある。(Source: [[@2025__arXiv__Dissecting the NVIDIA Blackwell Architecture with Microbenchmarks]]) ## 未解決の問い - Hopper(H100)の FP8 テンソルコアは HPC の精度要件を満たせるか。FP8 LU 分解 + IRGM の収束条件はどう変わるか - BF16 テンソルコア(Ampere 以降)は FP16 より広い指数範囲を持つ。HPC の密線形代数ソルバーに BF16-TC は有利か - テンソルコアの 4×4 MMA 演算単位は大規模行列の LU 分解においてどのようにタイリングされるか。高性能実装の詳細は - Hopper の Tensor Memory Accelerator(TMA)や Blackwell の FP4 テンソルコアは混合精度ソルバーの設計をどう変えるか - B300/GB300 で指数関数ユニットのスループットが 32 ops/clock/SM に倍増する場合、FlashAttention のボトルネック分布はどう変化するか（FA4 で指数関数がボトルネックに特定されたため） - Rubin の指数関数スループット向上(BF16/FP16 で 4 倍)は FA4 が特定した softmax ボトルネックをどの程度解消するか。テンソルコアスループット(K 次元倍増)と指数関数スループットの比が変わることで、FlashAttention の次世代設計(FA5 相当)はタイルサイズ・パイプライン構成をどう再最適化すべきか - Blackwell の FP4/FP6 Tensor Core を、HPC の数値安定性が必要な行列計算へ適用したとき、FP32 累積・スケーリング・反復精密化でどこまで誤差を抑えられるか。 ## 関連 - 隣接 concept: [[GPU最適化]] / [[混合精度訓練]] / [[混合精度反復求解法]] / [[カーネルフュージョン]] - ハードウェア: NVIDIA V100(Volta) / NVIDIA A100(Ampere) / NVIDIA H100(Hopper) - ハードウェア: [[NVIDIA Blackwell]] / [[RTX 5080]] - 関連エンティティ: [[MAGMA]] / [[Azzam Haidar]] / [[Jack Dongarra]] - 関連 MOC: [[分散深層学習 - MOC]] - [[FlashAttention]] — テンソルコアの深層学習アテンション応用の代表例 - [[NVIDIA Rubin GPU]] — 精度拡張とは独立な K 次元スループット倍増を導入した最新世代 - [[AI Chip Architectures]] — NVIDIA Tensor Core 6世代の発行モデル進化とAMD Matrix Core(MFMA)との対比 - [[CUDA]] — Tensor Core命令を包む warp/thread 階層のプログラミングモデル ## 出典 - [[@2018__SC__Harnessing GPU Tensor Cores for Fast FP16 Arithmetic to Speed up Mixed-Precision Iterative Refinement Solvers]] — SC 2018。Haidar ら。V100 テンソルコアを HPC 密線形代数ソルバーに初めて適用し、FP64 比 4× 高速化を実証 - [[@2024__arXiv__FlashAttention-3 - Fast and Accurate Attention with Asynchrony and Low-precision]] — FA3。FP8 テンソルコアを Hopper でアテンションに適用。ブロック量子化で数値誤差 2.6× 改善 - [[@2026__arXiv__FlashAttention-4 - Algorithm and Kernel Pipelining Co-Design for Asymmetric Hardware Scaling]] — FA4。Blackwell TMEM + 2-CTA MMA モードでアテンションカーネルを再設計 - [[@2026__NVIDIA Developer Blog__Inside NVIDIA Rubin GPU Architecture Powering the Era of Agentic AI]] — Rubin GPU。K 次元テンソルコアスループット倍増・第 3 世代 Transformer Engine(NVFP4 最大 50 PFLOPS)・指数関数スループット向上 - [[AI Chip Architectures]] — Tensor Core命令(mma.sync→wgmma.mma_async→tcgen05.mma)の発行モデル進化、AMD Matrix Core(MFMA)との比較 - [[@2025__arXiv__Dissecting the NVIDIA Blackwell Architecture with Microbenchmarks]] — GB203 の第5世代 Tensor Core、FP4/FP6、QMMA/OMMA、ILP とワープスケーリング