FP8 / NVFP4 の weight-only 推論で RTX 3090 が失っているものは無い

公開した量子化モデルに対して、RTX 5090 を 2 枚持つ利用者から「INT4 の量子化を使うと、Blackwell の NVFP4 / FP8 の利点を失うのではないか」という趣旨の質問を受けた。

筆者の回答は「weight-only (A16) の範囲では何も失わない」である。本記事では、その根拠を vLLM のコード、実行バイナリの SASS、RTX 3090 での実測で示す。あわせて、よく見かける「3090 には FP8 / FP4 のテンソルコアが無いので、FP8 / NVFP4 のモデルは遅い」という主張のうち、どこが正しくどこが誤りかを整理する。

結論を先に示す。比較するのは、実在する RTX 3090 と、FP8 / FP4 のテンソルコアを追加した仮想の 3090 である。それ以外の条件 (メモリ帯域、SM 数、クロック) は同じとする。

3090 が失っているのは最後の 1 行だけである。それは「weight-only の推論」ではなく、アクティベーションも量子化する別の量子化方式の話である。

仮想の 3090 は存在しないので、直接は測れない。代わりに次の 3 点を実測で確認した。

• 5090 (sm_120) 用にビルドされた Marlin の A16 カーネル 480 個は、すべて BF16 / FP16 の HMMA だけを発行する。 FP8 / FP4 の MMA 命令も、FP8 / FP4 からの変換命令も含まれていない (2.1 節)。
• decode (M=1) の Marlin は DRAM 帯域の 91〜94% を使い、命令発行枠は 13〜33% しか使っていない。 逆量子化の命令数が違う FP8 と INT8、NVFP4 と INT4 は、同じ実効帯域で動く (4.3 節)。
• prefill (M=8192) の Marlin は、同じクロックで BF16 の cuBLAS の 95〜97% の速度で動く。 逆量子化を含む weight-only のコストは全体で 3〜5% である (4.4 節)。

1. 「FP8 のモデル」「NVFP4 のモデル」は 2 つの意味で使われる

「FP8 / NVFP4 のモデル」という言い方は、次の 2 つを区別していない。

Hugging Face で「FP8」「NVFP4」と名前の付いたチェックポイントの多くは、下の 2 行 (W8A8 / W4A4) として作られている。Blackwell の FP8 / FP4 テンソルコアが速度に効くのは、この 2 行だけである。

3090 で W8A8 / W4A4 のチェックポイントを読むと、vLLM はアクティベーションを量子化せず、weight-only として実行する (3 節)。つまり 3090 で動いている「FP8 のモデル」は、常に W8A16 である。

2. テンソルコアは、2 つの入力が同じ系統の型でないと使えない

PTX の mma 命令が受け付ける入力型の組み合わせは、おおむね次のとおりである。

bf16 × e4m3 や bf16 × e2m1 の組み合わせは存在しない。 アクティベーションが BF16 のままなら、重みを BF16 に変換してから bf16 × bf16 の mma を発行するしかない。これは GPU の世代に関係なく成り立つ。A16 である限り、FP8 / FP4 のテンソルコアは搭載されていても使われない。

vLLM の Marlin カーネルもこのとおりに実装されている。BF16 のアクティベーションに対して発行する命令は次の 1 つである (csrc/libtorch_stable/quantization/marlin/marlin_mma.h:71、vLLM b22afe45a)。

"mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32 "

同じファイルには f32.e4m3.e4m3.f32 の mma もある (79 行・97 行)。これはアクティベーションの型 (a_type) が FP8 のとき、つまり W4A8-FP8 のときだけ使われる。sm_89 未満ではコンパイル時に除外される (marlin_template.h:283)。

2.1 実行バイナリの SASS

ソースだけでなく、実際に配布されているバイナリも確認した。対象は vLLM 0.29.1rc1.dev47+gdc36fcce9 のイメージに含まれる _C_stable_libtorch.abi3.so である。この中の Marlin の全カーネルを cuobjdump -sass で逆アセンブルし、MMA 命令の種類を数えた。3090 (sm_86) は sm_80 のバイナリを実行する。

