Tensor Cores

1. What is it?#

A hardware unit that performs a small matrix multiply-accumulate in one instruction, rather than a scalar one.

CUDA core:   d = a * b + c              1 FMA, 2 FLOPs
Tensor core: D = A × B + C              a 16×8×16 matrix MMA, 4,096 FLOPs
                                        in roughly the same time

That’s where the 10-30x gap between “CUDA core FLOPs” and “tensor FLOPs” comes from, and why using them is non-negotiable for any matmul-shaped work.

2. Why does it exist?#

Because neural networks are ~95% matrix multiplication, and a general-purpose FMA unit spends most of its energy on instruction fetch, decode, and register access rather than on arithmetic.

Fixing the operation to “multiply these small matrices and accumulate” lets the hardware:

• Amortize instruction overhead over 4,096 FLOPs instead of 2.
• Use a systolic/dataflow arrangement with short wires and high reuse.
• Deliver an order of magnitude more FLOPs per watt.

H100:  FP32 CUDA cores    67 TFLOP/s
       FP16 tensor cores 990 TFLOP/s        14.8x
       FP8  tensor cores 1979 TFLOP/s       29.5x

If your matmul isn’t using tensor cores, you’re getting ~3-7% of the chip.

3. Simple analogy#

A general-purpose lathe versus a purpose-built stamping press.

The lathe can make any shape, one feature at a time. The press makes one specific shape — but makes it complete, in one stroke, thousands of times faster.

Tensor cores are the press. You can only feed them one shape of work (a small matrix multiply), and you must arrange your data into exactly that shape — which is what all the “fragment” machinery in CUTLASS is doing.

4. Tiny example#

Verify tensor cores are being used:

import torch, time

def bench(dtype, n=8192):
    a = torch.randn(n, n, device='cuda', dtype=dtype)
    b = torch.randn(n, n, device='cuda', dtype=dtype)
    for _ in range(3): c = a @ b
    torch.cuda.synchronize(); t0 = time.perf_counter()
    for _ in range(10): c = a @ b
    torch.cuda.synchronize()
    dt = (time.perf_counter()-t0)/10
    print(f"{str(dtype):20s} {dt*1e3:7.1f} ms  {2*n**3/dt/1e12:7.1f} TFLOP/s")

torch.backends.cuda.matmul.allow_tf32 = False
bench(torch.float32)          # true FP32, no tensor cores
torch.backends.cuda.matmul.allow_tf32 = True
bench(torch.float32)          # TF32 tensor cores
bench(torch.float16)
bench(torch.bfloat16)

Typical A100:

torch.float32 (no TF32)   271.4 ms     4.0 TFLOP/s     ← CUDA cores only
torch.float32 (TF32)       31.2 ms    35.2 TFLOP/s     ← 8.8x from a config flag
torch.float16              10.5 ms   104.7 TFLOP/s
torch.bfloat16             10.6 ms   103.7 TFLOP/s

The TF32 line is free performance most people leave on the table. It’s enabled by default in recent PyTorch, but check — some code disables it for reproducibility.

5. Technical explanation#

The MMA instruction#

At the PTX level:

mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32  d, a, b, c;
                 ^^^^^^^^^^                 ^^^ ^^^ ^^^ ^^^
                 M=16 N=8 K=16              acc  A   B  acc
                                            fp32 fp16 fp16 fp32

Executed by a whole warp cooperatively: the 32 lanes each hold a few elements of A, B, and the accumulator, in a specific “fragment” layout defined by the hardware. The instruction operates on the warp’s collective registers.

This is why tensor core code is warp-level, not thread-level. You can’t reason about it one thread at a time.

Supported types by generation#

Requirements to actually use them#

1. DTYPE: FP16, BF16, TF32, FP8, or INT8. (FP32 without TF32 → no tensor cores)
2. SHAPES: dimensions should be multiples of 8 (FP16) or 16 (INT8/FP8).
           Ideally multiples of the tile size (64/128).
3. ALIGNMENT: pointers aligned to 16 bytes.
4. LAYOUT: the library handles fragment layout, but a bad input layout
           forces a transpose.
5. LIBRARY: cuBLAS/cuBLASLt/CUTLASS. Hand-written scalar loops never use them.

Failure to meet any of these silently falls back to CUDA cores — a 10x regression with no error message. This is a real production failure mode.

Writing MMA directly#

#include <mma.h>
using namespace nvcuda::wmma;

__global__ void wmma_gemm(const half* a, const half* b, float* c, int M, int N, int K) {
    fragment<matrix_a, 16, 16, 16, half, row_major> a_frag;
    fragment<matrix_b, 16, 16, 16, half, col_major> b_frag;
    fragment<accumulator, 16, 16, 16, float> acc_frag;

    fill_fragment(acc_frag, 0.0f);

    int warpM = (blockIdx.x * blockDim.x + threadIdx.x) / 32;
    int warpN = blockIdx.y * blockDim.y + threadIdx.y;

    for (int k = 0; k < K; k += 16) {
        load_matrix_sync(a_frag, a + warpM*16*K + k, K);
        load_matrix_sync(b_frag, b + k*N + warpN*16, N);
        mma_sync(acc_frag, a_frag, b_frag, acc_frag);
    }
    store_matrix_sync(c + warpM*16*N + warpN*16, acc_frag, N, mem_row_major);
}

