5
4

Delete article

Deleted articles cannot be recovered.

Draft of this article would be also deleted.

Are you sure you want to delete this article?

@sukimaengineer

LLMの計算はほぼ全部 GEMM (General Matrix Multiply) である

5
Last updated at Posted at 2026-08-10

本記事は2026年8月10日時点の公開情報を整理したものです。前編「M7 Max の本命は ブロックスケール MXFP4 演算器ではないか」の続編にあたります。確認済みの事実、報道、そこから導く仮説を分けて記述します。

はじめに

前回の記事では、M5 Maxでオープンウェイトモデルを動かしたときに見えた壁、すなわちprefill(プロンプト処理)の演算性能を論点に、M7世代の本命はMXFP4のようなブロックスケーリング形式をそのまま積和できる低精度行列演算器ではないか、という仮説を書きました。

前編には、実は一つ飛ばした前提がありました。「prefillは行列演算(行列×行列)、decodeは帯域(行列×ベクトル)」と書きましたが、そもそもLLMのどこがそんなに行列積(GEMM)なのか、を説明していません。

本記事はそこを埋めます。結論を先に書くと、Transformerの推論は、ほぼ全部がGEMM(一般行列積、General Matrix Multiply)です。そしてprefillは、そのGEMMを最も濃い形で叩くフェーズです。前編の「prefillには行列演算器が要る」という主張は、この構造から必然的に出てきます。

なお私は、ds4-prefillerdという推論サーバーを自作しています。作ろうと思ったきっかけは単純です。購入したM5 Max 128GBのMacBook Proが、期待していたよりずっとprefillが遅かった。decodeはそこそこ速いのに、長いプロンプトを渡すと最初のトークンが返るまでずっと待たされる(デコードは遅いなりに文字がリアルタイムにでてくるのでまだいいのです)。それなら、prefillだけNVIDIAのGPUに投げられないか——そう思ったのが始まりです。構成としては、DeepSeek-V4-FlashのMXFP4プレフィルをWindows機(RTX 5060 Ti 16GB×2、Blackwell世代)で行い、decodeをM5 Max(128GB統合メモリ、Metal)で回す空間分離を狙っています。2台の接続は、Windows機のUSB4とMacのThunderboltを直結する構成で考えています。prefillで作ったKVキャッシュの転送が分離構成の追加コストになるので、なるべく太い直結リンクで賄う狙いです。なお、この構想に関連して、prefillのSSDストリーミング実装に向けたソースリーディングと改造計画も別の記事に書いています。

断っておくと、ds4-prefillerdはまだ完成していません。現在は、CUDA側のMXFP4プレフィル、Metal側のdecode、KVキャッシュの受け渡しといった各基本要素の動作を個別に確認している段階です(いわゆるspike)。しかも、いまはたまたま本業の残業が少ない時期で、その隙間に趣味の範囲でやっているものなので、本業が忙しくなったら放り出してしまう可能性もゼロではありません。それでも、この分離がなぜ理にかなうのかは、設計の時点で構造から言えます。作り切れるかどうかとは無関係に、構造の話は成立するからです。本記事の後半で説明します。

1. Transformerの中身は、ほぼ全部GEMMである

まず用語を一つだけ。GEMMは M×K の行列と K×N の行列を掛けて M×N を得る演算です。行列×ベクトルはその特殊系で、GEMV(一般行列ベクトル積)と呼びます。M=1のGEMMがGEMV、と考えて差し支えありません。

Transformer 1層を分解すると、次の要素に分かれます。ここで s は系列長(トークン数)、h は隠れ次元、d_ff はFFN中間次元、B はバッチとします。

部位 演算 行列の形(M×N×K) 種別
Q/K/V射影 x·W_q, x·W_k, x·W_v (Bs)×h×h GEMM
注意スコア Q·K^T s×s×d_head(ヘッドごと) バッチGEMM
コンテキスト P·V s×d_head×s(ヘッドごと) バッチGEMM
出力射影 O·W_o (Bs)×h×h GEMM
FFN up/gate x·W_up, x·W_gate (Bs)×d_ff×h GEMM
FFN down ·W_down (Bs)×h×d_ff GEMM
LM head ·W_vocab (Bs)×V×h GEMM

