tensor core programming

**Tensor Core Programming** is **the specialized technique for utilizing dedicated matrix multiplication hardware units on NVIDIA GPUs (Volta, Turing, Ampere, Hopper) that perform 4×4×4 or 16×16×16 matrix operations in a single instruction — achieving 8-20× higher throughput than CUDA cores for mixed-precision matrix multiplication (FP16/BF16 inputs, FP32 accumulation) and enabling training and inference of large neural networks at unprecedented speeds**. **Tensor Core Architecture:** - **Compute Capability**: Volta (7.0) introduced Tensor Cores with 125 TFLOPS FP16; Ampere (8.0) added BF16, TF32, INT8, INT4 support with 312 TFLOPS; Hopper (9.0) provides FP8 support and 1000+ TFLOPS with sparsity; each SM contains 4 Tensor Cores (Ampere) or 4th-gen Tensor Cores (Hopper) - **Matrix Dimensions**: Volta/Turing perform 16×16×16 matrix multiply-accumulate (D = A×B + C); Ampere/Hopper support 16×8×16, 16×8×8 for different data types; operation completes in a single instruction across the warp (32 threads cooperatively compute the result) - **Data Types**: FP16 (half precision), BF16 (bfloat16), TF32 (TensorFloat-32, 19-bit format), FP8 (Hopper), INT8, INT4, and binary; accumulation typically in FP32 for numerical stability; TF32 provides FP32 range with reduced precision, enabling drop-in acceleration for FP32 code - **Throughput**: A100 delivers 312 TFLOPS FP16 Tensor Core vs 19.5 TFLOPS FP32 CUDA Core — 16× advantage; H100 delivers 1000+ TFLOPS FP8 with sparsity vs 60 TFLOPS FP32 — 16-20× advantage; Tensor Cores dominate training and inference performance **WMMA API (Warp-Level Matrix Multiply-Accumulate):** - **Fragment Declaration**: wmma::fragment a_frag; declares a fragment (distributed across warp threads) for a 16×16 matrix of half-precision elements; each thread holds a portion of the matrix - **Load Operation**: wmma::load_matrix_sync(a_frag, a_ptr, lda); cooperatively loads matrix from global/shared memory into fragment; all 32 threads in warp participate; lda is leading dimension (stride) of the matrix in memory - **Matrix Multiply-Accumulate**: wmma::mma_sync(c_frag, a_frag, b_frag, c_frag); performs D = A×B + C using Tensor Cores; single instruction computes 16×16×16 = 4096 multiply-add operations; result distributed across warp threads in c_frag - **Store Operation**: wmma::store_matrix_sync(c_ptr, c_frag, ldc, wmma::mem_row_major); cooperatively stores result from fragment to memory; all threads participate; supports row-major and column-major layouts **Optimization Techniques:** - **Tiling for Tensor Cores**: decompose large matrix multiplication into 16×16×16 tiles; outer loops iterate over tiles; inner loop loads tiles into fragments, performs mma_sync, accumulates results; similar to traditional tiling but aligned to Tensor Core dimensions - **Shared Memory Staging**: load tiles from global memory to shared memory with coalesced access; load fragments from shared memory to registers; enables efficient data reuse and avoids repeated global memory access; shared memory acts as software-managed cache - **Double Buffering**: overlap Tensor Core computation on current tile with loading next tile from global memory; requires two sets of fragments and shared memory buffers; hides memory latency behind computation; achieves 80-90% of peak Tensor Core throughput - **Warp Specialization**: assign different warps to different tasks (loading, computing, storing); producer warps load data into shared memory; consumer warps perform Tensor Core operations; maximizes throughput by overlapping memory and compute **Mixed Precision Training:** - **FP16/BF16 Forward Pass**: activations and weights stored in FP16/BF16; Tensor Core matrix multiplications use FP16/BF16 inputs with FP32 accumulation; 2× memory bandwidth reduction and 8-16× compute speedup vs FP32 - **FP32 Master Weights**: optimizer maintains FP32 copy of weights; updates computed in FP32 for numerical stability; updated weights cast to FP16/BF16 for next iteration; prevents underflow in small gradient updates - **Loss Scaling**: multiply loss by scale factor (1024-32768) before backward pass; scales gradients to prevent underflow in FP16 range; unscale gradients before optimizer step; dynamic loss scaling adjusts scale based on gradient overflow detection - **BF16 Advantages**: bfloat16 has same exponent range as FP32 (8 bits) but reduced mantissa (7 bits vs 23 bits); eliminates loss scaling requirement; better numerical stability than FP16; preferred for training on Ampere/Hopper **PTX-Level Programming:** - **MMA Instruction**: mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 {d0,d1,d2,d3}, {a0,a1}, {b0,b1}, {c0,c1,c2,c3}; direct PTX instruction for Tensor Core operation; provides fine-grained control over data layout and operation - **Asynchronous Copy**: cp.async.cg.shared.global [smem_addr], [gmem_addr], 16; asynchronously copies data from global to shared memory; overlaps copy with Tensor Core computation; critical for achieving peak performance - **Barrier Instructions**: cp.async.wait_group and __syncthreads() coordinate asynchronous copies with computation; ensures data is ready before Tensor Core operations begin **Performance Analysis:** - **Tensor Core Utilization**: nsight compute reports tensor_precision_fu_utilization; target >80% for compute-bound kernels; low utilization indicates insufficient parallelism, memory bottlenecks, or suboptimal tiling - **Memory Bandwidth**: Tensor Cores consume data at 312 TFLOPS × 2 bytes (FP16) / 2 (multiply-add) = 312 TB/s; far exceeds HBM bandwidth (1.9 TB/s on A100); requires aggressive data reuse through tiling and shared memory - **Arithmetic Intensity**: Tensor Core GEMM achieves 100-200 FLOPs per byte; traditional CUDA Core GEMM achieves 10-20 FLOPs per byte; higher arithmetic intensity enables better utilization of memory bandwidth Tensor Core programming is **the key to unlocking the full performance of modern NVIDIA GPUs — by mastering warp-level matrix operations, mixed-precision techniques, and memory optimization patterns, developers achieve 10-20× speedups over CUDA Core implementations, making Tensor Cores the foundation of all high-performance deep learning training and inference workloads**.