TF32 (Hybrid format for accelerating tensor computations)

TF32 is a compromise numerical format from NVIDIA. It takes a 10-bit mantissa from FP16 for precision and an 8-bit exponent from FP32 for a wide dynamic range. This allows automatic acceleration of neural network training without code changes, achieving performance close to half precision but without loss of stability.

FP16 (Compact representation of Floating-Point numbers)

The format is activated in deep learning software environments (PyTorch, TensorFlow) and frameworks like cuDNN. It is aimed at accelerating convolutional and fully connected layers, primarily on GPUs of Ampere (A100, RTX 3090), Hopper, and newer architectures. The greatest effect is achieved in computer vision tasks and transformers, where matrix multiplication operations are the bottleneck. In inference, the format is rarely used, with preference given to pure FP16 or INT8 for minimal latency and model size.

A standard inconvenience is related to the opacity of behavior: the cuDNN engine applies TF32 implicitly, which can slightly alter model convergence compared to reference FP32 without the developer’s knowledge. The problem of result reproducibility is especially critical during strict numerical comparison. Furthermore, in layers sensitive to the exponent range (overflow beyond ~3.4×10³⁸ is impossible), there is no speed increase, and accumulator addition operations still proceed in FP32, creating the illusion of full network conversion to reduced precision.

How TF32 works

The working principle is based on introducing a unique 19-bit tensor data type into the matrix multiplication flow. When Tensor Cores receive a command, the input operands (FP32) are converted: 10 bits are allocated to the mantissa, 8 to the exponent, and 1 sign bit. This is fundamentally different from classic FP16, where the exponent occupies only 5 bits and the range easily overflows. Compared to BFloat16, which also has an 8-bit exponent but a mantissa truncated to 7 bits, TF32 provides higher precision for intermediate products. Internal summation of these products is performed in a full 32-bit accumulator. The final result is converted back to IEEE FP32, which practically eliminates losses beyond the 10-bit mantissa. This approach, unlike the total conversion of the entire network to FP16, eliminates the need for manual gradient scaling (loss scaling) characteristic of mixed precision, as the dynamic range of FP32 is fully preserved.

Tensor (Multidimensional container for numerical data)

