GPU Tensor Core Programming is the technique of leveraging specialized matrix-multiply-and-accumulate hardware units in modern GPUs to achieve dramatic speedups for linear algebra operations — performing 4×4 or larger matrix operations in a single clock cycle with throughput exceeding 1 PFLOPS on high-end GPUs.

Tensor Core Architecture:

• Matrix Operation: each tensor core performs D = A × B + C where A and B are typically FP16/BF16/TF32/INT8 and C/D are FP32 — a single SM contains 4-16 tensor cores depending on GPU generation
• Throughput Progression: Volta: 125 TFLOPS (FP16); Ampere: 312 TFLOPS (FP16/BF16); Hopper: 989 TFLOPS (FP16) — tensor cores provide 8-16× throughput improvement over standard CUDA cores for supported operations
• Data Types Supported: FP16×FP16→FP16/FP32 (training), BF16×BF16→FP32 (training), TF32×TF32→FP32 (easy migration from FP32), INT8×INT8→INT32 (inference), FP8×FP8→FP16/FP32 (Hopper inference/training)
• Warp-Level Operation: tensor core operations execute at warp granularity — all 32 threads cooperatively provide matrix fragments and receive results

Programming Interfaces:

• WMMA API (Warp Matrix Multiply-Accumulate): C++ API with fragment types for A, B, C matrices — load_matrix_sync, mma_sync, and store_matrix_sync operations manage fragment data; supports 16×16×16 and other tile sizes
• MMA PTX Instructions: lower-level PTX assembly providing finer control over tensor core operations — mma.sync.aligned instruction specifies exact matrix dimensions and data types; used by library developers for maximum performance
• cuBLAS/cuDNN: high-level libraries automatically use tensor cores when input dimensions and data types are compatible — cuBLAS GEMM with FP16 inputs automatically dispatches to tensor cores; easiest adoption path
• Cutlass: NVIDIA template library for custom GEMM implementations using tensor cores — provides building blocks (tile iterators, warp-level MMA, epilogue fusion) for researchers needing custom matrix operation variants

Optimization Techniques:

• Tile Size Selection: matrix dimensions should be multiples of tensor core tile size (16 for WMMA) — padding to multiples achieves full tensor core utilization; odd dimensions waste partial tiles
• Memory Layout: column-major or row-major layout must match the fragment loading pattern — mismatched layout requires transpose operations that reduce effective throughput
• Epilogue Fusion: combining matrix multiply with subsequent element-wise operations (bias add, activation, scaling) in the same kernel avoids writing/reading intermediate results — improves memory efficiency by 2-3×
• Occupancy vs. Tile Size: larger tiles improve computation efficiency but reduce SM occupancy — optimal tile size balances tensor core utilization with memory latency hiding

Tensor cores represent the primary performance driver for modern AI workloads — understanding how to structure computations to leverage tensor cores is essential for achieving the published TFLOPS ratings of modern GPUs, as standard CUDA cores provide only a fraction of this throughput.