PaperlayerView exact revision on arXiv

Source-linked AI summary

NVIDIA Tensor Core Programmability, Performance & Precision

Stefano Markidis, Steven Wei Der Chien, Erwin Laure, Ivy Bo Peng, Jeffrey S. Vetter

arXiv:1803.04014v1cs.DCcs.PF

TL;DR

The paper addresses how to program NVIDIA Tensor Cores for HPC while assessing their performance and mixed-precision precision loss. It surveys interfaces, benchmarks matrix multiplication, and evaluates an error-reduction technique. Tensor Cores achieve substantial matrix-multiplication performance, while precision loss can be reduced through additional computation.

  • Problem

    The paper investigates Tensor Core programmability, performance, and precision loss because mixed-precision matrix multiplication may affect HPC applications.

  • Method

    The authors survey programming interfaces, benchmark large and batched matrix multiplication against CUDA-core operations, and measure precision loss using varying matrix sizes.

  • Results

    Tensor Cores achieved 83 Tflops/s with cuBLAS GEMM and 2.5×-12× performance increases for batched small matrix multiplications.

  • Takeaways & Limitations

    HPC applications using matrix multiplications can benefit from Tensor Cores, but the benefit involves decreased precision and potentially increased computation.

Abstract

from arXiv · show

The NVIDIA Volta GPU microarchitecture introduces a specialized unit, called "Tensor Core" that performs one matrix-multiply-and-accumulate on 4x4 matrices per clock cycle. The NVIDIA Tesla V100 accelerator, featuring the Volta microarchitecture, provides 640 Tensor Cores with a theoretical peak performance of 125 Tflops/s in mixed precision. In this paper, we investigate current approaches to program NVIDIA Tensor Cores, their performances and the precision loss due to computation in mixed precision. Currently, NVIDIA provides three different ways of programming matrix-multiply-and-accumulate on Tensor Cores: the CUDA Warp Matrix Multiply Accumulate (WMMA) API, CUTLASS, a templated library based on WMMA, and cuBLAS GEMM. After experimenting with different approaches, we found that NVIDIA Tensor Cores can deliver up to 83 Tflops/s in mixed precision on a Tesla V100 GPU, seven and three times the performance in single and half precision respectively. A WMMA implementation of batched GEMM reaches a performance of 4 Tflops/s. While precision loss due to matrix multiplication with half precision input might be critical in many HPC applications, it can be considerably reduced at the cost of increased computation. Our results indicate that HPC applications using matrix multiplications can strongly benefit from using of NVIDIA Tensor Cores.

I. INTRODUCTION

The paper examines how NVIDIA Tensor Cores can be programmed for HPC, their performance benefits, and precision loss from mixed-precision computation. It surveys programming interfaces, benchmarks matrix multiplication, quantifies precision loss, and proposes a mitigation technique.

  • Motivation: Specialized Tensor Cores target dense matrix computation in AI and deep-learning applications.The paper situates Tensor Cores within hardware developments driven by tensor workloads.
  • Tensor Core Context: An NVIDIA Tensor Core performs one 4×4 matrix-multiply-and-accumulate per GPU clock cycle using half-precision inputs and single-precision accumulation.This mixed-precision mode combines half-precision multiplication with single-precision accumulation.
  • Tensor Core Context: The Tesla V100 provides 640 Tensor Cores with a theoretical peak of 125 Tflops/s in mixed precision.The paper also notes that eight V100 GPUs could theoretically reach one Pflops/s in mixed precision.
  • Objectives: The study evaluates Tensor Core programmability, performance, and precision loss in HPC applications.Its scope includes programming interfaces, large and batched matrix multiplication, and mixed-precision accuracy.
  • Contributions: The authors survey programming interfaces, compare matrix multiplication with CUDA cores, and quantify precision loss across matrix sizes.The comparison is intended to measure performance boost, while the precision study varies matrix size.
  • Contributions: A proposed technique decreases precision loss in Tensor Core matrix multiplication at the cost of increased computation.The paper frames this technique as one of its main contributions.

II. RELATED WORK

