GPU Architecture · Optimization
A second kind of arithmetic unit that multiplies whole matrix tiles in one shot — what it does, what work it speeds up, and why feeding it is the hard part.
Until now, one thread has meant one lane doing scalar arithmetic (Lesson 2). But every recent NVIDIA SM carries a second kind of math unit beside those lanes — the Tensor Core — and it is the reason a modern GPU is quoted in hundreds of "AI TFLOPS." It does not make your existing kernels faster on its own. It accelerates exactly one shape of computation, and getting the speedup is a memory problem before it is a math problem. This lesson opens the optimization arc: taking the machine you now understand and making it fast.
A CUDA core multiplies two numbers; a Tensor Core multiplies two matrices. Its single operation is D = A×B + C on small tiles, with low-precision inputs accumulated in high precision. Anything you can express as dense matrix multiply — neural nets, convolutions, attention, dense linear algebra — rides it. Anything else does not. And because it is so fast, the whole game becomes keeping it fed from on-chip memory.
It is tempting to picture a Tensor Core inside a CUDA core. It isn't. Recall the SM from Lesson 2: four sub-partitions, each with ~32 FP32 lanes, a warp scheduler, and a register file. A Tensor Core sits beside those lanes — one per sub-partition, four per SM on current architectures — and it is a warp-level shared resource, not something a single thread owns.1
One SM sub-partition
Each does one scalar FMA per clock — a*b + c on single numbers. One thread per lane.
Does a whole matrix-multiply-accumulate tile per operation. Fed by the whole warp at once.
So you cannot think "each thread has a Tensor Core." A warp issues one Tensor Core instruction and all 32 lanes cooperate to feed a single tile through it. That warp-level cooperation is baked into the API you'll see below — every call ends in _sync. (This is also why AMD's equivalent, the Matrix Core, is wavefront-level, and why optimized matmul code on either vendor is organized around warps/wavefronts, not threads.)
A Tensor Core is not general-purpose. It cannot branch, load, add scalars, or compute a transcendental. It does exactly one thing, parameterized by tile shape and datatype — a matrix multiply-accumulate, "MMA":
The hardware works in fixed tile shapes — the canonical one you program against is 16×16×16 (a 16×16 tile of D, formed from a 16×16 slab of A times a 16×16 slab of B). To multiply larger matrices you tile them into 16×16 blocks and sweep, accumulating — which is exactly what the code does. One MMA replaces the inner loop of scalar FMAs; that is the whole source of the speedup.
The trick is mixed precision: inputs come in low precision, but the products accumulate into a wider format so error doesn't compound over a long dot product. Each GPU generation widened the menu, and each new type unlocked a different workload:2
| Generation | Input types added | Accumulate | What it unlocked |
|---|---|---|---|
| Volta | FP16 | FP32 | the first deep-learning training boom |
| Turing | INT8, INT4, INT1 | INT32 | quantized inference |
| Ampere | TF32, BF16, FP64, 2:4 sparsity | FP32 / FP64 | FP32 code sped up untouched; HPC on Tensor Cores |
| Hopper | FP8 (E4M3 / E5M2) | FP32 | large-language-model training & inference |
| Blackwell | FP4, FP6 (microscaling) | FP32 | your RTX 5050's 5th-gen cores; low-precision inference |
Two of these are worth pausing on:
The rule is simple and broad: anything you can express as dense matrix multiplication (GEMM). A surprising fraction of real computing reduces to it — which is the whole reason this hardware exists.
In a real network the Tensor Cores do the heavy layers and the ordinary lanes do everything between them — the activation functions, the residual adds, the layernorms. Both kinds of unit run; the art is overlapping them.
Here is where the optimization arc really begins. A Tensor Core is so fast that data movement, not arithmetic, becomes the bottleneck. Go back to the roofline from Lesson 4: adding Tensor Cores raises the compute ceiling enormously — but the memory bandwidth line doesn't move. So the machine balance shifts hard, and the arithmetic intensity you need to stay compute-bound shoots up.
If every 16×16 MMA fetched its inputs from DRAM, you'd be pinned to the bandwidth wall and the Tensor Cores would sit idle waiting. The fix is the two-sided tiling pattern from Lesson 6: stage tiles of A and B into shared memory once, then let many MMAs reuse them. Each byte crosses the memory bus once and feeds a whole row/column of tile products. That reuse — arithmetic intensity — is the only thing that converts Tensor Core peak into real throughput. This is why __syncthreads() and shared memory, which felt abstract in Lesson 6, are load-bearing here.
Newer generations attack the same feed problem in hardware: Hopper's wgmma reads operands directly from shared memory and runs asynchronously; Blackwell adds dedicated tensor memory (TMEM) to shorten the path further. All of it is in service of one goal — keep the tile inputs close so the Tensor Core never starves.
You will almost always stay at the top of this ladder. Descend only when you need more control than the level above gives you.
These dispatch to Tensor Cores automatically when your dtype and shape qualify. PyTorch calling cuBLAS/cuDNN is how the overwhelming majority of Tensor Core cycles are actually issued. You write a matmul; the library picks the hardware.3
nvcuda::wmma in CUDA C++Warp-cooperative fragments with load_matrix_sync / mma_sync / store_matrix_sync. The level where you can see the tile mechanism — and the level the code example uses.1
mma.sync / wgmmaMaximum control over exactly which instruction and layout, for squeezing the last few percent. This is what CUTLASS emits under the hood; you rarely write it by hand.
The companion example, 04-wmma-matmul.cu, computes the same matrix multiply two ways — scalar FP32 on the CUDA cores and 16×16 tiles on the Tensor Cores via WMMA — checks both against a CPU reference, and times them. The WMMA kernel is the whole idea in a dozen lines:
// One WARP owns one 16×16 output tile.
wmma::fragment<wmma::matrix_a, 16,16,16, half, wmma::row_major> aFrag;
wmma::fragment<wmma::matrix_b, 16,16,16, half, wmma::row_major> bFrag;
wmma::fragment<wmma::accumulator, 16,16,16, float> acc;
wmma::fill_fragment(acc, 0.0f); // C tile starts at zero
for (int k = 0; k < n; k += 16) {
wmma::load_matrix_sync(aFrag, A + rowTile*n + k, n); // stage a 16×16 A tile
wmma::load_matrix_sync(bFrag, B + k*n + colTile, n); // and a 16×16 B tile
wmma::mma_sync(acc, aFrag, bFrag, acc); // acc += aFrag × bFrag (one MMA)
}
wmma::store_matrix_sync(C + rowTile*n + colTile, acc, n, wmma::mem_row_major);Four things to notice, each a direct echo of the concepts above:
fragment types are opaque. You never index them element by element — the hardware decides how a tile's 256 elements spread across the warp's 32 lanes' registers. You address tiles, not scalars._sync. All 32 threads of the warp execute each one together; that's the warp-level cooperation the hardware requires, made explicit.half, the accumulator is float. Mixed precision, exactly as the datatype table said — and why the example's error is small FP16 rounding, not a bug.k loop is the "sweep." Each step does one MMA and accumulates; larger matrices are just more tiles. This inner loop of scalar FMAs in the naive kernel becomes a single mma_sync here.Build and run it on the 5050: nvcc -arch=sm_120 04-wmma-matmul.cu -o wmma && ./wmma. You'll see the Tensor Core version land several times the naive kernel's GFLOP/s while agreeing with the CPU reference. Then — the arc's real payoff — profile it: ncu --set full ./wmma and watch this kernel report as compute-bound on the Tensor pipe, where the naive one is limited by the CUDA-core FMA rate. This is a modest, un-tiled WMMA kernel; closing the gap to cuBLAS is a shared-memory feeding problem, which is exactly where the arc heads next.
From memory — one click locks each answer.
Mental modelA Tensor Core, relative to a CUDA core, is best described as —
Tensor Cores sit beside the CUDA-core lanes — one per sub-partition, four per SM — and are shared by the whole warp. A CUDA core multiplies scalars; a Tensor Core multiplies 16×16 tiles. Not inside a lane, not per-thread.
What acceleratesWhich of these will not be sped up by Tensor Cores?
Tensor Cores accelerate dense matrix multiply. FC layers, attention's matmuls, im2col/implicit-GEMM convolutions, and blocked factorizations all reduce to GEMM. An element-wise activation is not a matmul — it stays on the ordinary CUDA-core lanes.
The catchYou replace a kernel's inner loop with mma_sync but see little speedup. The most likely reason —
Tensor Cores are so fast that feeding them dominates. Fetching each tile straight from DRAM pins you to the bandwidth wall (Lesson 4). The fix is tiling (Lesson 6): stage A and B tiles into shared memory once and reuse them across many MMAs. FP32 accumulation is normal and correct; the API works fine in your own kernels.
CUDA C++ Programming Guide — Warp Matrix Functions (WMMA): fragment types, load/mma/store_matrix_sync, supported tile shapes and datatypes. Read it here. Then skim NVIDIA's "Programming Tensor Cores in CUDA" for the picture end to end: developer.nvidia.com.
The optimization arc is open. You now have a second engine and the one law that governs it: feed it or waste it. Next we make that law bite — a properly shared-memory-tiled matmul (WMMA fed from on-chip tiles, and a plain FP32 tiled version measured against your 320 GB/s) profiled under Nsight Compute, so you watch arithmetic intensity climb and the memory-stall reasons from Lesson 10 fall on your own code. Every physical lesson becomes a number you can move. Say the word when you're ready.
nvcuda::wmma API operates at warp scope on opaque fragment operands; load_matrix_sync, mma_sync, store_matrix_sync, fill_fragment; supported shapes (incl. 16×16×16) and input/accumulator type pairs. Tensor Cores are per-sub-partition SM units (one per sub-partition, four per SM on Turing and later). docs.nvidia.com/cuda/cuda-c-programming-guidewgmma) and Blackwell (FP4/FP6 microscaling, 5th-gen Tensor Cores) — NVIDIA Tensor Core overview: nvidia.com/tensor-cores. Illustrative throughput: A100 dense FP16 Tensor ≈ 312 TFLOPS vs FP32 ≈ 19.5 TFLOPS (~16×).