4.1 Streaming Multiprocessors & Tensor Core Evolution

4.1 Streaming Multiprocessors & Tensor Core Evolution

Core Architecture: The Streaming Multiprocessor (SM) is the foundational processing unit of modern NVIDIA GPU microarchitectures. Unlike general-purpose CPUs that optimize single-thread instruction latency through deep branch prediction and out-of-order execution, the SM is a massively parallel, throughput-optimized compute engine designed to execute thousands of concurrent threads via the Single Instruction, Multiple Threads (SIMT) execution model.

---

1. SM Microarchitecture & Sub-Core Partitioning

Every NVIDIA GPU ASIC is composed of multiple Graphics Processing Clusters (GPCs), Texture Processing Clusters (TPCs), and Streaming Multiprocessors (SMs). In data center architectures such as NVIDIA Ampere (GA100), Hopper (GH100), and Blackwell (GB200), each SM is architected into four distinct processing partitions (also called sub-cores or processing blocks).

┌─────────────────────────────────────────────────────────────────────────────────┐
│                     STREAMING MULTIPROCESSOR (SM) TOPOLOGY                      │
│                                                                                 │
│  ┌───────────────────────────────────────────────────────────────────────────┐  │
│  │         Configurable L1 Data Cache / Shared Memory (128 KB - 256 KB)       │  │
│  └───────────────────────────────────────────────────────────────────────────┘  │
│                                                                                 │
│  ┌──────────────────┐ ┌──────────────────┐ ┌──────────────────┐ ┌─────────────┐  │
│  │   SUB-CORE 0     │ │   SUB-CORE 1     │ │   SUB-CORE 2     │ │  SUB-CORE 3 │  │
│  │ ┌──────────────┐ │ │ ┌──────────────┐ │ │ ┌──────────────┐ │ │ ┌─────────┐ │  │
│  │ │Warp Scheduler│ │ │ │Warp Scheduler│ │ │ │Warp Scheduler│ │ │ │Warp Sched│ │  │
│  │ ├──────────────┤ │ │ ├──────────────┤ │ │ ├──────────────┤ │ │ ├─────────┤ │  │
│  │ │Dispatch Unit │ │ │ │Dispatch Unit │ │ │ │Dispatch Unit │ │ │ │Dispatch │ │  │
│  │ ├──────────────┤ │ │ ├──────────────┤ │ │ ├──────────────┤ │ │ ├─────────┤ │  │
│  │ │16K Reg File  │ │ │ │16K Reg File  │ │ │ │16K Reg File  │ │ │ │16K Reg  │ │  │
│  │ ├──────────────┤ │ │ ├──────────────┤ │ │ ├──────────────┤ │ │ ├─────────┤ │  │
│  │ │FP32 / INT32  │ │ │ │FP32 / INT32  │ │ │ │FP32 / INT32  │ │ │ │FP32/INT32│ │  │
│  │ ├──────────────┤ │ │ ├──────────────┤ │ │ ├──────────────┤ │ │ ├─────────┤ │  │
│  │ │FP64 Cores    │ │ │ │FP64 Cores    │ │ │ │FP64 Cores    │ │ │ │FP64 Cores│ │  │
│  │ ├──────────────┤ │ │ ├──────────────┤ │ │ ├──────────────┤ │ │ ├─────────┤ │  │
│  │ │Tensor Core   │ │ │ │Tensor Core   │ │ │ │Tensor Core   │ │ │ │Tensor C. │ │  │
│  │ ├──────────────┤ │ │ ├──────────────┤ │ │ ├──────────────┤ │ │ ├─────────┤ │  │
│  │ │LD/ST & SFUs  │ │ │ │LD/ST & SFUs  │ │ │ │LD/ST & SFUs  │ │ │ │LD/ST/SFU │ │  │
│  │ └──────────────┘ │ │ └──────────────┘ │ │ └──────────────┘ │ │ └─────────┘ │  │
│  └──────────────────┘ └──────────────────┘ └──────────────────┘ └─────────────┘  │
│                                                                                 │
│  ┌───────────────────────────────────────────────────────────────────────────┐  │
│  │      Tensor Memory Accelerator (TMA) & Async Transaction Barrier Unit      │  │
│  └───────────────────────────────────────────────────────────────────────────┘  │
└─────────────────────────────────────────────────────────────────────────────────┘

Sub-Core Functional Components

Each of the four partitions within an SM functions as an independent instruction processing pipeline containing dedicated execution resources:

1. Warp Scheduler & Instruction Dispatch Unit:
   • One warp scheduler per partition (4 warp schedulers per SM).
   • Each scheduler manages active warps assigned to its partition, issuing instructions for one ready warp every clock cycle.
   • Modern schedulers utilize advanced instruction-level scoreboarding to track operand dependencies and data hazards without stalling the entire SM pipeline.
2. Register File Slices:
   • The total SM register file consists of 65,536 32-bit registers (256 KB SRAM), partitioned into four 16,384-register slices (64 KB per sub-core).
   • Registers are allocated dynamically to active threads at thread block launch time. Register pressure (threads requiring excessive registers) reduces warp occupancy.
3. Primary Execution Pipelines (CUDA Cores):
   • FP32 Pipelines: Dedicated single-precision 32-bit floating-point units executing standard arithmetic, fused multiply-add (FMA), and comparison instructions.
   • INT32 Pipelines: Dedicated 32-bit integer arithmetic logic units running concurrently with FP32 pipelines, allowing simultaneous address calculation and floating-point math.
   • FP64 Pipelines: Double-precision 64-bit floating-point units engineered for scientific HPC simulations, running at a 1:2 rate relative to FP32 in enterprise silicon (e.g., A100/H100).
4. Special Function Units (SFUs) & Load/Store (LD/ST) Units:
   • SFUs: Specialized hardware approximating transcendental mathematical operations (sine, cosine, square root, reciprocal, exponential, and logarithmic functions) using quadratic interpolation tables.
   • LD/ST Units: Dedicated pipelines calculating source and destination memory addresses for global, local, and shared memory load/store operations.
5. Unified L1 Data Cache / Shared Memory:
   • A high-bandwidth, ultra-low latency SRAM pool shared across all four sub-cores within the SM (e.g., 192 KB on A100, 256 KB on H100).
   • Software can dynamically configure the allocation split between hardware-managed L1 caching and software-managed __shared__ memory.
6. Tensor Memory Accelerator (TMA) & Asynchronous Transfer:
   • Introduced in Hopper (GH100) and expanded in Blackwell (GB200), the TMA is a dedicated hardware engine that handles multi-dimensional tensor data movement directly between Global Memory (HBM) and Shared Memory.
   • TMA operates completely out-of-band without consuming SM register file capacity or tying up warp execution pipelines.

---

2. Generational Evolution of NVIDIA Tensor Cores

Traditional CUDA cores operate on scalar values, computing a single arithmetic operation per thread per cycle. In contrast, Tensor Cores are specialized multi-dimensional matrix processing units that execute entire matrix tile multiply-accumulate operations in hardware across all 32 threads in a warp.

                      TENSOR CORE GENERATIONAL TIMELINE

  [2017: Volta GV100]  ──► 1st Gen: FP16/FP32 Mixed Precision MMA (125 TFLOPS FP16)
           │
  [2018: Turing TU104] ──► 2nd Gen: INT8 (130 TOPS on T4), INT4 (260 TOPS), INT1
           │
  [2020: Ampere GA100] ──► 3rd Gen: TF32, BF16, FP64 TCs, 2:4 Sparsity (312 TFLOPS FP16 dense)
           │
  [2022: Hopper GH100] ──► 4th Gen: FP8 (E4M3/E5M2), DPX, TMA (1,979 TFLOPS FP8 dense)
           │
  [2024: Blackwell B200]─► 5th Gen: FP4 (E2M1), 2nd Gen Transformer Engine (~9 PFLOPS FP8)

Comprehensive Generational Matrix

---

3. Matrix Multiply-Accumulate (MMA) Mechanics

At the core of deep learning compute—including linear layers, convolutional filtering, and transformer self-attention projections—is the Matrix Multiply-Accumulate (MMA) operation.

$\mathbf{D} = \mathbf{A} \times \mathbf{B} + \mathbf{C}$

Where:

• $\mathbf{A}$ is an $M \times K$ matrix (input activations or batch tokens).
• $\mathbf{B}$ is a $K \times N$ matrix (learnable weight parameters).
• $\mathbf{C}$ is an $M \times N$ matrix (bias vector or intermediate accumulator).
• $\mathbf{D}$ is the resulting $M \times N$ accumulated output matrix.

      Matrix A (M x K)        Matrix B (K x N)           Matrix C (M x N)
     ┌───┬───┬───┬───┐       ┌───┬───┬───┬───┐          ┌───┬───┬───┬───┐
     │a00│a01│a02│a03│       │b00│b01│b02│b03│          │c00│c01│c02│c03│
   M │a10│a11│a12│a13│   x K │b10│b11│b12│b13│    +   M │c10│c11│c12│c13│
     │a20│a21│a22│a23│       │b20│b21│b22│b23│          │c20│c21│c22│c23│
     └───┴───┴───┴───┘       │b30│b31│b32│b33│          └───┴───┴───┴───┘
             K               └───┴───┴───┴───┘                  N
                                     N
                                     ║
                                     ▼
                      Result Matrix D = A x B + C (M x N)

