PaperlayerView exact revision on arXiv

Source-linked AI summary

FP64 Is All You Want, INT8 Is All You Need, FP4/6/8 Is All You Have

Pratyai Mazumder, Alexandru Calotoiu, Torsten Hoefler

arXiv:2609.37693v1cs.DC

TL;DR

FP64 throughput has declined while low-precision GPU formats have advanced, and existing Ozaki II variants had been described for particular formats. This paper optimizes scheme selection across formats and derives GEMM-count lower bounds; its Blackwell implementations run faster than native FP64 at the tested sizes, up to 83 times on B300.

  • Problem

    Ozaki II adaptations existed for FP8 and FP4, but schemes had not been given for FP6 or INT4 and prior variants addressed particular formats.

    Source: §I · p. 1

  • Method

    The paper formulates Ozaki II selection as an exact optimization of moduli and product forms for minimum GEMM count, with lower bounds over the family.

    Source: §I · p. 1, A. Contributions · p. 1

  • Results

    The method gives the first FP6 schemes and an FP8 scheme with 30 GEMMs versus the previous best of 42; Blackwell implementations run faster than native FP64 at tested sizes, up to 83 times on B300.

    Source: A. Contributions · p. 1, §IX · p. 11

  • Takeaways & Limitations

    A new low-precision format needs only its values and accumulator to be given a scheme.

    Source: §IX · p. 11

    Minimizing GEMM count minimizes time only asymptotically; at finite sizes, encoding, output, and reconstruction work can make a higher-count configuration faster.

    Source: A. Minimizing Q Minimizes Time Only… · p. 10

Abstract

from arXiv · show

Ozaki scheme II emulates FP64 matrix products with INT8 ones through residues modulo pairwise coprime moduli, and variants for FP8 and FP4 have followed. We treat these schemes as one family and pose the choice of a scheme as a combinatorial program that minimizes the number of low-precision GEMMs. Given, for each modulus, a finite set of ways to compute products modulo it from low-precision GEMMs, we find the choice of moduli and ways with the fewest GEMMs, for any format, accumulator and inner dimension, and derive lower bounds on the GEMM count over the whole family. Applied to the formats of current GPUs, the method gives the first FP6 schemes, an FP8 scheme with fewer GEMMs than any previous one, and an FP4 scheme that the bounds show needs the fewest GEMMs of any scheme in the family whose moduli lie in a stated range. Implemented on three Blackwell GPUs, the INT8, FP8 and FP4 schemes run faster than native FP64, up to 83x on B300.

I. INTRODUCTION

This paper formulates Ozaki II scheme selection as minimizing low-precision GEMM count, addressing the shift from declining FP64 throughput toward faster low-precision GPU formats. It introduces schemes for additional formats and reports Blackwell speedups over native FP64.

I. INTRODUCTION

  • FP64 throughput has declined on recent NVIDIA and AMD GPUs while peak throughput for lower-precision formats has continued to rise.Integer units stop at eight bits, and AI-focused GPUs have fastest formats below eight bits.Source: §I · p. 1
  • Earlier Ozaki II adaptations targeted FP8 and FP4, but no FP6 or INT4 scheme had been given, and prior low-precision schemes were described for particular formats.The cited FP8 and FP4 schemes split each residue into two values; AWE selected moduli with an integer program whose formulation and solver were not given.Source: §I · p. 1

A. Contributions

  • The selection program jointly chooses moduli, residue splits, and product-formation methods to minimize GEMM count for a given inner dimension.It solves exactly over candidates and derives lower bounds over ways of forming products.Source: §I · p. 1, A. Contributions · p. 1
  • The search finds the first FP6 schemes and an FP8 scheme using 30 GEMMs versus the previous best of 42, while its FP4 scheme uses 59 GEMMs.Bounds show no FP4 scheme in a stated modulus range needs fewer than 59 GEMMs; on B300, FP4 runs up to 83× faster than native FP64.Source: A. Contributions · pp. 1–2

A. Emulation Problem

  • Emulation converts FP64 operands to integers, forms their exact product on a low-precision unit, and rounds once at the end; lossless conversion returns fl(AB).When exact-product conditions hold, truncation and final rounding are the only errors, so constructions are compared by their low-precision GEMM counts.Source: A. Emulation Problem · p. 2
  • Shared power-of-two row and column scaling keeps converted entries below 2^53, but truncation can discard low-order bits and the integer product can exceed low-precision range.The scaling exponents are chosen per row of A and column of B; no low-precision unit represents integers of the resulting product magnitude directly.Source: A. Emulation Problem · p. 2