• 5090 用の A16 カーネルも、3090 用と同じ HMMA (BF16 / FP16 のテンソルコア) しか発行しない。
• FP8 のテンソルコア命令 QMMA が現れるのは、アクティベーションが FP8 のカーネルだけである。
• A16 のカーネルに含まれる変換命令は F2FP.BF16.PACK_AB (FP32 の累算結果を BF16 の出力に詰める命令) だけで、FP8 / FP4 の重みを変換する命令は無い。

同じ設定 (256 スレッド、W8A16 FP8) の 1 カーネルで比べると、命令の種類は同じで、数の差はコンパイラのスケジューリングの差である。

NVIDIA 自身の実装も同じである。vLLM が使う FlashInfer (0.6.18.post1) には、Blackwell 向けの BF16 × NVFP4 のカーネルがある。SM100 版 (dense_gemm_bf16_fp4_sm100.py) の説明には次のように書かれている。

The kernel decodes each 16-value E2M1 weight block and its E4M3 scale to BF16, then executes BF16 tcgen05 MMA with FP32 accumulation.

GeForce Blackwell (SM12x) 版 (dense_gemm_bf16_fp4_sm12x.py) は、cute.nvgpu.warp.MmaF16BF16Op (F16 / BF16 の MMA) を使っている。Blackwell 専用に書かれた W4A16 カーネルでも、FP4 のテンソルコアは使われていない。

3. A16 では、vLLM は 5090 でも 3090 と同じカーネルを選ぶ

以下は計測に使った vLLM 0.29.1rc1.dev47+gdc36fcce9 のコードである。

NVFP4 の weight-only (NVFP4A16)。 SM100 / SM103 (B200 など) では FlashInfer の CuTe-DSL カーネルを、それ以外の GPU では Marlin を使う (vllm/model_executor/kernels/linear/__init__.py:1081-1092)。5090 (SM120) は Marlin である。

elif linear_backend == "auto" and use_a16:
    ...
    # Weight-only: prefer FlashInfer CuTe-DSL W4A16 on SM100/103,
    # otherwise Marlin.
    ...
    if compute_capability in (100, 103) and cutedsl_ok:
        force_kernel = FlashInferCuteDslNvFp4W4A16LinearKernel
    else:
        force_kernel = MarlinNvFp4LinearKernel

FlashInfer のカーネルが BF16 の MMA を使うことは 2.1 節で確認した。

FP8 の weight-only (W8A16)。 候補は HummingFP8ScaledMMLinearKernel と MarlinFP8ScaledMMLinearKernel の 2 つだけである (__init__.py:475-479)。どちらも SM75 以降で動作する。Humming の内部は本記事では読んでいない。ただし 2 節の命令の制約により、A16 である以上は BF16 / FP16 の mma を使うしかない。

MoE。 fused MoE の Marlin カーネル (csrc/libtorch_stable/moe/marlin_moe_wna16/marlin_template.h) は、上記と同じ dequant.h と marlin_mma.h を include している。

したがって A16 のチェックポイントでは、3090 と 5090 は同じカーネルのコードを実行する。速度の差を生むのは、メモリ帯域・SM 数・クロック・キャッシュ容量など GPU 全般の性能であり、FP8 / FP4 のテンソルコアの有無ではない。

3090 で W8A8 / W4A4 のチェックポイント (compressed-tensors 形式) を読んだ場合は、次のようになる。

• W8A8 の FP8: カーネルを選ぶ前の段階で振り分けられる。W8A8 のスキームは compute capability 8.9 以上を要求する (compressed_tensors_w8a8_fp8.py:83-85)。そのため 3090 では W8A16 のスキームになり (compressed_tensors.py:840-857)、W8A16 と同じ Humming か Marlin で実行される。
• W4A4 の NVFP4: FP4 のカーネルの候補 (__init__.py:544-) を 3090 の上で順に判定させると、FP4 のテンソルコアを使う候補はすべて非対応を返し、最初に通るのは Marlin だった。

どちらも、アクティベーションは量子化されずに BF16 のまま計算される。

4. 逆量子化のコストの実測

A16 の GEMM では、重みを BF16 / FP16 に変換する処理 (逆量子化) が毎回入る。仮想の 3090 との差が生じうるのはここだけなので、そのコストを実測した。

4.1 Marlin の実装

Marlin は重みを VRAM 上で BF16 に展開しない。GEMM のたびに重みのタイルを共有メモリからレジスタに読み、レジスタ上で変換して mma に渡す。起動時の処理は並べ替え (repack) だけである。

