A tensor core is not a faster floating-point unit. It is a different kind of unit: instead of taking two scalars and producing one, it takes two small matrices and a running accumulator and produces a whole tile at once. That change of granularity is where the throughput comes from, and also where every difficulty comes from: the instruction has fixed shapes, wants its operands laid out a hardware-defined way across a whole warp, is fussy about alignment, and will silently decline to run at all. This piece walks the primitive, the programming interfaces, why mixed precision is really a statement about the accumulator, what each rung of the dtype ladder costs, and how to tell whether your kernel touches the pipe at all.
The MMA primitive — a matrix instruction, not a wider FMA
A tensor core implements matrix multiply-accumulate: D = A × B + C, where A, B, C and D are small dense tiles rather than scalars. One issue of the instruction performs every multiply and every add in that tile product. A classic FMA does one multiply-add per lane per cycle; an MMA retires hundreds, because the multiplier array and adder tree form a systolic block rather than independent lanes.
The consequence worth internalising is data reuse. Every element of A participates in many multiplies, and so does every element of B, so an operand is fetched once from the register file and used repeatedly inside the array. The instruction's own arithmetic intensity is high, which is why the pipe can be fed at all — and why anything that breaks the tile structure, like a skinny matrix, erodes the advantage before you write a line of code.
Why a warp cooperates on one tile
The most confusing property of the MMA instruction is that it is warp-scoped. It is not a per-thread operation that happens to be fast; all 32 lanes issue it together and collectively own one tile. The A, B and accumulator tiles live distributed across the lanes' registers, and no thread holds a complete row or column of anything.
That distribution is the fragment layout: architecture-defined and deliberately opaque. Hence the interface hands you load_matrix_sync and store_matrix_sync instead of letting you index fragment elements: the library knows the pattern for the target architecture and you do not, and code that hard-codes one layout breaks on the next generation. The _sync suffix is the other half — the whole warp must arrive, converged, before the instruction can issue. A divergent branch around an MMA is not a slowdown, it is undefined behaviour.
WMMA and MMA — two levels of interface
There are two ways to reach the matrix pipe directly. The WMMA API is the C++ one in the nvcuda::wmma namespace: declare fragment objects for A, B and the accumulator, fill them with load_matrix_sync, call mma_sync, write back with store_matrix_sync. It is portable and hides the layout entirely, at the cost of hiding the shape and scheduling choices that matter most for peak throughput.
Below it sit PTX mma.sync and, on newer parts, warpgroup-level instructions, where operand shape, operand source and synchronisation are explicit. Kernels chasing peak live here or in CUTLASS, which wraps them in templates that also handle tiling and pipelining. Most people should write neither: cuBLAS, cuDNN and framework backends already pick tuned kernels. Reach for WMMA to fuse a matmul into something bespoke, and for CUTLASS once you have measured the library kernel leaving throughput behind.
Instruction shape versus the tile your kernel chooses
Two things get called “the tile size,” and conflating them causes real confusion. The instruction shape — the M×N×K of one MMA — is fixed by the hardware and the operand dtype; you pick from a small architecture-specific menu and cannot invent a shape. The CTA and warp tile is yours: how much of the output one thread block computes, subdivided among its warps and built from many MMA instructions looping over K.
The kernel-chosen tile is the real tuning knob. Larger tiles reuse each operand more and cut trips to HBM, but burn registers and shared memory, limiting how many blocks an SM can host. Because the matrix pipe drains operands so fast, feeding it is the hard part: asynchronous copies, TMA, multi-buffered shared-memory stages and dedicated producer warps all exist to keep the MMA units busy, and each has its own article here.
Mixed precision — narrow operands, wide accumulator
“Mixed precision” in the tensor-core sense means one thing: the operands entering the multiplier are narrow and the accumulator is wider. A typical path multiplies fp16 or bf16 A and B tiles and accumulates the products into fp32. The products are never rounded back to the input format before being summed; they enter a wide accumulator and stay there for the whole K loop.
This asymmetry is the entire design. Narrow inputs buy throughput and halved memory traffic: fewer bits per operand means more operands per cycle and per byte of bandwidth. The wide accumulator buys correctness. This is a property of the instruction; the training recipes built on top of it (loss scaling, fp32 master weights, optimizer-state precision) are a separate concern.
Why the accumulator width is what protects the numerics
It looks backwards that you can be reckless with inputs but not with the accumulator, yet the two errors behave differently. Rounding an input to fp16 or bf16 is a one-time, bounded relative error on that element — small, and it does not grow. Accumulation error compounds over K, each partial sum inheriting the error of the one before it.
There is a sharper failure than drift. Once the running partial sum grows much larger than the next addend, the addend falls off the end of the mantissa and rounds away entirely: in a narrow accumulator, contributions stop being counted at all. That is systematic loss, not noise, and it worsens as K grows — which is why an fp16-in/fp16-accumulate path can look fine at a toy K and destroy a real reduction. Push the operand format down as far as the model tolerates; leave the accumulator alone.
The dtype ladder — what each step trades away
Each rung down the ladder halves the operand bits and gives something up. TF32 is the gentle first step for code that thinks it is doing fp32: keep fp32's exponent range, truncate the mantissa so values fit the matrix pipe, accumulate in fp32. fp16 has ten mantissa bits but a narrow exponent, so it is precise in range and overflows or flushes to zero outside it. bf16 chooses the opposite — fp32's range, far fewer mantissa bits — which is why it became the training default: range failures are catastrophic, precision loss is largely absorbed by the wide accumulator.
fp8 comes in two flavours that make the trade explicit: an E4M3-style encoding favouring precision, an E5M2-style one favouring range. At eight bits the dynamic range is so narrow that per-tensor or finer scaling factors become mandatory rather than optional; four-bit and integer paths push the same logic further.
Structured sparsity — the 2:4 pattern
Recent tensor cores exploit a structured sparsity pattern: within every aligned group of four values along the reduction dimension, at most two may be non-zero. The rigidity is the point: unstructured zeros land anywhere and cannot be skipped without irregular control, whereas under a fixed 2-of-4 rule the compressed operand plus a small metadata field naming which two positions survived lets the hardware select the matching elements of the other operand and skip the rest.
So the sparse operand is stored at roughly half size plus metadata, and the sparse MMA path does the tile in about half the math — an upper bound of roughly 2× on that instruction, illustrative rather than a delivered number. End-to-end gains are usually far smaller: real kernels are often limited by operand movement rather than matrix math, only one operand (in practice the weights) is pruned, and getting there requires prune-and-retrain to recover accuracy.
The rules that silently disable tensor cores
The failure mode that costs the most time is not a crash. The library selects a non-tensor kernel, the result is numerically fine, and you are several times slower with nothing in the log. The usual causes:
| Cause | Why it disqualifies the MMA path |
|---|---|
| fp32 inputs, TF32 not enabled | Strict fp32 has no matrix path; the flag (cuBLAS math mode, allow_tf32) must be on |
| Misaligned pointers or leading dimensions | Operand loads are wide vector accesses; below that alignment the tuned kernel is ineligible |
| K not a multiple of the instruction's K | The ragged tail needs padding or a fallback epilogue; a naive path may win instead |
| Mismatched dtypes | An accidental upcast in one operand pushes selection off the mixed-precision path |
| Unsupported layout or transpose combination | Not every row/column-major pairing has a tuned tensor kernel |
Most reduce to one habit: pad dimensions to friendly multiples, keep allocations aligned, and be deliberate about dtype at every boundary.
Diagnosing whether a kernel actually used them
Do not infer tensor-core usage from having asked for bf16. The direct evidence is a hardware counter: a kernel profiler reports tensor-pipe utilization alongside the other pipes, and a matmul-heavy kernel showing near-zero there is not using the matrix units. The kernel name is a fast secondary signal — library kernels usually encode the dtype path and instruction family in their symbol.
The cheapest structural check is arithmetic: compute the FLOPs the GEMM must perform, divide by measured kernel time, and compare against the vendor's dense and matrix-pipe rates for your part. Landing near the non-matrix rate strongly suggests a fallback. When a shape looks suspicious, pad it and re-measure — a large jump from padding alone is the signature of an eligibility rule you were missing. Verify per shape and dtype; do not assume it carries over.
Related: GPU Tensor Cores — Hardware Matrix Math for Deep Learning covers the GEMM data path built to keep this unit fed, and how to tell from a profiler whether it is.