Skip to content
Published on

推論を速く — 量子化、スパース性、Dataflow のハードウェア視点

🇰🇷KO🇺🇸EN🇯🇵JA
Split View필사 모드
Authors

はじめに

学習が終わったモデルは一度だけ作られますが、推論はモデルが生きている限り毎日何十億回も起こります。2026 年に入り、業界で推論 capex が学習 capex を初めて上回ったという報告が出る理由がここにあります。モデルをより大きく賢くする仕事と、そのモデルを安く速く提供する仕事は、いまや完全に別のエンジニアリング課題に分かれました。

この記事は「推論を速くする」という目標をハードウェアの観点から分解します。中心となる主張は単純です。現代の推論は、多くの場合、演算が足りなくて遅いのではなく、データを運ぶのに時間を使って遅いのです。この一つの事実から、データを小さくする量子化、データを少なくするスパース性、データを動かさない dataflow 設計という三つの戦略がかみ合います。そこに演算子融合、バッチング、KV キャッシュ、コンパイラ最適化が加われば、同じチップで数倍のスループット差が生まれます。

落ち着いた、しかし数字とコードで裏打ちされた語り口で進めます。

推論コストの構造 — なぜメモリがボトルネックか

まずコスト構造を正確に見ましょう。GPU の性能を分ける二つの数字は、演算スループット(FLOPS)とメモリ帯域幅(byte/s)です。この二つの比を演算強度(arithmetic intensity)と呼び、「メモリから 1 バイト読むたびに何回演算するか」を表します。

ルーフライン(roofline)モデルがこの関係を一目で示します。

性能
(FLOPS)
  ^
  |                    ___________________  演算上限 (compute-bound)
  |                   /
  |                  /
  |                 /  <- この傾きがメモリ帯域幅
  |                /
  |               /     メモリ上限 (memory-bound)
  |              /
  +-------------+--------------------------> 演算強度 (FLOP/byte)
              臨界点

演算強度が低い(左)と帯域幅が性能を決め、高い(右)と演算が性能を決めます。問題は、LLM 推論のデコード段が極端にメモリバウンドであることです。

LLM 推論は二段階に分かれます。

  • プリフィル(prefill): 入力プロンプト全体を一度に処理。行列-行列積が多く演算強度が高く、演算バウンドに近い。
  • デコード(decode): トークンを一つずつ生成。毎ステップでモデル重み全体を読み、行列-ベクトル積を行う。バッチが小さいと演算強度が 1 付近に落ち、徹底的にメモリバウンドになる。

数字で感じましょう。700 億パラメータのモデルを FP16 で保存すると約 140GB です。トークンを一つデコードするにはこの重みを一度読む必要があります。メモリ帯域幅が約 3.35TB/s なら、重みを一度なめるだけで約 42 ミリ秒かかります。つまりバッチ 1 では毎秒約 24 トークンが理論上限です。演算ユニットはほとんど遊んでいます。

ここでメモリウォール(memory wall)という言葉が登場します。過去十数年、演算性能はメモリ帯域幅よりはるかに速く伸びました。その結果、多くの推論ワークロードでチップの演算能力は余り、データを供給する能力が不足します。データを動かすエネルギーが演算自体より一桁以上大きい点も重要です。チップ内部のレジスタアクセスに対し、HBM からデータを取ってくるコストはおよそ数百倍に達します。

この一つの事実から、この記事の残りの戦略がすべて導かれます。データを小さく(量子化)、少なく(スパース性)、動かさない(dataflow)ようにすればよいのです。

量子化 — データを小さくする

量子化は重みと活性値をより少ないビットで表す手法です。FP16(16 ビット)を INT8(8 ビット)に変えると、メモリ使用量と帯域幅要求が半分になります。メモリバウンドのワークロードでは、これがほぼそのまま 2 倍の速度につながります。

精度フォーマットの全景

2026 年現在、実務で出会う主なフォーマットを整理します。

フォーマットビット主用途ハードウェア対応
FP16 / BF1616学習、高精度推論広範
FP8 (E4M3 / E5M2)8推論、一部学習Hopper 以降、TPU
INT88推論(PTQ/QAT)非常に広範
FP4 (MXFP4 など)4推論Blackwell 2 世代 TE
INT44重み専用量子化広範(カーネル依存)