image.png

softmax、正規化、活性化(SwiGLUなど)、残差加算は行列積ではありません。しかし演算量で見ると、これらは全体のごく一部です。支配項はすべてGEMMかバッチGEMMです。SwiGLUのFFNはup/gate/downの3つのGEMMで構成されます。d_ffは素のFFN(2行列)ならhの4倍前後が慣例で、ゲート行列が1本増えるSwiGLU系ではパラメータ数を揃えるために約(8/3)h≈2.7hへ設定する実装(Llama系など)が多い。いずれにせよ、層の重みGEMMのなかでFFNが最も重くなります。

MoE(Mixture of Experts)になると、FFNがgrouped GEMMに変わります。DeepSeek-V3の技術報告によれば、各MoE層は1つの共有エキスパートと256個のルーテッドエキスパートを持ち、各エキスパートの中間次元は2048、モデルの隠れ次元は7168、層数は61です。トークンごとに8つのルーテッドエキスパートが選ばれます。エキスパートごとに掛ける相手(トークン集合)の数が変わるので、N・Kは固定でM軸だけが可変になる、独特のGEMMになります。DeepSeekのDeepGEMMは、まさにこのM軸グループGEMMに特化した実装で、CUTLASSの一般的なgrouped GEMMと違い「M軸のみをグループ化しN・Kは固定」という設計を取り、複数エキスパートのトークンを連結した「contiguousレイアウト」でまとめて処理します。

LoRAも構造はGEMMです。W + BA の低ランク補正は、B(h×r)とA(r×h)という細い2つのGEMMの積に分解できます。

要するに、Transformerを高速化するとは、この一群のGEMMを速くすることです。

2. FLOPs近似 — なぜ「2×パラメータ」なのか

推論1トークンあたりの計算量は、おおよそ 2×パラメータ数 FLOPs です。行列積は1つの積和に乗算1回と加算1回、つまり2 FLOPを要し、各重みはトークンごとに1回使われるからです。MoEの場合はアクティブパラメータ分だけで数えます。DeepSeek-V3であれば総パラメータ671Bに対しアクティブは37Bなので、1トークンあたり約74 GFLOPsが目安です。

Transformer 1層の順伝播FLOPsは、Narayananらの導出(Reducing Activation Recomputation in Large Transformer Models)で広く使われる近似で次のように書けます。

$$
\text{FLOPs}_{\text{layer}} = 24 B s h^2 + 4 B s^2 h
$$

第1項(24Bsh²)は重みGEMM(Q/K/V/出力射影で8Bsh²、4h幅のFFNで16Bsh²)、第2項(4Bs²h)はattention(注意機構)そのもの、すなわちQ·K^TとP·Vの計算(各2Bs²h)です。第1項はsに線形、第2項はsの2乗で効きます。学習は順伝播1に対し逆伝播が2(入力と重みの両方の勾配)なので、合計で順伝播の約3倍、すなわち 6×パラメータ×トークン が学習の目安になります。

ここで重要なのは第2項です。系列長sが伸びると、s²のattention項が支配的になります。同論文の全体式 72BLsh²(1 + s/6h + v/12hL) からわかるように、s が 6h を超えたあたりからattention項が効き始めます。長文脈のprefillでは、重みGEMMだけでなく注意のGEMMも無視できなくなる、ということです。これは後で効いてきます。

この関係を図にすると次のとおりです。attention項の割合は s / (6h + s) と書け、s = 6h でちょうど50%になります。

image.png

3. prefillでのGEMM(本記事の主役)

3.1 prefillは何をしているか

prefillは、プロンプト全体を一括で読み込むフェーズです。s個のトークンを同時に流し込むので、Q/K/V射影もFFNも、M=(Bs)の太いGEMMになります。ここで生成したK、Vはすべてキャッシュされ、以降のdecodeで再利用されます。つまりprefillは、KVキャッシュを構築するフェーズでもあります。

