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:

D = A × B + C

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.

// Minimal WMMA sketch (Volta, fp16 inputs, fp32 accumulator)
#include <mma.h>
using namespace nvcuda::wmma;

fragment<matrix_a, 16, 16, 16, half, row_major> a_frag;
fragment<matrix_b, 16, 16, 16, half, col_major> b_frag;
fragment<accumulator, 16, 16, 16, float>         c_frag;

load_matrix_sync(a_frag, a_ptr, lda);
load_matrix_sync(b_frag, b_ptr, ldb);
fill_fragment(c_frag, 0.0f);
mma_sync(c_frag, a_frag, b_frag, c_frag);
store_matrix_sync(d_ptr, c_frag, ldd, mem_row_major);

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.

FP8, introduced in Hopper (H100) and formalised in the paper "FP8 Formats for Deep Learning" (Micikevicius et al., 2022, arXiv:2209.05433), splits into two sub-formats: E4M3 for activations and weights (more precision, less range) and E5M2 for gradients (more range, less precision). A well-tuned FP8 training run matches BF16 quality with roughly 2x the throughput.

Shape alignment requirements

Tensor Cores are unforgiving about tile shape. The supported tile sizes for the WMMA API on Volta and Turing are (M, N, K) = (16, 16, 16), (32, 8, 16), and (8, 32, 16). cuBLAS and cuDNN handle tiling internally, but there is a non-negotiable rule: matrix dimensions must be multiples of the tile K dimension.

For most transformer models the relevant constraint is that the hidden dimension and number of attention heads must be multiples of 8 (for BF16/FP16) or multiples of 16 to achieve peak occupancy. A hidden size of 4096 with 32 heads (128 per head) is clean. A hidden size of 4097 breaks alignment and falls back to CUDA cores.

The same logic applies to batch size. A batch size of 1 at inference time means K=1 for the weight matmul, which will not tile onto Tensor Cores at all. This is why batching inference requests is so important for throughput, and why vLLM's continuous batching was a significant practical advance.

Beyond shape, the operand tensors must be contiguous in memory in a layout the Tensor Core expects. Padding, non-contiguous views, or misaligned base pointers will silently fall back to a slower kernel.

How frameworks use them

In practice, users rarely call WMMA directly. The call stack is:

1. User code: torch.nn.Linear, torch.matmul, or an attention kernel.
2. PyTorch dispatcher selects a cuBLAS or custom CUDA kernel.
3. cuBLAS internally tiles the problem and dispatches WMMA or PTX mma.sync instructions.
4. At the warp level, each mma.sync instruction maps to one or more Tensor Core clock cycles.

PyTorch's torch.amp.autocast (formerly torch.cuda.amp.autocast) is the recommended way to enable FP16/BF16 Tensor Core paths. It wraps the forward pass and automatically casts eligible operations (linear layers, convolutions, attention) to the target dtype while keeping accumulation and sensitive operations in FP32.

with torch.amp.autocast(device_type="cuda", dtype=torch.bfloat16):
    output = model(input)

Without autocast, a model defined in FP32 will use FP32 CUDA cores even on an Ampere GPU, unless TF32 is enabled (which it is by default since PyTorch 1.7 for matmul, but not convolutions - that requires torch.backends.cudnn.allow_tf32 = True separately).

When it falls down

Shape misalignment. Odd hidden dimensions, non-power-of-two sequence lengths passed directly into matmul, or unpadded vocabulary projections all break the tile alignment requirement. The computation silently continues in CUDA core mode. The GPU Nsight profiler will show near-zero Tensor Core utilisation.

FP16 overflow during training. FP16's maximum representable value is about 65,504. Gradients can easily exceed this before the first LR warmup step completes. Without loss scaling (multiply the loss by a scale factor, divide gradients back before the optimiser step), the update becomes NaN. BF16 mostly eliminates this problem but still has lower precision than FP32 in the mantissa.

TF32 precision surprises. Code that was validated against FP32 arithmetic may produce numerically different results after a PyTorch upgrade that silently enabled TF32. Reproducibility tests comparing against saved FP32 checkpoints can silently fail. The fix is explicit: torch.backends.cuda.matmul.allow_tf32 = False in contexts where you need exact FP32 semantics.

Accumulator saturation in FP8. FP8 E4M3 has a maximum value of 448. Activations in later transformer layers, after residual accumulation, can exceed this. The typical mitigation is per-tensor scaling (a learned or dynamic scale factor stored alongside the tensor), but incorrect scaling leads to clipping artefacts that are hard to debug.

Not all GPUs have Tensor Cores. Pre-Volta consumer cards (GTX series before the RTX 20xx line) have no Tensor Cores. Code written for Tensor Cores silently degrades to CUDA cores on those devices. A model that benchmarks at 80 TFLOP/s effective throughput on an A100 may run at 10 TFLOP/s on an older V100 not because of bandwidth, but because the FP format path is unavailable.

Further reading

• Mixed Precision Training - Micikevicius et al., ICLR 2018
• FP8 Formats for Deep Learning - Micikevicius et al., 2022
• Programming Tensor Cores in CUDA 9 - NVIDIA Developer Blog
• PyTorch Automatic Mixed Precision - torch.amp documentation