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**. --- **AI Accelerator Architecture — Compute, Memory, and Interconnect.** Modern AI chips are purpose-built for matrix multiplication: a systolic array or tensor core computes thousands of multiply-accumulate (MAC) operations per cycle, fed by a memory hierarchy (registers → SRAM → HBM) connected through a network-on-chip (NoC) that determines whether the compute units starve or stay busy. The single metric that captures this interaction is the roofline model: peak performance (TFLOPS) vs memory bandwidth (TB/s), where the arithmetic intensity of the workload (FLOPs/byte) determines which resource limits throughput. **Tensor Cores — The Matrix Multiply Unit.** NVIDIA tensor cores perform 4$\times$4 matrix multiply-accumulate (D = A$\times$B + C) in a single clock cycle at mixed precision (FP16 inputs, FP32 accumulate). The H100 has 528 tensor cores across 132 SMs, delivering 989 TFLOPS at FP16 or 1,979 TFLOPS at FP8 — a 3$\times$ generational improvement over A100 (312 TFLOPS FP16). Programming tensor cores requires structuring data in tile-friendly layouts (16$\times$16 or 32$\times$8 fragments) via CUDA WMMA or MMA PTX instructions. Utilization typically reaches 60–80% in production training (compute-bound GEMM) but drops to 10–30% during inference decode (memory-bound, limited by KV cache reads). AMD CDNA3 Matrix Cores and Google TPU v5 MXUs provide equivalent functionality at comparable TFLOPS/W. **KV Cache and Inference Efficiency.** During autoregressive LLM inference, each generated token requires reading the full key-value cache of all prior tokens — creating a memory-bandwidth bottleneck where arithmetic intensity drops to 1–5 FLOPs/byte (far left of the roofline). A 70B-parameter model at sequence length 4096 stores 40 GB of KV cache in HBM; generating each token reads 40 GB at 3.35 TB/s = 12 ms latency per token — regardless of compute capacity. Solutions: PagedAttention (vLLM) eliminates KV cache fragmentation; multi-query attention (MQA/GQA) reduces KV size by 8$\times$; speculative decoding verifies 4–8 draft tokens per forward pass, increasing effective throughput 2–4$\times$; continuous batching (Orca) amortizes KV reads across multiple sequences in flight. **Network-on-Chip (NoC) for AI Accelerators.** The NoC connects hundreds of compute tiles (tensor cores, memory controllers, I/O ports) through a mesh, ring, or hierarchical topology — and its bisection bandwidth determines the maximum data rate for all-reduce operations during distributed training. An H100 has a 12$\times$11 crossbar connecting 132 SMs, 6 HBM3 stacks, and 18 NVLink ports. The total internal bandwidth exceeds 30 TB/s. For multi-chip training, NVLink 4.0 provides 900 GB/s chip-to-chip (18 links $\times$ 50 GB/s each) while PCIe 5.0 adds 128 GB/s for host communication. The NoC design determines whether the GPU can keep all tensor cores fed during a 2048-GPU training run where each iteration requires an all-reduce of 1–10 GB of gradients across the fabric. **Mixture of Experts (MoE) — Hardware Implications.** MoE models (GPT-4, Mixtral, Switch Transformer) activate only 2–8 experts per token out of 64–256 total, reducing compute by 10–30$\times$ relative to a dense model of equivalent capacity — but at the cost of massive memory footprint (every expert's weights must reside in HBM) and irregular memory access patterns that stress the NoC and memory controller. A Mixtral 8$\times$7B model has 46.7B total parameters but only 12.9B active per token; the challenge is that expert routing is data-dependent and unpredictable, causing load imbalance across GPU SMs and across nodes in distributed inference. Hardware solutions include expert parallelism (each GPU holds a subset of experts), capacity factors limiting expert overload, and all-to-all communication patterns that require high bisection bandwidth.