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

Flashcard·Easy·4.0 · 0·~30s·Asked atComet Ml Datarobot Kpmg·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 every cycle. 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 light up the tensor-core pipeline. For LLM serving this matters because attention and MLP forward passes are almost entirely matmuls, so tensor-core utilization is the primary lever on prefill throughput.

Your answer

**A tensor core** is a specialized arithmetic unit on each streaming multiprocessor (SM) of a modern NVIDIA GPU. Its job is to do one operation extremely 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 roughly one clock cycle. This matters because LLM forward passes are almost entirely matrix multiplications. Every attention computation, every MLP layer, every output projection is a big matmul that decomposes into many small tile matmuls. Tensor cores do those tile matmuls in parallel, where a traditional CUDA core would have to grind through one scalar multiply-add at a time. The practical effect is that reduced-precision matmuls (FP16, BF16, FP8, INT8) run on tensor cores at many times the speed of FP32 matmuls, which still run on CUDA cores. An H100 GPU is rated at about 67 TFLOPS for FP32 but around 1,979 TFLOPS for FP8 dense matmul, because the FP8 number is the tensor-core throughput. This is why inference engines push to use FP16 or FP8 weights and activations: it unlocks the tensor-core pipeline and dramatically increases tokens per second on compute-bound work. Tensor cores have been on every NVIDIA datacenter GPU since the V100 (introduced 2017). Each generation has expanded the precision tiers: V100 added FP16, A100 added BF16 and TF32, H100 added FP8, B200 added FP4 and FP6.

**A tensor core** is a specialized matrix-multiply-accumulate (MMA) unit embedded in every streaming multiprocessor of a modern NVIDIA GPU. It operates on small matrix tiles (typically `16 x 16` for FP16 or `64 x 32` for FP8) and computes a fused `D = A * B + C` result in a small number of cycles, in parallel across many tiles per SM. Compare this to a CUDA core, which performs one scalar fused-multiply-add per cycle. CUDA cores are general-purpose: they handle scalar ops, control flow, FP32 math, and anything that does not fit the tensor-core pattern. Tensor cores are narrow but extreme: they only do MMA on supported precisions, but they do an entire tile of MMA in one cycle. For workloads dominated by matmul, the throughput ratio between the two pipelines is huge. The LLM connection is direct. A transformer forward pass is essentially three big matmul stages per layer: QKV projection, attention, and MLP. Together these account for over 95% of the FLOPs in inference. Anything that runs on tensor cores runs many times faster than the same op on CUDA cores. Reduced-precision matmul is the gateway: weights and activations need to be in FP16, BF16, FP8, INT8, or (on B200) FP4 to feed the tensor cores; FP32 matmul falls back to CUDA cores. Generation history: - **V100 (2017)**: introduced tensor cores at FP16 with FP32 accumulate. - **A100 (2020)**: added BF16 and TF32; introduced sparsity support for 2:4 patterns. - **H100 (2022)**: added FP8 (E4M3 and E5M2 variants); roughly 6x A100 throughput on supported precisions. - **B200 (2024-2025)**: added FP4 and FP6 with finer-grained block scaling; FP4 throughput roughly 5x H100 FP8. Production inference engines (TensorRT-LLM, vLLM, SGLang) detect the available precision tiers at runtime and route weights through quantization paths (FP8 KV cache, FP8 weight-only quant, INT8 SmoothQuant) precisely to keep matmul work on the tensor-core pipeline.

TL;DR

A tensor core is a fused matrix-multiply-accumulate unit on each SM that does an entire small tile of MMA per cycle; 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 dramatically 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 are why modern LLM inference is possible on a single GPU. Without them, a Llama 70B forward pass would be many times slower and far less power-efficient. Understanding what tensor cores are, what they accelerate, and what they do not accelerate is a prerequisite for reasoning about any inference performance problem.

This deep dive covers the hardware: what tensor cores are physically, how they differ from CUDA cores, what precision tiers each GPU generation added, and how serving engines route LLM compute through them. The takeaway is that tensor-core utilization is the primary lever on prefill throughput, while decode is bottlenecked elsewhere.

What a tensor core is, physically

A modern NVIDIA GPU is a sea of streaming multiprocessors (SMs). The H100 has 132 SMs; the B200 has roughly 148 enabled SMs per die. 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. On both Hopper and Blackwell there are four tensor cores per SM; Blackwell's 5th-gen cores deliver higher throughput per core (notably native FP4) rather than doubling the count.

The tile shape is precision-dependent. FP16 matmul uses 16x16 tiles fed to a single tensor core; FP8 matmul uses larger 64x32 or 128x32 tiles via 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. The total arithmetic per cycle on the tensor-core path is the product of tile area and clock rate, which dwarfs what 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) uses asynchronous TMA loads plus warpgroup MMA (WGMMA) to keep tensor cores fed during attention, hitting ~75% of peak FP16 throughput.

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 per cycle. 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.