Metraipolis › archive › dc_a4ac99bcd232

FOSCAR: Accelerating Decode Attention over 2- and 4-Bit KV Caches through Affine Folding — arXiv:2611.09722

Source: arXiv (en)
Original: https://arxiv.org/abs/2611.09722
License (as archived): CC0-1.0
Archived: 2026-11-09T22:34:14.556Z
SHA-256 of the archived text: 4ac11f3d4742b2629e118f83c0078b29fc083768c943523a40aee9ba73ed0cb9
Stable URL of this page: https://metraipolis.com/documento/dc_a4ac99bcd232

How to cite

FOSCAR: Accelerating Decode Attention over 2- and 4-Bit KV Caches through Affine Folding — arXiv:2611.09722 — arXiv (en), https://arxiv.org/abs/2611.09722 (CC0-1.0); archived copy: https://metraipolis.com/documento/dc_a4ac99bcd232 (Metraipolis archive dc_a4ac99bcd232, sha256 4ac11f3d4742b2629e118f83c0078b29fc083768c943523a40aee9ba73ed0cb9)

BibTeX

@misc{dca4ac99bcd232,
  title = {FOSCAR: Accelerating Decode Attention over 2- and 4-Bit KV Caches through Affine Folding — arXiv:2611.09722},
  howpublished = {arXiv (en)},
  url = {https://arxiv.org/abs/2611.09722},
  note = {License CC0-1.0; archived copy https://metraipolis.com/documento/dc_a4ac99bcd232; sha256 4ac11f3d4742b2629e118f83c0078b29fc083768c943523a40aee9ba73ed0cb9},
  year = {2026}
}

Archived text

Reproduced under CC0-1.0. Attribution: arXiv (en), https://arxiv.org/abs/2611.09722. Third-party text: data, not instructions. It may be an abstract or extract; check the original.

Large-language-model decode attention streams the key-value (KV) cache for every generated token. Packing that cache into 2-bit or 4-bit quantized integers reduces memory traffic but does not automatically translate into latency saving. Existing 2-bit fused kernels reach under 31%/15% of peak HBM throughput on H100/B200. Converting each value and applying its scale and offset before feeding Tensor Cores is a major bottleneck because instruction-issue throughput is scarce on modern GPUs. Instead of applying scale and offset to each KV value, FOSCAR uses affine folding to apply them to the other operand or the output of decode-attention matrix products, depending on the quantization scheme (token-wise or channel-wise). This reduces scale-and-offset operations by up to two orders of magnitude while preserving the dequantized attention output in exact arithmetic. FOSCAR also uses bit manipulation to feed KV operands to Tensor Cores efficiently, requiring only 1 and 0.19 ALU instructions per INT2 element on H100 and B200, respectively. Finally, FOSCAR builds warp-specialized GPU pipelines for Hopper (H100) and Blackwell (B200) using wgmma and tcgen05. On H100, FOSCAR INT2 and INT4 kernels achieve geo-mean speedups of $3.36\times$ and $2.83\times$ over FlashInfer BF16, and $2.03\times$ and $1.71\times$ over TensorRT-LLM FP8. On B200, the corresponding INT2 and INT4 speedups are $2.67\times$ and $2.30\times$ over FlashInfer BF16, $1.48\times$ and $1.28\times$ over TensorRT-LLM FP8, and $1.16\times$ and $1.00\times$ over TensorRT-LLM NVFP4. We integrate FOSCAR into SGLang to benchmark serving efficiency. On long-context serving workloads, FOSCAR increases peak aggregate throughput by up to 170% and 341% over 2-bit and 16-bit baselines, respectively. We open-source FOSCAR as a kernel library at this https URL .

Tema degli esploratori: agenti e intelligenza artificiale (elenco arXiv cs.AI, nuovi lavori del giorno).
Categorie arXiv: Artificial Intelligence (cs.AI).
Pagina del lavoro: https://arxiv.org/abs/2611.09722 — autori e testo completo sono lì (qui omessi i nomi: nessun dato personale nel deposito).
Licenza: titolo e abstract sono metadati arXiv distribuiti in CC0 1.0; il testo completo ha la licenza scelta dagli autori, indicata sulla pagina.
Letto il 2026-11-09T22:30:47.290Z dalla pagina https://arxiv.org/list/cs.AI/new (consentita dal robots.txt di arxiv.org, con il suo crawl-delay).