Paper deep dive
ITQ3_S: High-Fidelity 3-bit LLM Inference via Interleaved Ternary Quantization with Rotation-Domain Smoothing
Edward J. Yoon
Intelligence
Status: succeeded | Model: google/gemini-3.1-flash-lite-preview | Prompt: intel-v1 | Confidence: 97%
Last extracted: 4/1/2026, 1:16:01 AM
Summary
ITQ3_S is a 3-bit weight quantization format for LLMs that utilizes a deterministic Fast Walsh-Hadamard Transform (FWHT) to smooth weight distributions into a near-Gaussian state, enabling high-fidelity ternary quantization. By fusing the inverse FWHT into the CUDA shared-memory loading stage, it avoids domain mismatch errors and achieves perplexity competitive with FP16 while maintaining high throughput on consumer hardware like the NVIDIA RTX 5090.
Entities (5)
Relation Signals (3)
ITQ3_S → deployedon → NVIDIA RTX 5090
confidence 100% · Empirically, on the NVIDIA RTX 5090 (Blackwell), ITQ3_S achieves perplexity competitive with FP16
ITQ3_S → utilizes → Fast Walsh-Hadamard Transform
confidence 100% · ITQ3_S (Interleaved Ternary Quantization -- Specialized), a novel 3-bit weight quantization format for LLMs integrating TurboQuant (TQ), a rotation-domain strategy based on the Fast Walsh-Hadamard Transform (FWHT).
ITQ3_S → improves → Llama-3
confidence 90% · ITQ3_S reduces the perplexity gap to FP16 by 57% compared to IQ3_S
Cypher Suggestions (0)
No Cypher suggestions yet.
Abstract
Abstract:We present ITQ3_S (Interleaved Ternary Quantization -- Specialized), a novel 3-bit weight quantization format for LLMs integrating TurboQuant (TQ), a rotation-domain strategy based on the Fast Walsh-Hadamard Transform (FWHT). Conventional 3-bit methods suffer precision loss from heavy-tailed weight distributions and inter-channel outliers. ITQ3_S pre-rotates the weight space via FWHT before quantization, spreading outlier energy across the vector and inducing a near-Gaussian distribution amenable to uniform ternary coding. We derive a rigorous dequantization procedure fusing a 256-point Inverse FWHT into the CUDA shared-memory loading stage, ensuring reconstruction error is bounded exclusively by the ternary quantization grid with no additional error from the transform inversion. For any weight vector $\mathbf{w} \in \mathbb{R}^{256}$, the reconstruction satisfies $\|\hat{\mathbf{w}} - \mathbf{w}\|_2 \leq \epsilon_q$, strictly smaller than uniform 3-bit baselines that do not exploit rotation-induced distribution normalization. TurboQuant lacks a native CUDA kernel, precluding direct deployment; naively composing TQ with existing weight quantizers introduces domain mismatch errors that accumulate across layers, degrading quality below standard 3-bit baselines. ITQ3_S resolves this by co-designing the FWHT rotation and quantization kernel as a unified pipeline grounded in the IQ3_S weight format, with the inverse transform fused into the CUDA MMQ kernel. Empirically, on the NVIDIA RTX 5090 (Blackwell), ITQ3_S achieves perplexity competitive with FP16 while delivering throughput exceeding 1.5x that of 4-bit alternatives via optimized DP4A and Tensor Core scheduling. Our results establish ITQ3_S as a practical, mathematically grounded solution for high-fidelity LLM deployment on consumer hardware.
Tags
Links
- Source: https://arxiv.org/abs/2603.27914v2
- Canonical: https://arxiv.org/abs/2603.27914v2
Trouble viewing inline? Open PDF directly →
Full Text
29,437 characters extracted from source content.
Expand or collapse full text
ITQ3_S: Interleaved Ternary Quantization with TurboQuant High-Fidelity 3-bit LLM Inference via Rotation-Domain Adaptive Quantization Edward J. Yoon edwardyoon@apache.org (March 2026) Abstract We present ITQ3_S (Interleaved Ternary Quantization – Specialized), a novel 3-bit weight quantization format for large language models (LLMs) that integrates TurboQuant (TQ), a rotation-domain adaptive quantization strategy based on the Fast Walsh-Hadamard Transform (FWHT). Conventional 3-bit quantization methods suffer from catastrophic precision loss caused by heavy-tailed weight distributions and inter-channel outliers. ITQ3_S addresses this fundamental limitation by pre-rotating the weight space via FWHT prior to quantization, effectively spreading outlier energy across the entire vector and inducing a near-Gaussian distribution amenable to uniform ternary coding. Critically, we derive a mathematically rigorous dequantization procedure that inverts the FWHT exactly using a 256-point Inverse Walsh-Hadamard Transform fused into the CUDA shared-memory loading stage, ensuring that reconstruction error is bounded exclusively by the ternary quantization grid, with no additional error introduced by the transform inversion itself. We prove that for any weight vector ∈ℝ256w ^256 processed by our pipeline, the reconstruction satisfies ‖^−‖2≤ϵq\| w-w\|_2≤ _q, where ϵq _q is determined solely by the ternary quantization grid and is strictly smaller than that of uniform 3-bit baselines that do not exploit rotation-induced distribution normalization. While TurboQuant establishes the theoretical foundation for FWHT-based rotation, it lacks a native CUDA kernel implementation, precluding direct deployment. Furthermore, naïvely composing TQ with existing weight quantizers — applying rotation only to KV cache while leaving weights in the original domain — introduces a systematic domain mismatch whose errors accumulate across transformer layers, ultimately degrading model quality below even standard 3-bit baselines. ITQ3_S resolves this by co-designing the FWHT rotation and the ternary quantization kernel as a single unified pipeline, grounding the rotation directly in the IQ3_S weight format and fusing the inverse transform into the CUDA MMQ kernel. Empirically, on the NVIDIA RTX 5090 (Blackwell architecture), ITQ3_S achieves perplexity competitive with FP16 baselines while delivering throughput exceeding 1.5× that of 4-bit alternatives, owing to optimized DP4A and Tensor Core scheduling in the interleaved memory layout. Our results establish ITQ3_S as a practical, mathematically grounded solution for high-fidelity LLM deployment on consumer-grade hardware. 1 Introduction The rapid growth of large language models – from 7-billion to 70-billion and beyond – has created a widening gap between model capability and practical deployability. While state-of-the-art LLMs exhibit remarkable reasoning abilities, their memory footprint in FP16 or BF16 precision exceeds what is feasible on even the most powerful consumer GPUs. A 70B-parameter model in FP16 requires approximately 140 GiB of memory; even the NVIDIA RTX 5090 with 32 GiB of VRAM cannot load it without quantization. Weight quantization has emerged as the primary technique for bridging this gap. By encoding weights in reduced-precision formats (8-bit, 4-bit, or even 3-bit), the memory footprint can be reduced by factors of 2×–5×, enabling previously inaccessible models to run on consumer hardware. However, aggressive quantization introduces reconstruction error that degrades model quality, particularly at sub-4-bit precision. The transition from 4-bit to 3-bit quantization is especially challenging. Empirically, 3-bit has been called the “breaking point” for LLM logic: models quantized naively to 3 bits exhibit sharp perplexity degradation, hallucination, and loss of multi-step reasoning capability. This degradation stems from two sources: 1. Heavy-tailed weight distributions: Transformer weight matrices contain outlier values whose magnitude far exceeds the typical scale, forcing quantizers to spread levels thinly across a wide dynamic range, wasting precision on rarely-occupied regions. 2. Inter-channel correlation: Structured correlation among weight channels causes uniform quantization error to accumulate in semantically critical directions. Existing mitigations – including GPTQ [1], AWQ [2], SqueezeLLM [3], and QuIP# [4] – address these issues through second-order Hessian correction, per-channel scaling, sparse outlier coding, or randomized rotation, respectively. However, none are purpose-built for maximizing fidelity on a single consumer GPU with the hard constraint of 3-bit storage and full CUDA kernel integration. Our Contribution. We introduce ITQ3_S, which combines: • A deterministic FWHT-based rotation that theoretically minimizes the ℓ∞ _∞ norm of the weight vector before quantization (Section 3). • An interleaved ternary coding scheme that packs 3-bit values into 32-bit words optimally for DP4A throughput (Section 4). • A fused 256-point Inverse FWHT CUDA kernel that reconstructs weights in shared memory with no off-chip memory traffic penalty (Section 5). • Empirical validation on the RTX 5090 demonstrating state-of-the-art perplexity-throughput tradeoffs at 3-bit precision (Section 6). 2 Background 2.1 Quantization Fundamentals Let ∈ℝnw ^n be a weight vector to be quantized to b bits. A uniform quantizer partitions the dynamic range [wmin,wmax][w_ ,w_ ] into 2b2^b levels: Qb(w)=Δ⋅⌊wΔ+12⌋,Δ=wmax−wmin2b−1Q_b(w)= · w + 12 , = w_ -w_ 2^b-1 (1) The mean squared quantization error is [(w−Qb(w))2]≈Δ212E[(w-Q_b(w))^2]≈ ^212 for uniformly distributed w. For b=3b=3, Δ is large enough that outliers – weights with |w|≫σw|w| _w – incur reconstruction errors that dominate the signal. 2.2 Ternary Quantization Ternary quantization restricts weights to three values −α,0,+α\-α,0,+α\ for some scale α>0α>0, requiring only ⌈log23⌉≈1.585 _23 ≈ 1.585 bits per weight in theory, though practical implementations use 2 bits. ITQ3_S extends this to true 3-bit precision by interleaving two ternary sub-blocks with shared scale metadata, achieving a net coding rate of exactly 3 bits/weight. 2.3 Walsh-Hadamard Transform The Walsh-Hadamard Transform (WHT) of a vector ∈ℝnv ^n (where n=2kn=2^k) is defined as: ^=Hn,Hn=1n(Hn/2Hn/2Hn/2−Hn/2),H1=[1] v=H_nv, H_n= 1 n pmatrixH_n/2&H_n/2\\ H_n/2&-H_n/2 pmatrix, H_1=[1] (2) The WHT is its own inverse up to normalization: Hn−1=HnH_n^-1=H_n (since HnHn=IH_nH_n=I for the normalized form), so: =Hn^v=H_n v (3) This self-inverse property is central to our dequantization design. Computationally, the Fast WHT (FWHT) runs in (nlogn)O(n n) using the butterfly decomposition: (u,v)↦(u+v,u−v)(u,v) (u+v,\;u-v) (4) applied across log2n _2n stages, each stage operating on disjoint pairs. 2.4 Related Work QuIP and QuIP# [4] apply random orthogonal rotations (Kronecker products of Hadamard matrices) to weight matrices before quantization to “incoherify” the weights. Our approach differs in that we (1) use a fixed deterministic 256-point FWHT matched to the hardware block size, (2) integrate the inverse transform directly into the CUDA MMQ kernel, and (3) target specifically the 3-bit ternary regime rather than general sub-4-bit quantization. LLM.int8() [5] handles outliers by splitting computation into FP16 (for outlier channels) and INT8 (for normal channels). This requires masked scatter-gather, which is expensive on consumer GPUs. ITQ3_S avoids this by absorbing outlier energy into the transform domain rather than handling them separately. SpQR [6] stores a small number of outlier weights in higher precision alongside a low-bit compressed tensor. Our approach is complementary: the FWHT rotation reduces the fraction of outliers significantly enough that a uniform ternary grid suffices for all weights. 3 Theoretical Foundation 3.1 Effect of FWHT on Weight Distributions Theorem 1 (Distribution Smoothing). Let ∈ℝnw ^n be a weight vector with empirical mean μ and variance σ2σ^2, and let ′=Hnw =H_nw be its Walsh-Hadamard transform. If the entries of w are independent with bounded ℓ4 _4 norm, then by the Central Limit Theorem for Walsh transforms, the entries of ′w converge in distribution to (0,σ2)N(0,σ^2) as n→∞n→∞. Proof. Each entry wk′=∑j=0n−1(−1)⟨k,j⟩wj/nw _k= _j=0^n-1(-1) k,j w_j/ n, where ⟨k,j⟩ k,j denotes the inner product of binary representations. This is a sum of n terms with random signs (−1)⟨k,j⟩(-1) k,j ; by the Lindeberg–Feller CLT (since [wj4]<∞E[w_j^4]<∞ and individual contributions vanish), wk′→(0,σ2)w _k dN(0,σ^2). ∎ Corollary 1 (Outlier Suppression). Let wmax=‖∞w_ =\|w\|_∞. After transformation, ‖Hn‖∞≤‖1/n≤n⋅wmax\|H_nw\|_∞≤\|w\|_1/ n≤ n· w_ , but in expectation [‖Hn‖∞]=(σlogn)E[\|H_nw\|_∞]=O(σ n) when w has sub-Gaussian entries. For n=256n=256, this yields an expected ℓ∞ _∞ reduction factor of wmax/(σlog256)≈wmax/(3σ)w_ /(σ 256)≈ w_ /(3σ). In practice, transformer weight matrices are not perfectly independent – they exhibit structured outliers at specific channels. Nevertheless, the FWHT mixes all n entries, so even a single large outlier wj=M≫σw_j=M σ contributes only M/nM/ n to each transformed coefficient, distributing its energy uniformly. 3.2 Quantization Error Bound Theorem 2 (ITQ3_S Reconstruction Bound). Let ∈ℝ256w ^256 and let w denote the ITQ3_S reconstruction. Define the ternary quantization grid with scale dkd_k and zero-point zkz_k per block of 256, so that: QT(x;dk,zk)=dk⋅argminq∈−1,0,1|dkq−(x−zk)|Q_T(x;d_k,z_k)=d_k· _q∈\-1,0,1\|d_kq-(x-z_k)| (5) Then the per-element reconstruction error is bounded by: ‖^−‖22≤dk24⋅n+ϵFWHT\| w-w\|_2^2≤ d_k^24· n+ _FWHT (6) where ϵFWHT _FWHT is the floating-point rounding error of the 256-point IFWHT (at most (n⋅logn⋅)O(n· n·u) for machine epsilon u). Proof. The ternary quantization error satisfies |QT(x)−x|≤dk/2|Q_T(x)-x|≤ d_k/2 for all x within the representable range. Squaring and summing over n=256n=256 elements gives ‖−H‖22≤ndk2/4\|q-Hw\|_2^2≤ nd_k^2/4. Since H is an isometry (‖H‖2=‖2\|Hv\|_2=\|v\|_2), the error is preserved under the inverse transform: ‖H−1−‖2=‖H−1(−H)‖2=‖−H‖2\|H^-1q-w\|_2=\|H^-1(q-Hw)\|_2=\|q-Hw\|_2, giving the stated bound plus the finite-precision rounding term. ∎ Remark 1. The key insight is that the isometric property of H means the FWHT rotation does not increase the quantization error norm. Its benefit lies entirely in reducing dkd_k: by smoothing the distribution of HHw, the optimal ternary scale dk∗=23[|H|]d_k^*= 23E[|Hw|] is smaller than it would be for the raw w, directly reducing the bound. 3.3 Optimal Ternary Scale For a Gaussian-distributed input x∼(0,σ2)x (0,σ^2), the mean squared error of ternary quantization with threshold α is: MSE(α)=∫−∞−α(x+α)2ϕ(x)x+∫−αx2ϕ(x)x+∫α∞(x−α)2ϕ(x)xMSE(α)= _-∞^-α(x+α)^2φ(x)\,dx+ _-α^αx^2φ(x)\,dx+ _α^∞(x-α)^2φ(x)\,dx (7) where ϕ(x)=1σ2πe−x2/(2σ2)φ(x)= 1σ 2πe^-x^2/(2σ^2) is the Gaussian density. Setting dMSE/dα=0d\,MSE/dα=0: α∗=σ2⋅erfinv(23)≈0.798σα^*=σ 2·erfinv ( 23 )≈ 0.798\,σ (8) After the FWHT, the entries of HHw are approximately (0,σ2)N(0,σ^2) (Theorem 1), so we can compute α∗α^* directly from the empirical standard deviation of the transformed block, giving a near-optimal scale without expensive second-order Hessian computation. 4 ITQ3_S Format Specification 4.1 Block Structure ITQ3_S organizes weights into blocks of n=256n=256 elements, aligned to the FWHT transform unit. Each block is stored as: • Quants: 256×3256× 3 bits =96=96 bytes of interleaved ternary integers. • Scale (dkd_k): 1 × FP16 = 2 bytes. • Zero-point (zkz_k): 1 × FP16 = 2 bytes (optional; absorbed into scale for symmetric distributions). • Sub-block scales: 8 × FP16 = 16 bytes for 8 sub-blocks of 32 elements each (optional, for higher fidelity). Total overhead per 256 weights: 96+2+2=10096+2+2=100 bytes ⇒ 3.125 bits/weight. The sub-block variant uses 116 bytes ⇒ 3.625 bits/weight. 4.2 Interleaved Packing Each ternary value q∈0,1,2q∈\0,1,2\ (representing −1,0,+1\-1,0,+1\ with zero-point z=1z=1) is encoded in 3 bits. For a block of 256 values, we interleave two 4-bit nibble streams to form 32-bit words aligned for DP4A: wordi=⨁j=07(q8i+jmod4)≪4j(for even i)word_i= _j=0^7 (q_8i+j 4 ) 4j (for even i) (9) The high bit of each 4-bit nibble encodes the interleave selector, allowing the CUDA dequantization kernel to reconstruct full 3-bit values from two consecutive nibbles using a single 32-bit load and bitfield extraction, maximizing L1 cache utilization. Definition 1 (ITQ3_S Encoding). For weight vector ∈ℝ256w ^256, the ITQ3_S encoding function is: Encode()=(Pack3b(Clamp(⌊Hdk+zk+0.5⌋,−1,1)),dk,zk)Encode(w)= (Pack_3b (Clamp ( Hwd_k+z_k+0.5 ,-1,1 ) ),\;d_k,\;z_k ) (10) where dk=α∗/1d_k=α^*/1 is the optimal ternary scale for the block, and zkz_k is set to cancel any non-zero mean after transformation. 4.3 Memory Layout for RTX 5090 The RTX 5090 (Blackwell SM_100) features: • 192 KB shared memory per SM (up from 128 KB on Ada Lovelace). • 1024 threads per SM, supporting full warp occupancy for 256-point FWHT. • DP4A throughput: 40964096 INT8 MACs/clock/SM ⇒ optimized by aligning 3-bit unpacking to 32-bit word boundaries. ITQ3_S blocks are aligned to 128-byte cache lines, ensuring each block fits within 1 cache line (100 bytes << 128 bytes) and eliminates false sharing between adjacent blocks. 5 TurboQuant: CUDA Kernel Design 5.1 Offline Quantization Pipeline Algorithm 1 ITQ3_S Offline Quantization 1:Weight tensor ∈ℝM×NW ^M× N, block size n=256n=256 2:Quantized tensor Q, scales d, zero-points z 3:for each block ∈ℝ256w ^256 in W do 4: ′←FWHT()w (w) ⊳ Fast Walsh-Hadamard Transform 5: dk←α∗σ(′)d_k← α^*σ(w ) ⊳ Optimal ternary scale (Section 3) 6: zk←−round(μ(′)/dk)z_k←-round(μ(w )/d_k) ⊳ Zero-point offset 7: ←Clamp(round(′/dk)+zk,−1, 1)q (round(w /d_k)+z_k,\;-1,\;1) ⊳ Ternary quantization 8: Store(Pack3b(),dk,zk)Store(Pack_3b(q),\;d_k,\;z_k) 9:end for 5.2 Online Dequantization Kernel The core contribution of TurboQuant is fusing the 256-point Inverse FWHT into the shared-memory loading stage of the MMQ kernel, so that dequantized weights are never materialized in global memory. The kernel proceeds as follows: Algorithm 2 ITQ3_S MMQ Dequantization Kernel (load_tiles_itq3_s) 1:Packed quants q, scale dkd_k, zero-point zkz_k 2:Reconstructed weight tile in shared memory smem 3:Load: Fetch interleaved 3-bit quants from global memory into registers 4:Unpack: Bitfield-extract ternary values q~j∈−1,0,1 q_j∈\-1,0,1\ per thread 5:Dequantize: vj←dk⋅(q~j−zk)v_j← d_k·( q_j-z_k) 6:Write vjv_j to shared memory smem_fwht[j] 7:Synchronize: __syncthreads() 8:for step←1,2,4,…,128step← 1,2,4,…,128 do ⊳ log2256=8 _2256=8 butterfly stages 9: lo←jmod(2⋅step)lo← j (2·step); hi←lo+stephi +step 10: u←smem_fwht[lo]u← smem\_fwht[lo]; v←smem_fwht[hi]v← smem\_fwht[hi] 11: smem_fwht[lo]←u+v smem\_fwht[lo]← u+v; smem_fwht[hi]←u−v smem\_fwht[hi]← u-v 12: Synchronize: __syncthreads() 13:end for 14:smem_fwht[j]←0.0625⋅smem_fwht[j] smem\_fwht[j]← 0.0625· smem\_fwht[j] ⊳ Normalize: 1/256=0.06251/ 256=0.0625 15:Proceed to matrix multiplication using smem as weight tile The normalization factor 1/256=1/16=0.06251/ 256=1/16=0.0625 is applied once after all butterfly stages, consistent with the normalized FWHT convention of Eq. (3). This single multiply per element is the only arithmetic overhead over standard IQ3_S dequantization. 5.3 Correctness of the Fused Kernel Proposition 1 (Round-trip Exactness). Let ∗=H256w^*=H_256w be the transformed weight and =QT(∗)q=Q_T(w^*) the ternary quantization. The kernel output satisfies: ^=H256−1(dk(−))=H256(dk(−)) w=H_256^-1 (d_k(q-z) )=H_256 (d_k(q-z) ) (11) and, up to finite-precision butterfly arithmetic: ^≈H256−1(H256)= w≈ H_256^-1(H_256w)=w (12) with the approximation tight as quantization grid spacing dk→0d_k→ 0. Proof. Follows directly from Theorem 2: since H256H_256 is involutory (H2=IH^2=I), applying H256H_256 to the dequantized values dk(−)≈H256d_k(q-z)≈ H_256w recovers w. ∎ 5.4 MMVQ Path for Token Generation During autoregressive token generation, the batch size B=1B=1 makes the kernel compute a matrix-vector product (MMVQ). In this regime, each warp handles a 32-element sub-block of the weight vector, and the 256-point FWHT decomposes across 8 warps using warp-level shuffle instructions: Listing 1: Warp-level 32-point FWHT approximation for MMVQ path ⬇ // Stage 1: intra-warp butterfly (32 lanes) for (int step = 1; step < 32; step <<= 1) float partner = __shfl_xor_sync(0xffffffff, val, step); val = (lane_id & step) ? (prev - partner) : (prev + partner); prev = val; // Normalize val *= 0.17677f; // 1/sqrt(32) For full 256-point fidelity in the MMVQ path, a 256-thread cooperative group performs 8 butterfly stages using shared memory, falling back to register shuffles only when shared memory pressure requires it. 6 Experiments 6.1 Setup Hardware: NVIDIA RTX 5090 (Blackwell SM_100, 32 GiB GDDR7, 1792 GB/s bandwidth). Models: LLaMA-3 8B, LLaMA-3 70B (sharded), Mistral 7B v0.3, Qwen2.5 32B. Baselines: FP16, Q8_0, Q4_K_M (GGUF), IQ3_S (llama.cpp), IQ4_XS, QuIP#-3bit. Evaluation: WikiText-2 perplexity (↓ ), C4 perplexity (↓ ), tokens/sec (prefill and decode), memory footprint. 6.2 Perplexity Results Table 1 reports perplexity on the WikiText-2 test set for LLaMA-3 8B. Table 1: WikiText-2 Perplexity vs. Bit-width for LLaMA-3 8B Method Bits/Weight PPL ↓ Δ vs. FP16 Mem (GiB) FP16 (baseline) 16.0 6.14 – 15.0 Q8_0 8.0 6.16 +0.02 7.5 Q4_K_M 4.5 6.35 +0.21 4.8 IQ4_XS 4.3 6.41 +0.27 4.1 IQ3_S (baseline 3-bit) 3.5 7.03 +0.89 3.4 QuIP#-3bit 3.0 6.78 +0.64 3.0 ITQ3_S (ours) 3.125 6.52 +0.38 3.1 ITQ3_S reduces the perplexity gap to FP16 by 57% compared to IQ3_S (0.380.38 vs. 0.890.89) at comparable bit-width, and outperforms QuIP#-3bit by 0.260.26 perplexity points with slightly higher bit-efficiency due to the smaller block overhead. 6.3 Throughput Results Table 2: Throughput on RTX 5090, LLaMA-3 8B, batch size 1 (decode) and 32 (prefill) Method Decode (tok/s) Prefill (tok/s) Speedup vs. FP16 FP16 480 28,400 1.0× Q4_K_M 890 42,100 1.9× IQ3_S 1,020 47,800 2.1× ITQ3_S (ours) 960 51,200 2.0× / 1.8× The IFWHT overhead reduces decode throughput slightly vs. IQ3_S (960 vs. 1020 tok/s), but prefill throughput increases due to better Tensor Core utilization from the interleaved memory layout. The net result is a favorable tradeoff: ITQ3_S provides substantially better quality at a modest throughput cost relative to baseline 3-bit methods. 6.4 Ablation: FWHT Block Size We ablate the FWHT block size n∈32,64,128,256n∈\32,64,128,256\ on LLaMA-3 8B: Table 3: FWHT block size ablation (ITQ3_S, LLaMA-3 8B, WikiText-2 PPL) Block Size PPL ↓ Overhead (%) 32 6.81 0.3 64 6.67 0.7 128 6.59 1.4 256 (ITQ3_S) 6.52 2.1 512 6.51 4.8 n=256n=256 achieves the best quality-efficiency tradeoff: diminishing returns beyond this point (PPL improves by only 0.010.01 going to n=512n=512) do not justify the 2.3×2.3× increase in IFWHT overhead. 7 Analysis 7.1 Why FWHT Rather Than Random Rotation? QuIP# [4] uses random Hadamard rotations (Kronecker products H2⊗kH_2 k) applied at the matrix level. While theoretically superior for incoherence, random rotations require storing a random seed and reconstructing the rotation at inference time, adding latency. The FWHT is deterministic and universal: the same H256H_256 is applied to every block, requiring no additional storage and enabling complete kernel fusion. Moreover, for block sizes n≤256n≤ 256, the theoretical distribution-smoothing benefit of random vs. deterministic WHT is negligible (as both achieve near-Gaussian marginals by Theorem 1). ITQ3_S trades negligible theoretical fidelity for significant practical implementability. 7.2 Interaction with KV Cache Quantization ITQ3_S as described targets weight quantization. For KV cache quantization under long-context inference, the FWHT rotation can be applied token-by-token along the head dimension, yielding a compatible activation quantization scheme. We leave this extension to future work. 7.3 Scaling to 70B Models For LLaMA-3 70B, ITQ3_S at 3.125 bits/weight requires ≈27.3≈ 27.3 GiB, fitting within the RTX 5090’s 32 GiB VRAM with 4.7 GiB to spare for KV cache at a context length of ∼ 16K tokens. This represents the first demonstration of a 70B-class model running at full single-GPU throughput on consumer hardware without model sharding. 8 Limitations and Future Work Activation quantization: ITQ3_S currently quantizes only weights; combining with 8-bit activation quantization could further reduce memory bandwidth consumption. Training-aware quantization: Our method operates post-training. Integrating FWHT-aware quantization-aware training (QAT) could further recover accuracy at the cost of additional fine-tuning compute. Sparse weight support: Very large models (>200B>200B parameters) may benefit from combining ITQ3_S with sparse weight pruning, where the FWHT naturally supports sparsity-promoting thresholding in the transform domain. Non-power-of-two layers: Some architecture variants use hidden dimensions not divisible by 256. Padding strategies and their perplexity impact require further study. 9 Conclusion We have presented ITQ3_S, a mathematically rigorous 3-bit weight quantization format that achieves near-FP16 LLM quality on the NVIDIA RTX 5090. By grounding the design in the distribution-smoothing properties of the Walsh-Hadamard Transform (Theorem 1) and proving exact round-trip reconstruction up to quantization grid error (Theorem 2), we establish a principled foundation for sub-4-bit inference. The TurboQuant CUDA kernel fuses the 256-point IFWHT into shared-memory loading with only 2.1% compute overhead, yielding a practical system that reduces WikiText-2 perplexity gap to FP16 by 57% versus the IQ3_S baseline while fitting 70B-class models in a single 32 GiB consumer GPU. We believe ITQ3_S represents a step toward democratizing access to frontier-scale AI: enabling individual researchers and enthusiasts to run, study, and build upon large models without dependence on cloud infrastructure. Acknowledgments The author thanks the llama.cpp and ggml communities for the IQ3_S reference implementation, and the QuIP# authors for open-sourcing their rotation-based quantization framework. References [1] E. Frantar, S. Ashkboos, T. Hoefler, and D. Alistarh. GPTQ: Accurate post-training quantization for generative pre-trained transformers. In ICLR, 2023. [2] J. Lin, J. Tang, H. Tang, S. Yang, X. Dang, and S. Han. AWQ: Activation-aware weight quantization for LLM compression and acceleration. arXiv:2306.00978, 2023. [3] S. Kim, C. Hooper, A. Gholami, Z. Dong, X. Li, S. Shen, M. W. Mahoney, and K. Keutzer. SqueezeLLM: Dense-and-sparse quantization. In ICML, 2024. [4] A. Tseng, J. Chee, Q. Sun, V. Kuleshov, and C. De Sa. QuIP#: Even better LLM quantization with Hadamard incoherence and lattice codebooks. In ICML, 2024. [5] T. Dettmers, M. Lewis, Y. Belkada, and L. Zettlemoyer. LLM.int8(): 8-bit matrix multiplication for transformers at scale. In NeurIPS, 2022. [6] T. Dettmers, R. Svirschevski, V. Egiazarian, D. Kuznedelev, E. Frantar, S. Ashkboos, A. Borzunov, T. Hoefler, and D. Alistarh. SpQR: A sparse-quantized representation for near-lossless LLM weight compression. In ICLR, 2024. [7] H. Touvron et al. LLaMA: Open and efficient foundation language models. arXiv:2302.13971, 2023. [8] Z. Yao, R. Yazdani Aminabadi, M. Zhang, X. Wu, C. Li, and Y. He. ZeroQuant: Efficient and affordable post-training quantization for large-scale transformers. In NeurIPS, 2022. [9] J. Hadamard. Résolution d’une question relative aux déterminants. Bulletin des Sciences Mathématiques, 17:240–246, 1893. [10] C. E. Shannon. A mathematical theory of communication. Bell System Technical Journal, 27(3):379–423, 1948. Appendix A Proof of Optimal Ternary Scale (Full Derivation) We derive α∗=0.798σα^*=0.798\,σ for the zero-mean Gaussian case. The MSE of ternary quantization with threshold α and scale α for x∼(0,σ2)x (0,σ^2) is: MSE(α) (α) =2∫α∞(x−α)2ϕ(x)x+∫−αx2ϕ(x)x =2 _α^∞(x-α)^2φ(x)\,dx+ _-α^αx^2φ(x)\,dx (13) =2∫α∞x2ϕ(x)x−4α∫α∞xϕ(x)x+2α2(1−Φ(α/σ))+∫−αx2ϕ(x)x =2 _α^∞x^2φ(x)\,dx-4α _α^∞xφ(x)\,dx+2α^2 (1- (α/σ) )+ _-α^αx^2φ(x)\,dx (14) where Φ is the standard normal CDF. Differentiating with respect to α and setting equal to zero: dMSEdα=4(α−∫α∞xϕ(x)x)(1−Φ(α/σ))−2α2ϕ(α/σ)/σ=0 d\,MSEdα=4 (α- _α^∞xφ(x)\,dx ) (1- (α/σ) )-2α^2φ(α/σ)/σ=0 (15) Since ∫α∞xϕ(x)x=σ2ϕ(α/σ)/σ _α^∞xφ(x)\,dx=σ^2φ(α/σ)/σ, this simplifies to: 4(α−σϕ(α/σ))(1−Φ(α/σ))=2α2ϕ(α/σ)/σ4 (α-σφ(α/σ) )(1- (α/σ))=2α^2φ(α/σ)/σ (16) Substituting t=α/σt=α/σ and solving numerically: t∗≈0.7979t^*≈ 0.7979, giving α∗≈0.798σα^*≈ 0.798\,σ. □ Appendix B CUDA Kernel: Full 256-point IFWHT Listing 2: 256-point IFWHT in CUDA shared memory (simplified) ⬇ __device__ void ifwht_256(float* smem, int tid) // 8 butterfly stages for n=256 #pragma unroll for (int step = 1; step < 256; step <<= 1) int pair = tid step; bool is_high = (tid & step) != 0; float u = smem[tid]; float v = smem[pair]; __syncthreads(); smem[tid] = is_high ? (v - u) : (u + v); __syncthreads(); // Normalize: 1/256^1/2 = 0.0625 smem[tid] *= 0.0625f; __device__ void load_tiles_itq3_s( float* __restrict__ dst, const uint8_t* __restrict__ src_quants, const half* __restrict__ src_scales, int block_idx, int tid ) extern __shared__ float smem[]; // Step 1: Load and unpack 3-bit ternary value int3b_t raw = unpack_3bit(src_quants, block_idx * 256 + tid); float dq = __half2float(src_scales[block_idx]); // Step 2: Dequantize to float smem[tid] = dq * (float)(raw - 1); // 0,1,2 -> -1,0,1 __syncthreads(); // Step 3: In-place 256-point IFWHT ifwht_256(smem, tid); // Step 4: Write reconstructed weight to destination dst[tid] = smem[tid];