Related work describes the rise of specialized tensor hardware for AI and the Volta architecture underlying NVIDIA Tensor Cores. It also highlights mixed precision as a throughput-oriented design choice with accuracy implications.

  • AI Hardware and Frameworks: AI and deep-learning frameworks increasingly represent tensor operations as computational-graph nodes.TensorFlow is given as a prominent example, alongside Caffe, Torch, and CNTK.
  • Specialized Hardware: Specialized alternatives include FPGAs, VPUs, TPUs, neural-network processors, and neuromorphic hardware for high-performance inference or tensor workloads.The related-work discussion contrasts several hardware approaches and their intended workloads.
  • Precision Trade-offs: Narrower numerical representations reduce power and energy footprints but also reduce computational accuracy.The passage connects lower-bit representations with reduced power consumption and precision.
  • Volta Architecture: The Volta GV100 GPU contains six GPCs and 16 GB HBM2 memory.The architecture description identifies these as major GV100 features.
  • Volta Architecture: The Tesla V100 uses 80 SMs, totaling 2,560 FP64 cores, 5,120 FP32 cores, and 640 Tensor Cores.Each SM contains Tensor Cores alongside FP64, FP32, INT32, and special-function resources.
  • Tensor Core Arithmetic: Volta Tensor Cores perform 64 half-input FMA operations per cycle, with half- or single-precision outputs.FMA uses one rounding operation rather than two, which produces a more accurate output.
  • Tensor Core Arithmetic: Tensor Cores provide 125 Tflops/s theoretical maximum performance, versus 31.4 Tflops/s in half precision and 15.7 Tflops/s in single precision.The same passage reports 7.8 Tflops/s for double precision.
  • Precision Trade-offs: Half-precision data requires half the memory bandwidth and footprint of single-precision data, enabling faster transfer.The passage identifies this reduction as a contributor to high throughput.

IV. PROGRAMMING NVIDIA TENSOR CORES

The paper presents Tensor Core programming through GEMM interfaces at different abstraction levels, beginning with direct WMMA access and its warp-based execution model.

  • GEMM Abstraction: Tensor Cores fundamentally perform GEMM, C = αAB + βC, by multiplying two matrices and accumulating into a third.The paper uses GEMM to illustrate Tensor Core programmability across interfaces.
  • WMMA API: CUDA 9 WMMA is the lowest-level direct interface for programming NVIDIA Tensor Cores.The paper notes that WMMA was a preview feature without backward-compatibility guarantees.
  • WMMA API: CUDA 9 WMMA computes 16×16 matrix multiply-and-accumulate operations with one CUDA warp of 32 threads.Although hardware performs 4×4 multiplications, WMMA exposes larger operations through overdecomposition.
  • WMMA API: A WMMA kernel declares input and accumulator fragments, initializes the accumulator, loads matrices, performs multiplication, and stores the result.The workflow uses fragment operations for register-resident inputs and accumulation.
  • WMMA API: WMMA requires one-dimensional tensor storage to be interpreted as either row-major or column-major.The layout must be declared when handling the underlying arrays.

A. Matrix Multiplication

Because direct WMMA supports fixed-size operations, the paper describes tiling and higher-level libraries for arbitrary-size matrix multiplication. CUTLASS and cuBLAS provide distinct library-based approaches.

  • Arbitrary-Size GEMM: Direct CUDA 9 WMMA performs fixed-size GEMM, while arbitrary-size matrix multiplication requires additional methods.The paper introduces tiling, CUTLASS, and cuBLAS as alternatives.
  • Tiled WMMA: Tiled WMMA partitions the result into 16×16 tiles and assigns a warp to compute each tile.Each output tile is formed by summing products of corresponding A and B tiles.
  • CUTLASS: CUTLASS is a CUDA C++ header-only templated library supporting GEMM in multiple precisions, WMMA-based wgemm, tiling, and software pipelining.Its software pipelining hides GPU memory latency.
  • cuBLAS: cuBLAS provides Tensor Core GEMM routines when its math mode is set to CUBLAS_tensorOp_MATH.The paper identifies cublasGemmEx() and cublasSgemm() as callable routines for this mode.

B. Batched Matrix Multiplications

