FP8 Mixed-Precision Serving in Production: Comparing E4M3 vs. E5M2 Formats, Per-Tensor vs. Block-Wise Dynamic Scaling, FP8 KV Cache, and Tensor Core GEMM Economics

Serving large language models at enterprise scale requires balancing computational throughput against high-bandwidth memory (HBM) capacity. While 16-bit floating-point formats (FP16 and BF16) remain the standard for model pre-training and fine-tuning, their memory footprint and arithmetic bandwidth create severe bottlenecks during high-concurrency production inference. Deploying models in INT8 or INT4 reduces memory pressure but often introduces non-trivial quantization latency and accuracy loss on sensitive attention layers.

Native 8-bit floating-point (FP8) execution has emerged as the production standard across modern accelerator architectures, including NVIDIA Hopper (H100, H200), Blackwell (B200), Ada Lovelace (L40S), and AMD Instinct (MI300X). FP8 matrix multiplication kernels deliver up to twice the raw theoretical TFLOPS of 16-bit operations while cutting parameter storage and Key-Value (KV) cache footprints in half.

Implementing FP8 serving in production requires navigating architectural trade-offs between numeric representation formats (E4M3 versus E5M2), quantization granularities (per-tensor static scaling, delayed dynamic scaling, and block-wise tile scaling), and memory cache quantization strategies.

The Mathematics of FP8: E4M3 vs. E5M2

Standardized by the Open Compute Project (OCP) Microscaling Formats specification and supported in hardware Tensor Cores, FP8 comprises two distinct 8-bit floating-point representations tailored for different stages of the deep learning pipeline:

E4M3 Bit Layout (High Precision, Narrow Dynamic Range):
[Sign: 1 bit] [Exponent: 4 bits (bias = 7)] [Mantissa: 3 bits]
Max Value: +/- 448 | Min Positive Normal: 2^-6 ≈ 0.015625 | Dynamic Range: ~4.8 orders of magnitude

E5M2 Bit Layout (Low Precision, Wide Dynamic Range):
[Sign: 1 bit] [Exponent: 5 bits (bias = 15)] [Mantissa: 2 bits]
Max Value: +/- 57344 | Min Positive Normal: 2^-14 ≈ 6.10e-5 | Dynamic Range: ~10.2 orders of magnitude

The mathematical formula for an FP8 value is:

$$x = (-1)^s \times 2^{e - \text{bias}} \times \left(1 + \sum_{i=1}^m b_i 2^{-i}\right)$$

where $s$ is the sign bit, $e$ is the unsigned integer exponent, $\text{bias}$ is the exponent bias (7 for E4M3, 15 for E5M2), and $m$ is the number of mantissa bits.

Core Architectural Differences

• E4M3 (1 sign, 4 exponent, 3 mantissa): Provides 3 bits of fraction precision with an exponent bias of 7, supporting values up to $\pm 448$. It does not encode positive or negative infinities (all exponent and mantissa 1s represent NaN). The unit roundoff is $2^{-4} = 0.0625$. It serves as the primary format for inference weights, forward-pass activations, and KV caches.
• E5M2 (1 sign, 5 exponent, 2 mantissa): Preserves the same 5-bit exponent field and bias (15) as IEEE 754 half-precision (FP16), supporting values up to $\pm 57344$ along with standard representations for positive/negative infinities and NaN. The unit roundoff is $2^{-3} = 0.125$. It is primarily utilized for backward-pass gradients during training and unnormalized attention score logits where values span many orders of magnitude.

According to research published in NVIDIA's FP8 Formats for Deep Learning and empirical evaluations across hardware accelerators (arXiv:2502.01070), E4M3 consistently outperforms E5M2 in quantization accuracy across LLM inference workloads. Because matrix multiplications in inference are bounded and normalized via LayerNorm/RMSNorm, preserving mantissa precision is critical for maintaining output distribution fidelity and minimizing perplexity drift.

Quantization Granularity and Scaling Topologies

Because FP8 has a narrow dynamic range compared to 16-bit formats, input tensors must be scaled before conversion to maximize the utilization of available representational bins. Three scaling topologies govern production inference engines:

1. Per-Tensor Static Scaling

In static post-training quantization (PTQ), a single scaling factor $s \in \mathbb{R}$ is precomputed for each weight tensor and activation channel using a calibration dataset:

$$X_{\text{FP8}} = \text{clip}\left(\left\lfloor \frac{X_{\text{FP32}}}{s} \right\rceil, -V_{\max}, V_{\max}\right)$$

$$s = \frac{\max(|X_{\text{calibration}}|)}{V_{\max}}$$

where $V_{\max} = 448$ for E4M3. During inference, Tensor Core GEMM operations execute directly without runtime scale reduction overhead. However, static scaling is susceptible to activation outlier degradation: if out-of-distribution user prompts generate activation spikes, clipping results in representation collapse.

2. Delayed Scaling vs. Just-in-Time Dynamic Scaling

In dynamic scaling, the scale factor is computed on live activations at runtime. The NVIDIA Transformer Engine library defines two operational strategies:

• Delayed Scaling (DelayedScaling): Instead of computing a costly full-tensor maximum reduction kernel before every GEMM, the serving runtime maintains a rolling history buffer of absolute maximum values $H = \{h_t\}_{t=1}^\ell$ over $\ell$ iterations (e.g., $\ell = 1024$). The scale is calculated as:

$$\rho(X) = \frac{\text{FP8MaxValue}}{2^m \cdot \max(H)}$$ where $m$ is an optional safety margin. This approach amortizes reduction latency to near zero but risks numerical overflow if sudden batch distribution shifts occur.

