Tensor Core Programming is the utilization of specialized matrix multiplication hardware on NVIDIA GPUs to achieve 10-20× higher throughput than CUDA cores — where Tensor Cores perform mixed-precision matrix operations (FP16/BF16 input, FP32 accumulation) at 312 TFLOPS on A100 and 989 TFLOPS on H100 compared to 19.5 TFLOPS and 67 TFLOPS for CUDA cores, accessed through WMMA (Warp Matrix Multiply-Accumulate) API or cuBLAS/cuDNN libraries that automatically utilize Tensor Cores, requiring specific matrix dimensions (multiples of 8 for FP16, 16 for INT8) and memory layouts (row-major or column-major with proper alignment) to achieve peak performance, enabling 5-15× faster training of large language models and 10-30× faster inference through INT8 quantization, making Tensor Core programming essential for AI workloads where matrix multiplication dominates (60-90% of compute) and proper utilization can reduce training time from weeks to days.

Tensor Core Capabilities:

• A100: 312 TFLOPS (FP16), 624 TFLOPS (BF16), 1248 TOPS (INT8), 2496 TOPS (INT4); 16× faster than CUDA cores
• H100: 989 TFLOPS (FP16), 1979 TFLOPS (BF16), 3958 TOPS (INT8); 3× faster than A100; FP8 support
• V100: 125 TFLOPS (FP16); first generation Tensor Cores; 8× faster than CUDA cores
• Supported Types: FP16, BF16, TF32, FP8 (H100), INT8, INT4; mixed precision with FP32 accumulation