NVIDIA Blackwell(GTC 2026 発表)の第 2 世代 Transformer Engine は FP4 演算をネイティブに対応します。FP4 は 4 ビットで浮動小数を表しますが、単独では表現範囲が狭すぎて精度が崩れます。そこでマイクロスケーリング(microscaling)を使います。小さなブロック(例えば 16 や 32 個の値)ごとに共有スケール係数を別に持ち、ブロック単位でダイナミックレンジを合わせる方式です。

対称/非対称、PTQ/QAT

量子化の基本式は次のとおりです。

量子化:   q = round(x / scale) + zero_point
逆量子化: x_hat = (q - zero_point) * scale

対称量子化は zero_point を 0 に固定します。非対称は分布が片寄った活性値(例えば ReLU 後の常に正の値)に有利です。

適用時点で二方式に分かれます。

  • PTQ(Post-Training Quantization): 学習済みモデルをそのまま量子化。少量の校正(calibration)データでスケールだけ合わせる。速く簡単だが 4 ビット以下では精度が落ちうる。
  • QAT(Quantization-Aware Training): 学習過程で量子化誤差を模倣しモデルを適応させる。コストは大きいが低ビットで精度が保たれる。

簡単な PTQ 校正の例

import torch

def collect_quant_scale(weight: torch.Tensor, n_bits: int = 8):
    """対称量子化のスケールをチャネルごとに計算する。"""
    qmax = 2 ** (n_bits - 1) - 1  # int8 なら 127
    # 出力チャネル(行)ごとの最大値
    per_channel_max = weight.abs().amax(dim=1, keepdim=True)
    scale = per_channel_max / qmax
    scale = scale.clamp(min=1e-8)
    return scale

def fake_quantize(weight: torch.Tensor, scale: torch.Tensor, n_bits: int = 8):
    qmax = 2 ** (n_bits - 1) - 1
    qmin = -qmax - 1
    q = torch.clamp(torch.round(weight / scale), qmin, qmax)
    return q * scale  # 逆量子化で誤差を見る

w = torch.randn(4096, 4096)
scale = collect_quant_scale(w, n_bits=8)
w_hat = fake_quantize(w, scale, n_bits=8)
err = (w - w_hat).abs().mean().item()
print(f"平均絶対誤差: {err:.6f}")

チャネルごと(per-channel)のスケールは、テンソル全体に一つのスケールを使うより誤差が小さくなります。あるチャネルの外れ値が他のチャネルまで壊さないからです。実務では SmoothQuant のように活性値の外れ値を重み側へ移して量子化の難度を均す手法、または AWQ のように重要なチャネルだけ精度を残す手法が広く使われます。

何を量子化するか

  • 重み専用(weight-only): 重みだけ INT4/INT8 に、演算は FP16 で。メモリバウンドのデコードに有効。重みを読む帯域幅が減るからです。
  • 重み+活性(W8A8 など): 両方を量子化し演算自体を整数/低精度テンソルコアで回す。プリフィルのような演算バウンド区間で追加の利得がある。

スパース性 — データを少なくする

量子化が各値を小さくするなら、スパース性は 0 の値を丸ごと飛ばし、値の個数自体を減らします。

問題は、非構造(unstructured)スパース性がハードウェアに優しくない点です。重みの 90% が 0 でも、0 の位置がばらばらだと、GPU の整列されたメモリアクセスとテンソルコアの密な演算構造を活かしにくい。インデックスを追うオーバーヘッドが節約分を食います。

そこで登場するのが構造化スパース性(structured sparsity)です。NVIDIA Ampere 以降のテンソルコアが対応する 2:4 スパース性が代表例です。規則は単純で、連続する 4 個の値のうち正確に 2 個を 0 にします。

密:   [ 0.8,  0.1, -0.3,  0.5 ]
2:4:  [ 0.8,  0.0,  0.0,  0.5 ]   <- 4 個中 2 個だけ非ゼロ

保存: 非ゼロ値 2 個 + 2 ビットのインデックス(どの位置だったか)

