Tensor Core

A Tensor Core is a specialized execution unit inside NVIDIA GPUs that computes a small matrix multiplication and accumulation, D = A x B + C, as a single hardware operation instead of as a long sequence of scalar fused multiply-add instructions. NVIDIA introduced the units with the Volta architecture in 2017 and has extended them in every data center generation since, mainly by adding narrower numeric formats. Because the matrix multiply at the heart of a neural network layer is exactly the operation these units implement, Tensor Cores rather than conventional CUDA cores now supply almost all of the peak arithmetic throughput available for deep learning training and inference on NVIDIA hardware. On the A100, for example, NVIDIA quotes 19.5 TFLOPS of standard FP32 against 312 TFLOPS of dense FP16 Tensor Core throughput ^[4].

The first implementation, in the Tesla V100, performed a 4x4x4 matrix multiply-accumulate per clock per unit with FP16 inputs and FP32 accumulation ^[1]. NVIDIA announced the part on 10 May 2017, rating it at 120 teraflops of deep learning performance ^[2]; the architecture whitepaper published shortly afterwards lists up to 125 Tensor TFLOPS at GPU Boost clock ^[1]. Nine years later, NVIDIA rates a single Rubin GPU at up to 50 petaFLOPS of NVFP4 inference ^[14]. The two headline numbers are not measured the same way, since they use different formats and different sparsity assumptions, but the roughly 400-fold gap between them tracks what changed: more units per chip, a wider operation per unit, and a data format that shrank from 16 bits to 4.

The term is NVIDIA's, but the idea is not unique to NVIDIA. AMD ships Matrix Cores in its CDNA compute units ^[29], and Google's TPUs are built around a matrix multiply unit called the MXU ^[31]. Google also uses the name "TensorCore" for the larger block that contains its MXUs, vector unit, and scalar unit ^[31], which is a different level of the hierarchy from an NVIDIA Tensor Core.

What a Tensor Core computes

General matrix-matrix multiplication (GEMM) computes D = A * B + C, where A is an m-by-k matrix, B is k-by-n, and C and D are m-by-n. The problem size is conventionally written m-by-n-by-k ^[4]. Tensor Cores implement small fixed shapes of this operation directly in hardware, and libraries build large GEMMs out of many such tiles.

In Volta, each Tensor Core performed 64 floating point FMA operations per clock, enough for one 4x4x4 mixed-precision matrix multiply per clock. Each streaming multiprocessor (SM) held eight of them, two per processing block, giving 512 FMA operations per SM per clock, and the full Tesla V100 contained 640 Tensor Cores ^[1]. The mixed-precision behavior matters as much as the throughput: the FP16 multiply produces a full-precision product, which is then accumulated in FP32 with the other partial products ^[1]. That arrangement is what makes reduced-precision matrix math usable for training, and it pairs with the software technique described in the 2017 mixed precision training paper by Micikevicius and colleagues, which keeps an FP32 master copy of the weights and applies loss scaling to preserve small gradient values ^[26].

Ampere redesigned the unit rather than just replicating it. Each A100 Tensor Core executes 256 FP16 FMA operations per clock, computing an 8x4x8 mixed-precision matrix multiply, and each SM holds four of them for 1,024 FP16 FMA operations per SM per clock. With 108 SMs the A100 has 432 Tensor Cores and 312 TFLOPS of dense FP16 throughput ^[4]. Hopper kept four Tensor Cores per SM, with 132 SMs and therefore 528 Tensor Cores on the H100 SXM5, and doubled the per-SM MMA rate again for equivalent data types ^[5].

Generational evolution

NVIDIA numbers the generations from Volta. Each one adds numeric formats, so the practical story of Tensor Cores is a steady descent in precision, although the sub-8-bit integer modes introduced on Turing later dropped out of NVIDIA's published data center data type lists.

Scroll sideways for more →

Volta supported only FP16 inputs. Against a Pascal-based Tesla P100, NVIDIA claimed up to 12x higher peak TFLOPS for training relative to standard FP32 and up to 6x for inference relative to standard FP16 ^[1].

Turing added integer modes for inference workloads that tolerate quantization. A TU102 GPU contains 576 Tensor Cores, and the INT8 mode runs at double the FP16 rate, 2,048 integer operations per SM per clock ^[3]. The Tesla T4 was rated at up to 130 TOPS of INT8 and 260 TOPS of INT4 ^[3]. NVIDIA's later A100 whitepaper credits Turing with a binary 1-bit mode as well ^[4]. Turing also put Tensor Cores into consumer GeForce parts for the first time, where they run Deep Learning Super Sampling ^[3].