<svg viewBox="0 0 760 470" xmlns="http://www.w3.org/2000/svg" font-family="-apple-system,Segoe UI,Roboto,sans-serif">
  <rect x="0" y="0" width="760" height="470" fill="#0d1117"/>
  <text x="380" y="28" fill="#e6edf3" font-size="21" font-weight="700" text-anchor="middle">Tensor Cores — Matrix Multiply Accelerators</text>
  <text x="380" y="48" fill="#8b98a5" font-size="12" text-anchor="middle">hardware units that compute D = A×B + C in one cycle for 4×4 or 16×16 matrix tiles</text>

  <!-- Core operation -->
  <rect x="30" y="65" width="700" height="155" rx="6" fill="#080d14" stroke="#233043" stroke-width="1"/>
  <text x="380" y="84" fill="#e6edf3" font-size="11" font-weight="600" text-anchor="middle">Tensor Core Operation: D = A × B + C (one warp, one cycle)</text>

  <!-- Matrix A -->
  <rect x="70" y="105" width="60" height="60" rx="3" fill="#1e3a5f" stroke="#60a5fa" stroke-width="1"/>
  <text x="100" y="125" fill="#93c5fd" font-size="9" text-anchor="middle">A</text>
  <text x="100" y="140" fill="#60a5fa" font-size="7" text-anchor="middle">16×16</text>
  <text x="100" y="155" fill="#6b7684" font-size="6.5" text-anchor="middle">FP16/BF16/FP8</text>

  <text x="145" y="138" fill="#e6edf3" font-size="14" text-anchor="middle">×</text>

  <!-- Matrix B -->
  <rect x="162" y="105" width="60" height="60" rx="3" fill="#14261f" stroke="#34d399" stroke-width="1"/>
  <text x="192" y="125" fill="#6ee7b7" font-size="9" text-anchor="middle">B</text>
  <text x="192" y="140" fill="#34d399" font-size="7" text-anchor="middle">16×16</text>
  <text x="192" y="155" fill="#6b7684" font-size="6.5" text-anchor="middle">FP16/BF16/FP8</text>

  <text x="237" y="138" fill="#e6edf3" font-size="14" text-anchor="middle">+</text>

  <!-- Matrix C -->
  <rect x="254" y="105" width="60" height="60" rx="3" fill="#1c1633" stroke="#a78bfa" stroke-width="1"/>
  <text x="284" y="125" fill="#c4b5fd" font-size="9" text-anchor="middle">C</text>
  <text x="284" y="140" fill="#a78bfa" font-size="7" text-anchor="middle">16×16</text>
  <text x="284" y="155" fill="#6b7684" font-size="6.5" text-anchor="middle">FP32 (accum)</text>

  <text x="332" y="138" fill="#e6edf3" font-size="14" text-anchor="middle">=</text>

  <!-- Matrix D -->
  <rect x="349" y="105" width="60" height="60" rx="3" fill="#2a1a0a" stroke="#f59e0b" stroke-width="1.2"/>
  <text x="379" y="125" fill="#fbbf24" font-size="9" text-anchor="middle">D</text>
  <text x="379" y="140" fill="#f59e0b" font-size="7" text-anchor="middle">16×16</text>
  <text x="379" y="155" fill="#6b7684" font-size="6.5" text-anchor="middle">FP32 result</text>

  <!-- Throughput numbers -->
  <rect x="440" y="100" width="270" height="100" rx="4" fill="#0d1117" stroke="#334155" stroke-width="0.5"/>
  <text x="575" y="118" fill="#e6edf3" font-size="9" font-weight="600" text-anchor="middle">H100 Tensor Core TFLOPS</text>
  <text x="460" y="138" fill="#f87171" font-size="8">FP32 (CUDA cores):</text><text x="610" y="138" fill="#f87171" font-size="8">67 TFLOPS</text>
  <text x="460" y="156" fill="#a78bfa" font-size="8">TF32 (Tensor Core):</text><text x="610" y="156" fill="#c4b5fd" font-size="8">495 TFLOPS</text>
  <text x="460" y="174" fill="#60a5fa" font-size="8">BF16 (Tensor Core):</text><text x="610" y="174" fill="#93c5fd" font-size="8">990 TFLOPS</text>
  <text x="460" y="192" fill="#34d399" font-size="8" font-weight="600">FP8 (Tensor Core):</text><text x="610" y="192" fill="#6ee7b7" font-size="8" font-weight="600">1,979 TFLOPS</text>

  <text x="380" y="210" fill="#f59e0b" font-size="8" text-anchor="middle">tensor cores deliver 15–30× throughput vs CUDA cores for matrix ops</text>

  <!-- How to use them -->
  <rect x="30" y="230" width="345" height="148" rx="6" fill="#080d14" stroke="#233043" stroke-width="1"/>
  <text x="202" y="250" fill="#e6edf3" font-size="10" font-weight="600" text-anchor="middle">Programming Tensor Cores</text>

  <text x="50" y="274" fill="#60a5fa" font-size="8.5" font-weight="600">cuBLAS / cuDNN:</text>
  <text x="50" y="290" fill="#8b98a5" font-size="8">automatic for GEMM / conv (just use the library)</text>

  <text x="50" y="310" fill="#34d399" font-size="8.5" font-weight="600">WMMA API (CUDA):</text>
  <text x="50" y="326" fill="#8b98a5" font-size="8">wmma::load, wmma::mma_sync, wmma::store</text>
  <text x="50" y="340" fill="#6b7684" font-size="7.5">warp-level: 32 threads cooperate on one tile</text>

  <text x="50" y="360" fill="#a78bfa" font-size="8.5" font-weight="600">Triton / torch.compile:</text>
  <text x="50" y="376" fill="#8b98a5" font-size="8">tl.dot() auto-uses tensor cores when shapes align</text>

  <!-- Requirements -->
  <rect x="390" y="230" width="340" height="148" rx="6" fill="#080d14" stroke="#233043" stroke-width="1"/>
  <text x="560" y="250" fill="#e6edf3" font-size="10" font-weight="600" text-anchor="middle">Requirements for Tensor Core Use</text>

  <text x="410" y="274" fill="#f59e0b" font-size="8.5" font-weight="600">Dimension alignment:</text>
  <text x="410" y="290" fill="#8b98a5" font-size="8">M, N, K must be multiples of 16 (or 8 for FP8)</text>

  <text x="410" y="310" fill="#f59e0b" font-size="8.5" font-weight="600">Data type:</text>
  <text x="410" y="326" fill="#8b98a5" font-size="8">inputs: FP16/BF16/FP8/INT8. accum: FP32/FP16</text>

  <text x="410" y="346" fill="#f59e0b" font-size="8.5" font-weight="600">Memory layout:</text>
  <text x="410" y="362" fill="#8b98a5" font-size="8">data must be in shared memory, properly padded</text>

  <text x="410" y="376" fill="#6b7684" font-size="7.5">if any condition not met → falls back to slow CUDA cores</text>

  <!-- Bottom -->
  <rect x="30" y="390" width="700" height="40" rx="6" fill="#080d14" stroke="#233043" stroke-width="1"/>
  <text x="380" y="410" fill="#8b98a5" font-size="8.5" text-anchor="middle">FlashAttention = clever SRAM tiling that keeps tensor cores fed. cuBLAS GEMM = 95% of peak tensor core utilization.</text>
  <text x="380" y="424" fill="#6b7684" font-size="7.5" text-anchor="middle">every modern AI workload is designed around tensor core tile sizes — they define the microarchitecture of ML compute</text>

  <text x="380" y="452" fill="#6b7684" font-size="11" text-anchor="middle">Tensor cores turned GPUs from graphics chips into AI supercomputers — they do 95% of LLM compute.</text>