Many HPC applications solve numerous small matrix multiplications in parallel, but standard batched GEMM support is not optimized for Tensor Cores.

  • B. Batched Matrix Multiplications: Many HPC applications rely on several small-size matrix multiplications executed in parallel.Examples include Nek5000 spectral-element CFD and Fast Multipole Method-accelerated FFT.
  • B. Batched Matrix Multiplications: BLAS GEMM routines are optimized for large matrix multiplications and perform poorly on small-size matrix multiplications.LIBXSMM and Intel MKL provide high-performance alternatives for small matrices.
  • B. Batched Matrix Multiplications: NVIDIA cuBLAS provides batched single-precision GEMM through cublasSgemmBatched(), but batched GEMM was not supported by Tensor Cores at the time.The paper therefore implements a simple batched GEMM to evaluate potential Tensor Core benefits.

V. PRECISION LOSS

Tensor Core mixed precision can introduce substantial rounding error because multiplication uses half-precision inputs, motivating residual-based precision refinement at additional computational and memory cost.

  • V. PRECISION LOSS: Half-precision inputs can cause large rounding errors that affect simulation accuracy, despite single-precision accumulation.Tensor Cores multiply half-precision entries and add them to a single-precision accumulator.
  • V. PRECISION LOSS: Traditional HPC applications are generally more sensitive to narrow-precision rounding errors than deep neural network training.The paper therefore characterizes mixed-precision effects in GEMM.
  • V. PRECISION LOSS: Half precision represents values over a limited range, with numbers above 65,504 mapped to infinity and overly small values mapped to zero.Its machine epsilon is 2^-10.
  • V. PRECISION LOSS: Half precision loses fractional accuracy as value ranges grow, providing only 1,024 representable values per power-of-two interval.Between 32,768 and 65,536, the accuracy is only ±32.
  • V. PRECISION LOSS: A half-precision residual matrix R stores the difference between a matrix before and after single-to-half rounding.For matrix A, the paper denotes the pre- and post-rounding matrices as Asingle and Ahalf.
  • V. PRECISION LOSS: Precision refinement decomposes AsingleBhalf into a residual product and a half-precision product, requiring one additional Tensor Core matrix multiplication.The method also requires memory to store RA.
  • V. PRECISION LOSS: Applying refinement to both matrices reduces precision loss using four Tensor Core matrix multiplications and additional memory for RA and RB.Because Bhalf is rounded from Bsingle, refining only A eliminates precision loss only partially.
  • V. PRECISION LOSS: The refinement method assumes that an original 32-bit value can be represented by two 16-bit numbers, subject to distribution error.Residual values record the portion lost during conversion.

VI. EXPERIMENTAL SET-UP

The experiments evaluate Tensor Core GEMM performance, precision loss, and refinement cost on a Tesla V100 using several CUDA and Tensor Core implementations.

  • VI. EXPERIMENTAL SET-UP: The test platform is a Tesla V100 accelerator connected to an Intel E5-2690v3 Haswell host, using CUDA 9.0 and compute capability 7.0.The experiments use CentOS Linux 7.4.1708 and specified CUDA compilation flags.
  • VI. EXPERIMENTAL SET-UP: Performance is measured for GEMM C = αAB + βC with α = 1.0 and β = 1.0, after single-precision inputs are rounded to half precision for Tensor Cores.Rounding time is excluded from the overall execution time.
  • VI. EXPERIMENTAL SET-UP: Square matrices of size N are evaluated using Tflops/s as the primary performance metric and O(N^3) operations for operation counts.CUDA events measure GPU kernel execution time.
  • VI. EXPERIMENTAL SET-UP: Results aggregate 5 to 100 tests using harmonic means of flops/s, while execution-time results use arithmetic means.Error bars are omitted when the error is below 1%.
  • VI. EXPERIMENTAL SET-UP: The study compares naive CUDA 9 WMMA, optimized CUTLASS configurations, and cuBLAS implementations.CUTLASS uses shared memory and software pipelining, with tiling configurations selected separately for each matrix size.
  • VI. EXPERIMENTAL SET-UP: For batched GEMM, the implementation uses square 16×16 matrices and 512 threads per block, with 16 matrix multiplications per block.Each 16×16 multiplication is assigned to one 32-thread warp.
  • VI. EXPERIMENTAL SET-UP: Precision loss is measured with the maximum absolute entry of e = Chalf − Csingle while varying square matrix size N.Random single-precision inputs from [-1,1] are converted to half precision before multiplication.
  • VI. EXPERIMENTAL SET-UP: The experiments implement one- and two-residual precision refinement and assess their cost using four pipelined GEMM operations.The implementation uses four cuBLAS calls, while optimized versions remain possible.

VII. RESULTS