FP8 (E4M3) を FP16 に変換するコードは次のとおりである (dequant.h:321-336、vLLM b22afe45a)。

template <>
__device__ inline void dequant<half2, vllm::kFE4M3fn.id(), true>(
    int q, half2* frag_b) {
  constexpr int FP8_EXPONENT = 4, FP16_EXPONENT = 5;
  constexpr int RIGHT_SHIFT = FP16_EXPONENT - FP8_EXPONENT;
  constexpr int MASK = 0x7F007F00;

  int Out1 = (q & 0x80008000) | ((q & MASK) >> RIGHT_SHIFT);
  q <<= 8;
  int Out2 = (q & 0x80008000) | ((q & MASK) >> RIGHT_SHIFT);

  frag_b[1] = *reinterpret_cast<const half2*>(&Out1);
  frag_b[0] = *reinterpret_cast<const half2*>(&Out2);
}

• 32 bit の整数 1 つに入った 4 要素を、and・シフト・or だけで FP16 のビット列にする。
• 指数のバイアスの差 (FP16 なら 2^8、BF16 なら 2^120) は、ロード時にスケールへ掛けておく (marlin_utils_fp8.py:35)。そのため変換のたびに乗算する必要は無い。
• E4M3 と E2M1 の値は、すべて FP16 / BF16 で正確に表せる。変換で丸めは起きない。
• dequant.h 内のアーキテクチャ分岐は __CUDA_ARCH__ >= 750 の 1 か所だけである (70 行)。sm_89 以降にある FP8 → FP16 のハードウェア変換命令は使われていない。これは 2.1 節の SASS とも一致する。

4.2 計測条件

NVFP4 の重みは vLLM のテスト用関数で作った乱数である。本記事は速度だけを扱い、精度は扱わない。FP8 の W8A16 は、vLLM では Humming が第一候補である (3 節)。本記事で測ったのは Marlin だけである。

4.3 decode (M が小さい領域)

M を振ったときの GEMM 1 回の時間 (μs) は次のとおりである。

• 同じビット幅の FP と INT (FP8 と INT8、NVFP4 と INT4) は、M=1 の実効帯域が 1.5% 以内で同じである。2 つは逆量子化の命令がまったく違う (次の表) ので、その違いは decode の速度に現れていない。
• Marlin は BF16 の cuBLAS よりも高い実効帯域で重みを読んでいる。

Nsight Compute で M=1 の各カーネルを測った結果は次のとおりである。

命令数はスレッド単位で、重みの要素数 (K × N) で割った値である。

• decode の時間を決めているのは DRAM 帯域である。 Marlin は帯域の 91〜94% を使っている。一方で命令発行枠は、最も命令の多い NVFP4 でも 32.5% しか使っていない。
• 逆量子化は要素あたり 1.8〜2.9 命令である。 FP8 は 1.81 命令で、4.1 節のソースから数えた値 (4 要素に 7 命令) とほぼ一致する。
• 仮想の 3090 がハードウェアの変換命令で逆量子化の命令を減らしたとしても、空いている発行枠が増えるだけである。時間を決めている DRAM 帯域は変わらない。

SM クロックを固定して、演算側の能力を落としたときの M=1 の実効帯域 (GB/s) は次のとおりである。

「固定なし」でクロックが形式ごとに違うのは、電力上限 280 W に達したためである。

• 1,200 MHz 以上では、Marlin の 4 形式の実効帯域は 1% 以内で変わらない。
• 900 MHz では 4bit の 2 形式が 15〜16% 落ちる。ただし逆量子化をしない cuBLAS の BF16 は、1,200 MHz で 16%、900 MHz で 36% 落ちる。クロックへの感度は逆量子化の有無では決まっていない。
• 900 MHz でも、NVFP4 と INT4 の差は 1.6%、FP8 と INT8 の差は 1.3% である。

4.4 prefill (M が大きい領域)

Nsight Compute で M=8192 を測った結果は次のとおりである。ncu がクロックを固定するため、全形式が 1.69 GHz で動いている。

• Marlin は、同じクロックで cuBLAS の 95〜97% の速度で動く。この 3〜5% には逆量子化のほか、Marlin と cuBLAS のタイル構成の違いなどもすべて含まれる。逆量子化だけのコストはこれより小さい。
• テンソルコアのパイプ使用率は、Marlin も cuBLAS とほぼ同じ水準である。

