AMD、RDNA 4向けFP8 GEMMで小形状はhipBLASLt比+23.2%

AMD ROCm Blogsは、Adel Johar、Jan Stephan、Istvan Kissの3氏によるFP8 GEMM最適化の記事を2026年10月23日に公開しました。Radeon RX 9070 XT（gfx1201）上で4種類のタイル構成をhipBLASLtと比較し、192×192タイルのConfig 4はM=N=K=8192で250.5 TFLOP/s（hipBLASLt比−4.1%）に達しています。

AMD ROCmは、AMDが提供するGPU計算向けのオープンなソフトウェア群です。

FP8で何が変わるか

記事は2部構成の第2部で、第1部でRDNA 3/4向けに作ったFP16 GEMMカーネルの最終版をFP8へ移植しています。TNメモリレイアウト、バッファリソース記述子、LDS上の深いKスライス、キャッシュ再利用型のブロックスケジュール、インターリーブ型プリフェッチfoldはそのまま引き継ぎ、変更点はデータ型と命令です。FP8 WMMAはgfx12でのみ使えるため、対象GPUは1種類です。

• 命令: __builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12 ( a , b , acc )
• オペランド: FP8 E4M3（1要素1バイト）。128ビットロードで16値を運び、elems_per_b128が8から16になります
• フラグメント: レーンあたり8バイト。組み込み関数はFP8ベクタ型ではなく32ビット整数2個を受け取るため、int [[ clang :: ext_vector_type ( 2 )]]で宣言します
• アキュムレータ: FP32のみ（FP16版なし）。出力Cはfloatで、hipBLASLtと同じ列優先レイアウトです
• スケール係数: 適用せず、生のFP8値を扱います

要素幅が半分になるため、同じ64 KBのLDS予算でより大きな出力タイルや深いKスライスを使えます。FP16で80 KBを要した256×256/K=64のタイルは、FP8では40 KBに収まります。

4つのタイル構成

構成	出力タイル	K	wave grid	LDS	備考
Config 1	128×96	64	2×2	17.5 KB	FP16版の型だけ変更
Config 2	256×256	64	4×4	40.0 KB	アキュムレータ16
Config 3	256×160	128	4×2	58.5 KB	アキュムレータ20、235 VGPR
Config 4	192×192	128	2×4	54.0 KB	アキュムレータ18、210 VGPR

Config 4はタイル形状とwave gridを手動でスイープして選んだ構成です。VGPR数はコンパイラの-Rpass-analysis=kernel-resource-usageによる値で、どちらもSIMDあたり4 waveです。

ベンチマーク結果

形状は70B級モデル（Llama-2/3-70B、Qwen2-72B）の隠れ次元K=8192を想定した4種類です。hipBLASLtはHIPBLAS_COMPUTE_32Fで呼び出しています。

形状	hipBLASLt	Config 1	Config 2	Config 3	Config 4
8192×8192×8192	260.8 TFLOP/s	−18.8%	−13.3%	−6.0%	250.5（−4.1%）
8192×28672×8192	258.3	−21.6%	−12.7%	−5.2%	246.8（−4.7%）
8192×10240×8192	260.5	−20.2%	−14.4%	−6.2%	246.6（−5.7%）
3072×4096×8192	177.9	+12.8%	+19.6%	239.2（+25.7%）	231.5（+23.2%）

M=N=K=8192でConfig 4は4.39 ms、hipBLASLtは4.22 msでした。小さい形状ではConfig 3がConfig 4より速く、最適なタイルは形状によって異なります。正確性はAとBの全要素を−1または+1にして、hipBLASLtとビット単位で一致することで確認しています。

試すための要件

• ROCm（計測はROCm 10.0）とhipBLASLt開発パッケージ
• RDNA 4 GPU（検証はRX 9070 XT、gfx1201）
• 付属リポジトリの実装と、第1部の内容の理解

このカーネルはスケール係数を持たないため、記事は本番の量子化推論カーネルを作るための基盤と位置付けています。