What is a Tensor Core?

A Tensor Core is a GPU core built to execute one matrix multiply-accumulate instruction across an entire tile of two matrices, rather than the single scalar multiply-add a regular CUDA Core performs per instruction. Packing more arithmetic behind one instruction fetch cuts the power spent on decode and control logic per useful FLOP, and that saved power gets spent instead on raw throughput. NVIDIA's own chief scientist, Bill Dally, has made this exact argument in talks on GPU architecture: since Tensor Cores arrived, they have been the only route to a modern NVIDIA GPU's advertised peak FLOPS, delivering on the order of 100x more throughput than the same chip's CUDA Cores.

Introduced with Volta

Tensor Cores first appeared in the Volta Streaming Multiprocessor generation, inside the V100, which NVIDIA's own architecture white paper for that chip frames as the turning point that made GPUs a serious training platform for large neural networks rather than just a fast vector processor repurposed for the job. Every NVIDIA data-center GPU released since has shipped some generation of Tensor Core, and each architecture generation has widened what a single instruction can do: Volta's Tensor Cores handled FP16 multiply-accumulate, Turing added INT8 and INT4 paths for inference, Ampere introduced the TF32 format and fine-grained structured sparsity, and Hopper's Tensor Cores added FP8 support alongside a dedicated Transformer Engine that switches precision per layer automatically. Blackwell pushes the same trend further with native FP4 and FP6 formats, aimed squarely at squeezing more inference throughput out of the same die area and power budget.

What one instruction computes

A concrete example makes the "whole matrix per instruction" claim precise. The SASS-level instruction HMMA16.16816.F32 computes D = A×B + C for four matrices, where C and D are usually the same physical memory. "HMMA" is short for half-precision matrix-multiply-accumulate: the inputs A and B are FP16, and the accumulation happens in FP32 to avoid losing precision across many additions. The three numbers packed into the instruction name, 16, 8, and 16, are the m, n, and k dimensions NVIDIA uses in its own documentation: the outer dimensions of A and B come first, then the shared inner dimension the multiplication reduces over. Multiply those out and one instruction issues 16 × 8 × 16 = 2,048 multiply-accumulate operations.

No single thread executes that instruction alone. It is issued once per warp, and all 32 threads in the warp cooperate to produce the result, coordinated by the SM's warp scheduler. Spread the 2,048 MACs across 32 threads and you get 64 MACs per thread for that one instruction, still far more than a CUDA Core manages per cycle. Because of this whole-matrix-per-instruction shape, it's more accurate to think of a Tensor Core as complex-instruction-set-computer (CISC) hardware than as a scaled-up scalar core, a framing computer architect David Patterson (who coined the terms CISC and RISC) has applied to the similar systolic-array design of Google's TPUs.

Getting there from CUDA C++

That SASS instruction is the end of a compilation chain, not something a programmer writes directly. A CUDA C++ call to a warp-level matrix intrinsic like wmma::mma_sync compiles down through PTX intermediate representation (a wmma.mma.sync instruction operating on 16×16 tiles) and finally to a handful of SASS HMMA instructions, each covering a smaller sub-tile of the original 16×16 multiplication. Reading that generated SASS is a good way to see the pattern directly: a 16×16-by-16×16 multiply typically lowers to four HMMA.1688.F32 instructions, each an outer product between an 8-wide slice of one matrix and an 8-tall slice of the other, accumulated together into the final result.

Hand-writing at this level gets impractical fast, and on Hopper and Blackwell it stops being optional in pure CUDA C++ altogether: getting near peak Tensor Core throughput on those architectures requires PTX-level intrinsics for both compute and memory movement. In practice, almost nobody hand-writes Tensor Core kernels for production use. Most code reaches Tensor Cores through a vendor library like cuBLAS, a fixed catalog of tuned linear-algebra routines, or a template toolkit like CUTLASS (and its underlying CuTe layout library), which builds custom Tensor Core kernels from composable pieces instead of calling a fixed function.

Few, but enormous

Tensor Cores are far less numerous than CUDA Cores on the same chip, and each one is far larger. An H100 SXM5 SM has only four Tensor Cores (one per warp scheduler) alongside well over a hundred CUDA Cores, yet the four Tensor Cores contribute the overwhelming majority of the chip's dense FLOPS. That ratio is the reason a spec sheet's raw "core count" is close to meaningless for AI workloads without knowing the CUDA-Core/Tensor-Core split behind it, and why the two are worth tracking separately rather than as one number; see Aquanode's GPU pages for that breakdown on H100, H200, and B200 instances.

What actually happens electrically inside a Tensor Core is not public. NVIDIA has not documented the internal microarchitecture, and it is widely assumed, by analogy to TPUs, to be some form of systolic array, but there is no real consensus in the independent microbenchmarking literature on that point, and the design is understood to differ across architecture generations rather than staying fixed.

If you are picking hardware for a Tensor-Core-bound workload like transformer training or inference, the practical question is less "how many Tensor Cores" and more "which precision and sparsity formats does this generation support, and does my kernel library (cuBLAS, CUTLASS, or a framework's fused kernels) actually hit them." Aquanode's H100, H200, and B200 listings carry the Tensor Core generation and supported precisions for exactly that comparison.