Tensor Core (Hardware matrix multiplication with accumulation)

Tensor Core is a specialized computational unit in NVIDIA GPUs that performs the operation D = A × B + C on 4×4 matrices in a single clock cycle. This is a hardware accelerator optimized for mixed precision computing: it multiplies matrices in FP16 format and accumulates the result in FP32, delivering high performance without significant loss of precision.

Tensor Cores have become the architectural foundation for the deep learning industry and scientific simulations. They are heavily used in training and inference of neural networks (CNN, RNN, transformers), where computations essentially reduce to tensor multiplications. In professional visualization, the technology is applied to rendering with built-in neural network-based denoising, and in scientific computing — for matrix decompositions, molecular dynamics modeling, and accelerating cuBLAS library tasks.

The main problem with Tensor Cores is precision limitation: native use of FP16 or INT8 often leads to gradient divergence or artifacts in model output. Developers have to apply mixed precision techniques such as loss scaling (multiplying the loss function by a coefficient to prevent gradient underflow in the FP16 mantissa). There is also structural asymmetry, where data format transformation operations (matrix unpacking/packing) can take longer than the actual computation on the Tensor Core, negating the performance gain.

How Tensor Core works

The operating principle fundamentally differs from scalar CUDA cores and vector SIMD extensions of x86 processors. While a classic CUDA Core processes one multiply-add operation per cycle (scalar FMA), and AVX-512 vector extensions multiply vectors element-wise, a Tensor Core performs synchronous multiplication of 4×4×4 submatrices as an MMA (Matrix Multiply-Accumulate) primitive. The warp hardware scheduler supplies 16 elements in FP16 format to inputs A and B, while matrix C and output accumulator D are stored in FP32. In a single clock cycle, the unit computes 64 multiplications and 64 additions (128 operations total) as a single indivisible HMMA (Half Matrix Multiply and Accumulate) instruction. In comparison with analogues: Google TPU tensor processors use a systolic array for pulsating data propagation without random register access, while AMD Matrix Cores accelerators in CDNA operate with Matrix FMA format using variable-size blocks. NVIDIA implements rigidly deterministic pipelining, where data between warps is synchronized through Shared Memory, and partitioning large matrices into 4×4 tiles allows scaling performance proportionally to the number of Tensor Cores in streaming multiprocessors SM.

Tensor Core functionality

1. MMA operation software model. Tensor Core performs the operation D = A × B + C on fixed-size matrices in FP16, BF16, TF32, INT8, or INT4 formats with accumulation in FP32 or INT32, providing fused multiply-add without intermediate rounding precision losses.
2. BF16 (Floating-point number format for machine learning)TF32 (Hybrid format for accelerating tensor computations)
3. Atomic matrix instruction. Each warp-level instruction, such as mma.sync, initiates the multiplication of matrix fragments stored in the thread register file and summation with the accumulator, forming one block of the result matrix per cycle.
4. Mixed precision support. Input operands A and B can be supplied in FP16, while accumulation C is performed in FP32, enabling accelerated half-precision input without degradation of the final result thanks to the 32-bit accumulator.
5. Asynchronous copy and synchronization. Programming via cp.async asynchronous instructions loads data from global memory directly into SM shared memory, bypassing registers, after which a synchronization barrier activates the Tensor Core for fragment multiplication.
6. Sparse matrix support. The hardware compression block selects non-zero values from the weight matrix at a 2:1 ratio, feeding only significant elements to the multipliers, which doubles effective throughput without changing the logical operand dimensions.
7. INT8 and INT4 multiplicative mode. Integer tensor cores process INT8 and INT4 matrices with 32-bit accumulation, achieving double or quadruple peak performance relative to FP16 due to reduced bit width and vector packing of operands.
8. TF32 input format. TF32 input data is interpreted as a 19-bit FP number with a 10-bit FP32 mantissa and an 8-bit BF16 exponent, preserving the FP32 dynamic range with internal precision close to FP16 without modifying training code.
9. No indexing overhead. Matrix multiplication inside Tensor Core does not require explicit element addressing: fragments are loaded into registers distributed across warp threads, and the core itself performs fully connected multiplication of rows and columns of a strictly defined block.
10. Hierarchical accumulator. Multiple FMAC cascades inside the core sum partial products in a distributed register file, after which the final fragment is moved back to the general-purpose registers of the threads for further reduction or writing to memory.
11. Parallel register banks. The Tensor Core register file is divided into banks, enabling simultaneous reading of rows A and columns B without conflicts; the engineering layout ensures that each multiplier receives unique operand pairs in each sub-cycle.
12. cuBLAS and CUTLASS library encapsulation. High-level primitives like cublasGemmEx and CUTLASS templates abstract register fragmentation, tile sizes, and pipeline stages, allowing the compiler to generate optimal interleaving of asynchronous copies and math instructions.
13. Multi-Instance GPU (MIG). In A100 and H100 architectures, Tensor Core computational blocks are isolated at the SM level within a GPU instance, guaranteeing predictable latency and dividing bandwidth without mutual interference from parallel workloads.
14. Automatic type casting in PyTorch. The torch.cuda.amp context manager activates a cast-down policy: forward passes of convolutions and fully connected layers execute in FP16 on Tensor Cores, while master weights and loss functions remain in FP32 for numerical stability.
15. Dynamic power management. Upon receiving a matrix multiply instruction, the SM local voltage controller forces an elevated V-Frequency domain exclusively for Tensor Core banks, not affecting integer and special function units, which smooths peak consumption.
16. Completion signal and Wavefront. The hardware scheduler launches the warp instruction mma.sync, and the Tensor Core executes it as an atomic wavefront: all warp threads pause until the multiplication is fully completed and the result is written to accumulator registers.
17. Fourth-generation streaming multiprocessors. Each SM of the Hopper architecture contains four Tensor Core sub-cores that can simultaneously process independent MMA operations from different warps, multiplicatively increasing utilization under mixed inference workloads.
18. Internal multiplication architecture. Matrix multipliers are built as a 16×16 systolic array, where each element per cycle passes a partial sum to its neighbor, implementing pipelined MAC without intermediate data read cycles from registers.
19. Computational fault tolerance. Built-in mechanisms for end-to-end parity checking of multipliers and accumulators detect single faults directly during matrix multiplication, initiating instruction replay without stopping the GPU core.
20. Programmable mantissa switching. Tensor Core microcode dynamically changes the mantissa length from 10 bits (TF32) to 7 bits (BF16) based on a flag in the instruction descriptor, reconfiguring multipliers without changing the structure of accumulator registers.
21. Direct write to shared memory. The accumulation result can be directly offloaded from registers to SM shared memory using the stmatrix instruction, combining fragments from all warp threads into a continuous tile without additional shuffle instructions.

