Coelacanth's Dream

FP8 と BF8 に対応する AMD GFX940/CDNA 3 と BF8 に対応する Intel Xe-HPC

AMD の Stanislav Mekhanoshin (rampitec) 氏により、GFX940/CDNA 3 で追加された FP8/BF8 フォーマット命令の対応を LLVM に追加するパッチが投稿された。パッチはすでに llvm/llvm-project の main ブランチに取り込まれている。

GFX940/CDNA 3 の FP8 フォーマット対応自体は AMD Financial Analyst Day にて発表されており、GFX940/CDNA 3 は AI学習性能が GFX90A/CDNA 2 の 8倍という比較において、用いるデータフォーマットは GFX940/CDNA 3 が FP8、GFX90A/CDNA 2 は FP16 となっていた。

FP8/BF8

FP8/BF8 フォーマットを用いた演算命令は MFMA (Matrix-Fused-Multiply-Add), SMFMAC 命令と miSIMDユニットで対応し、通常の SIMDユニット側では変換命令にのみ対応する。
FP8/BF8 フォーマットは 8-bit ということはわかるが、指数部 (range, exponent) と仮数部 (precision, mantissa) がそれぞれどうなっているかはパッチ内には記述されていない。
8-bit の浮動小数点フォーマットには NVIDIA H100 GPU が先に対応しており、そして指数部 5-bit 仮数部 2-bit の E5M2 と指数部 4-bit 仮数部 3-bit の E4M3 の 2種類に対応している。
GFX940/CDNA 3 における FP8/BF8 がそれらと同様かはまだ不明だが、仮にそうだとすると、精度を寄った E4M3 は FP8 に、ダイナミックレンジを重視した E5M2 は BF8 となるだろうか。

行列のレイアウトは MFMA 命令では 16x16x32, 32x32x16 に、SMFMAC 命令では 16x16x64, 32x32x32 に対応している。
アキュムレータと最終的な出力フォーマットは F32 のみに対応する。
命令名の末尾は _{BF8,FP8}_{BF8,FP8} からなる 4つのパターンがあるが、既存の命令名からそれぞれの入力フォーマットと考えられる。
MFMA, SMFMAC 命令ではこれまで入力される 2個の行列は同じフォーマットだったが、FP8BF8 に限っては異なる組み合わせに対応するのだろう。

 +  defm V_MFMA_F32_16X16X32_BF8_BF8 : MAIInst<"v_mfma_f32_16x16x32_bf8_bf8", "F32_I64_X32",    int_amdgcn_mfma_f32_16x16x32_bf8_bf8>;
 +  defm V_MFMA_F32_16X16X32_BF8_FP8 : MAIInst<"v_mfma_f32_16x16x32_bf8_fp8", "F32_I64_X32",    int_amdgcn_mfma_f32_16x16x32_bf8_fp8>;
 +  defm V_MFMA_F32_16X16X32_FP8_BF8 : MAIInst<"v_mfma_f32_16x16x32_fp8_bf8", "F32_I64_X32",    int_amdgcn_mfma_f32_16x16x32_fp8_bf8>;
 +  defm V_MFMA_F32_16X16X32_FP8_FP8 : MAIInst<"v_mfma_f32_16x16x32_fp8_fp8", "F32_I64_X32",    int_amdgcn_mfma_f32_16x16x32_fp8_fp8>;
 +  defm V_MFMA_F32_32X32X16_BF8_BF8 : MAIInst<"v_mfma_f32_32x32x16_bf8_bf8", "F32_I64_X16",    int_amdgcn_mfma_f32_32x32x16_bf8_bf8>;
 +  defm V_MFMA_F32_32X32X16_BF8_FP8 : MAIInst<"v_mfma_f32_32x32x16_bf8_fp8", "F32_I64_X16",    int_amdgcn_mfma_f32_32x32x16_bf8_fp8>;
 +  defm V_MFMA_F32_32X32X16_FP8_BF8 : MAIInst<"v_mfma_f32_32x32x16_fp8_bf8", "F32_I64_X16",    int_amdgcn_mfma_f32_32x32x16_fp8_bf8>;
 +  defm V_MFMA_F32_32X32X16_FP8_FP8 : MAIInst<"v_mfma_f32_32x32x16_fp8_fp8", "F32_I64_X16",    int_amdgcn_mfma_f32_32x32x16_fp8_fp8>;

