NVIDIA Tensor Cores

In short: NVIDIA Tensor Cores are dedicated warp-level units that perform a fused mixed-precision matrix-multiply-accumulate (D = A·B + C) in one hardware pass, multiplying small FP16/TF32/INT8 tiles while accumulating in FP32 — delivering an order-of-magnitude throughput gain over the standard FP32 CUDA-core pipeline for the matrix math at the heart of deep learning.

Why a matrix unit, and not just more FP32 ALUs?

A general-purpose CUDA core issues one fused-multiply-add (FMA) per lane per cycle. That is fine for element-wise work, but a neural-network layer is dominated by general matrix multiply (GEMM): an M×K matrix times a K×N matrix costs roughly 2·M·N·K floating-point operations. For a single 4096×4096×4096 GEMM that is ~137 GFLOP — and a transformer does thousands of these per forward pass. Scaling FP32 lanes to meet that demand burns area and power on precision the model does not need.

NVIDIA's answer, introduced with the Volta V100 in 2017 and refined every generation since, was to hard-wire a small systolic-style matrix engine into each sub-core. Instead of one scalar FMA, a Tensor Core retires a whole tile multiply per issue. On an A100 (Ampere), the FP16 Tensor Core path delivers ~312 TFLOP/s versus ~19.5 TFLOP/s for FP32 CUDA cores — roughly a 16× gap on the exact operation that matters most.

A Volta Tensor Core computes a 4×4×4 FP16 MMA — 64 multiply-adds — in a single operation. 128 of them per V100 yield 125 TFLOP/s of FP16 throughput, about 8× the chip's FP32 rate.

Mixed precision: the key idea

The defining trick is asymmetric precision. The multiply inputs A and B are stored in a compact format, but every partial product is summed into a 32-bit floating-point accumulator. This keeps the dot-product's dynamic range intact even when the operands are coarse. Multiplying two FP16 values is cheap and roughly quadratic in mantissa width to implement in silicon; accumulating in FP32 prevents the catastrophic rounding that would wreck a 4096-term sum done entirely in FP16.

Supported input formats have grown with each architecture:

• FP16 (1-5-10): the original DL workhorse; small exponent range needs loss scaling during training.
• BF16 (1-8-7): same 8-bit exponent as FP32, so it matches FP32's range and usually trains without loss scaling — at the cost of mantissa bits.
• TF32 (1-8-10): an Ampere-only internal format that takes FP32 inputs, truncates the mantissa to 10 bits for the multiply, and accumulates in FP32. It is a drop-in accelerator — your code still passes FP32 tensors.
• INT8 / INT4: integer MACs for quantized inference, accumulating into INT32; INT8 roughly doubles the throughput of FP16 again.
• FP8 (E4M3 / E5M2): Hopper and Blackwell add 8-bit float, pushing H100 to ~2 PFLOP/s (sparse) and enabling large-model training at half the memory bandwidth of BF16.

TF32 is the quiet default on Ampere+: cuBLAS and cuDNN route FP32 GEMMs through it unless you opt out. You get up to 8× speedups on FP32 code you never changed — with a small, usually negligible, accuracy change.

Warp-level MMA: how 32 threads cooperate

A Tensor Core is not addressable by a single thread. The hardware operates at warp granularity — all 32 threads of a warp jointly supply fragments of A, B and C, the unit performs the tile MMA, and the result fragments are distributed back across the warp's registers. CUDA exposes this through the nvcuda::wmma API (fragments plus load_matrix_sync, mma_sync, store_matrix_sync) and, for finer control, the lower-level mma.sync PTX instruction. Tile shapes are fixed, e.g. 16×16×16 for FP16 WMMA.

In practice almost nobody writes WMMA by hand. The CUTLASS template library and cuBLASLt compose these primitives into tiled, pipelined kernels that stage operands through shared memory, overlap global loads with compute, and keep the Tensor Cores fed. The engineering challenge is memory, not math: a Tensor Core can consume operands far faster than HBM can deliver them, so a good kernel is mostly a careful software pipeline around a tiny, extremely fast MAC engine.

Structured sparsity: a free 2×

From Ampere onward, Tensor Cores also exploit 2:4 structured sparsity — in every group of four weights, two are forced to zero. The hardware skips the zero multiplies and doubles effective throughput. Pruning to this pattern and fine-tuning typically recovers accuracy, which is why the A100's "sparse" FP16 figure of 624 TFLOP/s is double its dense rate.

Where this matters beyond training

Tensor Cores now drive inference economics as much as training. INT8 and FP8 quantized models halve or quarter memory traffic, letting a single H100 serve large language models that would otherwise need multiple GPUs. For hardware teams, the same mixed-precision MAC principle increasingly shows up in custom silicon — accelerators synthesized and verified with Synopsys and Cadence flows reuse exactly this narrow-multiply / wide-accumulate structure. If you are designing such datapaths rather than merely programming them, the RTL-to-GDS methodology taught by Avecas design services covers the physical implementation side.

Want to go from "what is a Tensor Core" to writing and profiling your own CUDA GEMM kernels? Our AI Hardware & GPU Computing track at CourseTron walks through WMMA, CUTLASS and precision trade-offs with hands-on labs.

To connect this to the software side, read our companion piece on mixed-precision training and how loss scaling keeps FP16 gradients stable, then see the GPU memory hierarchy to understand why feeding these units is the real engineering problem. For broader VLSI foundations, our partner ChipXpert VLSI training is a good next step.

FAQ

Do Tensor Cores reduce model accuracy?

Usually negligibly. Because products accumulate in FP32, the loss of precision is confined to the inputs. BF16 and TF32 preserve FP32's exponent range, so training converges almost identically; FP16 needs loss scaling, and aggressive INT8/FP8 inference needs calibration or quantization-aware training.

How are Tensor Cores different from CUDA cores?

CUDA cores are scalar/SIMT ALUs that do one FMA per lane. Tensor Cores are warp-level matrix engines that compute a whole tile multiply-accumulate per instruction, specialized for GEMM and convolution. Both exist on the same SM and are used together.

Do I have to write WMMA code to use them?

No. cuBLAS, cuDNN and framework back-ends (PyTorch, TensorFlow) dispatch to Tensor Cores automatically when shapes and dtypes qualify. You only touch WMMA, mma.sync or CUTLASS when writing custom high-performance kernels.

What is TF32 exactly?

A 19-bit internal math mode on Ampere+ with an 8-bit exponent and 10-bit mantissa. It accepts normal FP32 tensors, truncates the mantissa for the multiply, and accumulates in FP32 — giving most of FP16's speed with no code change.