B. Ozaki II

Ozaki II computes integer matrix products modulo several pairwise-coprime moduli, then uses the Chinese Remainder Theorem to recover the product. This section defines schemes as collections of modulus-specific forms whose ranks determine their low-precision GEMM cost.

B. Ozaki II

  • Ozaki II reduces integer operands modulo pairwise-coprime moduli, forms each modular product with a low-precision GEMM, and recovers the full product by the Chinese Remainder Theorem.The moduli’s product must exceed the range of the integer product, and the GEMM sums must be exact.Source: B. Ozaki II · p. 2
  • INT8 can represent every residue for moduli up to 256, whereas FP8 and FP4 hold only short runs of consecutive integers and need multi-GEMM forms for larger moduli.FP-based schemes represent each reduced operand as D0 + bD1, using three or four GEMMs per modulus.Source: B. Ozaki II · p. 2

A. Forms and Channels

  • A form of rank r uses residue-indexed low-precision tables and integer weights so that the weighted products equal tu modulo m.The form applies its tables entrywise to the integer matrices and issues r GEMMs; its rank is its GEMM cost.Source: A. Forms and Channels · p. 2
  • A scheme assigns a form to each modulus in a pairwise-coprime set, and its cost Q is the sum of their ranks.The CRT recovers the product once the moduli’s product is sufficiently large; the optimization seeks the least Q.Source: A. Forms and Channels · p. 2

B. Forms Built from Limbs

Limb-based forms represent residue classes as tuples of low-precision values and combine limb-plane products with integer weights. Schoolbook forms compute every pair of limbs, while channel admissibility and radius constrain which moduli can be used.

B. Forms Built from Limbs

  • A limb tuple represents an integer as a base-b sum of signed low-precision limbs, and a modulus is admissible when every residue class has such a representation.Each limb plane is formed from the corresponding limbs of all matrix entries, and the forms use fixed integer combinations of those planes.Source: B. Forms Built from Limbs · p. 3
  • A channel selects a representation of each residue class whose plane values lie in the format, minimizing the largest plane magnitude; that radius bounds admissible moduli by m ≤ |S_rm|^ℓm.For E2M1 at base 13, two limbs cover [−56, 56], and the schoolbook channel at modulus 169 has radius 8.Source: B. Forms Built from Limbs · p. 3
  • The schoolbook form multiplies every pair of limb planes as a GEMM and weights each cross-product by the corresponding power of the base.Products with equal total limb index form a group.Source: B. Forms Built from Limbs · p. 3

3 GEMMs

The section describes product forms at each modulus and the exactness constraints that determine which configurations work for a given inner dimension. Accumulator and epilogue bounds combine with the CRT reconstruction window to set Kmax.

4 GEMMs

  • At two limbs, the Toom form has rank 3; at three limbs, the subset form has rank 6 versus Toom’s 5, with plane values formed from limb combinations.Toom uses evaluations at {0, 1, ∞} for two limbs and {0, 1, −1, 2, ∞} for three; the subset form uses unit vectors and selected 0/1 vectors.Source: §4 · p. 3

C. Exactness Constraints and Kmax

  • CRT reconstructs an entry only within its window, so for FP64 operands the modulus product must exceed K · 2^107.For K ≥ 3, a decoder using only the residues and scalings needs a window within log2(K/(K − 2)) bits of this bound even to return a faithful rounding.Source: C. Exactness Constraints and Kmax · pp. 3–4

A GEMM sums K products of stored values. Write ρA

  • Accumulator exactness depends on the magnitudes and number of GEMMs summed together; schoolbook grouping changes that count, while other forms assign each GEMM its own accumulator.The accumulator must represent each product and sum without rounding; the radius bound is sufficient, not necessary, and a limb meeting only smaller products can exceed the general bound.Source: A GEMM sums K products of stored… · p. 4
  • Kmax is the minimum of the reconstruction and accumulation caps under reduce-then-combine, and also the composition cap under combine-then-reduce.The cap attaining that minimum sets Kmax; the window grows with K, while the other caps limit the number of terms summed by an accumulator.Source: A GEMM sums K products of stored… · p. 4

D. Scheme Selection as a Combinatorial Program

Scheme selection is formulated as a combinatorial program that minimizes the total GEMM count while satisfying the configuration’s constraints.

