Tensor Core
Definition
A Tensor Core is a dedicated hardware unit that performs a small matrix-multiply-accumulate (MMA) — on tiles (e.g. ) — in a handful of cycles, packing hundreds of multiply-adds into one issued instruction. This is where the ~15× throughput gap over general-purpose FP32 CUDA cores comes from, and it is the reason GEMM (hence transformers) runs fast on GPUs. An H100 SM has 4 (4th-gen), one per sub-partition.
Key math
Contrast a CUDA core (one FMA per lane per cycle) with a Tensor Core consuming a whole tile:
- Amortized instruction overhead — one instruction, hundreds of MACs.
- Amortized operand movement — inputs read once, reused across a systolic array.
- Low-precision inputs, high-precision accumulate — BF16/FP8 multiplicands, FP32 accumulator, which is what makes low-precision training numerically viable.
H100 SXM5 dense throughput (the precision ladder — each step ≈ 2×):
| Precision | Dense TFLOP/s | vs BF16 |
|---|---|---|
| TF32 | ~495 | 0.5× |
| BF16 / FP16 | ~990 | 1× |
| FP8 (E4M3/E5M2) | ~1,979 | 2× |
| INT8 | ~1,979 TOPS | 2× |
(All roughly double again with 2:4 structured sparsity — the marketing numbers.) Compare FP32 CUDA-core peak ~67 TFLOP/s → the ~15× gap. Blackwell adds FP4.
Access: via cuBLAS/cuBLASLt/CUTLASS (what PyTorch matmul dispatches to), the WMMA API, or inline mma.sync / Hopper’s wgmma (warp-group MMA over 128 threads); TMA streams tiles global→SMEM to keep them fed.
Why it matters
The precision ladder is the industry’s march down precision in LLM training/inference: every step down is ~2× compute (and ~2× less HBM/bandwidth traffic), paid for with numerical-stability engineering. But Tensor Cores consume data so fast that feeding them — not compute — is usually the bottleneck; lower precision raises the ridge point, pushing more ops into the memory-bound regime. Softmax/GELU run on SFUs/CUDA cores (far slower), which is why non-GEMM glue must be fused (FlashAttention).
Taught in
- 01-gpu-architecture-and-simt — §8 Tensor Cores, the MMA and the precision ladder.
- 03-performance-modeling-roofline — §5 tensor-core throughput by precision.
See also
- gemm — the workload Tensor Cores exist for
- mixed-precision-training — low-precision inputs, FP32 accumulate
- streaming-multiprocessor — 4 Tensor Cores per SM
- arithmetic-intensity — why feeding them is hard
- roofline-model — precision moves the ridge point
- gpu-systems-for-llms