Ampere added three formats and one structural feature. TF32 has an 8-bit exponent (the range of FP32) and a 10-bit mantissa (the precision of FP16), reads and writes standard FP32 data, and is intended to accelerate existing FP32 code without changes ^[4]. BF16 runs at the same rate as FP16. A new double-precision MMA instruction replaced eight DFMA instructions and lifted A100 FP64 to 19.5 TFLOPS ^[4]. Finally, fine-grained structured sparsity enforces a 2:4 pattern, two non-zero values in every four-entry vector, which lets the hardware skip the zeros and doubles Tensor Core throughput. NVIDIA's recipe trains dense, prunes to the 2:4 pattern, then fine-tunes the surviving weights; the whitepaper notes that using sparsity to accelerate training, as opposed to inference, remained an open research area ^[4].

Hopper and Ada Lovelace introduced FP8 in two encodings: E4M3 with four exponent and three mantissa bits, and E5M2 with five exponent and two mantissa bits ^[5]. NVIDIA, Arm, and Intel proposed the same pair as a standard interchange format in a joint paper in September 2022 ^[27]. NVIDIA states that a Hopper SM delivers twice the MMA rate of an A100 SM on equivalent data types and four times the rate using FP8 ^[5]. The H100 SXM specification lists 989 TFLOPS TF32, 1,979 TFLOPS FP16 and BF16, 3,958 TFLOPS FP8 and 3,958 TOPS INT8, all with sparsity, plus 67 TFLOPS on the FP64 Tensor Cores ^[6]. On the consumer side, the fourth-generation unit gave the GeForce RTX 4090 660.6 dense FP8 Tensor TFLOPS, or over 1.3 petaFLOPS with sparsity ^[7].

Hopper also shipped the Transformer Engine, a mix of hardware and software that analyzes the statistics of the output values the Tensor Cores produce at each layer of a transformer model and scales tensor data into the representable FP8 range using factors computed from those statistics, choosing between FP8 and 16-bit calculations layer by layer ^[5]. The open-source library of the same name exposes this to PyTorch, JAX, Megatron-LM, and other frameworks; FP8 support requires compute capability 8.9 or higher, and on Blackwell the library adds MXFP8 and NVFP4 ^[23].

Blackwell, announced on 18 March 2024, packs 208 billion transistors into two reticle-limited dies joined by a 10 TB/s interconnect and presented as one GPU ^[8][9]. Its fifth-generation Tensor Cores add 4-bit and 6-bit floating point and a technique NVIDIA calls micro-tensor scaling, in which a scale factor is shared across a small block of values rather than a whole tensor ^[9]. A GB200 NVL72 rack of 72 Blackwell GPUs is rated at 1,440 PFLOPS of NVFP4 with sparsity and 720 PFLOPS dense ^[10]; an 8-GPU DGX B200 is rated at 144 PFLOPS sparse and 72 PFLOPS dense FP4, the same figures NVIDIA publishes for the HGX B200 board ^[11][32]. Blackwell Ultra, the mid-generation refresh, raises dense NVFP4 from 10 to 15 petaFLOPS per GPU, doubles the special function unit throughput used by attention layers for 2x faster attention-layer compute, and carries 288 GB of HBM3e ^[12].

Number formats and block scaling

The formats below 8 bits are not simply narrower floats. NVFP4 stores each value as E2M1, one sign bit, two exponent bits, and one mantissa bit, covering roughly -6 to 6, then attaches an FP8 E4M3 scale factor to every block of 16 values and a single FP32 scale to the whole tensor ^[13]. The effective cost is about 4.5 bits per value, which NVIDIA describes as a 1.8x memory reduction against FP8 and 3.5x against FP16. On DeepSeek-R1-0528 the company reports 1 percent or less accuracy degradation against the original FP8 model on key language modeling tasks ^[13].

