🔢 mathematics

Accurate Models of NVIDIA Tensor Cores

This paper presents accurate software models that emulate the non-IEEE 754 compliant inner product behaviors of NVIDIA Tensor Cores across V100, A100, H100, and B200 GPUs for 8-, 16-, and 19-bit floating-point formats, addressing the challenge of non-reproducible results across different hardware generations.

Original authors: Faizan A. Khattak, Mantas Mikaitis

Published 2026-04-07

📖 5 min read🧠 Deep dive

CC BY 4.0

Original authors: Faizan A. Khattak, Mantas Mikaitis

Original paper licensed under CC BY 4.0 (http://creativecommons.org/licenses/by/4.0/). ✨ This is an AI-generated explanation of the paper below. It is not written or endorsed by the authors. For technical accuracy, refer to the original paper. Read full disclaimer

Imagine you are a master chef trying to bake the perfect cake (a complex calculation) for a massive banquet. In the past, you had to mix the ingredients by hand, slowly and carefully, using a standard recipe (standard math rules). But now, you have a fleet of super-fast, specialized mixing machines (NVIDIA's Tensor Cores) that can whip up thousands of cakes in seconds.

The problem? These machines are a bit quirky. They don't follow the standard recipe book perfectly. Sometimes they round up a pinch of sugar, sometimes they chop off a crumb of flour, and different generations of machines (from 2017 to 2024) do it all slightly differently. If you bake a cake on Machine A in 2017 and then try to replicate it on Machine B in 2024, the taste might be just slightly off. For scientists and AI developers, this "off taste" can be a disaster.

This paper is about reverse-engineering these quirky machines to build a perfect "digital twin" (a software model) that behaves exactly like the real hardware.

Here is the breakdown of their work using simple analogies:

1. The Mystery of the "Black Box"

For years, NVIDIA's Tensor Cores have been like black boxes. We knew they were fast, but we didn't know exactly how they handled the tiny details of math.

The Issue: If you ask a machine to add $1.0000001 + 1.0000002$ , a standard computer gives you a precise answer. A Tensor Core might say, "Eh, that's close enough to 2.0000003," or it might drop the last digit entirely.
The Consequence: This leads to "non-reproducible results." Two scientists running the same simulation on different GPUs might get slightly different answers, which is a nightmare for scientific research.

2. The Detective Work: "The Taste Test"

The authors (Faizan and Mantas) decided to become culinary detectives. They didn't just guess how the machines worked; they tested them.

The Strategy: They created thousands of "test vectors" (special ingredient combinations designed to trigger specific behaviors).
- Analogy: Imagine you have a mystery spice. You mix it with lemon juice. If it fizzes, it's baking soda. If it turns red, it's beet juice.
- They mixed numbers in specific ways to see how the GPU reacted. Did it drop a tiny bit of precision? Did it round up or down?
The Iteration: They built a software model (a simulation) and compared it to the real GPU.
- The Loop: "The real machine gave me result X. My model gave me result Y. They don't match! Let's tweak the model's settings (like changing how many 'guard bits' or extra safety margins it uses) and try again."
- They repeated this millions of times until their software model produced bit-for-bit identical results to the real hardware.

3. The "Generations" of Machines

The paper covers a wide range of GPUs, from the V100 (the "vintage model" from 2017) to the B200 (the "brand new luxury model" from 2024).

The Evolution: They discovered that while the machines got faster, they also changed their "personality."
- Example: The older V100 was like a chef who chopped off the last crumb of flour immediately. The newer H100 and B200 are like chefs who keep a few extra crumbs in a safety pocket before deciding whether to keep or discard them.
The Surprise: They found some hidden tricks. For instance, on the newest machines, when mixing certain low-precision numbers, the machine actually processes them in an "interleaved" pattern (like shuffling two decks of cards together) before adding them up. This was a secret feature no one knew about until they built their model.

4. The Result: The "MATLAB Recipe Book"

The authors didn't just write a report; they built a toolbox (software available on GitHub) called MATLAB Tensor Core.

What it does: It allows any scientist to simulate exactly how an NVIDIA GPU will behave without needing to own the expensive hardware.
Why it matters:
- For Developers: They can test their AI algorithms on a laptop to see if the results will be accurate before they buy a million-dollar supercomputer.
- For Scientists: They can finally trust that their simulations are reproducible. If they share their code, others can run it on the "digital twin" and get the exact same answer, regardless of which physical GPU they have.

5. The Big Picture

Think of this paper as the user manual for a car engine that the manufacturer never wrote.

Before this, drivers (scientists) had to guess how the engine (Tensor Core) handled fuel (data).
Now, thanks to this paper, we have a detailed blueprint. We know exactly how many gears it shifts, how it handles bumps in the road (rounding errors), and how to tune it for the best performance.

In short: The authors built a perfect digital clone of NVIDIA's super-fast math engines. This allows the world to understand, predict, and trust the results of the AI and scientific calculations running on these powerful chips.

1. Problem Statement

Matrix multiplication is fundamental to AI training, inference, and high-performance computing (HPC). Modern GPUs utilize specialized hardware units called Tensor Cores to accelerate these operations. However, these units present significant challenges for reproducibility and numerical analysis:

Non-Standard Behavior: Unlike standard IEEE 754 floating-point arithmetic, Tensor Cores often deviate from standard rounding, normalization, and accumulation rules to prioritize throughput over precision.
Vendor and Architecture Variability: Different NVIDIA GPU generations (Volta, Ampere, Hopper, Ada, Blackwell) and even different input formats on the same chip exhibit distinct numerical behaviors (e.g., varying block FMA sizes, extra alignment bits, and normalization points).
Lack of Documentation: These numerical features are rarely documented by vendors, leading to non-reproducible results across hardware generations.
Limitations of Existing Models: Previous attempts to model these cores relied on limited test vectors that often failed to distinguish between hardware models or required manual derivation of constants, lacking the robustness needed for rigorous scientific computing.

2. Methodology

The authors propose a rigorous, iterative methodology to reverse-engineer and model the inner product behavior of NVIDIA Tensor Cores. The process consists of three main stages:

A. Generalised Numerical Feature Testing (GNFT)

Concept: Uses carefully designed test vectors (constants) to trigger specific bit-level differences in hardware outputs.
Goal: To identify initial numerical features such as support for subnormal numbers, extra alignment bits ( $n_{eab}$ ), and block FMA sizes ( $N_{FMA}$ ).
Advancement: Unlike previous work that required manual theorem proving or constant derivation, this method uses generalized expressions applicable across different formats (fp8, fp16, bf16, tf19, etc.).

B. Input Space Search Method (ISSM)

Concept: Once a conceptual model is approximated via GNFT, the authors perform massive randomized testing to find discrepancies between the software model and actual GPU hardware.
Scale: They generate approximately $10^7$ random input vectors drawn from normal and uniform distributions, plus targeted tests for edge cases (subnormals, infinities, NaNs).
Refinement Loop: If the model output mismatches the GPU output, the authors analyze the failure cases to hypothesize missing features (e.g., specific truncation rules or alignment limits) and refine the model. This loop repeats until the model is bit-accurate.

C. Software Implementation

The models are implemented in MATLAB using the CPFloat library for custom floating-point arithmetic.
The toolbox (MATLAB Tensor Core v0.4.1) includes pre-configured models for specific GPUs and a generalized model allowing users to customize parameters (FMA size, alignment bits, rounding modes).

3. Key Contributions

Comprehensive Feature Extraction: The paper determines precise numerical features for 31 distinct Tensor Core configurations across 9 GPU models (V100, A100, A2, A30, L40S, Ada RTX 1000, H100, H200, B200). This includes support for 8-bit, 16-bit, 19-bit (tf19), and 64-bit formats.
Discovery of Novel Behaviors:
- Exponent Capping: Discovered that when accumulating subnormal products, the maximum exponent used for alignment is capped (e.g., at -133 for A100/H100) rather than using the actual product exponent, affecting precision in the subnormal range.
- Interleaved Accumulation: Identified that H100/H200/B200 Tensor Cores process fp8 inputs via mma.sync in an interleaved fashion (splitting 32-element vectors into two groups processed separately) before summing, a behavior not present in previous architectures.
- Denormalized Accumulation: Confirmed that products remain denormalized during the alignment and accumulation phases in many configurations, affecting how bits are truncated.
Bit-Accurate MATLAB Models: Developed a validated software suite that replicates hardware behavior exactly (bit-for-bit) for mixed-precision matrix multiplication, including handling of NaNs and Infinities.
Open-Source Toolbox: Released the MATLAB Tensor Core v0.4.1 toolbox on GitHub, enabling researchers to simulate these cores without physical hardware access.

4. Key Results and Findings

Architecture Differences:
- V100 (Volta): $N_{FMA}=4$ for fp16; 24-bit alignment; no extra alignment bits ( $n_{eab}=0$ ).
- A100/A2/A30 (Ampere): $N_{FMA}=8$ (fp16/bf16); 26-bit alignment with 1 extra bit ( $n_{eab}=1$ ); introduces exponent capping at -133.
- H100/H200 (Hopper): $N_{FMA}=16$ (fp16/bf16); 27-bit alignment with 2 extra bits ( $n_{eab}=2$ ); supports native fp8 via wgmma with $N_{FMA}=32$ .
- B200 (Blackwell): Similar to H100 but introduces 5th-gen Tensor Cores (tcgen05) with $N_{FMA}=32$ for fp8 and 23 fractional bits.
Validation: The models were tested against $10^7$ random inputs per configuration. The final models achieved 100% bit-identical results compared to the physical GPUs across all tested formats.
Application Example: The authors demonstrated the utility of these models by simulating multi-word arithmetic (emulating high-precision GEMM using low-precision Tensor Cores). They showed that error characteristics vary significantly between generations (e.g., V100 vs. B200) due to differences in accumulator precision and rounding modes.

5. Significance

Reproducibility in Scientific Computing: Provides a reliable way to reproduce numerical results across different GPU generations, which is critical for scientific applications where cross-platform consistency is required.
Algorithm Development: Enables mixed-precision algorithm developers to test and optimize algorithms (like iterative refinement) on accurate software models before deploying to expensive or unavailable hardware.
Standardization Impact: By documenting these non-standard behaviors, the work highlights the need for industry standards (like IEEE P3109) to address mixed-precision arithmetic, particularly for reduction operations.
Debugging and Verification: The toolbox allows developers to verify test vectors and debug numerical issues in AI and HPC applications without needing direct access to specific GPU hardware.

In summary, this paper bridges the gap between opaque hardware implementations and transparent software modeling, providing the first bit-accurate, open-source simulation suite for NVIDIA Tensor Cores across multiple generations and precision formats.