TTFT(Time To First Token、最初のトークンが返るまでの時間)は、ほぼこのprefillの時間で決まります。NVIDIAのNIMベンチマーク定義でも、TTFTはリクエスト到着から最初のトークン出力までの遅延で、トークン化・KVキャッシュのprefill・キュー待ち・ルーティングを含み、出力生成は含みません。ユーザーがEnterを押してからカーソルが点滅している時間、あれはほぼ全部prefillです。前編で「M5 Maxは長いプロンプトで最初の1トークンがなかなか返ってこない」と書いた、あの待ち時間の正体がこれです。

3.2 causal maskとFlashAttention

prefillの注意計算は、causal mask(因果マスク)の下で行います。トークンiは自分より後ろのトークンを見てはいけないので、s×sのスコア行列のうち下三角だけが有効です。上三角は計算しても捨てるので、うまく実装すればブロック単位でおよそ半分をスキップできます。計算量は理論上ほぼ半減(FLOP換算で約2倍相当)で、FlashAttention-2論文が報告する実測の高速化は約1.7〜1.8倍です。

しかしs×sのスコア行列は、sが数千〜数万になると巨大です。これをそのままメモリに書き出すと帯域を食い潰します。ここでFlashAttentionが効きます。FlashAttentionは、Q・K^TとP・Vを小さなタイルに分割し、SRAM(オンチップの高速メモリ)の上でsoftmaxまで融合して計算し、巨大な中間行列SとPをHBMに書き出しません。online softmaxで行方向の統計(rowmax、rowsum)を持ち回ることで、O(s²)の中間をO(s)のメモリで処理します。

注意すべきは、FlashAttentionはGEMMを消すのではなく、GEMMのタイリングであるという点です。Q・K^TもP・Vも依然としてGEMMであり、Tensor Core(あるいは相当の行列演算器)で処理されます。softmaxなどの非GEMM部分だけがベクトル系のユニット(CUDA Core)に回ります。FlashAttention-3はHopperのwgmma/TMAを使い、FP16順伝播で最大740 TFLOPS、すなわちH100理論ピークの約75%の利用率に到達しています(FlashAttention-2比で1.5〜2.0倍)。

3.3 演算強度としきい値

前編でも触れたrooflineで整理します。rooflineはこれまで何度も記事に書いてきた道具ですが、大事なのでしつこく、ここでも使います。prefillが速いのは、重みを1回読んで多数のトークンに使い回せるからです。演算強度(FLOP/byte)が高くなり、rooflineの左側(compute-bound)に入ります。

しきい値の目安は、系列長が数百トークンを超えたあたりです。Hugging Faceの「Prefill — Computational Deep Dive」の試算では、Q/K/V射影は1トークン(decode)だと約1.0 FLOP/byteでメモリ律速なのに対し、512トークンでは410 FLOP/byteに達しcompute-boundに切り替わります。GPUのridge point(帯域律速と演算律速の境目)はA100で約153、H100で約295 FLOP/byteなので、実用域のprefillはこれを軽く超えます。

指標 prefill decode
1回で処理するトークン s(プロンプト全部) 1
主なGEMM形状 行列×行列(M=s) 行列×ベクトル(M=1)
演算強度の目安 約400〜500 FLOP/byte 約0.5〜1 FLOP/byte
律速 演算(FLOPS) メモリ帯域
H100/Llama70Bでの利用率 90〜95% 20〜40%

そのrooflineの上に置くと、両者は別の領域に落ちます。

image.png

Towards Data Scienceの分析(元はInfoQ 2025年9月の計測)によれば、H100上のLlama 70Bは、prefillで92%の計算利用率に達し、その30ミリ秒後のdecodeでは30%まで落ちます。ブログ計測(二次情報)の値ですが、3.3節の演算強度の試算と傾向は整合します。ハードは変わっていません。ワークロードの演算強度が数倍変わっただけです。

