Tensor cores are dedicated units that compute a small matrix multiply-accumulate in one instruction, $D = A \times B + C$, on tiles of roughly 16×16. Introduced with the Volta architecture, they exist because dense matrix multiplication dominates deep learning and much of scientific computing, and a specialised unit does it far faster than the general-purpose lanes.

The speed comes from doing more arithmetic per instruction and from operating at reduced precision. Inputs are typically FP16, BF16, TF32, or INT8, while accumulation happens in FP32 to keep the error acceptable:

FP32 on regular cores:     1x     (baseline)
TF32 on tensor cores:      ~8x
FP16 on tensor cores:      ~16x

The precision trade is the catch. Tensor cores are a good fit for neural network training and inference, where reduced precision is tolerable and often harmless. They are a poor fit for anything needing full FP64, which is why HPC codes with strict accuracy requirements see no benefit and why FLOPS figures quoted for a GPU need the precision attached to mean anything.

Most code should reach them through a library rather than directly. cuBLAS, cuDNN, and CUTLASS all dispatch to tensor cores automatically when the data types and dimensions allow it, and they are better tuned than hand-written code will be.

Direct use goes through the WMMA API:

#include <mma.h>
using namespace nvcuda::wmma;
 
fragment<matrix_a, 16, 16, 16, half, row_major> a;
fragment<matrix_b, 16, 16, 16, half, col_major> b;
fragment<accumulator, 16, 16, 16, float> acc;
 
fill_fragment(acc, 0.0f);
load_matrix_sync(a, ptr_a, lda);
load_matrix_sync(b, ptr_b, ldb);
mma_sync(acc, a, b, acc);
store_matrix_sync(ptr_d, acc, ldd, mem_row_major);

A fragment is distributed across a whole warp, so every WMMA call is warp-wide and all 32 threads must reach it. Matrix dimensions also have to be multiples of the tile size, so real kernels spend most of their code on tiling and padding rather than on the multiply itself.