TF32 functionality

  1. TF32 data format. TF32 uses 19 bits: 1 sign bit, 8 exponent bits, and 10 mantissa bits. This word length equals that of FP16/BF16, but the exponent range fully matches FP32. This is the key architectural decision for preserving dynamic range without overflow risk.
  2. BF16 (Floating-point number format for machine learning)
  3. Hardware acceleration on tensor cores. Computations in TF32 are performed exclusively on the tensor cores of Ampere, Hopper, and Blackwell architectures. Regular CUDA FP32 cores do not support this format. The input operands A and B of matrix multiplication are rounded to TF32 directly during loading into the tensor core registers.
  4. CUDA (Parallel computing on the graphics processing unit)
  5. FP32 accumulation mode. Despite truncating the input data mantissa to 10 bits, the multiplication and addition operations inside the tensor core are accumulated in full 32-bit floating-point format. The result of the matrix accumulation is always issued as standard IEEE FP32, minimizing precision loss during reduction.
  6. Implicit conversion mechanism. Conversion from FP32 to TF32 does not require explicit function calls or data type changes in code. The hardware unit discards the lower 13 bits of the input FP32 number mantissa, leaving 10 bits. Rounding occurs by truncation without stochastic mechanisms.
  7. Compatibility with FP32 code. The developer can activate TF32 without modifying the model tensor structure. When enabling the flag torch.backends.cuda.matmul.allow_tf32 = True, all nn.Linear operations and matrix multiplications automatically switch to TF32 math, preserving the data storage interface as torch.float32.
  8. Multiplication throughput. Tensor cores with TF32 support provide a multiple performance increase compared to FP32 on CUDA cores. On the Ampere architecture (A100), peak TF32 performance reaches 156 TFLOPS, while standard FP32 is limited to 19.5 TFLOPS, yielding an eightfold acceleration for dense operations.
  9. TFLOPS (Computational performance meter)
  10. Energy efficiency of the format. Reducing the mantissa bit width decreases the power consumption of arithmetic units. Tensor cores require fewer transistors and cycles per operation compared to full FP32 multiplication, directly reducing TDP while preserving exponent precision.
  11. TDP (Project processor heat dissipation limit)
  12. Precision in block multiplication tasks. The inner GEMM loop uses TF32 only for reading source matrices. Since the accumulator remains FP32, the rounding error accumulates more slowly. For typical matrix sizes in convolutional networks and transformers, the final numerical error often does not exceed 0.1% relative to reference FP32.
  13. GEMM (Matrix multiplication using the row times column method)
  14. Comparison of mantissas with BF16. BF16 has 7 mantissa bits, TF32 has 10. This gives the TF32 format an advantage in representing weights close to zero and gradients. For algorithms critical to quantization noise at small values, TF32 demonstrates a lower variance of rounding error.
  15. Impact on training convergence. The 8-bit exponent is identical to FP32, so the value range from 1e-38 to 3e38 is fully preserved. This prevents gradient loss due to underflow or overflow in the early stages of training. The loss function curves when using TF32 are practically identical to FP32 for most architectures.
  16. Handling of exceptional values. Support for normalized and denormalized numbers, infinities, and NaN is implemented in TF32 at the FP32 level. The hardware correctly processes these states without requiring software exception emulation, which is critical for training pipeline stability.
  17. Result determinism. TF32 operations on tensor cores of the same microarchitecture are strictly deterministic. However, when transferring computations between GPU generations (for example, from Ampere to Hopper), the lower bits of the FP32 result may differ slightly due to changes in the microarchitecture of the accumulation units.
  18. Math library policies. cuBLAS and cuDNN enable TF32 mode by default, unless overridden by the environment variable NVIDIA_TF32_OVERRIDE = 0. Libraries use heuristics to select TF32 or FP32 depending on matrix size, giving preference to TF32 for heavy tensor contractions.
  19. Selective precision control. NVIDIA provides granular control via cublasSetMathMode with the CUBLAS_TF32_TENSOR_OP_MATH flag. It is possible to forcibly disable TF32 only for convolutions (cudnnSetConvolutionMathType with CUDNN_FMA_MATH), leaving acceleration for fully connected layers in the transformer.
  20. Verification of computational graphs. With active TF32, loss values and metrics can micro-change within ULP (Unit in the Last Place). Tools like torch.autograd.detect_anomaly do not flag these differences as errors, but for critical systems, gradient comparison via torch.allclose with a relative tolerance of 1e-4 is recommended.
  21. Optimization of memory bandwidth. Since TF32 is not a storage format, no reduction in video memory volume occurs. Weights remain in FP32. The gain is achieved solely by increasing the throughput of arithmetic units, while the HBM bus is loaded with FP32 data unchanged.
  22. HBM (3D stacked memory with silicon vias)
  23. Interaction with Automatic Mixed Precision. AMP and TF32 solve adjacent but different tasks. AMP (FP16/BF16) halves memory and traffic, while TF32 keeps memory full, accelerating only computations. Simultaneous operation is allowed: TF32 for functions not converted to AMP, and BF16 for the main graph.
  24. Support in PyTorch. PyTorch scripts control TF32 via torch.set_float32_matmul_precision. The high level enables TF32, highest enables only pure FP32. For DataLoader and logging, it is important to ensure that the allow_tf32 flag is set before the first forward pass of the model.
  25. Activation in TensorFlow and JAX. In TensorFlow, TF32 is enabled globally via tf.config.experimental.enable_tensor_float_32_execution. In JAX, control occurs through the jax_default_matmul_precision flags, where the value tensorfloat32 activates mantissa truncation for dimension sizes greater than 128.
  26. Limitations for small-dimension layers. When multiplying matrices with a small inner dimension (K < 128), the overhead of launching tensor cores can negate the benefit. The cuBLAS driver automatically switches such tasks to standard FP32 cores, ignoring the TF32 permission flag to preserve performance.
  27. TF32 in inference. Although the format was designed for training, inference also gains a latency advantage. When running large language models in TF32 on Hopper architecture GPUs, the token generation speed is significantly higher than in FP32, without degradation of the perplexity metric.
  28. Utilization in Triton Inference Server. The Triton server integrates TF32 via the ONNX Runtime or PyTorch backend parameters. The config.pbtxt configuration file allows setting the enable_tf32 optimization, making it possible to serve the model with increased throughput without refactoring the model code.
  29. Effect of TF32 for transformers. The key advantage of TF32 is revealed in Self-Attention and MLP operations, where large matrix multiplications dominate. Preserving the 8-bit exponent is critically important for Softmax stabilization, and mantissa truncation in Feed-Forward layers does not distort GELU activations.
  30. Recommendations for disabling. It is necessary to forcibly disable TF32 for numerically unstable algorithms, such as inversion of ill-conditioned matrices, high-order singular value decomposition, or calculation of the logarithm of the determinant for normalizing flows.