Scalar CUDA Cores vs. Warp-Synchronous Tensor Cores

• Scalar CUDA Core Execution: To multiply two matrices, each individual thread loads scalar values, executes a single scalar FMA instruction ($d = a \cdot b + c$), and writes to a register. A warp of 32 threads computes 32 multiply-adds (64 FLOPs) per cycle.
• Tensor Core Warp Execution: In a Tensor Core, the entire 32-thread warp collaborates as a unified execution entity. Threads pool their registers to hold tiles of matrices $\mathbf{A}$, $\mathbf{B}$, and $\mathbf{C}$ (e.g., $16 \times 16 \times 16$ tile dimensions). A single warp-level $16 \times 16 \times 16$ MMA instruction issues 4,096 multiply-accumulate operations (8,192 FLOPs), which the Tensor Core array retires over a handful of clock cycles rather than one — orders of magnitude more work per issued instruction than a scalar FMA.
• Precision Accumulation: Inputs $\mathbf{A}$ and $\mathbf{B}$ can be supplied in low precision (e.g., FP16, BF16, FP8, or INT8) to maximize throughput and minimize memory bandwidth, while accumulator $\mathbf{C}$ and output $\mathbf{D}$ maintain high precision (FP32) to prevent numerical rounding drift.

---

4. 2:4 Structured Sparsity

Introduced in the NVIDIA Ampere architecture and native to Hopper and Blackwell, 2:4 Structured Sparsity is a hardware-silicon feature that doubles matrix math throughput and reduces memory bandwidth overhead by exploiting network pruning.

The 2:4 Sparsity Rule

For every contiguous block of four horizontal values in a weight matrix, exactly two values must be zero (50% sparsity).

  Original Dense Weight Matrix Row (4 Elements):     [ 0.42 ,  0.00 , -0.85 ,  0.00 ]
                                                            │        │
                                                            ▼        ▼
  Hardware Compressed Weight Storage (2 Values):     [ 0.42 , -0.85 ]
  Hardware 2-bit Index Metadata (Indices 0 & 2):     [  00  ,   10   ]

  Dense Matrix (4x4 Weights)             Sparse Compressed Weights (2x4) + Metadata
  ┌──────┬──────┬──────┬──────┐          ┌──────┬──────┐    ┌──────┬──────┐
  │ W0,0 │  0   │ W0,2 │  0   │          │ W0,0 │ W0,2 │    │  00  │  10  │ (Indices 0, 2)
  ├──────┼──────┼──────┼──────┤  ───►    ├──────┼──────┤    ├──────┼──────┤
  │  0   │ W1,1 │  0   │ W1,3 │          │ W1,1 │ W1,3 │    │  01  │  11  │ (Indices 1, 3)
  ├──────┼──────┼──────┼──────┤          ├──────┼──────┤    ├──────┼──────┤
  │ W2,0 │  0   │  0   │ W2,3 │          │ W2,0 │ W2,3 │    │  00  │  11  │ (Indices 0, 3)
  └──────┴──────┴──────┴──────┘          └──────┴──────┘    └──────┴──────┘

How Sparse Tensor Cores Operate in Silicon

1. Pruning Phase: Weights are trained normally in dense format. A magnitude-based pruning algorithm identifies the smallest two weights in each 4-element window and zeros them out.
2. Fine-Tuning Phase: The network is fine-tuned for a small number of iterations to allow the remaining non-zero weights to compensate for the pruned values, completely recovering baseline model accuracy.
3. Memory Storage: The 50% non-zero weights are stored contiguously in memory alongside compact 2-bit metadata indices representing their original column positions within each 4-element group.
4. Sparse MMA Execution: The Sparse Tensor Core hardware reads the packed 50% weight data and uses the 2-bit metadata indices to fetch only the matching activation vectors from Shared Memory or registers. It executes the dot products at 2× the rate of dense operations, effectively halving the required arithmetic operations and doubling FLOPS throughput.