D. Scheme Selection as a Combinatorial Program

E. Solving the Program

The paper gives an exact search over a finite candidate pool and rank-based lower bounds to certify when a found scheme is least. Examples include matching configurations for paired FP8 formats and least-count certificates for FP4 and INT8.

E. Solving the Program

  • The search reduces each modulus to its least-rank eligible candidate, then uses dynamic programming and a cutoff-based prime partition to find a least-Q set meeting the CRT window.It handles any remaining shared prime above the cutoff by rerunning the search with either conflicting member removed.Source: E. Solving the Program · p. 5
  • At K = 214, E4M3 and E2M3 share the same one- and two-limb candidate pools, so each FP8 alphabet uses its 6-bit partner’s configurations.The alphabets agree only up to 60, below Λ/K = 1024; the pools match because their stored values lie within rmax(1) = 32.Source: E. Solving the Program · p. 5

F. Certifying the Least GEMM Count

  • At K = 214 over E2M1, no configuration on moduli up to 474 and primes up to 246 568 393 has Q < 59; the constructed configuration attains Q = 59.The certificate uses lower bounds on rank for every form, not only the candidates in the search pool.Source: F. Certifying the Least GEMM Count · pp. 5–6
  • For E2M1, every form modulo a prime above 225 has rank at least five, and a general bound gives Q ≥ log2(2KBABB + 1)/log2|S|.The prime-rank bound follows from NS(p) ≥ 3.Source: F. Certifying the Least GEMM Count · p. 6
  • At K = 214, INT8 configurations need at least 16 GEMMs under the reconstruction condition, and the reported configuration attains that bound.Source: F. Certifying the Least GEMM Count · p. 6
  • For E2M1, |S · S| = 37 yields lower bounds of 20 GEMMs for returning A′B′ and 12 for returning fl(A′B′).The bounds use more than 299 distinct products and more than 258 distinct roundings, respectively.Source: F. Certifying the Least GEMM Count · p. 6

IV. INSTANTIATION

The method searches configurations for current low-precision GPU formats at selected inner dimensions and compares their GEMM counts with prior schemes. The results include lower bounds and format-specific limits on exactness and GEMM-count growth.

IV. INSTANTIATION

  • The search compares the least Q for E4M3, E2M3, E2M1, INT4, and INT8 at four inner dimensions with prior schemes.The float formats use FP32 accumulation and the integer formats use INT32 accumulation; the epilogue shares each accumulator’s format.Source: §IV · p. 6, Table I · p. 6
  • From K = 2^14 to 2^17, the 6-bit formats’ least GEMM counts rise by up to 12, while the 4-bit counts rise by at most 2.At every plotted point in that range, the 6-bit formats are limited by the accumulation cap.Source: §IV · p. 6, Figure 6 · p. 6
  • At K = 2^14, the float configurations listed are exact under reduce-then-combine, while combine-then-reduce caps OzII-FP8 at K = 243.Figure 6 plots the least Q against K, with each point placed at the largest K its Q suffices for.Source: §IV · p. 6, Figure 6 · p. 6
  • At K = 2^14, the INT8 GEMM count is 16 and is minimal over every configuration by the counting bound, without any search.Source: §IV · p. 6

A. FP8 and FP6: E4M3 and E2M3

The instantiation results compare optimized configurations with prior schemes across FP8/FP6, FP4, and integer formats. The implementation runs configurations through a common library and uses slabs when a configuration’s exactness limit or memory footprint requires them.

A. FP8 and FP6: E4M3 and E2M3

  • At K = 2^14, E4M3 uses 30 GEMMs versus OzII-FP8’s 42; E4M3 and E2M3 take their 6-bit partners’ configurations and counts at that K.The 30-GEMM E4M3 configuration combines 15 one-limb channels with 5 two-limb channels.Source: A. FP8 and FP6: E4M3 and E2M3 · pp. 6–7

B. FP4: E2M1

  • E2M1’s Toom configuration takes 59 GEMMs, down from OzII-FP4’s 75; AWE reaches the same 59-GEMM count.The intermediate choices reduce the count to 71 with base 11, 68 with mixed limb counts, and 66 with a base per channel.Source: B. FP4: E2M1 · p. 7

