先日 AMD の Matt Arsenault 氏 によって公開された LLVM への gfx950 関連のプルリクエストを取り上げた。
その後、さらなるプルリクエストが公開され、gfx950 では FP6/FP4 フォーマットを入力に取る命令がサポートされることが明らかとなった。
先日の記事では gfx950 が Instinct MI325X 関連の GPU ID ではないかと推測したが、これにより Instinct MI350X の GPU ID だと考えられる。
- AMDGPU: Define v_mfma_f32_{16x16x128|32x32x64}_f8f6f4 instructions by arsenm · Pull Request #116723 · llvm/llvm-project
- AMDGPU: Add v_mfma_ld_scale_b32 for gfx950 by arsenm · Pull Request #116722 · llvm/llvm-project
- AMDGPU: Add V_CVT_PK_BF16_F32 for gfx950 by arsenm · Pull Request #116678 · llvm/llvm-project
- AMDGPU: Define v_mfma_f32_32x32x16_bf16 for gfx950 by arsenm · Pull Request #116679 · llvm/llvm-project
- AMDGPU: Handle gfx950 96/128-bit buffer_load_lds by arsenm · Pull Request #116681 · llvm/llvm-project
- AMDGPU: Handle gfx950 global_load_lds_* instructions by arsenm · Pull Request #116680 · llvm/llvm-project
FP6/FP4 {fp6_fp4}
gfx950 では V_MFMA_SCALE_F32_{16X16X128,32X32X64}_F8F6F4
命令と V_MFMA_F32_{16X16X128,32X32X64}_F8F6F4
命令をサポートする。
それらの命令では入力に FP8/FP6/FP4 フォーマットの行列を取り、FP32 フォーマットに出力する。
- AMDGPU: Define v_mfma_f32_{16x16x128|32x32x64}_f8f6f4 instructions by arsenm · Pull Request #116723 · llvm/llvm-project
- AMDGPU: Add v_mfma_ld_scale_b32 for gfx950 by arsenm · Pull Request #116722 · llvm/llvm-project
V_MFMA_SCALE_F32_{16X16X128,32X32X64}_F8F6F4
命令ではスケール係数を指定することができる。
FP8/FP6/FP4 フォーマットはダイナミックレンジが狭いため、オーバーフロー/アンダーフローを防ぎ、精度を維持するのにスケーリング処理が使われる。