How FP8 Quantization Works in Modern Inference Engines
Learn how FP8 quantization works in LLM inference, including E4M3 and E5M2 data formats, scaling granularities, and asynchronous tensor core datapaths.
Introduction
In modern large language model (LLM) serving, inference throughput is fundamentally bound by high-bandwidth memory (HBM) bandwidth during the auto-regressive decode phase and by raw compute throughput during long-context prefill operations. While 16-bit floating-point formats (FP16 and BF16) long served as the de facto standard for training and inference pipelines, the compute and memory footprints of dense and mixture-of-experts (MoE) models have made lower precision arithmetic indispensable. FP8 quantization has established itself as the dominant standard across accelerator architectures including NVIDIA Hopper (SM90a), Blackwell (SM100), and AMD Instinct MI300X.
Unlike integer quantization schemes such as INT8 or INT4—which impose uniform quantization step sizes that often collapse when confronting activation outliers—FP8 quantization preserves high dynamic range through exponential bit allocations. Operating with FP8 halves the memory traffic required for weights and activations, effectively doubling the operational intensity of memory-bound operations, while doubling theoretical Tensor Core math throughput relative to 16-bit baselines. Realizing these performance gains without inducing model degradation requires an understanding of representation formats, scaling granularities, and the asynchronous hardware datapaths executing low-precision matrix multiplications.
FP16 (16-bit):
[ S (1) | E E E E E (5) | M M M M M M M M M M (10) ]
Bias: 15. Dynamic Range: ~10^-5 to 6.5 x 10^4. Mantissa: High Precision.
FP8 E4M3 (8-bit):
[ S (1) | E E E E (4) | M M M (3) ]
Bias: 7. Dynamic Range: ~1.95 x 10^-3 to 448. Mantissa: Moderate Precision.
FP8 E5M2 (8-bit):
[ S (1) | E E E E E (5) | M M (2) ]
Bias: 15. Dynamic Range: ~1.52 x 10^-5 to 57344. Mantissa: Low Precision.
Numerical Precision: E4M3 vs E5M2 in FP8 Quantization
The Open Compute Project (OCP) specification defines two primary 8-bit floating-point encodings: E4M3 (1 sign bit, 4 exponent bits, 3 mantissa bits) and E5M2 (1 sign bit, 5 exponent bits, 2 mantissa bits). Each format represents an explicit trade-off between numerical precision (signal-to-quantization-noise ratio, or SQNR) and dynamic range.
The E4M3 Format
The E4M3 format reserves 4 bits for the exponent and 3 bits for the mantissa, utilizing an exponent bias of 7. The value of an E4M3 bit vector is governed by:
$$V = (-1)^S \times 2^{E - 7} \times \left(1 + \sum_{i=1}^{3} b_{3-i} 2^{-i}\right) \quad \text{for } 0 < E < 15$$
For subnormal values ($E = 0$):
$$V = (-1)^S \times 2^{-6} \times \left(\sum_{i=1}^{3} b_{3-i} 2^{-i}\right)$$
A critical architectural distinction of E4M3 relative to IEEE-754 standards is the absence of infinite ($\pm\infty$) representations. The bit pattern 0bS1111111 is explicitly assigned to represent NaN (Not a Number), while all other bit combinations represent normalized or subnormal real numbers. Reclaiming the exponent pattern 15 allows E4M3 to reach a maximum absolute value of:
$$V_{\max} = 1.111_2 \times 2^{15-7} = 1.875 \times 256 = 448$$
The minimum positive normalized value is $2^{-6} \approx 0.015625$, while the minimum positive subnormal value is $2^{-9} \approx 0.001953125$. Because E4M3 dedicates three bits to the mantissa, it retains adequate precision for forward-pass activations and weight matrices, where values cluster predictably around zero within bounded distributions.
The E5M2 Format
The E5M2 format replicates the 5-bit exponent field of IEEE-754 FP16, combined with an exponent bias of 15 and 2 bits of mantissa. It adheres to standard floating-point conventions: an exponent field of 0b11111 with a zero mantissa denotes $\pm\infty$, while a non-zero mantissa denotes NaN.
The maximum representable finite value in E5M2 is:
$$V_{\max} = 1.11_2 \times 2^{30-15} = 1.75 \times 32768 = 57344$$
The minimum normalized positive value is $2^{-14} \approx 6.10 \times 10^{-5}$, and the smallest subnormal value is $2^{-16} \approx 1.52 \times 10^{-5}$. While E5M2 possesses an expansive dynamic range capable of absorbing volatile gradient spikes during backward propagation, its 2-bit mantissa delivers only 4 discrete quantization bins between consecutive powers of two. Consequently, using E5M2 for inference weights or activations introduces substantial truncation noise, rendering E4M3 the standard choice for inference GEMM computations.
| Property | FP16 (IEEE-754) | FP8 E4M3 (OCP) | FP8 E5M2 (OCP) |
|---|---|---|---|
| Exponent Bits | 5 | 4 | 5 |
| Mantissa Bits | 10 | 3 | 2 |
| Exponent Bias | 15 | 7 | 15 |
| Maximum Real Value | 65504 | 448 | 57344 |
| Smallest Normal | $\approx 6.10 \times 10^{-5}$ | $\approx 1.56 \times 10^{-2}$ | $\approx 6.10 \times 10^{-5}$ |
| Smallest Subnormal | $\approx 5.96 \times 10^{-8}$ | $\approx 1.95 \times 10^{-3}$ | $\approx 1.52 \times 10^{-5}$ |
| Primary Domain | Baseline Compute | Forward Weights & Activations | Backward Gradients / Wide Dynamic Range |
Scaling Granularities: Per-Tensor, Per-Token, and Block-Level Scaling
Because the finite dynamic range of E4M3 ($[-448, 448]$) cannot encompass the unnormalized distribution of multi-layer activations, inputs must be scaled before casting. The quantization transformation applies a scalar scale factor $S$:
$$X_{\text{fp8}} = \text{clip}\left(\left\lfloor \frac{X}{S} \right\rceil, -448, 448\right)$$
During matrix multiplication, the original precision is mathematically recovered via the dual scale product:
$$C = (A_{\text{fp8}} B_{\text{fp8}}) \cdot (S_A \cdot S_B)$$
The choice of scaling granularity dictates both the numerical fidelity of the model and the computational complexity of the runtime execution path.
Per-Tensor: [ Matrix A ] * S_A (Scalar)
Per-Token: [ Row 0: Tokens ] * S_0
[ Row 1: Tokens ] * S_1
[ Row 2: Tokens ] * S_2
Block-Level: [ Block 0 (1x128) ] [ Block 1 (1x128) ]
Scale: s_0 Scale: s_1
1. Static Per-Tensor Scaling
In static per-tensor quantization, an invariant scalar $S \in \mathbb{R}$ is calibrated offline for each weight tensor and activation tensor:
$$S = \frac{\max(|X|)}{V_{\max}}$$
While this scheme minimizes serving latency by eliminating runtime reduction overheads, it suffers severe degradation when activation channels exhibit systemic outliers. In architectures such as Llama-3 and Mistral, isolated hidden dimensions can produce activations that exceed median channel amplitudes by up to $100\times$. Under a static per-tensor scale, these outlier features force $S$ to expand, collapsing the remaining $99%$ of the tensor's values into zero or near-zero quantization bins.
2. Dynamic Per-Token (Per-Row) Scaling
To decouple token representations from divergent intra-batch activation swings, dynamic per-token scaling computes the scale factor $S_i$ over each sequence token $i$ online during inference:
$$S_i = \frac{\max_{j} |X_{i, j}|}{448}$$
This dynamic reduction ensures that prompt tokens and decoding tokens occupy the full numerical precision of E4M3. However, dynamic per-token scaling introduces an explicit reduction pass prior to GEMM execution. In systems scaling to thousands of concurrent requests across high-throughput distributed engines—such as those explored in Disaggregated Inference Architecture: Paged KV-Cache Tiering, Chunked Prefill, and Asynchronous Overlap—these extra memory sweeps can saturate memory bandwidth if not cleanly fused into antecedent layer norms.
3. Block-Level Scaling (Microscaling / MX Formats)
Modern hardware architectures (such as NVIDIA Blackwell and the OCP Microscaling Formats standard) implement block-level scaling. Instead of a single scale factor covering an entire row, vectors along the inner contraction dimension $K$ are partitioned into fixed-width sub-blocks (typically $B = 32$ or $B = 128$ consecutive elements).
Each sub-block $k$ has its own localized scale factor $s_k$:
$$X_{\text{fp8}}^{(k)} = \text{clip}\left(\left\lfloor \frac{X^{(k)}}{s_k} \right\rceil, -V_{\max}, V_{\max}\right), \quad s_k = \frac{\max(|X^{(k)}|)}{V_{\max}}$$
By confining activation outliers to their immediate 32-element or 128-element blocks, block-level scaling prevents outlier corruption from spilling across unrelated channels, maintaining model perplexity within $0.1%$ of FP16 baselines without manual outlier clipping.
Hardware Acceleration: Tensor Core Datapaths and TMA Pipelines
FP8 arithmetic cannot be analyzed purely as a representation format; its primary value lies in dedicated silicon execution paths. In modern architectures such as NVIDIA Hopper (SM90a), matrix multiplication is decoupled from standard thread execution via hardware-level asynchronous pipelines.
+-----------------------------------------------------------------------+
| Global Memory (HBM3e) |
+-----------------------------------------------------------------------+
|
| Asynchronous Transfer (TMA)
v
+-----------------------------------------------------------------------+
| Shared Memory (SRAM) |
| +---------------------------------------------------------------+ |
| | Swizzled FP8 Tiles: A Tile [M, K] | B Tile [K, N] | |
| +---------------------------------------------------------------+ |
+-----------------------------------------------------------------------+
|
| Warpgroup Matrix Multiply (WGMMA)
v
+-----------------------------------------------------------------------+
| SM90a Tensor Cores (4 Warps = 128 Threads) |
| +---------------------------------------------------------------+ |
| | FP8 Input Pairs --> [ Math Engine ] --> FP32 Accumulators | |
| +---------------------------------------------------------------+ |
+-----------------------------------------------------------------------+
|
| Scale Application & Fused Epilogue
v
+-----------------------------------------------------------------------+
| Registers / Output Writeback (FP16 / FP8) |
+-----------------------------------------------------------------------+
Tensor Memory Accelerator (TMA)
Prior GPU generations forced individual CUDA threads within a warp to issue manual load instructions (LDG.E) to transfer matrix tiles from HBM to Shared Memory (SRAM). On SM90a, the Tensor Memory Accelerator (TMA) executes multi-dimensional tensor copies asynchronously via hardware descriptor tables. A single thread issues a tma.load instruction specifying coordinate offsets; the hardware unit handles address translation, stride calculation, boundary clamping, and out-of-bounds zero-padding natively.
Because FP8 elements occupy a single byte, loads achieve twice the spatial data density per memory cycle compared to FP16, saturating memory crossbars while minimizing execution queue occupancy on the streaming multiprocessors (SMs).
WGMMA Instructions and Bank Conflict Avoidance
Inside the SM, FP8 GEMM is driven by Warpgroup Matrix Multiply and Accumulate (wgmma.mma_async) instructions. Unlike legacy warp-level matrix instructions (wmma or mma.sync) where 32 threads coordinate in lockstep, wgmma orchestrates four contiguous warps (128 threads) working as a single cooperative warpgroup.
Crucially, Tensor Cores do not accumulate in FP8. Internally, the hardware loads 8-bit floating-point pairs, parses the exponent and mantissa fields, executes hardware-fused multiply-adds, and accumulates the intermediate result directly into FP32 registers:
$$D_{\text{fp32}} = \sum_{k=0}^{K-1} \left(A_{\text{fp8}}[m, k] \times B_{\text{fp8}}[k, n]\right) + C_{\text{fp32}}[m, n]$$
This accumulation precision eliminates the numerical drift and underflow cascades that would otherwise occur if intermediate partial products were kept in 8-bit storage across long reduction loops ($K \ge 4096$).
To feed 128 threads simultaneously without stall cycles, Shared Memory is partitioned into 32 banks, each 4 bytes wide. Loading unaligned 1-byte FP8 elements can trigger pathological 32-way bank conflicts. Hopper addresses this by pairing the TMA with hardware Shared Memory Swizzling (swizzle_128b). The hardware applies an XOR-based permutation on the shared memory address bits:
$$\text{Bank Index} = \left(\frac{\text{Byte Address}}{4}\right) \oplus \left(\frac{\text{Byte Address}}{128}\right) \pmod{32}$$
This spatial shuffling guarantees that contiguous threads read across distinct physical banks simultaneously, ensuring that high-throughput asynchronous execution pipelines remain completely compute-bound.
Writing a Block-Scaled FP8 GEMM Kernel in Triton
To understand how software interfaces with underlying FP8 silicon, consider a modern Triton implementation. The following kernel implements a block-scaled FP8 matrix multiplication ($C = A \cdot B$) where input tensors are stored in tl.float8e4nv (E4M3), scaling factors are passed per-block, and accumulation occurs in tl.float32.
import triton
import triton.language as tl
@triton.jit
def fp8_gemm_kernel(
# Pointers to Matrices
a_ptr, b_ptr, c_ptr,
# Pointers to Scale Factors (Per-Block)
scale_a_ptr, scale_b_ptr,
# Matrix Dimensions
M, N, K,
# Strides for Matrix Traversals
stride_am, stride_ak,
stride_bk, stride_bn,
stride_cm, stride_cn,
stride_scale_am, stride_scale_ak,
stride_scale_bk, stride_scale_bn,
# Meta-parameters
BLOCK_SIZE_M: tl.constexpr,
BLOCK_SIZE_N: tl.constexpr,
BLOCK_SIZE_K: tl.constexpr,
):
"""
High-performance block-scaled FP8 GEMM kernel.
Accumulates intermediate products in FP32, then dequantizes via
the dual scaling product: C = (A_fp8 * B_fp8) * (Scale_A * Scale_B)
"""
pid_m = tl.program_id(axis=0)
pid_n = tl.program_id(axis=1)
# Offset calculations for tile coordinates
offs_am = (pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)) % M
offs_bn = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % N
offs_k = tl.arange(0, BLOCK_SIZE_K)
# Base pointers for the current warpgroup block
a_ptrs = a_ptr + (offs_am[:, None] * stride_am + offs_k[None, :] * stride_ak)
b_ptrs = b_ptr + (offs_k[:, None] * stride_bk + offs_bn[None, :] * stride_bn)
scale_a_ptrs = scale_a_ptr + (offs_am[:, None] * stride_scale_am)
scale_b_ptrs = scale_b_ptr + (offs_bn[None, :] * stride_scale_bn)
# FP32 accumulator to maintain precision across the contraction dimension
accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
# Main contraction loop over K
for k_idx in range(0, tl.cdiv(K, BLOCK_SIZE_K)):
# Load FP8 tiles from SRAM / Global Memory
a = tl.load(a_ptrs, mask=offs_k[None, :] < (K - k_idx * BLOCK_SIZE_K), other=0.0)
b = tl.load(b_ptrs, mask=offs_k[:, None] < (K - k_idx * BLOCK_SIZE_K), other=0.0)
# Tensor Core invocation (wgmma): casting occurs in HW, accumulating to FP32
accumulator += tl.dot(a, b, out_dtype=tl.float32)
# Advance matrix pointers along K dimension
a_ptrs += BLOCK_SIZE_K * stride_ak
b_ptrs += BLOCK_SIZE_K * stride_bk
# Epilogue: Load scale factors and dequantize
scale_a = tl.load(scale_a_ptrs)
scale_b = tl.load(scale_b_ptrs)
# Combined scalar dequantization factor
combined_scale = scale_a * scale_b
output = accumulator * combined_scale
# Writeback converted result to global memory
offs_cm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
offs_cn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
c_ptrs = c_ptr + (offs_cm[:, None] * stride_cm + offs_cn[None, :] * stride_cn)
c_mask = (offs_cm[:, None] < M) & (offs_cn[None, :] < N)
tl.store(c_ptrs, output.to(tl.float16), mask=c_mask)
In this kernel, tl.dot(a, b, out_dtype=tl.float32) tells the Triton compiler to emit the native PTX wgmma instruction for the SM90a target. The software accumulator loop maintains floating-point stability across thousands of reduction steps, while the dequantization scale factors are multiplied in the epilogue stage directly before casting the final activation to the destination type (FP16 or BF16).
Systems Integration: Activation Quantization, Epilogues, and Overheads
Deploying FP8 quantization into production serving systems involves significantly more than executing isolated GEMMs. The surrounding framework must minimize memory round-trips and manage host-side dispatch overheads.
Standard Non-Fused Pipeline (Memory-Bound Overhead):
[RMSNorm] ---> Write FP16 ---> Read FP16 ---> [FP8 Quant] ---> Write FP8 ---> Read FP8 ---> [GEMM]
(HBM) (HBM) (HBM)
Fused Production Pipeline:
[RMSNorm + Dynamic FP8 Quantization] ---> Write FP8 Tile ---> [WGMMA GEMM Engine]
(SRAM) (Tensor Cores)
1. Fused Kernel Epilogues and Dynamic Quantization
In an un-optimized pipeline, converting activations to FP8 requires reading the FP16 output of a preceding layer from HBM, computing the max-absolute value reduction to find $S$, converting the elements to E4M3, and writing the quantized bytes back to HBM—only for the subsequent GEMM to read them back. This round-trip degrades operational intensity.
Production inference runtimes (such as TensorRT-LLM and vLLM) fuse dynamic quantization into the antecedent kernel's epilogue. For example, the RMSNorm or LayerNorm kernel calculates the normalized output, evaluates the row-level scaling factor via intra-warp shuffle primitives (__shfl_xor_sync), casts the normalized tokens directly to E4M3 in registers, and deposits the FP8 values directly into shared memory buffers or contiguous global memory for the next layer.
2. Runtime Scheduling and Memory Orchestration
When scaling inference across large-scale distributed clusters, host-side synchronization jitter can easily stall execution engines running low-latency FP8 kernels. Serving architectures often leverage actor-based orchestration to manage multi-stream workloads efficiently, avoiding runtime dispatch bubbles. System designs such as those documented in High-Throughput Deterministic Actor Runtimes: Cache-Conscious Architectures, Lock-Free Rings, and Kernel-Bypass I/O emphasize pinning worker threads and utilizing kernel-bypass synchronization to prevent CPU-side bottlenecks from starving ultra-fast GPU Tensor Cores.
3. Distributed Interconnect and Network Serialization
As models scale across multiple nodes via Tensor Parallelism (TP) or Pipeline Parallelism (PP), communication overhead across high-speed interconnects (NVLink and InfiniBand) can offset the compute gains of FP8. Modern serving engines resolve this by communicating directly in FP8.
Rather than gathering or scattering activations in FP16/BF16 and down-casting locally, engines serialize intermediate representations in their quantized 8-bit form directly over the wire. This effectively halves network bus bandwidth consumption during inter-node All-Gather and All-Reduce collectives. These networked serving strategies are particularly critical in large-scale disaggregated prefill-decode topologies, as explored in Disaggregated LLM Serving: Architecting Asynchronous Prefill-Decode Engines with Zero-Copy RDMA KV-Cache Migration.
Memory Footprint and Throughput Analysis
The mechanical advantages of FP8 can be quantified by analyzing memory bandwidth consumption and theoretical peak FLOPS on modern silicon. The roofline model defines the maximum achievable performance as:
$$\text{Attainable FLOPS} = \min\left(\text{Peak Compute FLOPS}, , \text{Operational Intensity} \times \text{Memory Bandwidth}\right)$$
Where operational intensity is measured in arithmetic operations per byte transferred:
$$I = \frac{\text{Floating Point Operations (FLOPs)}}{\text{Memory Access (Bytes)}}$$
For a standard GEMM operating on matrix dimensions $M$ (batch/token count), $N$ (hidden dimension), and $K$ (projection dimension):
$$\text{Arithmetic Work} = 2 \cdot M \cdot N \cdot K$$
$$\text{Data Transferred (FP16)} = 2 \cdot (M \cdot K + K \cdot N + M \cdot N) \text{ bytes}$$
$$\text{Data Transferred (FP8)} = 1 \cdot (M \cdot K + K \cdot N) + 2 \cdot (M \cdot N) \text{ bytes}$$
(assuming 16-bit accumulator write-back).
Operational Intensity vs Matrix Size (K=4096, N=4096, M variable):
Arithmetic Intensity (FLOPs / Byte)
^
| FP8 (2x Higher Ceiling)
| -------------------------------
| /
| / FP16 Baseline
| ----------------------------------
| /
| /
| /
+------------------------------------------------------------------>
0 (Decode: M=1) 128 512 (Prefill: M>=1024)
Token Context / Batch Size (M)
In memory-bound decoding regimes where $M = 1$, operational intensity drops significantly:
$$I_{\text{decode}} \approx \frac{2 \cdot 1 \cdot N \cdot K}{1 \cdot K \cdot N} = 2 \text{ FLOPs/Byte (for FP8)}$$
$$\text{Versus: } I_{\text{decode}} \approx \frac{2 \cdot 1 \cdot N \cdot K}{2 \cdot K \cdot N} = 1 \text{ FLOP/Byte (for FP16)}$$
Because memory bandwidth directly limits token generation speed during the decode phase, cutting the weight footprint from 2 bytes to 1 byte per parameter doubles generation throughput, independent of compute core frequency. In the compute-bound prefill phase ($M \ge 512$), FP8 doubles the theoretical execution ceiling on NVIDIA Hopper from 989 TFLOPS (BF16 Tensor Core limit) to 1,978 TFLOPS (FP8 Tensor Core limit), slashing time-to-first-token (TTFT) metrics for massive context windows.
Conclusion
FP8 quantization represents a fundamental evolution in deep learning systems architecture. By shifting from uniform integer quantization to structured, exponential representations (specifically E4M3 for forward execution and E5M2 for wide-dynamic-range gradients), modern inference engines circumvent the severe precision loss historically caused by activation outliers.
Realizing the full potential of FP8 requires a holistic systems approach: choosing the appropriate scaling granularity (whether dynamic per-token or block-level microscaling), pairing workloads with asynchronous hardware datapaths like TMA and WGMMA instructions on SM90a/SM100 silicon, and fusing quantization steps directly into adjacent layer norm epilogues. As context sizes and model parameters continue to expand toward the multi-trillion threshold, the hardware and software mechanics of FP8 serve as the foundational execution blueprint powering contemporary AI serving infrastructure.
References
- Open Compute Project. (2023). OCP 8-bit Floating Point Specification (OFP8) Revision 1.0. https://www.opencompute.org/documents/ocp-8-bit-floating-point-specification-ofp8-r1-0-2023-06-20-pdf
- NVIDIA Corporation. (2022). NVIDIA Hopper Architecture In-Depth. NVIDIA Developer Technical Blog. https://developer.nvidia.com/blog/nvidia-hopper-architecture-in-depth/
- Tillet, P., Kung, H. T., & Cox, D. (2019). Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations. GitHub Repository. https://github.com/triton-lang/triton