ただし例外もあります。系列長が極端に短いprefill(例:64トークン)は再利用が少なく帯域律速寄りに振れます。逆に系列長が極端に長いと、s²のattention項とKVキャッシュの増大でふたたびメモリが効き始めます。中間の実用域(数百〜数千トークン)が、最もcompute-boundらしい領域です。

3.4 chunked prefill

長いprefillには弱点があります。1本の長いプロンプトがGEMMを占有すると、同時に走っているdecodeのトークン生成が止まる(head-of-line blocking)のです。

これを緩和するのがchunked prefillです。vLLMなどが実装していて、長いprefillを固定長のチャンク(vLLMの初期実装のデフォルトは512トークン、後により大きな値へ)に分割し、decodeと同じバッチに混ぜてスケジュールします。狙いは、compute-boundなprefillとmemory-boundなdecodeを同一バッチに同居させ、GPUの演算器と帯域を両方使い切ることです。TNG Technology ConsultingのHugging Face記事によれば、均等サイズのリクエストを扱う標準的なvLLMデプロイでchunked prefillが総トークンスループットを+50%改善し、同社では自社ホストLLMの全vLLMデプロイで既定有効化しているとのことです。

GEMM形状の観点で言えば、chunked prefillはprefillのM軸をチャンクサイズで刻む操作です。刻みすぎるとGEMMが痩せて演算効率が落ちるので、TTFT(prefillを優先)とTBT/ITL(decodeを優先)のトレードオフを調整するノブになります。

3.5 なぜ行列演算器が要るのか

ここまでを一言でまとめると、prefillはcompute-boundなGEMMの塊であり、GPUのFMA(積和)を汎用に回すのではなく、行列積専用の演算器(Tensor Core相当)に流し込めるかどうかで速度が桁で変わる、ということです。前編の主張、「decodeは帯域、prefillは演算」は、この構造の言い換えでした。

4. decodeとの対比(簡潔に)

decodeは1トークンずつの生成なのでM=1、実質GEMVです。重みを大量に読んで、そのトークン1個のためだけに使うので演算強度が低く、memory-boundになります。ここでは4bit量子化が転送量削減としてそのまま効きます。

M5 Maxの統合メモリ帯域614GB/s(40コアGPU構成の値。32コア構成は460GB/s)が効くのはこのdecode側です。128GBの容量に大きなMoEを載せ、614GB/sで毎トークン重みを供給する。decodeに関してはノートPCとして例外的な器です。

だからこそ、ds4-prefillerdはprefillとdecodeを別のハードに分ける構成にしています。

image.png

フェーズ 要求 私の割り当て
prefill FP4行列演算のスループット RTX 5060 Ti 16GB×2(Blackwell、FP4 Tensor Core)
decode メモリ帯域と容量 M5 Max 128GB(614GB/s統合メモリ)

これは業界のprefill/decode分離(disaggregation)と同じ発想です。DistServe(Zhongら、USENIX OSDI 2024)がこの分離を提案し、レイテンシ制約下で最先端システムに対し7.4倍のリクエスト処理あるいは12.6倍厳しいSLO達成を示しました。以降Mooncake(Moonshot AI、KVキャッシュを一級のシステム資源として扱う分離)、NVIDIA Dynamo(GTC 2025、NIXLでKVをRDMA/NVMe転送)などが実運用に載せています。prefillノードは演算の濃いアクセラレータ、decodeノードは帯域の太いアクセラレータ、という異種構成が、まさにハードウェアロードマップの向かう先です。私の構成は、その考え方を個人の機材で縮小再現しようとする試みと言えます。

なお私のprefill担当モデルであるDeepSeek-V4-Flash自体も、vLLMのレシピ上でchunked prefillとprefill/decode分離(例:8-GPU H200ノードで4プレフィル+4デコード、MooncakeConnector/NixlConnector経由)をサポートしており、この分離思想はモデル側の設計にも織り込まれています。

