Sub-4-Bit Microscaling (NVFP4, MXFP4 & Walsh-Hadamard Rotation)

When dropping from 8-bit (FP8 E4M3) to 4-bit (W4A4) Tensor Core execution, traditional integer quantization (INT4-AWQ, INT4-GPTQ) hits a severe accuracy wall on multi-step reasoning and coding benchmarks. The culprit is Emergent Activation Outliers—extreme channel spikes (50x–100x larger than neighboring channels) that emerge dynamically after the SwiGLU (SiLU_and_Mul) activation immediately preceding down_proj.

Eluvix integrates hardware-accelerated OCP MXFP4 and NVIDIA Blackwell NVFP4 (E2M1) two-level microscaling paired with Orthogonal Walsh-Hadamard Outlier Rotation (QuaRot / SpinQuant / SVDQuant) and Quantization-Aware Switchyard Routing.

1. The `E2M1` 4-Bit Floating-Point Representation

Both OCP MXFP4 and NVIDIA NVFP4 standardize on the 4-bit E2M1 element format (1 sign bit, 2 exponent bits, 1 mantissa bit, exponent bias b = 1). Because sub-block scales are finite and pre-validated, E2M1 wastes zero bit patterns on Inf or NaN—all 16 bit patterns encode exact numerical values:

E2M1 Representable Magnitudes:
{ 0.0, 0.5, 1.0, 1.5, 2.0, 3.0, 4.0, 6.0 }  (and negative counterparts)

2. OCP MXFP4 (`E8M0`) vs. NVIDIA Blackwell NVFP4 (`FP8 E4M3` + `FP32`)

While both formats store E2M1 elements, their block scaling hierarchies differ fundamentally:

Why `NVFP4` Eliminates the `E8M0` Power-of-2 Sawtooth Cliff

In standard OCP MXFP4, the 8-bit E8M0 scale has zero mantissa bits—it can only represent pure powers of two (2^k). Whenever a 32-element block's maximum magnitude slightly exceeds 6.0 * 2^k (for example, 6.1), the E8M0 scale is forced to double to 2^(k+1), collapsing the effective dynamic range of the entire 32-element block into [0, 3.05] and wasting the 4.0 and 6.0 E2M1 steps.

NVIDIA Blackwell NVFP4 replaces E8M0 with an FP8 E4M3 Level-1 micro-scale (3 mantissa bits, providing 8 fractional sub-steps between every power of two) across a tighter 16-element micro-block, backed by a per-tensor FP32 Level-2 global scale, cutting block quantization Mean Squared Error (MSE) by ~2.4x inside the 5th-Generation Tensor Core (tcgen05.mma.kind::mxf4nvf4).

3. Orthogonal Walsh-Hadamard Outlier Rotation (`R1..R4`) & `SVDQuant`

Even with 16-element micro-blocks, a 100x activation outlier channel after SiLU_and_Mul will crush the 15 normal channels sharing its block. By exploiting orthogonal invariance:

Y = X * W^T = (X * Q) * (W * Q)^T   where Q * Q^T = I

QuaRot and SpinQuant multiply activations X and weights W by a normalized Walsh-Hadamard Matrix (H_d with entries ±1/sqrt(d)). This spreads a single 120.0 channel spike evenly across all d channels (damping peak amplitude by sqrt(8192) ≈ 90.5x) via a multiplier-free O(d log2 d) Fast Walsh-Hadamard Transform (FWHT) butterfly kernel:

R1 (Offline Residual Stream Rotation): Absorbed offline into QKV_proj, o_proj, up_gate_proj, and down_proj weights across RMSNorm (zero runtime latency).
R2 (Offline Per-Head Value/Output Rotation): Block-diagonal H_128 folded offline into W_V and W_O (zero runtime latency).
R3 (Online Post-RoPE Q/K Rotation): Fused into the RoPE kernel to enable 4-bit NVFP4 KV-Cache without attention sink degradation (<3 us).
R4 (Online Fused SiLU_and_Mul Epilogue): Fused O(d log2 d) register-level FWHT inside the SwiGLU kernel immediately before down_proj (<8 us), optionally paired with SVDQuant rank-32 FP16 LoRA residual branches (W = Q_hat + L1 * L2).

4. Llama-3.3-70B Accuracy & Throughput Benchmark

| Precision Format | Weight Memory | MMLU-Pro | MATH-500 (Reasoning) | LiveCodeBench | Relative Throughput | | :--- | :--- | :--- | :--- | :--- | :--- | | BF16 (Baseline) | 140.0 GB (TP=2+) | 68.9% | 76.4% | 54.2% | 1.00x | | FP8 (W8A8 E4M3) | 70.0 GB | 68.7% | 76.1% | 54.0% | 1.85x | | Weight-Only INT4-AWQ (W4A16) | 36.2 GB | 67.1% | 71.2% (-5.2%) | 50.8% | 2.10x (BW Only) | | Naive OCP MXFP4 (No Rotation) | 37.2 GB | 66.4% | 68.9% (-7.5%) | 49.1% | 3.35x | | Blackwell NVFP4 + Hadamard (W4A4) | 39.4 GB (Single GPU TP=1) | 68.6% | 76.1% (99.6% Rec) | 53.9% | 3.42x (Compute + BW) |

5. Quantization-Aware Switchyard Policy (`quantizationPolicy`)

Configure quantizationPolicy on the Switchyard Node or via the x-eluvix-quantization request header:

disabled: Standard routing across all quantized and unquantized upstreams.
prefer-nvfp4-hadamard: Prioritizes Blackwell NVFP4 / MXFP4 + Walsh-Hadamard workers (3.42x throughput at 99.6% reasoning accuracy retention).
lossless-reasoning: Automatically excludes naive/unrotated 4-bit upstreams (INT4-AWQ, INT4-GPTQ, or unrotated MXFP4) so complex math and coding prompts never suffer activation outlier collapse.

curl https://127.0.0.1:3588/api/v1/chat/completions \
  -H "Authorization: Bearer $ELUVIX_API_KEY" \
  -H "Content-Type: application/json" \
  -H "x-eluvix-quantization: lossless-reasoning" \
  -d '{
    "model": "auto",
    "messages": [{"role": "user", "content": "Derive the orthogonal invariance of Walsh-Hadamard rotation across RMSNorm..."}]
  }'

Responses include x-eluvix-quantization-format, x-eluvix-outlier-rotation, and x-eluvix-tensor-core-ptx telemetry headers.

Sub-4-Bit Microscaling (NVFP4, MXFP4 & Walsh-Hadamard Rotation)

1. The E2M1 4-Bit Floating-Point Representation

2. OCP MXFP4 (E8M0) vs. NVIDIA Blackwell NVFP4 (FP8 E4M3 + FP32)

Why NVFP4 Eliminates the E8M0 Power-of-2 Sawtooth Cliff