The competing open standard, OCP microscaling, shares an 8-bit power-of-two exponent across a block of 32 values and defines MXFP8, MXFP6 (in E3M2 and E2M3 variants), and MXFP4 (E2M1) ^[29]. Blackwell supports both families, and the distinction is visible in the PTX instruction set: matrix instructions take a .kind qualifier of kind::mxf8f6f4, kind::mxf4, or kind::mxf4nvf4, with a scale type of .ue8m0 for the microscaling formats and .ue4m3 for NVFP4, and a .scale_vec size that selects the block granularity ^[18]. The same instruction listings show the full low-precision operand set as .e4m3, .e5m2, .e3m2, .e2m3, and .e2m1 ^[18].

Programming model

Tensor Cores are not addressable by a single thread. A CUDA warp of 32 threads cooperates on each matrix operation, and the exposed abstraction has grown wider with each generation.

The original interface is the warp matrix functions API in the nvcuda::wmma namespace, added in CUDA 9 and available on devices of compute capability 7.0 and higher ^[17][19]. Programs declare fragment objects for matrix_a, matrix_b, and accumulator, then call load_matrix_sync, mma_sync, and store_matrix_sync; every thread in the warp must participate, and the mapping of matrix elements into fragment storage is deliberately unspecified so NVIDIA can change it between architectures ^[17]. Supported warp-level tiles include 16x16x16, 32x8x16, and 8x32x16 for FP16 and 8-bit integer inputs, 16x16x8 for TF32, 8x8x4 for FP64, and 8x8x32 or 8x8x128 for the experimental sub-byte and binary modes ^[17]. Below this sit the PTX mma.sync instructions and their sparse mma.sp variants ^[18].

Hopper widened the unit of cooperation from a warp to a warpgroup, defined as four contiguous warps whose first warp has a rank that is a multiple of four. Its asynchronous wgmma.mma_async instruction issues shapes such as m64n256k16, far larger than anything a single warp can hold, and pairs with the Tensor Memory Accelerator for bulk data movement ^[18][5]. Blackwell goes further and gives the fifth-generation Tensor Cores their own dedicated on-chip memory: on sm_100a the Tensor Memory is a 512-column by 128-lane array of 32-bit cells per CTA, explicitly allocated and freed by the kernel through tcgen05.alloc and related instructions ^[18].

Most developers never touch any of this. cuBLAS uses Tensor Cores automatically wherever possible unless a pedantic compute mode is selected, with CUBLAS_DEFAULT_MATH, CUBLAS_PEDANTIC_MATH, and CUBLAS_TF32_TENSOR_OP_MATH controlling the tradeoff and the NVIDIA_TF32_OVERRIDE environment variable able to override the library defaults ^[20]. cuDNN exposes an equivalent math-type setting for convolutions ^[19]. CUTLASS, NVIDIA's open-source template library, has provided CUDA C++ abstractions for GEMM since 2017 and covers Volta through Blackwell, including NVFP4, the OCP MX formats, 4-bit integers, and 1-bit data; version 3.0 introduced the CuTe layout algebra and version 4 added Python-native kernel authoring ^[22]. Triton reaches Tensor Cores through tl.dot, whose input_precision argument defaults to "tf32" on NVIDIA devices and is ignored when the device has no Tensor Cores or the inputs are not FP32 ^[24]. PyTorch exposes the same choice through torch.backends.cuda.matmul.allow_tf32, which has defaulted to False since PyTorch 1.12, and the newer fp32_precision settings that accept "ieee" or "tf32" ^[25].

Getting close to peak

Published peak numbers are ceilings that real kernels approach only under specific conditions. An early independent study by Markidis and colleagues measured up to 83 Tflops/s of mixed precision on a V100 against the 125 Tflops/s peak, and only about 4 Tflops/s for a WMMA-based batched GEMM, while noting that the precision loss from half-precision inputs can matter for scientific codes and can be reduced at the cost of extra computation ^[28].

NVIDIA's own performance guide sets out the practical constraints. Matrix dimensions should be aligned to multiples of 16 bytes, which works out to multiples of 8 elements for FP16 (64 on A100) and 16 elements for INT8 (128 on A100); cuBLAS 11.0 and cuDNN 7.6.3 and later will use Tensor Cores regardless, but efficiency still favors aligned sizes ^[21]. Tile quantization wastes work when a dimension is not divisible by the tile size, and wave quantization wastes it when the number of tiles is not a multiple of the SM count, which can roughly halve achieved GFLOPS ^[21]. Above all, a GEMM only benefits if it is math-bound rather than memory-bound: arithmetic intensity is the FLOP count divided by the bytes accessed, which for 16-bit data reduces to (mnk) divided by (mk + nk + m*n), and matrix-vector products, whose arithmetic intensity is less than 1, are always memory-limited ^[21]. This is why high-bandwidth memory capacity and bandwidth, kernel fusion, and techniques such as FlashAttention matter as much to delivered performance as the Tensor Core peak does.