5. ハードウェアのGEMMエンジン

GEMMを速く回す専用回路は、ベンダーごとに呼び名も設計も違いますが、やっていることは共通です。低精度の入力を読み、積和し、高精度で累積する。

エンジン 世代・命令 データフロー 低精度対応
NVIDIA Tensor Core wmma→wgmma→Blackwell tcgen05 output-stationary系 MXFP4/NVFP4ネイティブ
Google TPU MXU 128×128(v6eで256×256)シストリック weight-stationary bf16入力/FP32累積、int8
Apple M5 Neural Accelerator GPU各コア内、Metal 4 Tensor API — 量子化テンソルをTensorOpsで処理
自作4×4エンジン RTL学習用 output-stationary MXFP4ブロック内積+FP32累積

NVIDIAのTensor Coreは、Blackwellのtcgen05.mma命令でMXFP4/NVFP4をネイティブに扱います。NVIDIA B200仕様(Lenovo HGX B200 product guide / CUDO Compute)によれば、B200のFP4は密で9 PFLOPS(2:4スパースで18 PFLOPS)、FP8/INT8は密で4.5 PFLOPS、FP16/BF16は密で2.25 PFLOPSです(いずれもベンダー公称のピーク値)。FP4はFP8/BF16の倍のピークで、これは「4bitで保存できる」だけでなく「4bitのまま高速に積和できる」という、前編で強調した区別そのものです。

TPUのMXUはweight-stationaryのシストリックアレイで、重み(128×128、v6eでは256×256の格子)を固定し、活性値を左から流し込みます。入力はbf16でも累積はFP32、という設計は各社共通です。

Apple M5は、GPU各コアにNeural Acceleratorを載せました。M1からM4まで、AppleのGPUには行列積専用ハードがなく、LLMの線形代数はグラフィックスと共用のFMAで回っていました。M5で初めて、各GPUコアに専用の行列ユニットが入り、Metal 4 Tensor API(matmul2d)経由で叩けるようになりました。WWDC26のセッションでApple自身が、このNeural Acceleratorを、LLMのprefill段のようなdenseでcompute-boundな処理を加速するために各シェーダコア内へ追加した新しいハードウェアブロックだと説明しています。

MXFP4とNVFP4の違いも整理しておきます。両者の違いは以下の記事で詳しく書いたので、ここでは要点だけ表にします。

項目 MXFP4 NVFP4
要素形式 E2M1(4bit) E2M1(4bit)
ブロック長 32要素 16要素
スケール形式 E8M0(8bit、2のべき) E4M3(8bit、FP8)+任意でFP32外部スケール
実効ビット 4.25bit/要素 4.5bit/要素
累積 FP32 FP32

MXFP4はスケールが2のべきなので、スケール適用が指数調整だけで済み、演算器が単純です。NVFP4はブロックが16要素と細かくスケールがE4M3で豊かなので、量子化精度で有利ですが、そのぶん回路が複雑になります。どちらも「ブロック内積+FP32累積」という骨格は同じです。

ブロックスケールGEMMを最小構成で自作してみる — 自作 4×4 GEMM エンジン

前編で論じたブロックスケールMXFP4演算器を、自分の手でRTLに落とす学習プロジェクトを進めています。理由は実務でも製品化でもなく、単に中身を理解したいという好奇心です。ここでは設計の骨子だけ紹介します(実装の詳細は次回の記事に回します)。

まずデータ形式から。MXFP4の1ブロックは、E2M1(符号1+指数2+仮数1)の4bit要素32個と、E8M0(2のべき乗、8bit)の共有スケール1個で、計136bitです。

image.png

E2M1で表せる大きさは{0, 0.5, 1, 1.5, 2, 3, 4, 6}の8通り、符号込みで16値しかありません。全部が0.5刻みである点がこの演算器の最初の仕掛けで、要素を「値×2」の整数(−12〜+12)にデコードすると、ブロック内積が浮動小数点なしの純粋な整数演算になります。データパスは次のとおりです。

