TL;DR: A tensor core is a matrix-multiply-accumulate unit: one instruction takes a small tile of A and B (for bf16 on Hopper, 64 × 16 by 16 × N per warpgroup) and accumulates the product into a tile of C in fp32. It delivers about 4,096 FLOP per SM per cycle for dense bf16 against 256 for the fp32 lanes. To reach it, the kernel must express its work as a matmul with dimensions that are multiples of the tile, operands in bf16, fp16, tf32, fp8 or int8 at 16-byte alignment, and data staged in registers or shared memory in the layout the instruction expects. Elementwise code, odd shapes and fp32 inputs never touch it.
How to approach it
Define it as an instruction, not a component: say what the instruction consumes, produces and accumulates in. Give the throughput ratio against the CUDA cores, derived from the datasheet peak. Then list the conditions a kernel must meet, because that is what an infra engineer actually needs: how to know whether a given layer is running on tensor cores and what to change if not. Close with the profiler metric.
A strong answer
A typical situation: a team quotes 989 TFLOPS from the datasheet, measures 60, and concludes the card is faulty. The kernel was running on the fp32 lanes the whole time, because one dimension was not a multiple of 16, and nothing anywhere reported an error.
Tensor Cores and Matrix Units exist because a matmul has structure the scalar datapath cannot exploit. In a scalar FMA, each lane fetches two operands from registers and produces one result: three register accesses per two FLOPs. In a 16 × 16 × 16 tile MMA, the unit loads 512 input values and performs 8,192 FLOPs on them, reusing each input 16 times inside the unit without touching the register file again. That reuse is where the speed comes from; the tensor core is a hardware version of register tiling.
The throughput ratio is derivable from the datasheet:
H100 SXM dense bf16 peak = 989 TFLOPS; 132 SMs; ~1.83 GHz implied clock
FLOP per SM per cycle on tensor cores = 989e12 ÷ 132 ÷ 1.83e9 ≈ 4,096
FLOP per SM per cycle on fp32 lanes = 128 lanes × 2 = 256
ratio = 16x
fp8 doubles it again: 1,979 TFLOPS dense, 8,192 FLOP per SM per cycle
sanity: 4,096 FLOP per cycle over 4 tensor cores per SM is 512 bf16 FMAs per tensor core per
cycle; a 16 × 8 × 16 warp-level mma is 2,048 FMAs, so it retires in a few cycles
What the instruction looks like. On Ampere and later, the warp-level primitive is mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32: a warp cooperatively holds fragments of a 16 × 16 A tile and a 16 × 8 B tile in its registers, and the instruction accumulates a 16 × 8 fp32 C tile. On Hopper the preferred form is wgmma.mma_async issued by a warpgroup of four warps, with M fixed at 64, N from 8 to 256, K of 16 for bf16, and the B operand (and optionally A) read straight from shared memory in a swizzled layout described by a matrix descriptor. Blackwell moves the accumulator into a dedicated tensor memory with tcgen05.mma. In every generation, the operands must be laid out exactly as the instruction expects, which is why ldmatrix, swizzled shared memory and the CUTLASS layout machinery exist.
What a kernel must satisfy to use one:
| Requirement | Why | What breaks it |
|---|---|---|
| The work is a matmul or batched matmul | the unit only does D = A × B + C on tiles | elementwise ops, reductions, softmax on their own |
| Input dtype is bf16, fp16, tf32, fp8 or int8 (fp4 on Blackwell) | there is no fp32 tensor-core path except tf32, which truncates the mantissa | a model left in fp32; torch.backends.cuda.matmul.allow_tf32 = False |
| M, N, K are multiples of the tile (8 or 16) and pointers are 16-byte aligned | fragments are fixed size; misaligned loads fall back | a vocab of 32,001, a hidden size padded to an odd multiple, a sliced tensor with a non-aligned offset |
| Operands staged in shared memory or registers in the right layout | the instruction reads fragments, not arbitrary addresses | naive kernels that read global memory per element |
| Accumulate in fp32 | the hardware accumulates wider than it multiplies | asking for fp16 accumulation to save registers and getting overflow in long K |
The consequence that matters for an infra engineer: a model's linear layers get tensor cores automatically through cuBLAS or CUTLASS when the dtype and shapes cooperate. Attention gets them through FlashAttention Internals, which reformulates QK^T and PV as tile matmuls. Everything else (layernorm, activation functions, rotary embeddings, sampling) runs on the fp32 lanes and is memory-bound regardless.
A worked check on a decode-time linear layer shows why "on a tensor core" is not the same as "fast":
layer: [M tokens × 8,192] × [8,192 × 8,192] in bf16, M = 8 (a small decode batch)
FLOPs = 2 × 8 × 8,192 × 8,192 = 1.07 GFLOP
bytes = weights 8,192 × 8,192 × 2 B = 134 MB (activations negligible)
time at 3.35 TB/s = 134e6 ÷ 3.35e12 = 40 µs
achieved = 1.07e9 ÷ 40e-6 = 27 TFLOPS = 2.7% of peak
sanity: the kernel is running on tensor cores and is still at 2.7% because M = 8 gives an
arithmetic intensity of 8 FLOP/B, far below the 295 FLOP/B ridge; the tensor core is
idle waiting on HBM, not misused
So the conditions above are necessary and not sufficient. Tensor cores reach their peak only when the tile is fed faster than it computes, which needs both an arithmetic intensity above the ridge and a kernel that keeps loads ahead of MMAs (double buffering, TMA on Hopper, warp specialization).
How to tell, in a profile: Nsight Compute reports sm__pipe_tensor_cycles_active.avg.pct_of_peak_sustained_active (or the "Tensor (All)" row under Compute Workload Analysis). Near zero on a linear layer means the dtype, shape or alignment sent it to the fp32 path; sm__inst_executed_pipe_tensor confirms whether any MMA instructions were issued. In PyTorch, the quicker test is the kernel name in a trace: ampere_bf16_s16816gemm, sm90_xmma_gemm or a cutlass name means tensor cores; a gemv or sgemm name does not.
Decision: keep weights and activations in bf16 or fp8, pad dimensions to multiples of 16 (vocab sizes in particular), keep tensors contiguous and aligned, and confirm with the pipe metric. The reversal condition: a matmul small or awkward enough that the fp32 lanes finish before the tile machinery is even set up, and the framework picks a non-tensor-core kernel on purpose. That is the right call, and the way to tell it apart from an accidental fallback is PROF_PIPE_TENSOR_ACTIVE near zero on a kernel whose shapes say it should be high. The Roofline Model says whether the tensor cores could have helped at all.
What interviewers probe next
- "Does tf32 use tensor cores?" Yes: fp32 inputs are rounded to a 10-bit mantissa and run through the tensor-core path at 8x the fp32-lane rate; it is on by default for convolutions in PyTorch and off for matmuls unless you enable it.
- "What is the sparsity figure NVIDIA quotes?" 2:4 structured sparsity: if two of every four weights are zero, the hardware skips them and the datasheet number doubles. Dense peaks are what a normal GEMM sees.
- "Why fp32 accumulation?" Summing 8,192 products of bf16 values in bf16 would lose most of the low bits; the accumulator has 23 mantissa bits and the round happens once at the end.
- "What is TMA?" Hopper's Tensor Memory Accelerator: a single thread issues a descriptor-based bulk copy of a tile from global to shared memory in the swizzled layout wgmma wants, freeing the warps to compute.
Common mistakes
- Describing a tensor core as "a faster core" rather than a tile MMA instruction with shape and layout requirements.
- Assuming a bf16 model is on tensor cores everywhere, then finding an lm_head with vocabulary 32,001 running on the fp32 path.
- Quoting the sparse peak (1,979 TFLOPS bf16 with sparsity) as the dense one.
- Concluding from a low tensor-pipe utilization that the kernel is not using tensor cores, when it is memory-bound at small M.
Key takeaways
- A tensor core executes D = A × B + C on fixed tiles (m16n8k16 per warp; m64nNk16 per warpgroup on Hopper), accumulating in fp32.
- Dense bf16: ~ 4,096 FLOP per SM per cycle, 16x the fp32 lanes; fp8 doubles it.
- Requirements: matmul-shaped work, bf16/fp16/tf32/fp8/int8 inputs, dimensions in multiples of 16, 16-byte alignment, operands staged in the expected layout.
- Being on the tensor core is not being at peak; intensity above ~ 295 FLOP/B on H100 is also required.