C. INT4 and INT8

  • INT4 configurations issue far more GEMMs than INT8, and no tested Blackwell instruction multiplies 4-bit integers faster than 8-bit ones.Accordingly, only INT8 configurations are evaluated for Blackwell GPUs.Source: C. INT4 and INT8 · p. 7
  • At K = 2^14, INT8 needs 16 one-limb GEMMs, the minimum for any INT8 configuration; the selected moduli give Kmax = 151 647.The selected moduli reach 239, compared with 256 for OzII-INT8, whose Kmax is 131 071.Source: C. INT4 and INT8 · p. 7

V. IMPLEMENTATION

  • The CUDA library accepts forms as tables, builds tables for standard form families, checks exactness conditions, and refuses K above κrec.It runs configurations from Table I and Figure 6, except INT4 and AWE.Source: §V · p. 7

A. Structure

  • A call scales and reduces operands, issues the channel GEMMs, combines their products in an epilogue, then reconstructs and rounds the result.Only the GEMMs use the low-precision unit; the other four phases are custom vector CUDA kernels.Source: A. Structure · p. 7

B. Slabs

  • Beyond Kmax, accumulation-limited calls can split K into slabs, each with its own accumulation and epilogue; callers choose Kt.Slabbing also serves calls whose operand planes do not fit in memory.Source: B. Slabs · p. 7

C. The GEMMs

The implementations use format- and GPU-specific GEMM paths, with custom kernels handling reduction, epilogue, and reconstruction. Across the tested configurations, calls beat native FP64 wherever both were timed, while the float configurations also outperform the vendor emulation on B300 in nearly all reported cases.

C. The GEMMs

  • E4M3 and E2M1 products use CUTLASS, E4M3 uses cuBLASLt, and the 6-bit formats use CUTLASS because cuBLASLt has no type for them.INT8 uses CUTLASS on B200 and cuBLASLt on B300 and RTX PRO 6000.Source: C. The GEMMs · p. 7
  • E2M1 uses the block-scaled instruction because the dense instruction runs 4-bit operands at the 8-bit rate, while block scaling reaches the 4-bit rate.Every block scale is 1, so products are unchanged and the partial-sum bound holds in any addition order.Source: C. The GEMMs · p. 7
  • CUTLASS output tiles are banded by columns; tuning band width shortens B200’s INT8 call at eight tiles wide.The tile order is set per GPU, and separately for INT8 through CUTLASS’s swizzle parameter.Source: C. The GEMMs · p. 7

D. Around the GEMMs

  • The epilogue reduces each accumulator before combining, an order required to guarantee exactness for every float configuration in Table I.One kernel computes the weighted sum in the accumulator’s format; reconstruction then forms the CRT sum and rounds once.Source: D. Around the GEMMs · pp. 7–8

VI. EVALUATION

  • The reported benchmark setup lists the GPUs, software, methods, and run-time measurement protocol.The evaluation also reports the methods’ throughput as a share of format peak.Source: Table II · p. 8, B. Throughput · p. 8
  • The speed comparisons pair ours, native FP64, and vendor emulation in one process per GPU and configuration, interleaving calls and checking results against native first.Prior-scheme bars and released implementations are measured in separate processes.Source: §VI · p. 8

A. Speedup

  • Every configuration is faster than native FP64 wherever both were timed; on B300, float configurations beat the vendor emulation in dynamic mode at every size.In pinned mode, they do so everywhere except E4M3 at 2^15; the comparisons are against calls in each method’s own process.Source: A. Speedup · p. 8

B. Throughput

  • Native FP64 reaches at least three quarters of its own peak, while the profiled E4M3 and E2M1 GEMMs reach at least half their format’s datasheet peak, except E2M1 on B300.That exception is below one third against an FP4 peak derived from a rack figure; a cuBLASLt path reaches 48% at 2^13 and 52% at 2^14.Source: B. Throughput · p. 8

C. Pipeline Breakdown

The pipeline profile shows that GEMMs dominate calls at the largest profiled size, making GEMM count a useful runtime proxy there but not at smaller sizes. Comparisons with released implementations separate scheme choice from implementation, while the paper frames modulus selection as a cost optimization within the residue-number-system family.

C. Pipeline Breakdown

  • At 2^15, GEMMs take at least three quarters of every profiled call; at 2^13 on B200, they take 49–56%, so GEMM count alone does not order configurations there.The four surrounding phases scale with n^2, while GEMMs scale with n^3.Source: C. Pipeline Breakdown · p. 8