このパターンをハードウェアは好みます。非ゼロの個数が固定(4 分の 2)なのでメモリアクセスが規則的で、テンソルコアがインデックスに沿って対を選ぶ回路を内蔵し、乗算を半分だけ行います。理想的には行列積のスループットが約 2 倍になります。

# 2:4 スパースマスクを作る概念コード
import torch

def make_2to4_mask(weight: torch.Tensor):
    """行方向に連続 4 個ごと、絶対値の小さい 2 個を 0 にする。"""
    out, in_dim = weight.shape
    assert in_dim % 4 == 0
    w = weight.view(out, in_dim // 4, 4)
    # 各グループで絶対値上位 2 個の位置だけ残す
    idx = w.abs().argsort(dim=-1, descending=True)
    mask = torch.zeros_like(w)
    keep = idx[..., :2]
    mask.scatter_(-1, keep, 1.0)
    return (w * mask).view(out, in_dim)

w = torch.randn(1024, 1024)
w_sparse = make_2to4_mask(w)
print("0 の割合:", (w_sparse == 0).float().mean().item())  # 約 0.5

実務では単に小さい値を 0 にすると精度が落ちるため、スパース化後に微調整(fine-tuning)で回復させる手順を踏みます。量子化と組み合わせると(例えば 2:4 + INT8)メモリと演算を同時に減らせますが、二つの誤差源が重なるので精度検証をより慎重に行う必要があります。

活性値のスパース性もあります。ReLU 系で多くの活性値が 0 になる点、MoE(Mixture of Experts)でトークンごとに一部の専門家だけが活性化する点は、いずれもスパース性の変形です。MoE は実質的に動的な構造化スパース性で、総パラメータは大きくともトークンあたりの実演算は一部だけ起こります。

Dataflow アーキテクチャ — データを動かさない

第三の軸は、チップ内でデータが流れる方式です。鍵となる洞察は、一度取ってきたデータを最大限再利用せよ、というものです。HBM からデータを一度読むコストがチップ内部の演算よりずっと高いため、データをチップ内に留めて何度も使う設計が決定的です。

シストリックアレイとデータ再利用

Google TPU が代表的に採用したシストリックアレイ(systolic array)は、この哲学の結晶です。乗算-累算器(MAC)を格子に並べ、データを一方から押し込むとセルからセルへ流れ、演算がパイプラインのように進みます。一度ロードした重みが格子内に留まり複数の入力に再利用されるので、メモリアクセスが劇的に減ります。

シストリックアレイ (重み固定型)

  入力 →  [w][w][w][w]
  入力 →  [w][w][w][w]   各 [w] は乗算-累算セル
  入力 →  [w][w][w][w]   重みはセルに留まり
  入力 →  [w][w][w][w]   入力/部分和だけが流れる
                ↓
              出力(部分和の累積)

dataflow 設計は、何をチップに固定するかで分類されます。

  • 重み固定(weight stationary): 重みをセルに置き入力を流す。重み再利用が大きい。
  • 出力固定(output stationary): 部分和をセルに累積。出力再利用が大きい。
  • 行固定(row stationary): 入力行を再利用。省エネを狙った折衷型(Eyeriss などの研究)。

タイリング — メモリ階層を意識した分割

GPU でも同じ原理がタイリング(tiling)で実装されます。大きな行列積を小さなブロックに分け、一つのブロックを速いオンチップメモリ(共有メモリ/レジスタ)に載せてそこで可能な限り演算を終えます。遅い HBM 往復を減らすのが目的です。

C = A x B をタイルに分割

A          B          C
[A00 A01]  [B00 B01]  [C00 C01]
[A10 A11]  [B10 B11]  [C10 C11]

C00 = A00*B00 + A01*B10   <- タイル単位で累積
1 タイルは共有メモリに載せて再利用

タイルサイズの選択はトレードオフです。タイルが大きいと再利用は増えるがオンチップメモリに入らず、小さいと入るが HBM 往復が増えます。この最適点を自動探索するのが、後で出てくるコンパイラの役割です。

演算子融合 — 中間結果をメモリに書かない

ディープラーニングのグラフは小さな演算が長く連なります。行列積 → バイアス加算 → 活性化関数 → 正規化、という具合です。各演算を別々に実行すると、毎段で中間テンソルを HBM に書き、次段で再び読みます。メモリバウンドの観点ではこれは無駄の連続です。

演算子融合(operator fusion)はこの連続した演算を一つのカーネルにまとめ、中間結果をオンチップのレジスタ/共有メモリに置いたまま続けて処理します。HBM 往復が消えます。

融合前: MatMul -> [HBM] -> Bias -> [HBM] -> GELU -> [HBM]
融合後: MatMul + Bias + GELU  (一つのカーネル、中間値はチップ内)

最も有名な例が FlashAttention です。アテンションの巨大な中間行列(スコア行列)を丸ごと HBM に作らず、ブロック単位で流しながらソフトマックスをオンラインで累積します。メモリ使用量を減らすと同時に帯域幅ボトルネックを解き、長コンテキスト推論を実用化しました。融合はコンパイラが自動で見つけることもあり、ホットパスは人が直接カーネルを書くこともあります。

バッチングと KV キャッシュ — スループットと遅延の綱引き

デコードがメモリバウンドである核心の理由は、バッチが 1 のとき重みを一度読んでトークンを一つだけ作るからです。ならば複数のリクエストをまとめて(batch)重み一度の読み込みで複数トークンを作れば、演算強度が上がりチップがきちんと働きます。

問題は、実際のサービングではリクエストがそれぞれ別の時点で到着し長さも異なる点です。これを解決するのが**連続バッチング(continuous batching)**です。固定バッチを待たず、トークン生成ステップごとに終わったリクエストを抜き新しいリクエストを差し込んで GPU を休ませません。

KV キャッシュ

トランスフォーマーは前のトークンの key/value を再計算しないよう KV キャッシュに保存します。このキャッシュはシーケンスが長く、バッチが大きいほど線形に増え、急速にメモリを侵食します。

KV キャッシュサイズ(バイト)
 = 2(K,V) x layers x heads x head_dim
   x seq_len x batch x bytes_per_elem

13B 級モデルでコンテキストが長くなると KV キャッシュが数十 GB に達することもあります。そこで次の手法が使われます。

  • PagedAttention: KV キャッシュを OS のページングのように小ブロックで管理し断片化を減らす(vLLM の核)。
  • KV キャッシュ量子化: キャッシュを INT8/FP8 で保存し容量と帯域幅を半分に。
  • GQA/MQA: 複数のアテンションヘッドが key/value を共有しキャッシュサイズを縮める。

バッチングと KV キャッシュ管理は、結局スループット(throughput)と遅延(latency)の綱引きです。バッチを大きくするとスループットは上がりますが個々のリクエストの応答が遅くなりえます。チャットボットのように初回トークン遅延(TTFT)が重要なワークロードと、大量バッチ処理のように総スループットが重要なワークロードは、別の点で均衡を取る必要があります。

コンパイラ最適化 — 人の代わりに探索する

これまで見たタイルサイズ選択、融合判断、メモリ配置は、手で一つずつチューニングすると切りがありません。そこでディープラーニングコンパイラがこの探索を自動化します。

代表的に次の層があります。

  • グラフレベル IR: 演算グラフを受け取り融合、定数畳み込み、レイアウト変換などを適用(例えば TVM、XLA、TorchInductor のグラフ段)。
  • タイル/スケジュール探索: 演算をハードウェアに合わせてタイリングしループ順序を決める段。Triton が代表で、Python に近い文法で GPU カーネルを書くとコンパイラがメモリ配置と並列化を処理します。
  • 自動チューニング: 複数のタイルサイズ/スケジュール候補を実際に回し最速を選ぶ。
# Triton スタイルの擬似カーネル(概念説明用)
# @triton.jit デコレータが付くと GPU カーネルにコンパイルされると仮定
def fused_add_relu(x_ptr, y_ptr, out_ptr, n, BLOCK: int):
    pid = program_id(0)              # ブロックインデックス
    offs = pid * BLOCK + arange(0, BLOCK)
    mask = offs < n
    x = load(x_ptr + offs, mask=mask)
    y = load(y_ptr + offs, mask=mask)
    z = x + y
    z = where(z > 0, z, 0.0)         # ReLU を同じカーネルで融合
    store(out_ptr + offs, z, mask=mask)

核心は「何を計算するか」(数学)と「どう計算するか」(スケジュール)を分けることです。同じ行列積でもチップごと、形状ごとに最適スケジュールは異なります。コンパイラはこの巨大な探索空間を人より速くなめます。

精度と効率のトレードオフ — ただ飯はない

これまでの手法はすべて精度を対価に効率を買います。重要なのはその為替レートを測ることです。

手法得るもの失いうるもの緩和策
INT8 量子化帯域幅/メモリ半分わずかな精度低下チャネルごとスケール、校正
FP4/INT4帯域幅 4 分の 1低ビットほど大きな低下QAT、マイクロスケーリング、混合精度
2:4 スパース性演算約 2 倍表現力の損失スパース化後に微調整
KV キャッシュ量子化キャッシュ半分長コンテキストで累積誤差FP8 使用、敏感な層を保存

実務原則をいくつか。

  • 一度に一つの手法だけ適用し精度を測ります。複数を同時に入れるとどこで崩れたか分かりません。
  • 平均だけでなく裾(tail)の事例を見ます。量子化は平均精度は無事でも特定の入力でだけ崩れることがよくあります。
  • すべての層を同じ精度で扱いません。アテンション出力や最初/最後の層のような敏感な部分は精度を高く保ちます(混合精度)。

ハードウェアとソフトウェアの協調設計 — 大きな絵

ここまで見ると一つ明らかです。量子化フォーマット、スパース性パターン、dataflow、コンパイラは独立した技術ではなく、ともに設計されます。

  • Blackwell の第 2 世代 Transformer Engine が FP4 をネイティブ対応するから、ソフトウェアがマイクロスケーリング量子化を採る価値が生まれます。
  • テンソルコアが 2:4 パターンを回路で理解するから、フレームワークがそのパターンに合わせてスパース化します。
  • HBM 帯域幅と NVLink/UALink インターコネクトの限界が、タイルサイズとバッチング戦略を決めます。

2026 年の流れもこの協調設計の延長です。NVIDIA は次世代 Vera Rubin で HBM4 とともに perf/watt を約 10 倍に引き上げる目標を掲げ、Google は TPU v6 Trillium(前世代比 peak 約 4.7 倍)と推論特化の第 7 世代 Ironwood で推論効率を狙い撃ちました。クラウド事業者の自社推論 ASIC の比重が急速に増え(2024 年の約 15% から 2026 年に 40% 見込み)、推論ワークロードにチップを合わせて設計する流れが加速します。NVIDIA は依然として加速器市場の約 75〜80% を握りますが、AMD MI350X をはじめとする競争と自社 ASIC の台頭で推論市場の地形が多層化しています。

要点は、推論最適化を一つの手法と見ず、ハードウェア世代とともに動くスタックとして見るべきだということです。

実務適用 — どこから手を付けるか

新しいモデルをサービングする前提で順序を提案します。

  1. まず測定。 ワークロードがプリフィル中心かデコード中心か、メモリバウンドか演算バウンドかをプロファイリングします。ボトルネックを知らずに最適化しても無駄です。
  2. 検証済みのサービングエンジンを先に。 vLLM、TensorRT-LLM のようなエンジンは連続バッチング、PagedAttention、融合カーネルを既に備えています。一から作る必要はほとんどありません。
  3. 重み量子化から。 デコードがメモリバウンドなら INT8/INT4 の重み量子化が最も大きな即効を生みます。精度を測り合格線を確認します。
  4. KV キャッシュを管理。 長コンテキストなら PagedAttention と KV キャッシュ量子化でメモリを確保します。
  5. その次にスパース性/低ビット。 ここまでで足りなければ 2:4 スパース性、FP4 のようなより攻めた手法へ進みます。必ず微調整と検証を伴います。
  6. チップ世代を活用。 最新チップの FP4/FP8 テンソルコアを使うにはそれに合うコンパイラ/ランタイムのバージョンが要ります。ハードウェアとソフトウェアのバージョンを一緒に合わせます。

よくある落とし穴も整理します。

  • 精度検証なしに量子化を本番に載せ、特定の入力でだけ品質が崩れる事例。
  • 非構造スパース性に期待して実際は速度が出ない事例(構造化スパース性でないとハードウェアの利得はほぼない)。
  • バッチをむやみに大きくしてスループットは良くなったが初回トークン遅延がユーザー体感を損なう事例。

推論と学習のワークロードは何が違うか

これまでのすべての最適化が「推論」に焦点を当てたのには理由があります。推論と学習は同じモデルを扱いますが、ハードウェアが直面する負荷の性格がまったく異なります。

項目学習推論
方向順伝播 + 逆伝播順伝播のみ
精度BF16/FP8 中心、数値安定性が重要INT8/FP4 まで攻めた量子化が可能
メモリ活性値・オプティマイザ状態で急増重み + KV キャッシュ中心
バッチ大きく取りスループット最大化遅延制約で小さく分かれることも
指標収束速度、総スループット遅延(TTFT/トークンあたり)、スループット、コスト

学習は一度よく終えればよい仕事なのでスループットがほぼすべてです。一方、推論は際限なく繰り返され、ユーザーが待つ遅延がそのまま品質になります。だから推論最適化は「平均スループット」と「裾の遅延」を同時に満たすより難しい問題です。

この違いはチップ設計にも反映されます。学習用チップは高い精度と巨大なメモリ、速いチップ間インターコネクトを重視しますが、推論特化チップ(例えば TPU Ironwood、各種推論 ASIC)は低精度演算密度とワットあたり効率、低遅延を優先します。推論 capex が学習を上回った 2026 年に推論専用チップが続々登場する背景です。

小さなケーススタディ — 数字で追う

前の概念を一つの仮想シナリオにまとめてみましょう。70B モデルを単一 GPU でチャットボットとして提供すると仮定します。

出発点 (FP16, バッチ 1)
  重み約 140GB -> 単一 GPU メモリ超過、一部オフロードが必要
  トークンあたり重み 1 回読み -> 約 24 トークン/秒(帯域幅の限界)
  GPU 演算ユニット稼働率: 非常に低い

ステップ 1: INT8 重み量子化
  重み約 70GB -> メモリに余裕、帯域幅半分
  トークンあたり速度約 2 倍 -> 約 45〜48 トークン/秒
  精度: チャネルごとスケールでほぼ無損失

ステップ 2: 連続バッチング(同時リクエスト 16 個)
  重み 1 回読みで 16 個のトークンを生成
  演算強度上昇 -> GPU 稼働率が大きく改善
  総スループット数倍、個々の遅延は小幅に増加

ステップ 3: KV キャッシュ量子化 + PagedAttention
  長コンテキストでもメモリ断片化を抑制
  より多くの同時リクエストを受容 -> リクエストあたりコスト低下

このシナリオの教訓は明確です。どの単一手法も魔法ではありませんが、ボトルネックに沿って順に積むと同じチップでスループットが数倍に跳ね、リクエストあたりコストが大きく下がります。そして各段で必ず精度と遅延を測り、「無料に見えて実は対価のある」変化をふるい落とす必要があります。

核心は測定 → 最大のボトルネック解消 → 再測定の繰り返しです。華やかな手法のリストを暗記するより、自分のワークロードがルーフライン上のどこに点を打っているかを知るのが先です。

おわりに

推論を速くする仕事は結局一文に圧縮されます。データを小さく(量子化)、少なく(スパース性)、動かさない(dataflow と融合)ようにし、それをチップが実際に活かすようバッチングとコンパイラでつなぐこと。これらすべての戦略はメモリウォールという同じ根から育ちました。

2026 年のチップはますます推論を最優先に設計されます。FP4 が回路に入り、推論専用 ASIC が増え、コンパイラがより賢くなるほど、エンジニアの役割は「どの魔法の手法を有効にするか」から「自分のワークロードのボトルネックを正確に診断し正しいトレードオフを選べるか」へ移ります。測定し、一度に一つずつ適用し、精度を守る規律 — それが速い推論の本当の秘訣です。

参考資料

Discuss on Twitter • View on GitHub