The WMMA API is the readable entry point. It gets you maybe 40-50% of cuBLAS. Production kernels use raw mma.sync PTX or CUTLASS, with software pipelining, async copy, and careful shared memory swizzling.

Write this once to understand fragments. Then use CUTLASS.

2:4 structured sparsity#

Ampere+ can skip zeros in a specific pattern: 2 of every 4 consecutive weights must be zero.

Dense:   [a b c d e f g h]
Sparse:  [a 0 c 0 0 f 0 h]   ← exactly 2 nonzero per group of 4
         stored as values [a c f h] + 2-bit indices
→ 2x tensor core throughput, 2x less weight memory

The catch: the model must be trained or fine-tuned to satisfy the pattern with acceptable quality loss. Adoption has been limited because the quality cost is real and the retraining requirement is a barrier. [EMERGING] rather than [ESTABLISHED] for LLM inference — and note that vendor “peak FLOPs with sparsity” numbers assume it.

Accumulation precision#

FP16 inputs → FP32 accumulator (always)
BF16 inputs → FP32 accumulator
FP8  inputs → FP32 accumulator (or FP16 on some paths)
INT8 inputs → INT32 accumulator

The hardware forces this, correctly. A K=4096 reduction accumulated in FP16 would lose unacceptable precision.

6. Under the hood#

Verify usage in a profile:

ncu --metrics sm__inst_executed_pipe_tensor.sum,\
sm__pipe_tensor_cycles_active.avg.pct_of_peak_sustained_active ./bench

If sm__inst_executed_pipe_tensor.sum is 0 for a matmul kernel, tensor cores are not being used. The kernel name is also a giveaway:

ampere_h16816gemm_128x128_...       ← h = half, 16816 = m16n8k16. Tensor cores. ✓
sgemm_128x128_nn                    ← s = single precision. CUDA cores. ✗

7. Performance implications#

Path                                    A100 TFLOP/s
FP32 CUDA cores                          19.5
TF32 tensor cores                       156     (8x)
FP16/BF16 tensor cores                  312     (16x)
INT8 tensor cores                       624     (32x)
FP16 with 2:4 sparsity                  624     (32x)

H100:
BF16                                    990
FP8                                    1979

For LLM inference: prefill is compute-bound and gets the full benefit. Decode is memory-bound and gets almost none from tensor cores — but gets the benefit from the precision reduction that enables them (fewer bytes). Same change, different mechanism, different phase.

8. Production implications#

• Verify tensor core usage for prefill kernels. A silent fallback is a 10x regression.
• Enable TF32 (torch.backends.cuda.matmul.allow_tf32 = True) for any remaining FP32 math.
• Pad dimensions to multiples of 8/16. This is why vocabulary sizes are padded to values like 128256 rather than 128000.
• Use FP8 on Hopper if quality permits — it’s a 2x compute and 2x memory win.
• Don’t cite sparse peak FLOPs unless you’re actually using 2:4 sparsity.
• Check alignment when slicing tensors — a misaligned pointer can disable tensor core paths.

9. Common mistakes#

Running FP32 without TF32. 8x slower for free.

Dimensions not multiples of 8. Silent fallback or padding waste.

Assuming tensor cores help decode. They help via precision (bytes), not via FLOPs.

Hand-writing scalar matmul loops. Never touches tensor cores.

Quoting sparse TFLOPs in capacity planning.

Not verifying in the profile. The failure is silent.

10. Hands-on exercise#

A. Measure the gap. Run the benchmark in section 4. Record all four numbers in numbers.md. Compute the ratio of FP16 tensor to FP32 CUDA-core throughput. Does it match the spec sheet?

B. Shape sensitivity. Benchmark (M, K) @ (K, N) with K=N=4096 and M ∈ {4093, 4096, 4100}. Explain any differences.

C. Write WMMA. Implement the WMMA kernel in section 5. Verify correctness. Benchmark against cuBLAS. What fraction do you achieve?

D. Verify in a profile. Profile an LLM prefill and confirm tensor core instructions are being executed. Then force FP32 (disable TF32) and confirm they aren’t.

E. FP8 (Hopper only). Benchmark FP8 vs BF16 matmul. Measure the speedup for prefill-shaped and decode-shaped problems. Explain the difference.

11. Interview questions#

1. What is a tensor core and why is it 15-30x faster than CUDA cores for matmul?
2. What conditions must be met to use tensor cores?
3. Why is the accumulator always higher precision than the inputs?
4. Why do tensor cores help prefill much more than decode?
5. What is TF32 and why is it “free”?
6. What is 2:4 sparsity and why isn’t it widely used for LLMs?
7. How would you verify that a kernel is using tensor cores?

12. Further reading#

• [REFERENCE] CUDA C++ Programming Guide, “Warp Matrix Functions” (WMMA)
• [REFERENCE] PTX ISA documentation, mma.sync and wgmma
• [ESTABLISHED] CUTLASS repository and documentation
• [FUNDAMENTAL] NVIDIA architecture whitepapers, tensor core sections
• Next: 12 — Profiling CUDA