Precision is the other tradeoff. PyTorch's documentation reports TF32 matrix multiplication running roughly 7x faster than FP32 on an A100 while producing relative error about two orders of magnitude larger when both are compared against double precision ^[25]. Sparsity numbers deserve similar care: the headline sparse figures assume the 2:4 pattern and the prune-and-fine-tune recipe, and vendor spec sheets flag them explicitly, with NVIDIA's DGX B200 page noting that dense performance is half the sparse specification shown ^[11].

Comparable hardware

AMD calls its equivalent units Matrix Cores and drives them with the MFMA (matrix fused multiply add) instruction family, reachable from HIP through compiler builtins such as __builtin_amdgcn_mfma_f32_16x16x4f32 ^[30]. CDNA 3 added the two OCP FP8 encodings; CDNA 4 added MXFP8, MXFP6, and MXFP4 and doubled the execution resources for 16-bit and narrower types ^[29]. The Instinct MI355X has 1,024 Matrix Cores and peaks at 2.5 PF dense FP16, 5.0 PF dense FP8, and 10 PF for MXFP6 or MXFP4 ^[29]. AMD also moved TF32 out of hardware in CDNA 4, supporting it through BF16-based software emulation, and halved matrix FP64 throughput relative to MI300X ^[29].

Google's TPU takes a different structural approach. Its MXU is a systolic array of multiply-accumulators, 128x128 on versions before v6e and 256x256 on v6e and TPU7x, and Google's documentation states that each MXU performs 16K multiply-accumulate operations per cycle ^[31]. Rather than embedding many small matrix units across many SMs, the TPU concentrates the matrix math into a small number of large arrays, with a vector unit in the same TensorCore handling work such as activations and softmax and a scalar unit handling control flow and address calculation ^[31].

Recent developments

By mid-2026 the Blackwell generation was in broad deployment and its successor was ramping into production. NVIDIA's Tensor Core page lists the same data type set for Blackwell and Rubin: NVFP4, FP64, TF32, BF16, FP16, FP8 and FP6, and INT8, with Rubin rated at up to 50 petaFLOPS of NVFP4 inference per GPU ^[14]. The Vera Rubin NVL72 specification, marked preliminary, gives 3,600 PFLOPS of NVFP4 inference and 1,260 PFLOPS of dense FP8 and FP6 training across 72 Rubin GPUs and 36 Vera CPUs, with 20.7 TB of HBM4 ^[15]. NVIDIA said on 21 July 2026 that Vera Rubin NVL72 production was ramping with racks running at CoreWeave, Google Cloud, Microsoft Azure, and Oracle Cloud Infrastructure, and cited a CoreWeave DeepSeek-R1 measurement of 10x more tokens per second per megawatt than Grace Blackwell NVL72 ^[16].

The direction of travel across vendors is consistent: narrower elements, finer-grained scaling blocks, larger hardware matrix shapes, and more of the surrounding data movement handled asynchronously by dedicated hardware. Some formats have quietly disappeared along the way. NVIDIA's published data type lists for Hopper and Blackwell stop at INT8, with no sign of the INT4 and binary modes added on Turing ^[6][14], and AMD moved TF32 out of hardware in CDNA 4 ^[29].

See also

• CUDA
• CUTLASS
• Mixed precision training
• NVIDIA Hopper
• NVIDIA Blackwell
• Tensor Processing Unit

References