D. Performance Against Prior Implementations

  • Against GEMMul8, our intervals are above parity on B300 and RTX PRO 6000 and on B200 at 2^15; below that B200 size, they include parity.The released FP4 code takes 1.00–1.03 times the time of the paper’s Triton implementation of the same scheme.Source: D. Performance Against Prior… · p. 8, Figure 10 · p. 10

A. Residue Number Systems

  • Unlike prior residue-number-system work that takes every modulus at one word size, this method minimizes the cost of a coprime modulus set meeting the product-range condition.In this family, the alphabet determines how many GEMMs a modulus costs.Source: A. Residue Number Systems · p. 9

B. The Ozaki Schemes

  • Prior Ozaki I ports choose slice counts per call from operand range, inner dimension, and desired accuracy; Ozaki II fixes the count before seeing operands.Earlier FP4 work fixed base 13 and two limbs per channel; AWE allows a base per modulus and reaches 59 GEMMs for E2M1.Source: B. The Ozaki Schemes · pp. 9–10
  • The paper’s per-alphabet scheme selection also differs from Ozaki II’s use in complex GEMM, where INT8 moduli and a three-product identity handle real and imaginary parts.The three-product method is identified as 3M.Source: B. The Ozaki Schemes · p. 10

C. Implementations

GEMM count is a proxy for run time only when GEMMs dominate; the non-GEMM work can change which configuration is faster, especially when dimensions do not all grow.

C. Implementations

  • The implementation line around GEMMs includes fused residue-path kernels, a one-launch FP4 path, engineered residue-plane conversion, and a scale-invariant fast-mode formula.These approaches improve the work surrounding GEMMs rather than changing the scheme’s GEMM count.Source: C. Implementations · p. 10

A. Minimizing Q Minimizes Time Only Asymptotically

  • When non-GEMM work is substantial, a configuration with a higher GEMM count can still be faster, because configurations differ in modulus count and written planes as well as GEMM count.As problem size grows, GEMMs take a larger share of the call on every GPU.Source: A. Minimizing Q Minimizes Time Only… · p. 10

B. Asymptotically Optimal Only When All Dimensions Grow

  • When K is fixed and M and N grow, or when M and N are fixed and K grows, GEMMs do not dominate the call.In the tall-A/fat-B case, epilogue and reconstruction remain comparable to GEMM work; in the fat-A/tall-B case, Reduce remains comparable.Source: B. Asymptotically Optimal Only When… · p. 10

C. Not “True” Emulation of DGEMM

The timed configurations use lossy truncation, while exact emulation is possible for bounded row and column ranges at a higher GEMM count. Exactness also depends on accumulator behavior, and truncation can make the scheme less accurate than native multiplication on skewed rows.

C. Not “True” Emulation of DGEMM

  • The timed configurations begin with one lossy truncation, so they do not return fl(AB) for every input.For operands whose rows and columns span at most 2d, keeping 53+d bits removes the loss, but requires a wider CRT window and more GEMMs.Source: C. Not “True” Emulation of DGEMM · p. 10
  • At E2M1 and K = 214, the exact scheme for operands with dynamic range 216 raises the count from 59 to 78 GEMMs while widening the CRT window by 32 bits.For those operands, it returns fl(AB); at 264 it needs 134 GEMMs.Source: C. Not “True” Emulation of DGEMM · p. 11

D. Exactness May Rely on Undocumented Accumulation

  • FP4-to-FP8 exactness requires integer additions in [−224, 224] without FP32 rounding, but the B200 check did not establish this for cuBLASLt, B300, or RTX PRO 6000.The paper reports a bitwise B200 check through CUTLASS at E4M3 and K = 214; B200 and B300 tcgen05 accumulation behavior is undocumented in the cited ISA.Source: D. Exactness May Rely on Undocumented… · p. 11

E. May Be Less Accurate Than Native on Skewed Rows

  • For skewed rows or columns, truncation can make the scheme less accurate than native multiplication, and its fixed GEMM count cannot add work for such calls.The stated condition is a largest entry more than (K−2)/2 times the row’s mean magnitude when the opposite operand’s column entries have about equal magnitudes.Source: E. May Be Less Accurate Than Native on… · p. 11

IX. CONCLUSION

  • On three Blackwell GPUs, the schemes run faster than native FP64 at the timed sizes, reaching up to 83 times faster on B300.The conclusion also reports that the method finds schemes by minimizing GEMM count subject to an exact product.Source: §IX · p. 11
Loading 2609.37693v1…