What is a tensor core and why does it dominate matmul throughput on GPUs?

Flashcard·Easy·4.0 · 0·~30s·Asked atComet Ml Datarobot·Relevant atNVIDIA

Attempt it

A tensor core is a specialized matrix-multiply-accumulate (MMA) unit on each streaming multiprocessor that consumes a small matrix tile (e.g. 16 x 16 in FP16 or 64 x 32 in FP8) and produces a fused D = A * B + C result in a few cycles. Unlike CUDA cores, which do one scalar fused-multiply-add per cycle, tensor cores do an entire tile worth of MMAs in parallel. That is why FP16, BF16, and FP8 matmul throughput on H100 / B200 is many times the FP32 throughput on the same die: FP32 still uses CUDA cores, while reduced-precision matmuls run on the tensor cores. For LLM serving this matters because attention and MLP forward passes are almost entirely matmuls, so tensor-core utilization is the main factor in prefill throughput.

Your answer

**A tensor core** is a specialized arithmetic unit on each streaming multiprocessor (SM) of a modern NVIDIA GPU. It does one operation very fast: a small matrix-multiply-accumulate, or MMA. Given three small matrices A, B, and C, the tensor core computes `D = A * B + C` in a few clock cycles. This matters because LLM forward passes are almost entirely matrix multiplications. Every attention computation and every MLP layer is a big matmul that splits into many small tile matmuls. Tensor cores do those tiles in parallel, where a CUDA core does one scalar multiply-add at a time. So reduced-precision matmuls (FP16, BF16, FP8, INT8) run on tensor cores many times faster than FP32 matmuls, which run on CUDA cores. An H100 SXM is rated at about 67 TFLOPS for FP32 but about 1,979 TFLOPS for dense FP8 matmul. This is why inference engines use FP16 or FP8 weights and activations: it greatly increases tokens per second on compute-bound work. Tensor cores have been on every NVIDIA datacenter GPU since the V100 (2017), which introduced them with FP16. A100 added BF16 and TF32, H100 added FP8, and B200 added FP4 and FP6.

**A tensor core** is a specialized matrix-multiply-accumulate (MMA) unit in every streaming multiprocessor (SM) of a modern NVIDIA GPU. It computes a fused `D = A * B + C` on small matrix tiles (typically `16 x 16` for FP16 or `64 x 32` for FP8) in a few cycles. A CUDA core performs one scalar fused-multiply-add per cycle and handles general work: scalar ops, control flow, FP32 math. Tensor cores only do MMA in supported precisions, but finish a whole tile in a few cycles, so for matmul-heavy work their throughput is far higher. A transformer forward pass is mostly matmuls (QKV projection, attention, MLP), over 95% of inference FLOPs. Weights and activations must be in FP16, BF16, FP8, INT8, or (on B200) FP4 to use tensor cores; FP32 matmul falls back to CUDA cores. Generation history: - **V100 (2017)**: tensor cores at FP16 with FP32 accumulate. - **A100 (2020)**: added BF16, TF32 and 2:4 structured sparsity. - **H100 (2022)**: added FP8 (E4M3, E5M2). Dense BF16 is about 3x A100 (989 vs 312 TFLOPS); dense FP8 is about 6x A100's dense BF16. - **B200 (2024-2025)**: added FP4 and FP6 with block scaling; dense NVFP4 (about 10 PFLOPS) is roughly 5x H100 SXM dense FP8. TensorRT-LLM, for example, supports FP8 on H100 and later and FP4 on B200. FP8 weights and activations use the faster FP8 path; FP8 KV cache and INT4 weight-only quantization mainly cut memory traffic.

TL;DR

A tensor core is a fused matrix-multiply-accumulate unit on each SM that does an entire small tile of MMA in a few cycles; it makes FP16/BF16/FP8 matmul many times faster than FP32 on CUDA cores.

Memory aid

Easy to grasp

