Reverse-Engineering Nvidia B200's Tensor Core Down to the Bit
Engineers reverse-engineered Nvidia B200's tcgen05.mma tensor core, releasing a bit-exact open-source model revealing its hidden fixed-point accumulation.
A team of engineers built a software model of Nvidia Blackwell's tcgen05.mma tensor-core instruction that matches real B200 hardware output bit-for-bit. Since Nvidia's PTX documentation only describes the instruction's existence and descriptor format, without detailing operand layout, special-value handling, or rounding behavior, the team probed the actual silicon on a real B200, comparing raw output bits against candidate models until every case matched exactly.
A naive IEEE-style round-to-nearest-even accumulation model failed on 144 of 2,000 test cases; even a more careful two-stage model still missed 12. The mismatches clustered around cancellation, sub-ULP contributions, and denormal boundaries, revealing that the hardware isn't summing terms with per-step floating-point rounding at all. Instead, it truncates inputs into the TF32 container, decodes them as integer mantissas and exponents, and accumulates everything in a single shared fixed-point window carrying 25 fractional bits below the group's largest exponent, only truncating once at the alignment step.
The resulting open-source model now covers 46 hardware-validated datapaths spanning dense and sparse, floating-point and integer, and block-scaled mixed-format modes. For compiler and kernel engineers, this gives a ground-truth reference for tensor-core edge-case behavior, letting code target actual silicon semantics instead of documented abstractions when chasing peak utilization.