Warp Matrix Multiply and Accumulate (WMMA) Architecture

The Physics of NVIDIA Tensor Cores

Standard CUDA cores execute scalar math (one FMA per clock). Tensor Cores execute entire matrix multiplications per clock cycle at the warp level (32 threads). In an Hopper (H100) or Blackwell (B200) GPU, a single Tensor Core computes $D = A imes B + C$ where $A$ and $B$ are $16 imes 16$ matrices.

WMMA Fragment Geometry

All 32 threads in a warp cooperate to hold the matrix fragments in their register files:

// CUDA PTX WMMA instruction format accessed via Rust FFI
// wmma.mma.sync.aligned.m16n16k16.row.col.f32.f16.f16.f32
// Computes 16x16x16 GEMM chunk in a single hardware cycle!

Accessing Tensor Cores from Rust directly or through cuBLAS achieves up to 989 TFLOPS of FP16/FP8 compute on H100 SXM5, compared to ~67 TFLOPS on standard CUDA FP32 cores.