</svg>

WMMA API:

• Fragment Types: matrix_a, matrix_b, accumulator; represent 16×16 matrix tiles; stored in registers
• Load: wmma::load_matrix_sync(); loads tile from memory to fragment; requires aligned access
• Multiply-Accumulate: wmma::mma_sync(); performs D = A × B + C; single instruction; 16×16×16 operation
• Store: wmma::store_matrix_sync(); stores result to memory; coalesced access required

Matrix Dimensions:

• Tile Sizes: 16×16×16 (FP16/BF16), 8×8×32 (INT8), 8×8×128 (INT4); fixed by hardware
• Matrix Sizes: must be multiples of tile size; pad if necessary; M×K × K×N = M×N
• Alignment: 128-byte alignment for optimal performance; use cudaMalloc for automatic alignment
• Layouts: row-major or column-major; specify in load/store; affects memory access patterns

Programming Model:

• Warp-Level: Tensor Cores operate at warp level; all 32 threads cooperate; implicit synchronization
• Fragment Distribution: matrix fragments distributed across warp; each thread holds portion
• Accumulation: accumulator fragment accumulates results; FP32 for precision; multiple MMA operations
• Synchronization: implicit in wmma operations; no explicit __syncthreads() needed within warp

Matrix Multiplication Example:

// Declare fragments
wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> a_frag;
wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> b_frag;
wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;

// Initialize accumulator
wmma::fill_fragment(c_frag, 0.0f);

// Loop over K dimension
for (int k = 0; k < K; k += 16) {
    wmma::load_matrix_sync(a_frag, A + k, K);
    wmma::load_matrix_sync(b_frag, B + k * N, N);
    wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
}

// Store result
wmma::store_matrix_sync(C, c_frag, N, wmma::mem_row_major);

Performance Optimization:

• Tile Size: use largest supported tile (16×16×16); maximizes Tensor Core utilization
• Loop Unrolling: unroll K-dimension loop; reduces overhead; 10-20% speedup
• Shared Memory: stage data in shared memory; reduces global memory accesses; 2-5× speedup
• Multiple Accumulators: use multiple accumulator fragments; increases ILP; 20-40% speedup

Mixed Precision:

• FP16 Input: half-precision input; 2× memory bandwidth vs FP32; 312 TFLOPS on A100
• FP32 Accumulation: full-precision accumulation; maintains accuracy; prevents overflow
• BF16: bfloat16 format; same exponent range as FP32; better for training; 624 TFLOPS on A100
• TF32: TensorFloat-32; automatic on A100; 156 TFLOPS; no code changes; 8× faster than FP32

cuBLAS Integration:

• Automatic: cuBLAS automatically uses Tensor Cores; no code changes; cublasGemmEx() for mixed precision
• Performance: 80-95% of peak Tensor Core performance; highly optimized; 10-20 TFLOPS on A100
• Batched: cublasGemmStridedBatchedEx() for multiple matrices; amortizes overhead; 90-95% efficiency
• Tuning: use cublasSetMathMode(CUBLAS_TENSOR_OP_MATH); enables Tensor Cores explicitly

cuDNN Integration:

• Convolution: cudnnConvolutionForward() uses Tensor Cores; 10-20× faster than CUDA cores
• RNN: cudnnRNNForward() uses Tensor Cores for matrix operations; 5-15× speedup
• Attention: cudnnMultiHeadAttnForward() optimized for Tensor Cores; 10-30× faster
• Automatic: cuDNN automatically selects Tensor Core algorithms; no code changes