1. ^NVIDIA, "NVIDIA Tesla V100 GPU Architecture" whitepaper, WP-08608-001_v1.1, 2017. images.nvidia.com/...a-architecture-whitepaper.pdf
2. ^NVIDIA Newsroom, "NVIDIA Launches Revolutionary Volta GPU Platform, Fueling Next Era of AI and High Performance Computing," 10 May 2017. nvidianews.nvidia.com/...igh-performance-computing
3. ^NVIDIA, "NVIDIA Turing GPU Architecture" whitepaper, WP-09183-001_v01, 2018. images.nvidia.com/...g-Architecture-Whitepaper.pdf
4. ^NVIDIA, "NVIDIA A100 Tensor Core GPU Architecture" whitepaper, 2020. images.nvidia.com/...e-architecture-whitepaper.pdf
5. ^NVIDIA Technical Blog, "NVIDIA Hopper Architecture In-Depth." developer.nvidia.com/...pper-architecture-in-depth
6. ^NVIDIA, "NVIDIA H100 Tensor Core GPU" product specifications. nvidia.com/...h100
7. ^NVIDIA, "NVIDIA Ada GPU Architecture" whitepaper, v2.02. images.nvidia.com/...nvidia-ada-gpu-architecture.pdf
8. ^NVIDIA Newsroom, "NVIDIA Blackwell Platform Arrives to Power a New Era of Computing," 18 March 2024. nvidianews.nvidia.com/...er-a-new-era-of-computing
9. ^NVIDIA, "NVIDIA Blackwell Architecture." nvidia.com/...blackwell-architecture
10. ^NVIDIA, "NVIDIA GB200 NVL72" product specifications. nvidia.com/...gb200-nvl72
11. ^NVIDIA, "NVIDIA DGX B200" product specifications. nvidia.com/...dgx-b200
12. ^NVIDIA Technical Blog, "Inside NVIDIA Blackwell Ultra: The Chip Powering the AI Factory Era." developer.nvidia.com/...owering-the-ai-factory-era
13. ^NVIDIA Technical Blog, "Introducing NVFP4 for Efficient and Accurate Low-Precision Inference." developer.nvidia.com/...te-low-precision-inference
14. ^NVIDIA, "Tensor Cores: Unprecedented Acceleration for Generative AI." nvidia.com/...tensor-cores
15. ^NVIDIA, "NVIDIA Vera Rubin NVL72" product specifications. nvidia.com/...vera-rubin-nvl72
16. ^NVIDIA Blog, "NVIDIA Vera Rubin Driving Performance Per Watt, Lowest Token Cost for Partners Worldwide," 21 July 2026. blogs.nvidia.com/...vera-rubin
17. ^NVIDIA, "CUDA C++ Programming Guide," section on Warp Matrix Functions (archived version 12.4.0). docs.nvidia.com/...cuda-c-programming-guide
18. ^NVIDIA, "Parallel Thread Execution ISA" (PTX), sections on warp-level, warpgroup-level, and fifth-generation Tensor Core instructions. docs.nvidia.com/...parallel-thread-execution
19. ^NVIDIA Technical Blog, "Programming Tensor Cores in CUDA 9." developer.nvidia.com/...amming-tensor-cores-cuda-9
20. ^NVIDIA, "cuBLAS Documentation," math modes and `cublasMath_t`. docs.nvidia.com/...cublas
21. ^NVIDIA, "Matrix Multiplication Background User's Guide," Deep Learning Performance documentation. docs.nvidia.com/...rformance-matrix-multiplication
22. ^NVIDIA, CUTLASS repository README. github.com/...cutlass
23. ^NVIDIA, Transformer Engine repository README. github.com/...TransformerEngine
24. ^Triton documentation, `triton.language.dot`. triton-lang.org/...triton.language.dot
25. ^PyTorch documentation, "CUDA semantics," TensorFloat-32 section. docs.pytorch.org/...cuda
26. ^P. Micikevicius et al., "Mixed Precision Training," arXiv:1710.03740, October 2017 (ICLR 2018). arxiv.org/...1710.03740
27. ^P. Micikevicius et al., "FP8 Formats for Deep Learning," arXiv:2209.05433, September 2022. arxiv.org/...2209.05433
28. ^S. Markidis, S. W. D. Chien, E. Laure, I. B. Peng, J. S. Vetter, "NVIDIA Tensor Core Programmability, Performance & Precision," arXiv:1803.04014, 2018. arxiv.org/...1803.04014
29. ^AMD, "Introducing AMD CDNA 4 Architecture" whitepaper, 2025. amd.com/...amd-cdna-4-architecture-whitepaper.pdf
30. ^AMD ROCm Blogs, "AMD Matrix Cores." rocm.blogs.amd.com/...README
31. ^Google Cloud, "TPU architecture," system architecture documentation. docs.cloud.google.com/...system-architecture-tpu-vm
32. ^NVIDIA, "NVIDIA HGX Platform" specifications. nvidia.com/...hgx