The experiments compare GEMM performance with and without Tensor Cores across naive WMMA, CUTLASS, and cuBLAS implementations.

  • VII. RESULTS: GEMM performance is presented for CUDA cores in single and half precision and for Tensor Cores using WMMA, CUTLASS, and cuBLAS.White bars represent CUDA-core results, while grey bars represent Tensor Core results.

A. Performance

Tensor Cores substantially accelerate large and batched GEMM, but performance depends on the programming approach and memory management. cuBLAS reaches the highest reported large-GEMM performance, while WMMA batched GEMM outperforms cuBLAS batched sgemm across tested batch sizes.

  • Large GEMM: 83 Tflops/s is the maximum reported Tensor Core GEMM performance, achieved with cuBLAS GEMM at N = 8,192.This is approximately 74% of the theoretical Tensor Core performance and about 6× and 3× full single- and half-precision GEMM performance.
  • Large GEMM: CUTLASS outperforms cuBLAS GEMM on Tensor Cores for N = 16,384, likely because configurable tiling enables performance tuning.
  • Large GEMM: Naive CUDA 9 WMMA provides no performance improvement over CUDA-core sgemm and is outperformed by half-precision hgemm.Adding CUDA shared memory raises performance to about five times that of the naive implementation for N = 8,192.
  • Batched GEMM: 4 Tflops/s is the measured peak for WMMA batched GEMM using 16×16 matrix multiplies.The comparison is against cuBLAS batched sgemm running on CUDA cores.
  • Batched GEMM: 2.5× to 12× is the batched-GEMM performance range over cuBLAS batched sgemm as batch size varies.cuBLAS batched sgemm cannot run beyond 131,072 multiplications because available Tesla V100 memory is insufficient.

B. Precision and refinement

Half-precision inputs introduce matrix-multiplication error that grows with matrix size, while precision refinement reduces error by trading additional computation for accuracy.

  • Error growth: Rounding error increases quadratically with matrix size N because calculating one matrix element requires O(N^2) multiplications and summations.
  • Refinement: 30% is the error decrease observed for refinement using only RA at N = 8,192.This limited reduction is attributed to the two matrix norms being approximately equal.
  • Refinement: A factor of ten is the error reduction achieved by refinement using both RA and RB at N = 8,192.
  • Input dependence: 35× is the measured error decrease for random A and B values between ±16 at N = 4,096 when using both A and B refinement.The maximum error decreases from 8.32 without refinement to 0.24 with refinement.
  • Cost–accuracy trade-off: At N = 8,192, RA-only refinement reduces precision loss by approximately 30% at 2.25× computational cost, while RA-and-RB refinement yields approximately 10× lower error at 5× cost.The RA-and-RB refinement cost remains approximately 25% below GEMM execution without Tensor Cores, and the implementation is not optimized.

VIII. DISCUSSION AND CONCLUSIONS

The paper finds that NVIDIA Tensor Cores offer substantial GEMM acceleration through several programming approaches, while mixed-precision error remains a practical HPC concern that can be reduced with extra computation. Their applicability depends on workload structure and further accuracy evaluation.

  • Discussion and conclusions: 6× is the reported GEMM performance boost for large matrices, while 2.5×–12× is the boost for many small matrices processed in parallel.
  • Applicability: Many HPC applications can directly use Tensor Cores for large or parallel small matrix multiplications, whereas particle-based codes may require algorithm reformulation.The paper calls for deeper study of mixed-precision effects on accuracy in large HPC simulations.
  • Programmability: Three programming interfaces are evaluated: CUDA 9 WMMA, CUTLASS, and cuBLAS.WMMA directly accesses Tensor Cores, CUTLASS provides tunable tiling, and cuBLAS enables Tensor Core math mode.
  • Performance: 83 Tflops/s is the maximum measured performance, while naive WMMA provides no improvement and shared-memory use yields a 5× improvement over CUDA-core sgemm.The maximum occurs at N = 8,192; batched GEMM improves over cuBLAS batched sgemm by 2.5×–12×.
  • Precision: Mixed-precision error increases with input magnitude and matrix size because half-precision inputs are rounded and matrix entries require O(N^2) operations.
  • Precision: Precision refinement reduces loss from single-to-half conversion at the cost of increased computation.The paper identifies further refinement methods as a future possibility.
Loading 1803.04014v1…