Imagine a kitchen with two kinds of cooks. One does single tasks: chop one onion, dice one carrot, mix one spoon of sauce. The other operates a giant industrial wok that processes a whole tray of ingredients in one motion: an entire wave of stir-fry done at once. The first cook is the general-purpose math unit on the chip; the wok is the dedicated math accelerator built specifically to do grid by grid multiplication and adding. The wok only works on that one recipe, but for that recipe it is much faster than ten regular cooks combined. Modern LLM math is almost entirely that one recipe, mixed at lower-precision settings the wok is best at. The result: a GPU advertised at 'X teraflops' on FP32 hits many times X on FP16 or FP8 because those numbers come from the wok, not the regular line.

Key concepts

Matrix Multiplication GPU Hardware Mixed Precision FLOPS

Concept explanation~2 min read

Everything you need to truly understand this topic: intuition, mechanics, step by step explanation, code, formulas, and worked example. Click to expand.

Tensor cores do most of the matmul work in LLM inference on NVIDIA GPUs. Without them, a Llama 70B forward pass would be many times slower and far less power-efficient. Knowing what tensor cores accelerate, and what they do not, is needed to reason about any inference performance problem.

Tensor cores differ from CUDA cores in what they compute and how fast. Each GPU generation added precisions, and serving engines choose precisions to keep LLM compute on them. Tensor-core utilization is the main factor in prefill throughput, while decode is limited by memory bandwidth.

What a tensor core is, physically

A modern NVIDIA GPU is built from many streaming multiprocessors (SMs); the H100 SXM has 132. Each SM contains a mix of execution units: a register file, a warp scheduler, several CUDA cores (FP32 / INT32 ALUs), a few tensor cores, special function units for transcendentals, and a shared memory bank.

A tensor core is a specialized MMA execution unit. Physically it is a small array of multipliers and adders wired to consume two operand tiles (A and B) and an accumulator tile (C) and produce a result tile (D = A*B + C) in a few cycles. Hopper and Blackwell both have four tensor cores per SM. Blackwell's fifth-generation tensor cores deliver higher throughput per core, including native FP4 (NVFP4 and MXFP4), rather than a higher count.

The tile shape depends on precision. FP16 matmul uses 16x16 tiles fed to a single tensor core. FP8 matmul uses larger 64x32 or 128x32 tiles through warpgroup-level instructions. FP4 on Blackwell uses block-scaled tiles with explicit scale operands every 16 elements.

The distinction from CUDA cores matters at the instruction level. CUDA cores execute SASS instructions like FFMA (32-bit fused multiply-add): one scalar op per thread per cycle. Tensor cores execute MMA instructions like mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32: one tile of MMA per warp per few cycles. Tensor-core throughput is the work per tile times the number of tiles finished per second, which is far more than the CUDA-core path can do.

Precision tiers across generations

Why this dominates LLM matmul throughput

What does NOT saturate the tensor cores

Production realities: how serving engines exploit tensor cores

Situations where this technique stops working.

2–4 min · Everything important, quickly.

Real products, models, and research that use this idea.

An NVIDIA H100 SXM is rated at ~989 TFLOPS BF16 and ~1,979 TFLOPS FP8 dense matmul, versus only ~67 TFLOPS FP32 (the FP32 number runs on CUDA cores).
FlashAttention-3 (Hopper-specific) overlaps asynchronous TMA loads with warpgroup MMA (WGMMA) during attention, reaching up to 840 TFLOPs/s in BF16 (85% utilization) and 1.3 PFLOPs/s in FP8 on H100.

What an interviewer would ask next. Try answering before peeking at the approach.

QWhy is H100 FP32 throughput so much lower than its FP8 throughput?

FP32 matmul runs on CUDA cores, which do one scalar FMA per cycle. FP8 matmul runs on tensor cores, which do an entire tile of MMA in a few cycles. The ratio is the architectural difference between the two pipelines, not a quantization quality / loss tradeoff.

2 more follow-ups an interviewer would ask next. Sign in to reveal them.

Red flags & common mistakes

The phrases that signal junior thinking. Click to expand.

Most common mistake

Thinking tensor cores are general-purpose ALUs. They only accelerate fused MMA on small matrix tiles in supported precisions (FP16, BF16, FP8, INT8); everything else runs on CUDA cores.

60 second bullets to scan on the way to the call.

Define a tensor core as a hardware MMA unit on the SM.
Contrast it with a CUDA core (scalar FMA, general-purpose).

Primary sources. Browse if you want the original framing.