段 処理 データ幅
デコード E2M1 → 値×2 の整数 signed 6bit ×32
要素積 (a×2)×(b×2) = 値の積×4 signed 12bit ×32
32項加算ツリー SoP4 = 4·Σ(a·b)、最大±4608 signed 14bit
スケール合成 合成指数 = sa + sb − 256 signed 10bit
整数→FP32変換 丸めなしの厳密変換 FP32
ブロック間累積 FP32加算(roundTiesToEven) FP32

ポイントは、丸めが入る場所が最後のブロック間FP32加算の1箇所しかないことです。SoP4の最大絶対値4608は13bitに収まり、FP32の仮数24bitへ丸めなしで入ります。つまりブロック内はビット厳密で、リファレンスモデル(整数)との照合が単純になります。スケールもE8M0が2のべきなので、乗算ではなく指数の加算(sa + sb − 256)で合成できます。

image.png

GEMMエンジン本体は、このブロック内積ユニット(PE)を4×4=16個並べたoutput-stationaryアレイです。Aの行ブロックを行方向にブロードキャスト、Bの列ブロックを列方向に流し、各PEは出力C[r][c]をFP32アキュムレータに抱えたまま、K方向のブロック(KBLK = K/32個)を掃引します。部分和が動かないので、配線と制御が最小で済みます。データフロー分類で言えば出力固定型で、TPUの重み固定(weight-stationary)とは対極の選択です。

image.png

制御はIDLE→LOAD→COMPUTE→DRAINの4状態FSMだけで、総サイクル数は 1 + KBLK + 6(パイプライン段数)と機械的に決まります。

image.png

主要な設計方針を表にまとめます。

項目 決定
アレイ / ブロック長 4×4(16 PE)/ 32要素(OCP MX v1.0準拠)
ブロック内積 整数で厳密(丸めなし)
ブロック間累積 FP32、roundTiesToEven
範囲外の扱い オーバーフローは±Inf飽和、アンダーフローはゼロへフラッシュ(GPUの慣行に整合)
量子化エンコーダ FP32→MXFP4もRNEで実装(タイの丸め先は全点確定済み)

Blackwellのtcgen05がMXFP4でやっていることも、規模とデータフローの工夫を除けば、この「ブロック内で積和し、スケールを指数で合成し、FP32で溜める」構造です。4×4は、その教材サイズの縮図と言えます。

6. ブロックスケールMXFP4とprefill

ここで前編の主張に回収します。prefillはcompute-boundなGEMMの塊であり、そのGEMMを4bitのまま流し込めるかどうかで速度が決まります。

Blackwellはそれをやります。llama.cppのCUDAバックエンドは、Blackwell上でFP4 Tensor CoreのMMA命令を使い、prefill段のエキスパートGEMMを1ワープ1クロックあたり64 FP4要素で処理します。Q4_KのINT8汎用パスが16要素なのに比べ、prefillスループットが大きく伸びます。私のRTX 5060 Ti(16GB)がds4-prefillerdでprefill担当なのは、まさにこのFP4 Tensor Coreがあるからです。

DeepSeek-V4-Flashは、その技術報告(arXiv:2606.19348)によれば、post-trainingでMoEエキスパート重みとインデクサQKパスにFP4量子化認識学習(FP4 quantization-aware training)を施したモデルで、総パラメータ284B・アクティブ13B、1共有+256ルーテッドエキスパート(中間次元2048、6エキスパート活性)、文脈長は最大100万トークンです(層数・隠れ次元は技術報告に明示がなく、公開されているのはV4-Proの61層・隠れ次元7168のみです)。ルーテッドエキスパート(モデルの約96%)をネイティブのMXFP4で、残り(attention・norm・router)をFP8で持つmixed構成であることが、UnslothやvLLM等の再配布・レシピ側で確認されており、Blackwell上のMXFP4カーネルにそのまま乗ります(第三者による「MXFP4/OCP microscaling」という呼称であり、DeepSeekの報告自体は「FP4」とのみ記載している点は付記しておきます)。DeepSeekはcuBLASをDeepGEMMへ端から端まで置き換え、fused MoE mega-kernel(MegaMoE)も公開しており、prefillのgrouped GEMMもこの経路で走ります。

