GPU and Floating Point Number

Transcript

1. GPU and Floating Point Number 2025/04/01 Mikiya Michishita(@aztecher)
2. Floating Point Number GPU における演算精度 Blackwell Numerics Table of Contents
3. 想定読者 GPU の演算精度やそれに紐づく高速化の知見があまりないが興味のある方 TensorCore、Transformer Engine について知見があまりないが興味のある方 FP8 / FP6 / FP4 などの低精度演算について興味のある方 本資料によって獲得できる知見 Floating Point についての基本知識 混合精度、MMA、FMA についての簡単な理解 TensorCore やTransformer Engine などの簡単な理解 NVIDIA Hopper, Brackwell などの最新GPU で利用可能な低精度演算の基礎知識 想定読者と本資料によって獲得できる知見
4. 理解を優先するために一部適切ではない表記、表現がある可能性があります。 著者の理解が乏しく、正解でない記述がある可能性があります。 より正確な情報を拾いたい場合はスライド末尾に掲載する参考資料の一覧から一次情報を 拾うことを推奨します。 免責事項
5. Floating Point Number
6. Floating-Point Number （浮動小数点数）は実数をコンピュータで処理するために有限桁 の近似値として扱う方式 この方式はIEEE 754 にて標準化されている IEEE Standard for Floating-Point Arithmetic 現時点で一番新しいのは2019 年に改定されたIEEE 754-2019 ほぼ全てのモダンなシステムが利用している浮動小数点数の表現方式 本資料ではIEEE 754 の詳細の説明や解説は行わないので興味のある方は各自調べてね Floating Point Number とは?
7. Name Common Name Radix Sign bit fraction bits Exponent Bits Exponent Bias binary16 Half precision 2 1 10 5 (-14 ~ 15) 15 binary32 Single precision 2 1 23 8 (-126 ~ 127) 127 binary64 Double precision 2 1 52 11 (-1022 ~ 1023) 1023 IEEE 754 にはその他の形式も定義されている。上記は今回の説明のための最低限を抜粋。 Floating Point Number (IEEE 754)
8. Common Name Radix Sign bit fraction bits Exponent Bits Exponent Bias Single precision 2 1 23 8 (-126 ~ 127) 127 Si gn(1bi t) 値の正負を表現 E xponent(8bi t) 指数部の桁数を表現 Fracti on(23bi t) 小数点以下の数値を表現 正規化前提で の 部分は仮定 Fraction (23bit) Exponent (8bit) Sign (1bit) 32bit（SIngle Precision） Floating Point Format (binary32)
9. 1. まずは10 進数を2 進数に変換する: 2. 2 進数を の形式に変換する（正規化）: 3. binary32 形式に割り当てる Sign は （正の数） Exponent バイアス値を考慮すると 2 進数に変換して、 になる Fraction Fraction は先頭の を暗黙に仮定し、小数点以下の2bit 表現をそのまま割り当てるた め、 になる 例題: をbinary32 形式で表現してみる
10. Exponent bit が8bit なので、素直に考えると指数部のレンジは -127 ~ 128 一方で、定義では -126 ~ 127 。これは特殊な値の表現に利用されているため。 1. Exponent が全て1bit (128 のときのとき) Value Bit representation Sign Exponent Fraction 0111 1111 1000 0...0 +(0) 11111111 0 1111 1111 1000 0...0 -(1) 11111111 0 NaN X111 1111 1XXX X...X 任意 11111111 0 以外任意 特別な値の表現 (1)
11. 2. Exponent が全て0bit (-127 のとき) 指数部が全て0 ビットの場合は、0 、-0 、非正規化数という値を表現する この場合は、仮数部の暗黙の先頭1bit を想定せず、 として表現する例外を適用 ゼロの表現は以下の通り、指数部と仮数部が全て の場合で規定される Value Bit representation Sing Exponent Fraction 0 0000 0000 0000 0...0 +(0) 0000 0000 0 -0 1000 0000 0000 0...0 -(1) 0000 0000 0 特別な値の表現 (2)
12. 2. Exponent が全て0bit (-127 のとき) 指数部が 、 仮数部が 以外の場合は非正規化数（ の数値）を表現する 指数部の計算は、 ではなく として考える例外が適用されるので注意 正規化数から0 に向けて減っていくと、非正規化数の表現範囲に切り替わり、0 までなめら かに減っていくイメージ 特別な値の表現 (3)
13. 以下の表は、正規化数と非正規化数の切り替わりを具体的な10 進数の値で表現したもの 一番上の値は正規化数の中での最小値（正の値） 真ん中の値は非正規化数の中での最大値（正の値） 一番下の値は非正規化数の中での最小値（正の値） Value (10 進数) Bit representation Sign Exponent Fraction 0000 0000 1000 00...0 0 0000 0001 000...0 0000 0000 0111 11...1 0 0000 0000 111...1 0000 0000 0000 00...1 0 0000 0000 000..01 特別な値の表現（4 ）
14. 浮動小数点数の表現範囲（ダイナミックレンジ）と精度 指数部のbit 数が大きいほど、表現できる最大値、最小値の幅が広い 仮数部のbit 数が大きいほど、精度（少数点数の解像度）高く演算が可能 浮動小数点数を表現するbit 長さが長くなるほど、より多くの演算リソースやメモリ容量な どが必要になり、計算速度自体も遅くなるなど、計算コストが上がる 計算速度の高速化、演算リソース効率化の側面ではbit 長さを短くしたいが、それをするこ とでダイナミックレンジや精度に割り当てられるbit 数が減るため数値表現や精度に影響が でるというトレードオフがある 機械学習の分野では低精度を利用し演算速度を上げる技術が利用されるケースが多い 浮動小数点数の表現範囲と精度と計算コスト
15. GPU における演算精度
16. Precision Presented by Used by BFloat16 Google Brain NVIDIA , AMD, Google TPU, ... TF32 NVIDIA NVIDIA GPU (Ampare+) FP8 Open Compute Project* NVIDIA , AMD, Gaudi FP6 Open Compute Project NVIDIA(Blackwell+), AMD(MI355X+) FP4 Open Compute Project NVIDIA(Blackwell+), AMD(MI355X+) *FP8 はOpen Compute Project で仕様が策定されたが、IEEE が異なる使用を検討中のよう TF32 という形式を理解する上で、Mixed-Precision Arithmetric とFMA 、及びTensorCore の議論が欠かせないため、まずはこれらの説明を行う。 GPU のカタログで見かける演算精度
17. Floating Point の演算形式の一つで、異なる数値精度を利用する演算 最小限の情報を高精度のフォーマットで保持しつつ、低精度フォーマットで演算を狙っ ていく手法 精度を低下を抑えつつ、計算速度の向上やメモリフットプリントの低下を実現できる 主に機械学習の文脈で見られる手法 Mixed-Precision Arithmetric （混合精度演算）
18. MMA : Matrix Multiply- Add 行列の積を求めた後、行列和を求める演算 行列演算: ML アプリケーションでは数多く発生する種類の 演算形式 FMA: Fused Multiply-Add 積和演算を1 命令で行う（積算の結果を浮 動小数点数に丸めずに積和まで１命令で 実行する）ことで、最終演算結果の誤差 を小さくする技術 スカラー演算: IEEE 754-2008 で標準化されている MMA and FMA
19. NVIDIA のTensorCore は混合精度FMA を実行する演算器であり、そのためのフォーマット としてTensorFloat が定義されている 混合精度FMA では、低精度積 + 高精度Accumlation がよく利用される 実際はこの混合精度FMA を大量に並列実行し、MMA を実現している。 詳細はAppendix に追記（2025.05 ） 各演算（Multiply・Accumulation ）でサポートする演算精度は世代ごとに異なる NVIDIA TensorCore
20. TensorCore の世代ごとの乗算演算精度のサポート一覧 Gen(Arch) FP64 TF32 FP16 BF16 FP8 INT8 INT4 INT1 1 (Volta) - - ✓ - - - - - 2 (Turing) - - ✓ - - ✓ ✓ ✓ 3 (Ampare) ✓ ✓ ✓ ✓ - ✓ ✓ ✓ 4 (Hopper) ✓ ✓ ✓ ✓ ✓ ✓ - - TensorCore (Multiply) ref: https://ja.wikipedia.org/wiki/ Tensor コア
21. TensorCore の世代ごとの加算演算精度のサポート一覧 Gen(Arch) FP64 FP32 FP16 INT32 1 (Volta) - ✓ ✓ - 2 (Turing) - ✓ ✓ - 3 (Ampare) ✓ ✓ ✓ ✓ 4 (Hopper) ✓ ✓ ✓ ✓ TensorCore (Accumlation) ref: https://ja.wikipedia.org/wiki/ Tensor コア
22. TensorCore で実現できる混合精度FMA の一部をリストアップする Input Muliplication Accumulation Output Arch FP32 TF32 FP32 FP32 Ampare+ FP16 FP16 FP32 FP32 Volta+ BF16 BF16 FP32 FP32 Ampare+ FP8 FP8 FP16 FP16 Hopper+ 基本的には低精度積を実行し、その結果を高精度Accumulation して出力を得る Precisions of TensorCore FMA
23. 4✕4 正方行列A、B、C によるFMA、 における各要素の計算式 FMA の行列演算において、行列和部分の精度が出力精度に大きく影響する 積部分は低精度に落とすことで演算の効率化を行い、和部分（Accumulation ）では高精 度にすることで最終結果の誤差を大幅に低減する Mixed Precision における演算精度
24. 2020 年にTensorCore 向けに定義された独自のFloating Point の表現形式 AI ワークロードの精度要件として十分なFP16 と同等の10bit の仮数部、FP32 と同じ表現範 囲をサポートする8bit 指数部を持ち、符号部1bit と合わせて計19bit で、32bit に収まる。 AI ワークロードにより特化したFloatingPoint の表現形式となっている。 FP32 と比較して精度はかけるものの、計算が高速になる 23bit 8bit 10bit 8bit 10bit 5bit TensorFloat32 FP16 FP32 TF32 （TensorFloat-32 ） Format
25. TensorCore における混合精度FMA モードの一つ FP32 Matrix 同士の演算時、TF32 に変換した状態で乗算した結果に対してFP32 で加算 入力、出力をFP32 形式で行うことにより、ユーザアプリケーションでのコード変更が不要 FP32 Matrix FP32 Matrix Format to TF32 and Multiply FP32 Accum FP32 Matrix 23bit 8bit 10bit 8bit 10bit 5bit TensorFloat32 FP16 FP32 TF32 TensorCore
26. 2018 年にGoogle Brain が提唱した独自のFloating Point の表現形式 Google のTPU で採用され、後に様々なアクセラレータでサポートされる 仮数部に7bit 、指数部にFP32 と同じく8bit を割り当てる形になっており、精度は下がるも のの、FP32 と同様の表現範囲をサポートする 23bit 8bit 7bit 8bit 10bit 5bit BFloat16 FP16 FP32 BF16 (BFloat16) Format
27. Original (Decimal) FP32 ( Value) BFloat16 ( Value) BF16 Error FP16 ( Value) FP16 Error 1.0005 1.0005 1.000000 0.000500 1.000977 0.000477 1.0010 1.0010 1.000000 0.001000 1.000977 0.000023 1.0e+10 1.0e+10 1.0e+10 0.000000 inf inf 1.0e-05 1.0e-05 1.0e-05 0.000000 9.999848e- 06 1.51968e- 10 BF16 とFP16 の違いによる差
28. Ampare 以降、BF16 フォーマットがTensorCore に追加され、BF16 TensorCore の命令は FP16 と同様のスループットになった。 FP16 / BF16 Matrix FP16 / BF16 Matrix FP16 / BF16 Matrix FP32 Accum FP32 Matrix FP16/BF16 TensorCore
29. 従来のFloating Point は必ずしも機械学習のようなワークロードにマッチしていなかった これらのワークロードに適するようなFloating Point を再定義して利用するようになった 計算精度はある程度低くても良いため、仮数部は削る 表現可能数自体はある程度必要になるため、指数部分は比較的重要 このレベルでの演算最適化とハードウェアレベルの最適化によって高速化する世界 Insights about New Floating Point Format
30. Hopper では右のようなブロックを4 つと以下を まとめたものが1 つのSM である L1 Instructi on C ache Tensor Memor y Accelerato r 256 KB L1 Data C ache / S hare d M emor y Tex GB100 は144 SMs、H100 SX M5 は132 S Ms、 PCIe は114 SMs INT32 x16 FP32 x32 FP64 x16 Tensor Core 4Gen Register File Dispatch Unit Warp Scheduler L1 Instruction Cache SFU LD / ST x8 TensorCore Hardware Overview (in Hopper) ref: https://developer.nvidia.com/blog/nvidia-hopper-architecture-in-depth/
31. FP8 の話に進む前にScaling の議論を行っておく必要がある AI モデルの大規模化に伴い、メモリ使用量と計算速度の最適化が更に重要になり、FP8 な どの低精度フォーマットを利用した混合精度演算が現実的な手段として利用される 一方でこれらのフォーマットは表現可能な数値範囲が狭く、精度も悪くなるため、直接演 算した結果オーバーフロー、アンダーフローする可能性が高い問題がある Scaling は演算結果を低精度フォーマットで表現できる範囲に収める仕組み Scaling 自体はFP8 やFP4 に特有の話ではなく、低精度演算一般に対して考慮すべき話ではある が、精度が低い分より顕著な問題として発生する。 Mixed-Precision and Scaling
32. DNN の学習では4 種類のTensor が登場する activati ons / activati on gradi e nts wei ghts / wei ght gradients activati on gradi ents 以外はFP16 レンジに収ま ることが多いが、activatio n grad ients はFP16 のレンジより小さい値になるケースが多い FP16 レンジに収まるように"shift" する操作が Scaling であり、その"shi ft" 値をScaling Factor と呼ぶ FP16 Representable range Become zero in FP16 FP16 denorms log2(magnitude) 15 -24 -14 -1 1 Percentage of all activation gradient values Scale ("shift") Scaling on DNN (ex. FP16 Scaling) ref: https://developer.nvidia.com/blog/mixed-precision-training-deep-neural- networks/
33. Constant Scaling Factor 固定値でScaling Factor を規定する方式 Neural Network の構造、フレームワーク、ミニバッチサイズなどに依存するため、適切 な値を見つけるために試行錯誤が発生する 勾配の統計値を取得できている場合はそれをもとに決定することもできる Dynamic Scaling Factor 大きなScaling Factor を初期値とし、各トレーニングイテレーションでスケーリングファ クターを動的に増減させる方式 Scaling Factor の決定方式
34. Hopper 以降にTensorCore でサポートされるようになった演算精度 特徴としてFormat が2 種類存在する Format: E5M2, E4M3 (E = Exponent = 指数, M = Mantissa(Fraction) = 仮数) Accumulator: FP16, FP32 FP8 Matrix FP8 Matrix Multiply FP32 or FP16 Accum bias/ act/ ... 2bit 23bit 8bit 5bit 10bit 5bit FP8 E5M2 FP16 FP32 3bit 4bit FP8 E4M3 FP32 / FP16 / BF8 / FP8 Matrix convert FP8 Format
35. Training でFP8 を利用する場合、2 種類のFormat を両方利用することが推奨される forward: 比較的精度が求められるため、E4M3 の利用が望ましい backward: 比較的ダイナミックレンジが求められるため、E5M2 の利用が望ましい 当然この例から外れるケースは存在する FP8 を利用した混合精度演算を行う上でもScaling が非常に重要になる forward / backword でScaling Factor がそれぞれ変わる FP16 のようなGlobal に単一のScaling Factor では不十分 per-tensor scaling という方式の導入 FP8 on Training
36. FP8 ではTensor （weights / activation ）毎にScaling Factor を持つ必要がある 各Tensor 個別にScaling してFP8 に変換する方針を採用 FP8 Range FP8 FP8 FP8 FP8 Range FP8 per-tensor scaling ref: https://w w w.nvidia.com/ja-jp/on-demand/session/gtc24-s62457/
37. 生成されるテンソルの最大絶対値（amax ）に基づいscaling factor を決定する方式 コンセプトは単純だが実装が難しい 中間の高精度演算結果をメモリに保持したり、そこから演算を複数回実施したり かなりのオーバーヘッドが発生しFP8 を利用するメリットを削ってしまう Input FP8 Output Scale Find Maixmum Compute High Precision Output Operator FP8 Scaling Strategy: Just-In-Time Scaling ref: https://w w w.nvidia.com/ja-jp/on-demand/session/gtc24-s62457/
38. 過去ある回数の反復で観測されたamax に基づいてscaling factor を決定する方式 特別なオーバーヘッドはなくFP8 の利点を最大限発揮できるが、過去の履歴を保存する必 要がある FP8 Operator Scaling Factor Current Window Recipe amax history new amax Input FP8 Output FP8 Scaling strategy: Delayed Scalling ref: https://w w w.nvidia.com/ja-jp/on-demand/session/gtc24-s62457/
39. Delayed Scaling ではScaling Factor を決定する"Recipe" という概念があった これを具体的に実装したものが Transformer Engine （TE ） TE はTransformer のbuilding block 向けのFP8 のRecipe のOSS 実装 Delayed Scaling における複雑なScaling Factor の決定をこのライブラリが提供するAPI で 簡単に利用できる ＊ 2023 年度段階では、Hopper + TE でのFP8 の演算サポート箇所は一部の演算箇所にとどま り（GEMM 演算）エンドツーエンドでみるとまだまだ高精度な演算で賄っている箇所がある Transformer Engine
40. 2023 年12 月時点でもよく利用されている混合精度スキームはFP16-FP32 かBF16-FP32 実際FP16 は数値範囲が狭いためモデル学習が不安定になりがち。 BF16-FP32 のケースが多い（Megatron-Turing NLG-530B, Bloom-175B, Gopher ） 前述したとおり、FP8 の適用範囲は限定的 Weight の更新や勾配同期はより高い精度を利用 Microsoft Research の論文でFP8 演算を学習の広いフェーズで行い、BF16 の精度 と同等のまま、メモリコストや通信コストを抑える手法が提案されています Delayed Scaling に関しても不安定性などについて議論があるようです。FP8 training を支える技術1 がFP8 Training の日本語記事として非常によくまとまっているため、合 わせて拝読することをおすすめします。 Actual Case Study of Mixed-Precision Training
41. Brackwell Numerics
42. Blackwell では更に新しい低精度の演算のサポートが行われるようになった MXFP8, MXFP6, MXFP4 背景にはBlock Format とそれをベースとした新しいScaling Factor の導入がある このBlock Format はOCP で定義されているMicroscaling Format と同一と思われる 具体的に実装するに当たり、ハードウェアやソフトウェアにさらなる進化が加わったよう 5th Generation Tensor Core 2nd Generation Transformer Engine 現時点（2025/03/29 時点）で詳細な情報は見つけられていない。 Brackwell Numerics
43. OCP からMX のSpec がVersion 1.0 が2023/09/07 に公開されている 開発者は、AMD、Arm、Intel、Meta、Microsoft、NVIDIA、Qualcomm Microscaling (MX)-compliant data format を定義する MXFP8、MXFP6、MXFP4、MXINT4 Binary Encoding を定義する FP6、FP4、INT8 Blackwell のサポートするBlock Format が正しくこのOCP の仕様に準拠しているかは微妙で はあるが、現時点で公開されているものがこれのみであること、NVIDIA の方が貢献されてい ることから、このOCP のFormat に基本的には準拠しているとして以降の議論を進める MX - Microscaling Format
44. MX- compli ant format は以下の3 つの要素によ って特徴づけられる 1. Scale ( ) data t ype / encodi ng 2. Private element ( ) data t ype / encodi ng 3. Scali ng block si ze ( ) 全てのBlock の要素 は同じデータ型をもち、 scali ng factor は elem 全体に共有 各block (scale + - elems) は bi ts にエ ンコードされる X （shared scale） P1 （element） P2 （element） P3 （element） Pk （element） ・・・ w bits d bits k scalar elements Microscaling (MX)
45. Format Name Element Data Type Element Bits (d) Scaling Block Size (k) Scale Data Type Scale Bits (w) MXFP8 FP8 (E5M2) 8 32 E8M0 8 MXFP8 FP8 (E4M3) 8 32 E8M0 8 MXFP6 FP6 (E3M2) 6 32 E8M0 8 MXFP6 FP6 (E2M3) 6 32 E8M0 8 MXFP4 FP4 (E2M1) 4 32 E8M0 8 MXINT8 INT8 8 32 E8M0 8 MX-compliant Formats
46. E8M0 Exponent bias 127 Supported exponent range -127~127 Inf inities N/A NaN Zeros N/A E8M0 は「従来のバイアスされたFloat32 の指数部の符号なし表現」と規定されている Exponent に8bit 利用し、底については明確な言及がないがおそらく だろう NaN encoding のみ予約されており、bias は 、レンジは ~ E8M0 format of Scale Data Type
47. Scale ・・・ 8 bits 8 bits 32 scalar elements E4 M3 E4 M3 Scale ・・・ 8 bits 8 bits 32 scalar elements MXFP8 FP8 E4M3 MXFP8 INT8 Scale ・・・ 8 bits 8 bits 32 scalar elements E5 M2 E5 M2 MXFP8 FP8 E5M2 INT8 INT8 MX-Compliant format w/ FP and INT
48. Hopper 世代からあったFP8 のMX-compliant format 表現 (Inference / Training) Data format : E4M3, E5M2 Metadata format : E8M0 BlockSi ze: 32 E8M0 E5M2 E5M2 E5M2 E5M2 ・・・ 8 bits 8 bits 32 scalar elements E8M0 E4M3 E4M3 E4M3 E4M3 ・・・ 8 bits 8 bits 32 scalar elements E5M2 E4M3 MXFP8 (E5M2 / E4M3)
49. FP4 が何かしらの理由で利用できないが、FP8 よりメモリフットプリントを抑えたいケース Data format : E2M3, E3M2 Metadata format : E8M0 BlockSi ze: 32 E8M0 E3M2 E3M2 E3M2 E3M2 ・・・ 8 bits 6 bits 32 scalar elements E8M0 E2M3 E2M3 E2M3 E2M3 ・・・ 8 bits 6 bits 32 scalar elements E3M2 E2M3 MXFP6 (E3M2 / E2M3)
50. 現状最も高速に計算できるフォーマット（Inference ） Data format : E2M1 Metadata format : E8M0 BlockSi ze: 32 E8M0 E2M1 E2M1 E2M1 E2M1 ・・・ 8 bits 4 bits 32 scalar elements E2M1 MXFP4 (E2M1)
51. FP4 になると、表現可能な数値がリストアップできる 指数部 仮数部 仮数部 FP4 values
52. elements の各要素は、それぞれのFor mat で 表現可能な数値を持つ。 の場合は、1 block あたり32 個の値 Scale は8bi t で表現できる値（E8M0 ）であり、 この値をblock 毎に乗じることでScale させる Scale S ・・・ w bits d bits 32 scalar elements E M S E M 1.5 3.0 2.0 ・・・ 16 （scale） 8 bits 24 48 32 ・・・ k scalar elements MX-Compliant Example
53. where , はベクター と のそれぞれの block scale , はベクター と それぞれの 番目 の要素 32 scalar elements PA 1 PA 32 ・・・ XA 8 bits PA 2 32 scalar elements PB 1 PB 32 ・・・ XB 8 bits PB 2 PA 1 PA 32 ・・・ PA 2 PB 1 PB 32 ・・・ PB 2 C 32 scalar elements 32 scalar elements Dot Product of MX-Compliant Format Vector
54. 一般に大きなベクトルを計算する場合、M X 準拠 のブロックサイズ のサブベクトルに分割し、 各サブベクトル毎にDot Product を計算した結 果を総和することで得られる ベクトルは計算前にその長さがブロックサイズ の倍数になるように調整（Padd ing ）されるこ とが前提になる B1 C k scalar elements B2 k scalar elements ・・・ Bm A1 A2 ・・・ Am n (k x m) MX-compliant vectors n (k x m) MX-compliant vectors General Dot Product
55. -length vector V: を の形式に変換したい。 → Block Scale と、Element を算出する必要がある Block Scale: = を算出 が2 のべき乗なら とおく そうでない場合、 より小さい中で最大の2 のべき乗を とおく MX-compliant format のdata type における最大の2 のべき乗の値を とする で算出する Vector 中の最大値を求め、変換先の表現可能な最大値で割る操作 2 のべき乗で表現してスケーリングを計算し、E8M0 で表現可能な整数にしている Conversion from Vector to MX format (1)
56. Elements: で各要素のScale 値を計算する を、変換先のデータフォーマットで量子化する Scaling Factor が計算できているので、その値で単純にスケールしたのちに量子化することで 表現可能な範囲の数値に落とし込む Conversion from Vector to MX format (2)
57. MX-Compliant Format では、単純な数値表現以外にScale 値という追加のメモリフットプ リントが存在する Scale を共有する値（Scaling block size ）が大きいと、全体のBlock 数が減るためこの 追加のメモリフットプリントは小さくなるが、その分繊細なScaling ができなくなるため overflow/underflow の可能性が増加する。 逆に小さいと、繊細なScaling ができるようになりoverflow/underflow の可能性が低減す る一方、Block 数が増えるため、その分のScale 値が必要になりメモリフットプリントが増 加する Insight about MX-Compliant Format
58. 基本的な考えかたや概念はOCP のMicroscaling Formats と同様 FP4 などの演算をサポートするためにBlock Format の対応が行われている FP4 のMetadata Format （OCP Spec で言うScale Data Type ）にE4M3 が！ Floating Point なScaling Factor に対応したい狙いか？ FP4 のBlockSize に16 が！ Scaling の粒度をより詳細にしたい狙いか？ Format Data Format Metadata Format BlockSize FP8 E4M3, E5M2 E8 32 FP6 E2M3, E3M2 E8 32 FP4 E2M1 E8, E4M3 32, 16 Block Format on Blackwell GPU
59. Per-Tensor Scaling と比較して、Block Scaling はより繊細にScaling を実行する分、オー バーヘッドも余計にかかってしまう。 5th Gen TensorCore、2nd Gen Transformer Engine の登場 Tensor Core にはdata とmetadata を入力として渡す scaling-factor の処理はTensorCore 内部で実行する 2nd Gen TE は 5th Gen Tensor Core とTensorRT-LLM やNemo を組み合わせ、LLM やMoE モデルの推論やトレーニングをより加速させるように実装されている 具体的にどのような実装、仕組みになっているかの細部については、現時点（2025/03/29 時 点）でこれ以上の詳細な情報は見つけられていない。 Block Format and Hardware changes
60. GPU がサポートする浮動小数点演算は、混合精度の利用や低精度演算の利用が処理の高速 化に大きく寄与しており、この方式から基本逸脱せずに進化を遂げてきている。 FP8 の低精度演算にもまだまだ課題はある一方、ハードウェアとしてFP4 をサポートし始め るなど、まだまだ低精度演算の領域は進化中 今後、どうなるだろうか？ FP8/FP4 の演算の適用範囲を広げる研究やScaling に関する研究が増えそう。 NCCL レベルのレイヤーでも、これに対応する実装が入っていくかも？ これ以上の低精度は難しそうじゃない？どうなんだろうか？ 今後の性能向上の指針がどのような方向に向かっていくかが楽しみですね まとめ
61. Appendix. A ではTensorCore の動作を検証した論文であるModeling Deep Learning Accelerator Enabled GPUs を引用し、TensorCore について深堀りする この論文では、詳細な設計や動作原理が公開されていないNVIDIA TensorCore について 様々な評価によって推察していくという論文である したがって、以降の内容もあくまで推測であるという点には注意されたい この論文の内容に従うため、以降の話はVolta TensorCore という前提で話す Volta は初めてTensorCore が搭載されたNVIDIA GPU であり研究も盛んに行われている TensorCore の基礎を知る上で良い題材になると思われる Appendix.A: DeepDive into TensorCore (Volta)
62. TensorCore を利用してGEMM （GEneral Matrix Multiply ）演算するためのインターフェ イスとしてWMMA （Warp Matrix Multiply Accumulation ） と呼ばれるC++API が提供 されている 例えばMMA 演算を行う命令として、wmma::mma_sync 命令がある このPTX 命令は合計16 のHMMA 命令（4SETx4STEP ）に変換され処理される この処理を追ってみるとFMA 演算で構成されていることがわかる TensorCore MMA
63. 4 8 4 8 4 4 8 4 8 4 8 4 8 4 8 8 16 16 16 x 16 Matrix A A B C A B C A B C A B C Octet 0 Octet 1 Octet 2 Octet 3 h g f e d c b E B F C G D H h g f d c b B F C G D H h g d c C G D H h d D H SET 1 SET 2 SET 3 SET 4 Thread Group 0 Thread Group 4 e[2:3] x A e[2:3] x E e[0:1] x E e[0:1] x A a[2:3] x A a[2:3] x E a[0:1] x E a[0:1] x A f[2:3] x B f[2:3] x F f[0:1] x F f[0:1] x B b[2:3] x B b[2:3] x F b[0:1] x F b[0:1] x B g[2:3] x C g[2:3] x G g[0:1] x G g[0:1] x C c[2:3] x C c[2:3] x G c[0:1] x G c[0:1] x C h[2:3] x D h[2:3] x H h[0:1] x H h[0:1] x D d[2:3] x D d[2:3] x H d[0:1] x H d[0:1] x D STEP0 ~ STEP4 (SET 1) STEP0 ~ STEP4 (SET 2) STEP0 ~ STEP4 (SET 3) 4 x 4 A a Tensor Core Octet0 Thread Group 0 Thread Group 4 FEDP FEDP FEDP FEDP FEDP FEDP FEDP FEDP Accumulator Buffer Accumulator Buffer Matrix A Buffer Matrix A Buffer MUX Tensor Core Octet3 Octet2 Octet1 TG0/Thread0-3 = Lane 0-3 TG1 TG5 TG2 TG6 TG3 TG7 TG1/Thread0-3 =Lane4-7 TG5/Thread0-3 =Lane20-23 TG2/Thread0-3 =Lane8-11 TG5/Thread0-3 =Lane24-27 TG3/Thread0-3 =Lane12-15 TG7/Thread0-3 =Lane28-31 TG4/Thread0-3 =Lane 16-19 (Part of) Sub-Core Matrix B Buffer 1Lane/Thread 1Warpあたり32Threadあり、1Warpあたり 2xTensorCoreをActivationする。 1TensorCoreあたり2Octetあり、1Octetあたり 8Threadずつ存在している。ThreadGroupあた り4Threadの対応になる。 Lane Id = Thread index それぞれのThreadがRegisterFileに持っている 担当分のデータ（Fragment）をTensorCoreの 内部BufferにFetchする TensorCoreのBufferにFetchされたデータから 各FEDPに対してデータが供給される Other Execution Units Operand bus 3（Matrix C） Operand bus 2（Matrix B） Operand bus 1（Matrix A） TG0/Thread0-3 = Lane 0-3 TG0/Thread0-3 = Lane 0-3 TG4/Thread0-3 =Lane 16-19 TG4/Thread0-3 =Lane 16-19 FEDP （Four-Element Dot Product） Unit Pipeline Registers 16 x 16 Matrix B 16 x 16 Matrix C 16 x 16 Matrix D TG0/Thread0-3 = Lane 0-3 Writeback TG4/Thread0-3 =Lane 16-19 TG1/Thread0-3 =Lane4-7 TG5/Thread0-3 =Lane20-23 TG5/Thread0-3 =Lane24-27 TG2/Thread0-3 =Lane8-11 TG7/Thread0-3 =Lane28-31 TG3/Thread0-3 =Lane12-15 STEP0 ~ STEP4 (SET 4) RegisterFile（ 512 x 32 Thread x 32 bit = 64 kB） General-purpose registers accessible by Thread in Warp TensorCore MMA overview
64. 例として16x16x16 の行列演算を考える Octet という概念を仮に導入する。各Octe t で最 終結果のうち8x8 の領域を担当する Octet0 では、オペランド行列 は8x16, は 16x8, は8x8 の領域の演算になる 各オペランドの結果は最終結果の8x8 領域 演算は4SET に別れる。理解のためにオペランド を4x4 の行列毎に名前をつけて整理する 16 16 16 x 16 Matrix A A B C A B C A B C A B C Octet 0 Octet 1 Octet 2 Octet 3 h g f e d c b E B F C G h g f d c b B F C G D H h g d c C G D H h d D H SET 1 SET 2 SET 3 SET 4 4 x 4 A a 16 x 16 Matrix B 16 x 16 Matrix C 16 x 16 Matrix D TensorCore MMA overview - 16x16x16 Matrix (1)
65. 各SET は4x8 行列毎に異なるThread Gro up で処 理する（仮単位、以降"TG " と略す） TG あたり4Thread Octet あたり2TG 各SET はそれぞれのTG で4STEP に渡って演算が 行われ、全体として16STE P でOcte t あたり8x8 行列の演算結果が得られる このSTEP は具体的にSA SS 命令としては観測でき るが、1STEP 1C YCLE であることに注意 4 8 4 8 4 4 8 4 8 4 8 4 8 4 8 8 h g f e d c b E B F C G D H h g f d c b B F C G D H h g d c C G D H h d D H SET 1 SET 2 SET 3 SET 4 Thread Group 0 Thread Group 4 e[2:3] x A e[2:3] x E e[0:1] x E e[0:1] x A a[2:3] x A a[2:3] x E a[0:1] x E a[0:1] x A f[2:3] x B f[2:3] x F f[0:1] x F f[0:1] x B b[2:3] x B b[2:3] x F b[0:1] x F b[0:1] x B g[2:3] x C g[2:3] x G g[0:1] x G g[0:1] x C c[2:3] x C c[2:3] x G c[0:1] x G c[0:1] x C h[2:3] x D h[2:3] x H h[0:1] x H h[0:1] x D d[2:3] x D d[2:3] x H d[0:1] x H d[0:1] x D STEP0 ~ STEP4 (SET 1) STEP0 ~ STEP4 (SET 2) STEP0 ~ STEP4 (SET 3) 4 x 4 A a STEP0 ~ STEP4 (SET 4) TensorCore MMA overview - 16x16x16 Matrix (2)
66. Tensor Core Octet0 Thread Group 0 Thread Group 4 FEDP FEDP FEDP FEDP FEDP FEDP FEDP FEDP Accumulator Buffer Accumulator Buffer Matrix A Buffer Matrix A Buffer MUX Tensor Core Octet3 Octet2 Octet1 TG0/Thread0-3 = Lane 0-3 TG1 TG5 TG2 TG6 TG3 TG7 TG1/Thread0-3 =Lane4-7 TG5/Thread0-3 =Lane20-23 TG2/Thread0-3 =Lane8-11 TG5/Thread0-3 =Lane24-27 TG3/Thread0-3 =Lane12-15 TG7/Thread0-3 =Lane28-31 TG4/Thread0-3 =Lane 16-19 (Part of) Sub-Core Matrix B Buffer 1Lane/Thread 1Warpあたり32Threadあり、1Warpあたり 2xTensorCoreをActivationする。 1TensorCoreあたり2Octetあり、1Octetあたり 8Threadずつ存在している。ThreadGroupあた り4Threadの対応になる。 Lane Id = Thread index それぞれのThreadがRegisterFileに持っている 担当分のデータ（Fragment）をTensorCoreの 内部BufferにFetchする TensorCoreのBufferにFetchされたデータから 各FEDPに対してデータが供給される Other Execution Units Operand bus 3（Matrix C） Operand bus 2（Matrix B） Operand bus 1（Matrix A） TG0/Thread0-3 = Lane 0-3 TG0/Thread0-3 = Lane 0-3 TG4/Thread0-3 =Lane 16-19 TG4/Thread0-3 =Lane 16-19 FEDP （Four-Element Dot Product） Unit Pipeline Registers TG0/Thread0-3 = Lane 0-3 Writeback TG4/Thread0-3 =Lane 16-19 TG1/Thread0-3 =Lane4-7 TG5/Thread0-3 =Lane20-23 TG5/Thread0-3 =Lane24-27 TG2/Thread0-3 =Lane8-11 TG7/Thread0-3 =Lane28-31 TG3/Thread0-3 =Lane12-15 RegisterFile（ 512 x 32 Thread x 32 bit = 64 kB） General-purpose registers accessible by Thread in Warp TensorCore MMA overview - 16x16x16 Matrix (3)
67. 1Warp が2 つのTensorCore を利用（それぞれのTensorCore に2 つのOctet が存在） Octet 毎に8 つのFEDP ユニット（TG 毎に4 つ）があり、TensorCore のバッファにフェッチ されたデータから各FEDP ユニットにデータが供給される FEDP ユニットは4 要素の乗算と累積を4 ステージで実行する（=4cycle 、ではない） Thread がRegisterFile （通常のGPU レジスタファイル）に持っている担当分のデータ （Fragment ）が、TensorCore の内部バッファにフェッチされる 1Warp あたり32Thread で構成され、それぞれのThread に対してLane が存在するイメージ 図の通り、行列A , C に関してはThreadGroup 毎に分したバッファにフェッチされるが、行 列B は共有バッファにフェッチされるよう TensorCore MMA overview - 16x16x16 Matrix (4)
68. オペランド行列B が共有バッファに置かれる理由 基本的にOctet 単位（=2xTG 単位) で演算が進む 前図でも表現したが、改めてOcte t 毎の演算を見 ると右テーブルのような形になっていそう（合計 4SET ある演算の一部） 各STEP の演算で同一TensorCore 内の2 つのTG が参照するオペランド行列B の要素は同じになっ ている SET STEP TG X TG X+4 1 0 a[0:1] A e[0:1] A 1 a[2:3] A e[2:3] A 2 a[0:1] E e[0:1] E 3 a[2:3] E e[2:3] E 2 0 b[0:1] B f[0:1] B 1 b[2:3] B f[2:3] B 2 b[0:1] F f[0:1] F 3 b[2:3] F f[2:3] F TensorCore MMA overview - 16x16x16 Matrix (5)
69. 実際の数値演算を担うのはFE DP ユニット 入力は8 要素あり、2 要素ずつ乗算したものの累 積し、パイプラインレジスタ（ここには累積計算 に使うオペランド、行列 の要素）のデータとも 累積して最終出力を得る そのため、4x4 行列の演算における単一要素の演 算を一度に実行するイメージ（以下） は次に同一領域を演算する際の加算項として利用 FEDP （Four-Element Dot Product） Unit Pipeline Registers TensorCore MMA overview - 16x16x16 Matrix (6)
70. Mixed-Precision Training of Deep Neural Network Train With Mixed Precision FP8 training を支える技術 1 Using FP8 with Transformer Engine What's New in Transformer Engine and FP8 Training OCP Microscaling Formats (MX) Specif ication Version 1.0 NVIDIA Blackwell Architecture Technical Brief Blackwell Numerics for AI FP8-LM: Training FP8 Large Language Model References
71. NVIDIA A100 Tensor Core GPU Architecture NVIDIA H100 Tensor Core GPU Architecture NVIDIA Hopper Architecture In-Depth Numerical Behavior of NVIDIA Tensor Cores Modeling Deep Learning Accelerator Enabled GPUs References