変換命令は F32 <-> FP8,BF8 の変換に対応し、パックド命令にも対応している。
それ以外のフォーマット、F16, BF16, F64 との変換には対応していない。

 +let SubtargetPredicate = HasFP8Insts, mayRaiseFPException = 0,
 +    SchedRW = [WriteFloatCvt] in {
 +  let Constraints = "$vdst = $vdst_in", DisableEncoding = "$vdst_in" in {
 +    defm V_CVT_PK_FP8_F32 : VOP3Inst<"v_cvt_pk_fp8_f32", VOP3_CVT_PK_F8_F32_Profile>;
 +    defm V_CVT_PK_BF8_F32 : VOP3Inst<"v_cvt_pk_bf8_f32", VOP3_CVT_PK_F8_F32_Profile>;
 +  }
 +
 +  // These instructions have non-standard use of op_sel. In particular they are
 +  // using op_sel bits 2 and 3 while only having two sources. Therefore dummy
 +  // src2 is used to hold the op_sel value.
 +  let Constraints = "$vdst = $src2", DisableEncoding = "$src2" in {
 +    defm V_CVT_SR_FP8_F32 : VOP3Inst<"v_cvt_sr_fp8_f32", VOP3_CVT_SR_F8_F32_Profile>;
 +    defm V_CVT_SR_BF8_F32 : VOP3Inst<"v_cvt_sr_bf8_f32", VOP3_CVT_SR_F8_F32_Profile>;
 +  }

Intel Xe-HPC

intel/intel-graphics-compiler によれば、Intel GPU も Xe-HPC アーキテクチャBF8 に対応している。
Xe-HPC A0 が対応するフォーマットとして QF (Quarter Float) も以下引用部にあるが、他の部分の記述から QF のフォーマット名は BF8 にリネームされたことがわかる。
Xe-HPC のリビジョン更新と共に QFBF8 に変えた理由は明かされていないが、他ベンダーにおける 8-bit 浮動小数点フォーマットと統一するためという理由が想像できる。
2022/07/07 付で、AI 向けの IPU (Intelligence Processing Unit) を設計している Graphcore より、AMD と Qualcomm と共に 8-bit 浮動小数点フォーマットの標準化を提案したというリリースが出されており、このことが関係しているのではないかと思われる。

Xe-HPC では現状 BF8 のみの対応とし、FP8 の記述はない。
変換命令については、AMD GFX940/CDNA 3 では F32 <-> FP8,BF8 のみだったが、Intel Xe-HPC では F16 (HF) <-> BF8 (QF) のみの対応関係となっている。

     GED_DATA_TYPE_nf,      ///< 11, TGL, XE.HP, XE.HPG
     GED_DATA_TYPE_bf,      ///< XE.HP, XE.HPG, XE.HPC.A, XE.HPC
     GED_DATA_TYPE_qf,      ///< XE.HPC.A
     GED_DATA_TYPE_bf8,     ///< XE.HPC
     GED_DATA_TYPE_tf32,    ///< XE.HPC
 // [Deprecated. Use bf8 version instead]
 // Note: qf is renamed to bf8. Those cvt builtins are the same as
 //       those cvt builtins between bf8 and hf (qf is deprecated)
 //       Once all apps/tests move to use bf8 cvt, those can be deleted
 //
 // quarter float <--> half float conversion
 //    qf : no igc type for qf yet. Use char as *opaque* type for it.
 //
 // hf -> qf conversion builtins (rte rounding mode)
 char   __builtin_IB_hftoqf_1 (half   a) __attribute__((const));
 char2  __builtin_IB_hftoqf_2 (half2  a) __attribute__((const));
 char3  __builtin_IB_hftoqf_3 (half3  a) __attribute__((const));
 char4  __builtin_IB_hftoqf_4 (half4  a) __attribute__((const));
 char8  __builtin_IB_hftoqf_8 (half8  a) __attribute__((const));
 char16 __builtin_IB_hftoqf_16(half16 a) __attribute__((const));
	
 // qf -> hf conversion builtins (precise conversion)
 half   __builtin_IB_qftohf_1 (char   a) __attribute__((const));
 half2  __builtin_IB_qftohf_2 (char2  a) __attribute__((const));
 half3  __builtin_IB_qftohf_3 (char3  a) __attribute__((const));
 half4  __builtin_IB_qftohf_4 (char4  a) __attribute__((const));
 half8  __builtin_IB_qftohf_8 (char8  a) __attribute__((const));
 half16 __builtin_IB_qftohf_16(char16 a) __attribute__((const));

参考リンク