Comparisons

  • TF32 vs FP32. TF32 uses a 19-bit mantissa, identical to FP16, and an 8-bit exponent from FP32, creating a format with single-precision dynamic range but reduced structural precision. In A100 tensor cores, TF32 provides an eightfold performance increase compared to classic FP32 multiplication, however, result accumulation still occurs in full FP32, minimizing rounding errors.
  • TF32 vs BF16. The key difference lies in the mantissa size: BF16 has a 7-bit mantissa, whereas TF32 retains 10 bits after discarding the lower bits of FP32. This gives TF32 higher precision in representing the fractional part of a number with an identical 8-bit exponent range. BF16 requires explicit data type casting in code, while TF32 is activated at the math kernel level transparently for the developer.
  • TF32 vs FP16. Standard FP16 has a narrow dynamic range (5-bit exponent), which often requires gradient scaling to prevent overflow or signal vanishing. TF32 lacks this drawback thanks to the full 8-bit exponent inherited from FP32. In terms of computational throughput on the Ampere architecture, tensor cores with TF32 operate eight times faster than standard FP32 CUDA cores, but slower than specialized FP16 operations.
  • TF32 vs INT8. These formats are optimized for fundamentally different processing stages: TF32 is a floating-point format aimed at precise forward propagation in precision-sensitive layers, whereas INT8 is an integer format for inference with maximum energy efficiency. TF32 requires no model calibration or quantization procedures, preserving the original weight distribution, unlike INT8, where monitoring for exceeding the dynamic range is necessary.
  • TF32 vs FP64. Double precision is irreplaceable in scientific simulations requiring 53 mantissa bits to exclude the accumulation of catastrophic error in long loops, whereas TF32 with its truncated mantissa is oriented towards iterative, noise-tolerant stochastic processes of deep learning. The FP64 hardware units on A100 accelerators are physically isolated from the tensor cores using TF32 and have a throughput sixteen times lower than the performance of TensorFloat-32 matrix operations.

OS and driver support

TF32 is implemented exclusively at the level of the NVIDIA driver (starting from version 450) and the cuDNN/cuBLAS libraries included in CUDA Toolkit 11.0 and newer, therefore support does not depend on the specific operating system and works wherever the specified driver can be installed (Windows, Linux, WSL2). TF32 activation occurs via a flag in code (torch.backends.cuda.matmul.allow_tf32 = True in PyTorch or tf32 in TensorFlow), which forces the cuBLAS kernels to automatically use the truncated 19-bit mantissa format for input matrix multiplications and a 32-bit accumulator without explicit data type changes by the user.

Computational safety

The safety mechanism of TF32 is built on deterministic precision control through preserving the FP32 exponent range (8 bits) and truncating the mantissa to 10 bits when loading operands, while accumulation of the sum always proceeds in a full 32-bit format with a 23-bit mantissa, preventing the loss of significant digits during reduction. In risky scenarios (for example, gradient explosions in transformers), the developer can selectively disable TF32 for critical layers by setting the precision policy in the torch.autocast context manager, returning to standard FP32 for the backward pass, while the forward pass continues to use the accelerated mode.

Logging

Diagnostics and logging of TF32 usage are carried out through the NVIDIA Nsight Systems and Nsight Compute profilers, which mark kernels executed with accelerated arithmetic using special tags in the trace, displaying reduced tensor core utilization due to non-multiple matrix dimensions. At the framework level, the cuBLAS library exposes counters through the NVIDIA Management Library (NVML) API, allowing monitoring environments like DCGM to log tensor operation activation events, and in PyTorch, when enabling verbose mode (torch._C._set_cublas_allow_tf32(True)), the checksum of computations is verified against reference FP32, recording discrepancies to the system log for precision auditing.

Architecture limitations

TF32 is fundamentally limited only to matrix multiplication and convolution operations on tensor cores of Ampere, Ada Lovelace, and Hopper architectures, not working on GPUs of the Turing series and below due to the lack of hardware support for mantissa truncation in the input pipeline; meanwhile, standard arithmetic operations (addition, nonlinear activation functions) and operations on individual scalars are always performed in pure FP32. A software limitation is the mandatory alignment of input tensor dimensions to a boundary of 16 elements to activate the tensor execution path, as well as the inability to apply TF32 for sparse matrices requiring separate support for compression formats (structured sparsity 2:4).

History and development

The TensorFloat-32 format was first introduced by NVIDIA in May 2020 together with the Ampere microarchitecture (A100 GPU) as a compromise between FP16 speed and FP32 precision without the need to rewrite model code. With the release of Hopper (H100) in 2022, support for dynamic precision switching within a single instruction stream without pipeline stall was added, and with the appearance of Blackwell (B200) in 2024, transformer engines began to apply TF32 micro-scaling at the level of individual matrix blocks, dynamically adjusting the mantissa to the value distribution in quantized attentions, thereby expanding the applicability of the format to inference of large language models.

Latest publications