TL;DR
FOSCARは、2/4ビットKVキャッシュのデコードAttentionで、値ごとのスケール・オフセット適用がTensor Coreへのオペランド供給を圧迫する点を狙う。アフィン折り畳みでスケール/オフセットを行列積の入力側または出力側へ移し、ビット操作でKVオペランドを直接構築する。H100/B200向けカーネルでFlashInfer BF16比3.36×/2.67×（INT2）などの高速化を報告。評価は限定的で、他GPUや精度面は提示本文だけでは未確認。
既存手法との違い
原論文はFlashInfer、TensorRT-LLM、FlashAttention-2/3/4、BitDecoding、KIVI、Kitty、SAW-INT4、OSCARを比較。B=128、KV=32Kで、KVキャッシュ読出しの実効HBMスループットを各GPUピーク比で測定している。BF16 FlashInferはH100/B200とも90%超。B200のTensorRT-LLM NVFP4は4ビット群で72.4%。整数4ビットではBitDecodingがH100 70.0%/B200 36.1%、そのINT2は30.7%/14.5%。FOSCARはINT4/INT2でH100 91.3%/60.6%、B200 68.6%/44.4%。カーネル遅延ではH100 INT2/INT4がFlashInfer BF16比3.36×/2.83×、TensorRT-LLM FP8比2.03×/1.71×。B200はFlashInfer BF16比2.67×/2.30×、TensorRT-LLM FP8比1.48×/1.28×、NVFP4比1.16×/1.00×。SGLang統合の長文脈サービングではピーク集約スループットが2ビット比最大170%、16ビット比最大341%増。既存手法が値ごとに再構成するのに対し、FOSCARはアフィン演算を行列積側へ移し、パック整数からTensor Coreオペランドを直接作る点が違い。
使える条件
対象は自己回帰LLMのデコードAttention。非対称量子化で各グループがスケールsとオフセットzを共有する2/4ビットKVキャッシュを想定し、K/Vそれぞれでtoken-wise/channel-wiseを独立に選べる。H100（wgmma、BF16オペランド）とB200（tcgen05、Q/PはE4M3、INT4 KVはE4M3、INT2 KVはE2M1）向けカーネル。評価はB=128、KV=32Kのスループットと、SGLang統合の長文脈サービング。グループサイズ128・KVヘッドあたり4クエリヘッドの命令需要例がある。実装はオープンソースのカーネルライブラリとして公開。他GPU、他量子化方式、prefill、学習、精度評価は提示本文からは確認できない。
限界と未検証の範囲
提示本文の評価はH100/B200のデコードAttentionとSGLang長文脈サービングに限定される。B200のINT4はTensorRT-LLM NVFP4比1.00×で、常に大幅高速とは限らない。性能は量子化方式・グループサイズ・ヘッド数・パイプラインの遅い段に依存する。acc-foldは量子化グループ境界を越える部分積をそのまま累積できず、グループ単位のスケール適用が必要。厳密算術での等価性を主張するが、提示本文にはモデル精度や数値誤差の実測はない。他GPU世代、他モデル/データ、prefill、学習での有効性は確認できない。
解説
ねえ智也くん、この「FOSCAR」ってブログ、低ビットKVキャッシュのデコードAttentionを速くする話らしいけど、そもそもなんで2ビットとかにしても速くならないの？ メモリ通信量は減るんでしょ？
そう。KVキャッシュを2/4ビットにパックすればHBMから読むバイト数は減る。でも既存の2ビットカーネルはH100/B200でピークHBMスループットの31%/15%未満しか出てない。理由は、Tensor Coreに渡す前に各値へスケールとオフセットを適用する処理が命令発行スロットを食うから。ビット幅を下げるとメモリ転送時間は減るのに、この値ごとの処理コストは変わらないから、ボトルネックがメモリから命令発行に移るんだ。
へえ！ じゃあ速くするにはその「値ごとのスケール・オフセット」をなんとかしないといけないんだ。
FOSCARはアフィン折り畳みってのを使う。KVの各値にスケールとオフセットをかける代わりに、行列積のもう一方の入力側か出力側に移す。QK⊤はDで縮約するから、token-wiseのKなら部分QK結果にスケールをかけて、オフセット×クエリ和を足す（acc-fold）。channel-wiseのKならQ側にスケールを折り畳む（operand-fold）。PVはLで縮約するから、token-wiseのVはP側にoperand-fold、channel-wiseのVはacc-fold。これでスケール・オフセット演算が最大2桁減る。厳密算術では、各KV値をデクオンタイズしてから行列積するのと同じ出力になる。
ビット操作でTensor Coreオペランドを直接作るってのもあるよね？
Hopperではマンティッサ・インジェクションで、qに0x4300をORしてBF16のq+128を作る。BlackwellではINT2のビットパターンをそのままE2M1の下位ビットに置くとq/2になる。INT4はE4M3でq·2^-9。これでINT→FP変換の命令も減らせる。表1だとHopperは1.00 ALU/elem、BlackwellはINT4 0.38、INT2 0.19。
なるほどね！ で、どれくらい速くなるの？
H100だとINT2/INT4がFlashInfer BF16比3.36×/2.83×、TensorRT-LLM FP8比2.03×/1.71×。B200だとFlashInfer BF16比2.67×/2.30×、TensorRT-LLM FP8比1.48×/1.28×、NVFP4比1.16×/1.00×。SGLangに統合した長文脈サービングではピーク集約スループットが2ビット比最大170%、16ビット比最大341%増。
すごいじゃん！ でも、なんか注意点もあるんでしょ？
評価はH100/B200のデコードAttentionとSGLang長文脈サービングに限られてる。B200のINT4はNVFP4比1.00×で、常に大幅高速とは限らない。acc-foldは量子化グループ境界を越えて部分積をそのまま累積できないから、グループごとにスケールを適用する必要がある。それに、提示本文にはモデル精度や数値誤差の実測はない。他GPU、prefill、学習での有効性は確認できない。
ふーん、じゃあFOSCARって「速いけど、まだ全部の舞台で活躍するには修行が必要」って感じだね。まるで智也くんみたい！
…俺は修行中じゃない。あと、勝手に話をまとめるな。