CPUs — few cores, deep pipelines

• 8-64 powerful cores
• Optimized for serial execution
• Complex branch prediction, out-of-order execution
• Great for: OS, databases, control flow

GPUs — thousands of cores, wide parallelism

• 18,944 CUDA cores (B200)
• Optimized for doing the same thing to lots of data
• Simple cores, massive throughput
• Great for: matrix math, neural networks, rendering

Each SM contains

• 128 CUDA cores (FP32) — general-purpose math
• 4 Tensor cores — matrix multiply acceleration
• 256 KB register file (64K × 32-bit)
• 228 KB shared memory / L1 cache
• Warp schedulers (select which threads run)
• Special function units (sin, cos, sqrt)

// The SM is a mini-computer
// with its own memory hierarchy

SM {
  cuda_cores:     128 FP32
  tensor_cores:   4
  registers:      256 KB
  shared_memory:  228 KB
  l1_cache:       shared with above
  warp_schedulers: 4
  max_threads:    2,048
  max_warps:      64
}

// H100: 132 SMs = 16,896 CUDA cores
// B200: 148 SMs = 18,944 CUDA cores

The execution hierarchy

• Thread: A single execution unit running your kernel function
• Warp: 32 threads executing in lockstep (SIMT — Single Instruction, Multiple Threads)
• Thread Block: A group of threads (up to 1,024) that share an SM’s shared memory and can synchronize
• Grid: All thread blocks in a kernel launch — the full computation

// CUDA programming model
// CPU (host) launches kernel on GPU (device)

// Define a kernel
__global__ void add(float* a, float* b,
                    float* c, int n) {
    int i = blockIdx.x * blockDim.x
            + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}

// Launch: 256 threads per block
// ceil(n/256) blocks in the grid
add<<<(n+255)/256, 256>>>(a, b, c, n);

// GPU distributes blocks to SMs
// Each SM runs multiple blocks
// Each block = warps of 32 threads

• Explicitly managed by the programmer (not automatic like L2)
• 228 KB per SM on H100 (up from 164 KB on A100)
• Shared across all threads in a block
• ~100× faster than HBM
• Threads within a block can cooperate via shared memory

// Without FlashAttention:
// Attention matrix lives in HBM
Q, K, V = load_from_hbm()
S = Q @ K.T       // [seq, seq] in HBM!
P = softmax(S)    // [seq, seq] in HBM!
O = P @ V         // read from HBM again
// 128K context → 128K² = 16B elements
// = 32 GB just for attention!

// With FlashAttention:
// Tile Q, K, V into SRAM-sized chunks
for tile_q in tiles(Q):
  for tile_k, tile_v in tiles(K, V):
    // Compute attention in SRAM
    // Never materialize full matrix
    update_output(tile_q, tile_k, tile_v)
// 100x less HBM traffic

• Generating each token requires loading the entire model from HBM
• A 70B model in FP16 = 140 GB of weights
• At H100’s 3.35 TB/s, loading 140 GB takes ~42 ms
• This sets a hard floor on latency — no amount of compute can help
• The GPU’s tensor cores sit mostly idle during decode, waiting for data

Strategies to fight the wall

• Batching: Amortize weight loads across multiple sequences
• Quantization: FP16 → FP8 → FP4 reduces bytes loaded
• KV cache optimization: MLA, GQA reduce cache memory
• Speculative decoding: Generate multiple tokens per weight load
• More HBM bandwidth: B200 at 8 TB/s, Rubin at 22 TB/s

CUDA Cores

• General-purpose floating-point units
• One scalar operation per core per cycle: a * b + c
• 128 per SM on Hopper/Blackwell
• Good for: element-wise ops, non-matrix workloads

Tensor Cores

• Specialized matrix multiply-accumulate (MMA) units
• Process entire small matrices in a single instruction
• 4 per SM (592 total on B200)
• Good for: matrix multiplication, convolution, attention — the core of neural networks