電力上限 280 W のままクロックを固定せずに測った場合は、次のとおりである (M=8192、2 回の計測)。

• cuBLAS の 41.3〜41.9 TFLOPS/GHz は、3090 の公称値 (FP32 累算の BF16 で 71 TFLOPS、1,695 MHz のとき。41.9 TFLOPS/GHz) の 99% である。
• 電力上限があると、Marlin (特に 8bit) はクロックが cuBLAS より 1 割前後下がる。そのため素の TFLOPS では 6〜15% の差になる。これは逆量子化とメモリ読み出しが電力を使うためで、電力上限のある GPU では実際のコストである。ただし A16 である以上、FP8 / FP4 のテンソルコアの有無でこの電力は変わらない。

5. 3090 が実際に失っているもの

失っているのは、W8A8 / W4A4 を選ぶ権利 である。

• prefill と大きいバッチ: GEMM が演算で上限に達する領域では、FP8 のテンソルコアは同じ GPU の BF16 に対して演算上限が一般に 2 倍、FP4 で 4 倍である。4.4 節のとおり、A16 の Marlin は BF16 のテンソルコアの上限の近くまで出ている。これより速くするには、アクティベーションも量子化するしかない。
• decode (バッチ小): メモリ帯域が上限なので、A8 / A4 にしても速くならない。アクティベーションを量子化する処理が増える分、遅くなる場合もある。vLLM の Marlin は W4A8-FP8 を SM89 と SM12x にだけ許可しており、エラーメッセージに理由として「他の GPU では Marlin W4A16 より遅い」と書いている (marlin.cu:419、vLLM b22afe45a)。
• 品質: W8A8 / W4A4 では、重みの誤差にアクティベーションの丸め誤差が加わる。本記事では W4A4 の KLD を測っていない。

したがって「Blackwell なら NVFP4 で速い」は、同じモデルを別の GPU で動かした比較ではない。prefill の速度とアクティベーションの誤差を交換する、別の量子化方式を選んだ結果である。 weight-only のチェックポイントを使う限り、この選択肢は Blackwell でも使われていない。

6. 想定される反論と回答

6.1 「vLLM 自身が、3090 では性能が落ちると警告している」

3090 で FP8 / NVFP4 のチェックポイントを読み、Marlin が選ばれると、vLLM は次の警告を出す (kernels/linear/nvfp4/marlin.py:34-39、FP8 は marlin_utils_fp8.py:110-115)。

Your GPU does not have native support for FP4 computation but FP4 quantization is being used. Weight-only FP4 compression will be used leveraging the Marlin kernel. This may degrade performance for compute-heavy workloads.

この警告の内容は正しい。ただし、比べている相手は W4A4 を FP4 のテンソルコアで実行する場合であり、対象は “compute-heavy workloads” (prefill) に限られている。5 節と同じことを述べているだけで、「weight-only が 3090 で遅い」という主張の根拠にはならない。

なお、この警告は MarlinNvFp4LinearKernel.process_weights_after_loading の中で無条件に出力される。3 節のとおり、NVFP4A16 は SM100 / SM103 以外のすべての GPU で Marlin を使う。したがってコード上は、5090 で NVFP4A16 のチェックポイントを読んだ場合も「Your GPU does not have native support for FP4 computation」と表示される。この警告は、GPU の能力ではなく選ばれたカーネルに対して出るものである。

6.2 「Blackwell には FP8 / FP4 → BF16 の変換命令があるので、逆量子化が速い」

4.1 節と 2.1 節のとおり、Marlin はその命令を使っていない。使ったとしても、減らせるのは要素あたり 1.8〜2.9 命令の逆量子化の一部だけである。decode では命令発行枠の 67% 以上が空いており、時間を決めているのは DRAM 帯域である (4.3 節)。prefill では、逆量子化を含む weight-only のコスト全体が 3〜5% である (4.4 節)。

6.3 「NVFP4 は Blackwell で BF16 の 4 倍速い」

これは W4A4 の GEMM の演算上限の話である。weight-only のチェックポイントは、どの GPU でもこの経路を通らない (2 節)。NVIDIA が Blackwell 専用に書いた FlashInfer の W4A16 カーネルも、BF16 の MMA を使っている (2.1 節)。