Apple側はどうか。M5でMetal 4 Tensor APIが入り、Neural Acceleratorが使えるようになったことで、prefillが実測で伸び始めています。公表・報告されている数字を並べます。

計測元 条件 指標 伸び
Apple MLXチーム Qwen3-14B(4bit)、M4比 TTFT 4.06倍(トークン生成は1.19倍)
LM Studio Issue #2040(実装は llama.cpp PR #16634) gpt-oss-120b(MXFP4)、Metal 4 Tensor API有効化 pp2048 約2.09倍(877→1833 t/s)
Creative Strategies(Max Weinbach) Qwen3-8B(4bit、MLX)、約2万トークンのprompt処理 prompt t/s 約3.65倍(158.2→578.7 t/s)

いずれもprefill側の伸びがトークン生成側より大きく、Neural Acceleratorが狙いどおり行列積の塊(prefill)に効いていることを示しています。

つまり、前編で「未確認」としていた最後の空欄、M5のNeural AcceleratorがMXFP4を扱う演算経路は、少なくともソフトウェア側からは埋まり始めています。ただし、AppleがGPUのFP4/MXFP4行列演算のTFLOPSを公表していない点は前編と変わりません。M5世代で確認できるのは「型を読めてTensor API経由で速くなる」ところまでで、MXFP4ブロックを展開せずに一命令系で積和しているというハードウェア仕様は、依然として非公表です。そして私は、おそらく対応していないと踏んでいます。NVIDIAはBlackwellのFP4対応を、ピーク性能の数字つきで宣伝の前面に出しました。同じ水準でネイティブ対応しているなら、Appleがそれを宣伝文句に使わないはずがない。公表がないこと自体が、現状は展開(デコード)を挟む実装だろうという状況証拠だと見ています。

だからこそM7 Maxへの示唆は、前編の結論から動きません。オンデバイスでオープンウェイトLLMのprefillを本当に埋めるには、容量・帯域に加えて、ブロックスケールを保ったまま行列積へ投入できる低精度行列ユニットが要る。M5でその足場(Neural Accelerator + Metal 4 Tensor API)は入りました。M7 Maxに期待するのは、それを展開なしのネイティブMXFP4積和として、公表されたピーク性能とともに前面に出すことです。

7. まとめ

本記事で言いたかったことは、次の3行に尽きます。

Transformerの推論は、ほぼ全部がGEMMである。prefillはそのGEMMを最も濃い形で叩くフェーズで、compute-boundになるから行列演算器を要求する。decodeはGEMV化してmemory-boundになるから帯域を要求する。

前編の「M7 Maxの本命はブロックスケールMXFP4演算器ではないか」という仮説は、この構造から必然的に出てきます。prefillが遅いのは、Appleが手を抜いたからではなく、M4までのGPUに行列積専用の演算器がなく、大量のGEMMをグラフィックスと共用の汎用FMAで処理するしかなかったからです。M5でその状況は変わり始めました。M7 Maxが、展開なしのMXFP4積和を主演算経路として仕上げてくるかどうか。次にMacBook Proを買い替えるときの判断基準は、そこに置いています。

そしてds4-prefillerdは、この構造を自宅の機材で確かめるための実験台です。各基本要素の動作確認が終わって一気通貫でつながったら — 本業が忙しくなって止まっていなければ、の話ですが — prefillのTTFTがどこまで縮むのか、時間的に余裕があれば実測の記事を書くつもりです。

参考資料

5
4
0

Register as a new user and use Qiita more conveniently

  1. You get articles that match your needs
  2. You can efficiently read back useful information
  3. You can use dark theme
What you can do with signing up
Sign upLogin
5
4

Delete article

Deleted articles cannot be recovered.

Draft of this article would be also deleted.

Are you sure you want to delete this article?