FOSCAR boosts LLM decode attention 3.4× using 2‑bit KV caches

FOSCAR, a new kernel library released on arXiv, applies an affine‑folding technique to 2‑bit and 4‑bit key‑value (KV) caches used in large‑language‑model (LLM) decode attention. By eliminating per‑element scale‑and‑offset work, it cuts latency and raises throughput, a step that matters for any service that generates text token‑by‑token.

---

What is FOSCAR and how does it work?

During autoregressive decoding, each token triggers a read of the KV cache stored in high‑bandwidth memory (HBM). Low‑bit quantization stores each cached value as a small integer $q$ together with a group‑wise scale $s$ and offset $z$ so that the original floating‑point value is reconstructed as $x = s\,q + z$. Existing kernels dequantize every element before feeding it to Tensor Cores, which consumes many instruction‑issue slots.

FOSCAR flips that pattern. Instead of applying $s$ and $z$ to every element, it "folds" the affine transformation into the matrix‑multiply itself. For token‑wise quantization, the kernel scales the partial $QK$ result and adds the offset multiplied by the query sum; for channel‑wise quantization, the query is scaled before the $QK$ product and the offset is added later. The same idea applies to the $PV$ multiplication. The result is mathematically identical to full dequantization but requires only a handful of ALU instructions per element (as low as 0.19 on the B200). The paper describes the bit‑level manipulation needed to pack INT2/INT4 values directly into Tensor Core operands.

---

Why did previous low‑bit kernels lag behind peak HBM throughput?

Even though 2‑bit and 4‑bit caches shrink the number of bytes streamed from HBM, the decode‑attention latency often stayed dominated by on‑chip work. The arXiv preprint notes that prior fused kernels reached under 31 % (H100) and 15 % (B200) of peak HBM bandwidth. The bottleneck was the instruction‑issue bandwidth required to unpack each integer, multiply by its scale, and add the offset before the Tensor Core could start the matrix multiply. As the bit‑width drops, memory traffic falls but the per‑element arithmetic cost remains, shifting the limiting factor from memory to instruction dispatch.

---

What performance gains does FOSCAR deliver?

On NVIDIA Hopper (H100) GPUs, FOSCAR’s INT2 kernel achieved a geometric‑mean speedup of 3.36× over the FlashInfer BF16 baseline and 2.03× over TensorRT‑LLM FP8. The INT4 variant posted 2.83× and 1.71× speedups respectively. On the newer Blackwell (B200) silicon, the gains were 2.67× (INT2) and 2.30× (INT4) versus FlashInfer BF16, with smaller but still positive improvements over TensorRT‑LLM FP8. In terms of effective KV‑access throughput, FOSCAR reached 91.3 % of peak HBM bandwidth for INT4 on H100, compared with sub‑30 % figures for earlier INT2 kernels.

When integrated into the SGLang serving engine, FOSCAR boosted peak aggregate throughput on long‑context workloads by up to 170 % over a 2‑bit baseline and 341 % over a 16‑bit baseline. These numbers demonstrate that the reduction in instruction‑issue pressure translates directly into higher end‑to‑end serving capacity.

---

How can developers integrate FOSCAR into their LLM serving stacks?

The authors have open‑sourced the kernel library on GitHub (https://github.com/FutureMLS-Lab/FOSCAR). It provides CUDA kernels for both Hopper and Blackwell, complete with warp‑specialized pipelines that exploit wgmma and tcgen05 instructions. To use it, developers replace the standard decode‑attention kernel in their inference framework (e.g., FlashAttention, TensorRT‑LLM) with the FOSCAR version and ensure that the KV cache is quantized to INT2 or INT4 with the required group‑wise scale/offset layout.

The paper also describes a modest integration effort with SGLang, showing that the library can be dropped into an existing serving engine without major architectural changes. Because the kernels preserve exact arithmetic, model accuracy remains unchanged, as confirmed by the authors’ evaluation.

---

What relevance does this have for Grafy users?

Grafy’s Code studio lets developers spin up a CUDA‑enabled notebook or container directly in the browser. By cloning the FOSCAR repository inside a Grafy Code environment, you can compile the kernels, run the benchmark suite, and experiment with your own KV‑cache quantization pipeline—all without leaving the platform. Researchers who prefer to read the full paper can upload the PDF to Reader and ask follow‑up questions to the AI assistant, making it easy to explore the math behind affine folding.

For creators building AI‑driven applications on Grafy, the speedups mean lower latency when generating long texts, which can improve user experience in chatbots, story generators, or real‑time transcription tools.

---

Frequently asked questions

Does FOSCAR require specific GPU architectures?

FOSCAR is designed for NVIDIA Hopper (H100) and Blackwell (B200) GPUs because it leverages architecture‑specific instructions such as wgmma and tcgen05. The kernels compile on other recent NVIDIA GPUs but will not achieve the same speedups, as the warp‑specialized pipelines are tuned for those two generations.

Can FOSCAR be used with 8‑bit or higher‑precision caches?

The current implementation focuses on INT2 and INT4 quantization, where the instruction‑issue bottleneck is most severe. While the affine‑folding principle could be adapted to higher‑precision formats, the performance benefit would be smaller because dequantization overhead is less dominant at 8 bits or above.

Is the affine folding technique applicable to other matrix‑multiply workloads?

Yes. Affine folding works whenever a linear transformation (scale and offset) is applied element‑wise to one operand of a matrix multiplication. This includes certain vision models, recommendation systems, or any workload that uses per‑group quantization. Adapting the method would require custom kernels that respect the target hardware’s instruction set.

---

Source: Arxiv AI Recent — FOSCAR: Accelerating Decode Attention over 2- and 4-Bit KV Caches through Affine Folding