Skip to content

Quantization & Hardware Acceleration: Deep Engineering Realities

Numerics: The Bit-Level Reality

The transition from FP32 to lower precision formats fundamentally trades dynamic range and precision for memory bandwidth and FLOPs.

FP32, FP16, and BF16

  • FP32 (IEEE 754): 1 sign bit, 8 exponent bits, 23 mantissa bits. The baseline. Memory bandwidth bound for inference.
  • FP16: 1 sign, 5 exponent, 10 mantissa. Offers higher throughput via Tensor Cores (e.g., A100 HMMA instructions) but suffers from limited dynamic range (max value ~65504). Overflow/underflow is a constant threat in activations.
  • BF16: 1 sign, 8 exponent, 7 mantissa. Retains FP32's dynamic range, eliminating gradient overflow in training and activation overflow in inference. The lost precision in the mantissa behaves as a regularizer during training, but for inference, it acts as a low-cost FP32 drop-in.

FP8 (E4M3 and E5M2)

Hopper (H100) introduced native FP8 support to double throughput over FP16. - E4M3: 1 sign, 4 exponent, 3 mantissa. Narrower dynamic range but higher precision. Used primarily for weight and activation forward passes where value distributions are tightly constrained. - E5M2: 1 sign, 5 exponent, 2 mantissa. Matches FP16's dynamic range but with severe precision loss. Used for gradient communication during training where dynamic range prevents underflow.

Post-Training Quantization (PTQ) vs. Quantization-Aware Training (QAT)

PTQ

Quantizes a pre-trained FP16/BF16 model to INT8/INT4 without retraining. Relies on calibration data to determine scale and zero-point parameters. $$ X_q = \text{round}(X / S) + Z $$ PTQ is fast but suffers from outlier clipping, especially in large language models where activation outliers emerge systematically beyond 6.7B parameters.

QAT

Simulates quantization during training using Straight-Through Estimators (STE). The forward pass uses quantized weights (dequantized on the fly), while the backward pass uses full-precision gradients to update high-precision shadow weights. QAT allows the model to learn a weight distribution robust to the discrete quantization grid, yielding higher accuracy at ultra-low precisions (e.g., INT4 or lower), but requires massive compute equivalent to fine-tuning.

Weight-Only vs. Weight-and-Activation Quantization

Weight-Only Quantization (AWQ, GPTQ)

Targets the HBM bandwidth bottleneck during autoregressive generation (batch size = 1). Since weights must be loaded from HBM to SRAM for every token, reducing weight size is paramount. - GPTQ: Uses second-order information (inverse Hessian) to compensate for quantization error in the remaining unquantized weights. Iteratively quantizes weights column by column. - AWQ (Activation-Aware Weight Quantization): Observes that a small fraction (e.g., 1%) of salient weights drastically impacts performance. These salient weights correspond to large activation channels. AWQ scales up these salient channels and scales down the corresponding activation channels, preserving the FP16 semantics while pushing the weights into a friendlier quantization grid.

Weight-and-Activation Quantization (SmoothQuant)

Required for high-batch-size inference where compute, not just memory bandwidth, becomes the bottleneck. W8A8 (INT8 weights and activations) enables INT8 Tensor Core utilization. - SmoothQuant: Addresses the problem of massive activation outliers in LLMs. It mathematically shifts the quantization difficulty from activations (which are dynamic and hard to quantize) to weights (which are static). By applying a per-channel scaling factor \(s\), it smooths out activation magnitude: \(Y = (X \cdot \text{diag}(s)^{-1}) \cdot (\text{diag}(s) \cdot W)\).

The Hardware Reality: Memory Throughput vs. Compute in INT4

When running an INT4 weight-only model (like AWQ/GPTQ) on A100/H100, the GPU does not execute INT4 matrix multiplications. Hardware INT4 Tensor Cores are sparse or deprecated for this use case.

The Execution Flow: 1. HBM to SRAM: 32-bit registers pack eight INT4 weights. This is where the 4x memory bandwidth savings occurs compared to FP16. In autoregressive decoding, loading weights from HBM dominates the latency equation. 2. SRAM to Registers (Dequantization): Inside the streaming multiprocessor (SM), a custom CUDA kernel unpacks the INT4 weights. Bitwise operations (shifts and masks) extract the 4-bit values. 3. Casting: The INT4 values are cast to FP16 (or BF16) using the pre-computed scale and zero-point parameters loaded alongside the weights. 4. Compute: The actual matrix multiplication occurs on FP16 Tensor Cores (e.g., using mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16).

The Trade-off: We trade ALUs (for shifting, masking, and FP16 scaling) to save HBM read cycles. Because memory bandwidth is orders of magnitude slower than ALU execution in batch-size-1 inference, the SM is entirely stalled waiting for weights anyway. The unpacking math hides perfectly behind the memory latency.

Speculative Decoding Internals

Speculative decoding circumvents the memory bandwidth bottleneck of the target model by guessing multiple tokens sequentially with a smaller "draft" model, then verifying them in parallel with the target model.

  1. Draft Generation: A small draft model (e.g., 1B params) auto-regressively generates a sequence of \(K\) tokens. Because it fits in cache or takes minimal HBM bandwidth, this is fast.
  2. Target Verification: The large target model (e.g., 70B params) runs a single forward pass on the \(K\) draft tokens. Instead of \(K\) separate memory fetches for the weights, the target model's weights are loaded from HBM exactly once to process all \(K\) tokens in parallel (essentially changing the batch size from 1 to \(K\)).
  3. Matching & Rejection Sampling: The target model produces a probability distribution for each of the \(K\) positions. The hardware executes a modified rejection sampling kernel:
  4. If the target's probability for a draft token exceeds a uniform random threshold, it is accepted.
  5. The streak breaks at the first rejected token. The target model's distribution at the failure point is used to sample the corrected token.
  6. Speedup: If the draft model's acceptance rate \(\alpha\) is high enough, the expected number of generated tokens per target model forward pass is \(> 1\). The wall-clock latency per token plummets, as HBM bandwidth is amortized across multiple accepted tokens.