gpu tensor core programming

**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.**