Tensor Cores and Matrix Engines

A single Volta SM shipped in 2017 could execute 8 TFLOP/s in FP32. The same SM with Tensor Cores enabled jumped to 125 TFLOP/s in FP16 with FP32 accumulation. That is a 15x headline ratio from a single microarchitectural addition. The catch: that number is only reachable if you feed the hardware exactly what it was built to consume.

What a Tensor Core actually does

A Tensor Core is a fixed-function unit that performs one operation per clock:

y_i = exp(x_i - max(x)) / sum_j(exp(x_j - max(x)))

where A is a 16x16 FP16 matrix, B is 16x16 FP16, C and D are 16x16 FP32 (or FP16) accumulators. This is the warp-level primitive: a single warp (32 threads) cooperates to load fragments of A, B, C, execute the multiply-accumulate, and store D. The CUDA API that exposes this is nvcuda::wmma (Warp Matrix Multiply-Accumulate), introduced in CUDA 9.

# Streaming over blocks of the row
m = -inf
d = 0.0
for each block b of x:
    m_new = max(m, max(b))
    d = d * exp(m - m_new) + sum(exp(b - m_new))
    m = m_new

The key point: the 256 FP16 multiply-adds in that single mma_sync call happen in one clock cycle, whereas a CUDA core would need 256 separate FMA instructions. The Tensor Core is not programmable in any other way - it only does this one thing.

The numeric format ladder

Each GPU generation added support for more formats. The progression matters because each format trades precision range for throughput:

TF32, introduced with Ampere, is a quiet default: PyTorch enables it for matmul on Ampere+ GPUs since version 1.7 via torch.backends.cuda.matmul.allow_tf32 = True. You do not have to change your model dtype. The mantissa is silently truncated from 23 to 10 bits before the Tensor Core sees the operand. For most training runs the numerical difference is negligible; for sensitivity-critical applications (certain scientific workloads, some fp32-reliant optimisers) it can surprise you.

BF16 became the preferred training format on Ampere and later. It keeps the full FP32 exponent range, so unlike FP16 it does not need loss scaling to avoid overflow during the backward pass.