INT8 Quantization:

• Throughput: 1248 TOPS on A100; 4× faster than FP16; 2496 TOPS on H100
• Accuracy: 1-2% accuracy loss typical; acceptable for inference; calibration required
• Quantization: convert FP32 weights to INT8; scale factors for each layer; TensorRT automates
• Deployment: 10-30× faster inference; 4× less memory; enables larger batch sizes

FP8 (H100):

• E4M3: 4-bit exponent, 3-bit mantissa; for forward pass; 1979 TFLOPS on H100
• E5M2: 5-bit exponent, 2-bit mantissa; for gradients; wider range; 1979 TFLOPS
• Transformer Engine: automatic FP8 training; maintains FP16 accuracy; 2× faster than FP16
• Scaling: per-tensor or per-channel scaling; maintains accuracy; automatic in frameworks

Memory Considerations:

• Bandwidth: Tensor Cores consume 2-4× more bandwidth than CUDA cores; memory-bound at small sizes
• Tiling: use shared memory tiling; reduces global memory accesses; 5-20× speedup
• Prefetching: overlap memory transfers with compute; async copy; 20-50% speedup
• Alignment: 128-byte alignment critical; misalignment causes 2-10× slowdown

Occupancy:

• Register Usage: WMMA uses 256-512 registers per warp; limits occupancy; 50-75% typical
• Shared Memory: tiling requires 32-64KB per block; limits occupancy; balance with registers
• Block Size: 128-256 threads optimal; 4-8 warps per block; maximizes Tensor Core utilization
• SM Utilization: 80-100% SM utilization achievable; proper launch configuration critical

Performance Metrics:

• TFLOPS: measure achieved TFLOPS; compare to peak (312 on A100, 989 on H100); target 50-80%
• Memory Bandwidth: measure bandwidth utilization; 80-100% for large matrices; memory-bound for small
• Occupancy: 50-75% typical; limited by register usage; acceptable for Tensor Core workloads
• Efficiency: TFLOPS / peak TFLOPS; 50-80% achievable with optimization; 80-95% with cuBLAS

Common Pitfalls:

• Wrong Dimensions: matrix dimensions not multiples of tile size; pad matrices; 10-50% overhead
• Misalignment: unaligned memory access; 2-10× slowdown; use cudaMalloc or align manually
• Wrong Layout: row-major vs column-major mismatch; incorrect results or slowdown; specify correctly
• Insufficient Occupancy: too many registers; limits active warps; reduce register usage or increase block size

Frameworks Integration:

• PyTorch: automatic Tensor Core usage with torch.cuda.amp; mixed precision training; 2-3× speedup
• TensorFlow: automatic mixed precision with tf.keras.mixed_precision; 2-3× speedup
• JAX: automatic with jax.default_matmul_precision('high'); 2-3× speedup
• TensorRT: automatic INT8 quantization; 10-30× inference speedup; calibration required

Use Cases:

• Training: large language models, vision transformers; 5-15× faster with Tensor Cores; weeks to days
• Inference: real-time inference with INT8; 10-30× faster; enables larger batch sizes
• Scientific Computing: matrix-heavy workloads; molecular dynamics, climate modeling; 10-20× speedup
• Recommendation Systems: embedding lookups and matrix operations; 5-15× speedup

Best Practices:

• Use Libraries: cuBLAS, cuDNN, TensorRT; 80-95% of peak; highly optimized; easier than custom kernels
• Mixed Precision: FP16/BF16 for compute, FP32 for accumulation; 2× speedup; maintains accuracy
• Proper Dimensions: ensure matrix dimensions are multiples of tile size; pad if necessary
• Profile: use Nsight Compute; verify Tensor Core utilization; target 50-80% of peak

Tensor Core Programming represents the key to AI performance on NVIDIA GPUs — by utilizing specialized matrix multiplication hardware through WMMA API or cuBLAS/cuDNN libraries, developers achieve 10-20× higher throughput (312 TFLOPS on A100, 989 TFLOPS on H100) compared to CUDA cores, enabling 5-15× faster training and 10-30× faster inference through INT8 quantization, making Tensor Core programming essential for AI workloads where proper utilization can reduce training time from weeks to days.');