FuTransformer: FP8 prefill and INT4 decode for Llama-style inference on Hopper

What it claims

FuTransformer is a set of Hopper CUDA kernels with separate prefill and decode paths for a Llama-style block. The abstract adds several results the summary omits. Prefill reaches 1.29 PFLOP/s on one H100. INT4 decode is 31–51% faster than the authors' own FP8 decode path. The rotated INT4 KV cache adds 0.20 WikiText-2 perplexity on top of the weight-quantization loss. The kernels match the quantized reference computation to within 0.001 perplexity, so the accuracy loss comes from quantization, not kernel error.

How it works

Prefill builds on CODA's GEMM-plus-epilogue design for training. It uses persistent FP8 WGMMA GEMMs with 128×256×128 CTA tiles, a four-stage TMA pipeline, one producer warpgroup and two consumer warpgroups. Epilogue policies are chosen at compile time and add four inference-specific operations:

• KV-cache writes. The QKV epilogue applies RoPE and writes K and V straight into cache slots. This removes the dense K/V tensors and the separate scatter kernel.
• Folded RMSNorm gain. The fixed gain is folded into the next projection's weights before quantization. The row factor is applied in the epilogue, using sums of squares collected by the residual-producing epilogue.
• Delayed FP8 scaling. The gate/up epilogue writes FP8 directly, using the maximum from the previous invocation. The paper puts this at MI bytes, against 5MI for writing BF16 and quantizing it separately.
• Packed prefill. Variable-length requests run without padding, using a prefix-sum row map. Packed and padded inputs get separately compiled kernels.

Decode cannot use Hopper's tensor cores directly at low precision. Hopper has no INT4 MMA, and its FP8 WGMMA shapes are too large for a few query rows. FuTransformer therefore expands INT4 values into 16-bit mma.sync.m16n8k16 operands.

For weights, it uses one symmetric scale per output channel and 128-input group. Nibbles sit in B-fragment order. Each nibble is placed into the mantissa of BF16 128.0, and one packed FMA then subtracts 136. Group scales are applied after accumulation.

Two kernel families cover decode. A register-streaming kernel handles batch 1. A skinny kernel with split-K handles batches 2 to 32.

The KV cache keeps a BF16 primary copy plus an INT4 mirror. The mirror has one FP32 scale (amax/7) per token and KV head. Keys are stored in fragment order and expanded through a 256-entry shared-memory lookup table that maps each byte to two FP16 values. V stays in channel order and is transposed through shared memory. A block-diagonal Walsh–Hadamard rotation is applied to each cache row when it is written, and to the query and output once per step. Its cost does not grow with context length. Split-KV handles long contexts.

Evidence

All results are on one H100 80GB SXM5 with CUDA 13.0 and PyTorch 2.14.1+cu130.

• Against CODA, gate/up plus SwiGLU at the 8B shape: the BF16 path is 1.13× faster than the released CODA kernel. The FP8 path is 2.02–2.18× faster than BF16 CODA. That second comparison mixes precisions, and the paper labels it as such.
• Decode, 8B and 70B shapes, batch 1–32: INT4 decode is 2.1–3.3× faster than the compiled, CUDA-graphed BF16 PyTorch block. It is 1.3–6.3× faster than the same block using PyTorch's INT4 kernel. That kernel is built for one query row and is slower than PyTorch's own BF16 block at batch 32.
• Shapes: the paper's table gives 1B, 8B, 70B and 405B layer shapes, from 60.82M to 3187.67M projection weights per layer. For the 8B shape, weights per layer are 436 MB in BF16, 218 MB in E4M3 and 109 MB in INT4.

Limits

• Decode kernels cover only 1 ≤ M ≤ 32.
• The INT4 cache is an additional copy alongside the BF16 cache. It does not replace it.
• The authors list conditions that delayed FP8 scaling needs: a valid initial scale, a defined scale for all-zero tensors, and separate scale state per concurrent request group. They say production use also needs saturation counters, because a sudden change in value range can make the previous scale unsuitable.
• Accuracy is reported only as WikiText-2 perplexity on Llama-3.1-8B.
• The fetched text stops at the experimental setup. It gives no code location, licence or install steps.