• Introduced in Hopper H100, enhanced in Blackwell
• Dynamically manages precision per-layer during training
• Monitors statistics (max values) per tensor
• Automatically picks optimal scaling factors for FP8
• No manual tuning needed — just enable it

// The Transformer Engine workflow

// 1. Forward pass in FP8
for layer in model:
    // Monitor activation statistics
    stats = observe(activations)

    // Compute scaling factor
    scale = compute_scale(stats.max_val)

    // Scale to FP8 range
    x_fp8 = quantize(x, scale)

    // Matrix multiply in FP8
    // Accumulate in FP32
    out = tensor_core_mma(x_fp8, w_fp8)

    // Dequantize output
    out = dequantize(out, scale)

// All automatic — no manual tuning

• Arithmetic Intensity = FLOPs per byte of memory accessed
• If intensity is low: memory-bound (GPU starved for data)
• If intensity is high: compute-bound (GPU at full throttle)

H100 Ridge Point

• 990 TFLOPS ÷ 3.35 TB/s = ~295 FLOP/byte
• Operations below 295 FLOP/byte are memory-bound
• LLM decode: ~0.5-1 FLOP/byte — deeply memory-bound
• LLM prefill (batch): 100-1000+ FLOP/byte — compute-bound

• B200 is a dual-die package — two 104B-transistor dies
• Connected via 10 TB/s NV-HBI chip-to-chip link
• Each die: ~800 mm² (near reticle limit)
• Appears as a single GPU to software
• This is how NVIDIA gets 208B transistors — can’t fit on one die

Manufacturing cost breakdown

Rubin GPU (R200) — H2 2026

• 336 billion transistors (2 dies, TSMC 3nm)
• 288 GB HBM4 per GPU
• 22 TB/s memory bandwidth (2.8× Blackwell)
• 50 PFLOPS FP4 inference per GPU
• 5× inference, 3.5× training vs Blackwell
• 10× reduction in inference token cost

Vera CPU

• 227 billion transistors
• Custom Arm “Olympus” cores, 88 cores / 176 threads
• CPU↔GPU link: 1.8 TB/s (2× Grace’s 900 GB/s)

MI300X (current, CDNA 3)

• 192 GB HBM3 (2.4× H100’s 80 GB)
• 5.3 TB/s memory bandwidth (1.6× H100)
• ~750W TDP
• Price: $10,000-$15,000 (vs $25,000-$40,000 for H100)

MI350X (2025, CDNA 4)

• 288 GB HBM3E, 8 TB/s bandwidth
• TSMC 3nm, native FP4/FP6
• Up to 18-20 PFLOPS FP4
• 2.7× tokens/second vs MI325X

Progress (as of 2026)

• ROCm 7.1+ with PyTorch 2.9 support
• vLLM: ROCm is now a “first-class platform”
• 93% of vLLM AMD test groups passing
• Dedicated CI pipeline went live Dec 2025
• Pre-built Docker images (no more build-from-source)
• FP8/FP4 quantization, DeepSeek V3 support

• SMs are mini-computers — 128 CUDA cores + 4 tensor cores + 228 KB shared memory, replicated 132-148 times
• Memory bandwidth, not FLOPS, is the bottleneck for LLM inference. The roofline model shows decode is deeply memory-bound
• Tensor cores deliver 60× the throughput of CUDA cores on matrix ops. Five generations have evolved from FP16 to FP4
• FlashAttention exploits the memory hierarchy to make long-context practical

• Blackwell B200: 208B transistors, 192 GB HBM3e, 8 TB/s bandwidth, dual-die design. Costs ~$6,400 to make, sells for ~$35,000
• Vera Rubin (H2 2026): 336B transistors, HBM4, 22 TB/s. 5× inference over Blackwell
• AMD is the value play — MI300X wins on memory capacity and price, but ROCm software gap limits adoption
• Google TPUs are the real alternative — full vertical stack, no NVIDIA dependency