Comparisons

• Tensor Core vs CUDA Core. Tensor Core performs mixed-precision matrix multiplication with accumulation in one cycle, processing 4×4 matrix fragments. CUDA Core implements a scalar model: one multiply and one add per instruction. The throughput difference reaches an order of magnitude: one Tensor Core replaces dozens of CUDA Cores in convolution and linear algebra operations, radically reducing energy consumption per unit of computation.
• Tensor Core vs Systolic Array (TPU). Both units multiply matrices in a systolic rhythm, but the architectures differ. Tensor Core acts as a primitive inside a streaming multiprocessor, controlled by the CUDA instruction stream. Google TPU is implemented as a problem-oriented accelerator with an explicit systolic array, where data flows between elementary cells without register file access at each step.
• Tensor Core vs x86 AVX-512 VNNI. VNNI supplements AVX-512 with an instruction for multiplying two int8 vectors with accumulation in int32. This accelerates inference but preserves the vector paradigm with a horizontal operation over short registers. Tensor Core performs two-dimensional matrix multiplication directly in mixed FP16/FP32 precision, providing an order of magnitude greater operation density per square millimeter of die area in neural network training.
• Tensor Core vs Apple AMX. The Apple AMX matrix coprocessor is integrated into the processor cores and accelerates matrix multiplication instructions through a hidden control interface inaccessible to direct programming. NVIDIA Tensor Core, in contrast, is explicitly programmable via the CUDA API and cuBLAS library. This provides direct control over computation scheduling, asynchronous data copying, and kernel function interleaving to overlap memory latencies.
• Tensor Core vs AMD Matrix Cores (CDNA). AMD Matrix Cores accelerators in the CDNA architecture operate on matrix blocks and target FP64 and FP32 workloads for high-performance computing. NVIDIA Tensor Core has historically emphasized mixed FP16/BF16/FP8 precision with FP32 accumulation for accelerating AI training and inference. Both approaches are now converging, but the CUDA ecosystem provides more mature low-level optimizations for sparse tensors and Transformer Engine.

OS and driver support

Access to tensor core operations is provided exclusively through the proprietary CUDA stack and deep learning libraries, so operating system support is determined by NVIDIA driver compatibility: Linux features open kernel module drivers with maximum computational profile performance, Windows Server and client editions have full integration via WDDM, and virtualized guest environments gain access to tensor cores through NVIDIA vGPU technology with a mandatory license.

Computation isolation and integrity

Matrix operation security is based on hardware context separation through MIG, where each GPU instance is allocated an isolated set of tensor cores with its own caches and memory controllers, while PCIe traffic encryption and confidential computing via secure enclaves prevent data interception on the path from CPU to tensor core registers, with hardware thread scheduling activation eliminating mutual influence of tenant tasks.

Operation profiling and diagnostics

Tensor core event registration is performed by Nsight Compute and CUPTI tools, which use hardware counters to capture matrix multiplier utilization, shared memory access intensity, and granularity down to individual HMMA/IMMA instructions without significant performance impact, while system monitoring is supplemented by the NVML API collecting telemetry on temperature, clock frequencies, and computational errors from tensor clusters in real time.

Limitations

Tensor core usage imposes strict requirements on matrix dimensions and alignment, dependent on the architecture, with automatic adjustment of non-multiple dimensions via padding in cuBLAS and cuDNN, while computational precision is limited to FP16/BF16/TF32/FP8/INT8/FP4 formats, and maximum theoretical performance is achieved only when all four tensor cores within a streaming multiprocessor are fully utilized and the SRAM buffer memory is fully loaded.

Architecture evolution and generations

The development history began with Volta, where cores performing fused matrix multiply with accumulation in a single cycle first appeared, then Turing added INT8 and INT4 support, Ampere introduced TF32 and structural sparsity with a twofold performance increase, Hopper implemented FP8 with dynamic scaling and an asynchronous TMA engine, and Blackwell introduced fifth-generation tensor core microservices with FP4 precision and high-reliability error correction.