6.4 「実際に 5090 で NVFP4 のモデルを動かすと、3090 よりずっと速い」

この比較では、2 つの要因が混ざっている。

1. GPU 全般の性能。メモリ帯域だけでも 5090 は 1,792 GB/s、3090 は 936 GB/s である (公称値)。
2. 量子化方式。5090 は W4A4、3090 は W4A16 として実行している。

FP8 / FP4 のテンソルコアの効果を取り出すには、同じ GPU で W4A16 と W4A4 を比べる必要がある。GPU をまたいで比べるなら、同じ A16 のチェックポイントを両方で動かす必要がある。

6.5 「3090 の FP8 はエミュレーションなので精度が落ちる」

4.1 節のとおり、FP8 / FP4 から FP16 / BF16 への変換は正確で、丸めは起きない。3090 の Marlin は、5090 が同じ A16 のチェックポイントに対して実行するのと同じカーネル (2.1 節の sm_120 版) である。精度の点で不利なのは、むしろアクティベーションも丸める W8A8 / W4A4 の方である。

6.6 「W4A4 の品質の劣化は小さい。Blackwell なら A4 を使うべきだ」

本記事の主張の範囲外である。W4A4 を選ぶかどうかは、prefill の速度とアクティベーションの誤差の交換であり、KLD を同じ測定方法で測って判断するものである。筆者は本記事でその比較をしていない。本記事が主張しているのは、weight-only を選んだ時点で、Blackwell の FP8 / FP4 テンソルコアはどの GPU でも使われていない、という点だけである。

6.7 「クロックを下げると 4bit の decode が遅くなる。逆量子化が効いている証拠だ」

4.3 節のとおり、900 MHz まで下げると 4bit の Marlin は 15〜16% 遅くなる。しかし逆量子化をしない cuBLAS の BF16 は、同じ条件で 36% 遅くなる。1,200 MHz では Marlin は変わらず、cuBLAS は 16% 遅い。低クロックで遅くなる度合いは、逆量子化の有無では説明できない。また、同じクロックで FP と INT (逆量子化の命令が違う) を比べると、差はどのクロックでも 1.6% 以内である。

6.8 「測ったのは 3090 だけで、5090 では違うかもしれない」

5090 での計測はしていない。確認できているのは次の 2 点である。

• 5090 が実行する sm_120 のバイナリは、3090 の sm_80 と同じ種類の命令 (HMMA と整数演算) だけで構成されている (2.1 節)。
• 5090 の公称値 (170 SM、ブーストクロック 2.41 GHz、1,792 GB/s) で計算すると、命令発行枠の上限は 3090 の約 2.9 倍、メモリ帯域は約 1.9 倍である。4.3 節の NVFP4 の命令数 (要素あたり 3.66) で帯域を使い切った場合、発行枠の使用率は約 22% になる。同じ式に 3090 の実測の帯域使用率とクロックを入れると 32% になり、ncu の実測値 32.5% と一致する。5090 の方が、演算側の余裕は大きい。

まとめ

• FP8 / FP4 のテンソルコアの命令は、2 つの入力がどちらも FP8 / FP4 系の型であることを要求する。BF16 × FP8 の命令は存在しない。5090 用の Marlin の A16 カーネル 480 個も、HMMA (BF16 / FP16) だけを発行する。NVIDIA が Blackwell 用に書いた FlashInfer の W4A16 カーネルも、BF16 の MMA を使う。
• vLLM は A16 の場合、3090 と 5090 で同じカーネル (NVFP4A16 は Marlin、FP8 は Humming か Marlin) を選ぶ。
• 3090 の実測では、decode の Marlin は DRAM 帯域の 91〜94% を使い、命令発行枠は 13〜33% しか使っていない。逆量子化の命令数が違う FP と INT は、同じ実効帯域で動く。prefill では、同じクロックで cuBLAS の BF16 の 95〜97% の速度で動く。
• 3090 が失っているのは W8A8 / W4A4 を選ぶ権利、つまり prefill の演算上限 (FP8 で 2 倍、FP4 で 4 倍) である。それはアクティベーションの誤差と交換する別の量子化方式であり、weight-only の推論の話ではない。
• vLLM の「This may degrade performance」の警告は、W4A4 と比べた prefill の話である。コード上は、NVFP4A16 なら 5090 でも同じ警告が出る。