• Just-in-Time Current Scaling (Float8CurrentScaling): Computes the exact absolute maximum $\max(|X_t|)$ of the active tensor immediately prior to the GEMM kernel launch. While adding a lightweight reduction pass, it eliminates overflow risks during high-variance serving traffic.

3. Block-Wise and Tile Scaling (DeepSeek-V3 / FP8 Tiling)

To address the cross-channel outlier problem without falling back to high-overhead per-token per-channel scaling, modern architectures such as DeepSeek-V3 introduce fine-grained block-wise FP8 quantization.

In this layout, weights and activations are partitioned into localized 2D sub-matrices (e.g., $128 \times 128$ GEMM tiles or $1 \times 128$ row blocks). A dedicated FP32 scale factor is assigned to each individual block:

$$X_{\text{block}} = \text{clip}\left(\left\lfloor \frac{X_{\text{tile}}}{s_{\text{tile}}} \right\rceil, -448, 448\right), \quad s_{\text{tile}} = \frac{\max(|X_{\text{tile}}|)}{448}$$

This architecture isolates outlier channels within their local tile, preventing a single activation spike from degrading the precision of neighboring features across the entire tensor.

FP8 KV Cache Architecture and Serving Latency

During the autoregressive decoding phase of LLM inference, throughput is bounded by GPU memory bandwidth rather than compute. For long-context workloads (e.g., 32k to 128k context windows), the Key-Value (KV) cache consumes the majority of active HBM.

Storing Key and Value states in E4M3 cuts KV cache memory consumption by 50% relative to standard 16-bit representations (BF16/FP16):

$$\text{Memory}_{\text{KV}} = 2 \times B \times L \times N_{\text{heads}} \times D_{\text{head}} \times \text{BytesPerElement}$$

By halving $\text{BytesPerElement}$ from 2 bytes (BF16) to 1 byte (FP8), serving systems achieve two major operational advantages:

• Doubled Concurrent Batch Capacity: Deployments can accommodate up to $2\times$ more active concurrent requests before triggering cache preemption or host memory swapping.
• Reduced Memory Bandwidth Saturation: Each decoding step requires reading half as many bytes across the memory bus, directly lowering Inter-Token Latency (ITL).

Empirical evaluations published by the vLLM engineering team on production models (including Llama 3.1 8B/70B) demonstrate that FP8 KV cache yields a 14.9% increase in output throughput and a 14.8% reduction in median ITL under high concurrency loads, with negligible degradation on standard Needle-in-a-Haystack retrieval benchmarks.

When combined with Multi-Head Latent Attention (MLA) architectures in engines like SGLang and vLLM, FP8 KV caching expands effective token capacity per 8-GPU node by nearly an order of magnitude.

Framework Implementations and Kernel Topologies

Production serving engines leverage specialized CUTLASS, FlashInfer, and Triton kernels to execute W8A8 (8-bit weights, 8-bit activations) matrix multiplications on hardware Tensor Cores:

• vLLM (v0.7+): Integrates dynamic per-tensor quantization, ModelOpt checkpoints, and Compressed-Tensors formats. It utilizes CUTLASS FP8 GEMM, FlashInfer, and Marlin-FP8 kernels, supporting both E4M3 and E5M2 KV cache formats.
• TensorRT-LLM: Implements static calibration and specialized FP8 GEMM plugins built atop cuBLASLt and custom CUDA kernels, offering optimized execution pipelines for pre-quantized NGC checkpoints.
• SGLang: Features native support for DeepSeek-V3 block-wise FP8 GEMM kernels (FlashMLA, CutlassMLA) and dynamic W8A8 routing, maximizing throughput across large Mixture-of-Experts (MoE) deployments.
• Transformer Engine: NVIDIA's reference library providing DelayedScaling, Float8CurrentScaling, and BlockwiseScaling recipes with automatic format swizzling and layout transformation for GEMM inputs.

On NVIDIA Hopper architecture, FP8 Tensor Cores leverage asynchronous copy operations (cp.async) and warp-specialized pipelines to stream FP8 tiles directly from global memory into shared memory (SRAM), overlapping memory loads with Tensor Core computation.

Engineering Trade-Offs and Best Practices

To deploy FP8 mixed-precision serving in production without compromising model fidelity, engineering teams should apply three operational practices:

1. Selective Layer Precision Retention: The first embedding layer, the final linear projection layer (LM head), and normalization layers (LayerNorm/RMSNorm) exhibit extreme sensitivity to quantization noise. Keeping these components in BF16/FP16 while running all intermediate GEMM projections in W8A8 FP8 eliminates the majority of downstream perplexity degradation.
2. Dynamic Scaling for Multi-Turn Chat Workloads: For production APIs with variable input lengths and multi-turn conversations, use dynamic per-tensor or block-wise scaling rather than static PTQ calibrations to prevent unexpected clipping on out-of-domain prompts.
3. Hardware-Aligned Batch Sizing: FP8 Tensor Cores require matrix dimensions (M, N, K) to be multiples of 16 (and ideally 64 or 128) for optimal warp tiling. Ensuring sequence padding and batch aggregation align with these boundaries prevents kernel fallback to slower unaligned execution paths.

Sources

• Open Compute Project (OCP) Microscaling Formats (MX) Specification v1.0
• NVIDIA: FP8 Formats for Deep Learning (arXiv:2209.05433)
• An Investigation of FP8 Across Accelerators for LLM Inference (arXiv:2502.01070)
• DeepSeek-V3 Technical Report: Architecture and FP8 Mixed Precision Training (arXiv:2412.19437)
• NVIDIA Transformer Engine Documentation: FP8 Primer and Scaling Strategies
• vLLM Engineering: The State of FP8 KV-Cache and Attention Quantization
• SGLang Documentation: DeepSeek Models and FP8 Optimization
• Red Hat Developer: vLLM Brings FP8 Inference to the Open Source Community