GPU Architecture and Optimization

Concept Definition: A GPU Is Not a "Big CPU"

GPUs follow a completely different design philosophy from CPUs — CPUs optimize for low latency (a single thread running very fast), while GPUs optimize for high throughput (tens of thousands of threads in parallel). Understanding GPU architecture is the foundation for understanding all inference optimization — the memory hierarchy, Roofline, kernel fusion, and weight quantization are all built on this architecture.

Three key insights for understanding GPUs:

1. A GPU is a many-core architecture — a single thread is slower than a CPU core, but tens of thousands of threads in parallel deliver far higher aggregate throughput;
2. GPUs trade latency for bandwidth — a single memory access takes 200-400 cycles (vs. ~50 on a CPU), but bandwidth reaches 3-8 TB/s;
3. GPU compute comes from Tensor Cores — not ordinary ALUs, but matrix-computation units (each clock cycle processes an 8×8×8 matrix multiply-accumulate).

LLM inference engineers must understand GPUs — otherwise they can't explain "why quantization gives 4× speedup" or "why a bigger batch can multiply throughput 30×."

1. GPU Architecture

From Die to SM to Warp to Thread

┌─────────────────────────────────────────────────────────┐
│  GPU (e.g. H100)                                        │
│  ┌──────────────────────────────────────────────┐       │
│  │  SM 0   SM 1   SM 2  ... SM 131 (132 SMs)    │       │
│  │  ┌──────────────┐                            │       │
│  │  │Inside an SM: │                            │       │
│  │  │- 4 blocks    │                            │       │
│  │  │- 64 warps    │                            │       │
│  │  │- 2048 threads│                            │       │
│  │  │- Tensor Core │                            │       │
│  │  │- 228KB SRAM  │                            │       │
│  │  └──────────────┘                            │       │
│  └──────────────────────────────────────────────┘       │
│  ┌──────────┐  ┌──────────┐                            │
│  │ L2 Cache │  │ HBM      │                            │
│  │ 50 MB    │  │ 80 GB    │                            │
│  └──────────┘  └──────────┘                            │
└─────────────────────────────────────────────────────────┘

Level	Count (H100)	Resources	Role
GPU	1	All SMs + L2 + HBM	Overall scheduling
SM	132	4 warp schedulers + Tensor Core + SRAM	Streaming multiprocessor; one SM runs one or more blocks
Block	Software-defined	Multiple warps + shared memory	Intra-block synchronization and SRAM sharing
Warp	32 threads	Single instruction, multiple data (SIMT)	The GPU's minimum execution unit; 32 threads execute in lockstep
Thread	Single execution unit	Registers	Computes one element

The Warp Is the Minimum Execution Unit

A GPU cannot schedule "a single thread" — all 32 threads of a warp execute the same instruction in lockstep. If a branch (if/else) sends the 32 threads down different paths, all threads execute all branches, and the mask picks the results afterward — this is warp divergence, a performance killer. Minimize intra-warp branching when writing kernels.

Key SM Resources

Each SM (H100) has:

• 64 warp schedulers (can issue 4 instructions per clock cycle);
• 4 Tensor Core units (each computes a 16×8×16 matmul per clock cycle);
• 228 KB shared memory (configurable into different SRAM/L1 ratios);
• 256 KB L1 cache (includes the 228 KB above);
• 65,536 32-bit registers (per SM).

2. Tensor Core: The Source of Compute

Tensor Cores (introduced in Volta) are the GPU's matrix-computation units — tens of times faster than ordinary CUDA Cores.

Compute Capability

A Tensor Core computes D = A × B + C in one clock cycle (A is m×k, B is k×n, C/D are m×n matrices). By precision:

Precision	Tensor Core Shape	FLOPS per Op
FP64	8×8×4	256 FLOPS
FP32 (TF32)	8×8×8	512 FLOPS
FP16/BF16	8×8×16	1024 FLOPS
INT8	8×8×16	2048 OPS
FP8 (H100+)	8×8×32	4096 FLOPS
FP4 (B200+)	32×32×32	32768 FLOPS

Why Tensor Cores Are Tens of Times Faster Than Ordinary CUDA Cores

An ordinary CUDA Core is a scalar ALU (1-2 FLOPS per clock). A Tensor Core is a matrix unit (1024 FLOPS per clock) — the gap comes from data reuse: a traditional CUDA Core computing C[i,j] = ΣA[i,k]·B[k,j] reads each element of A and B many times; a Tensor Core computing an 8×8×16 matrix reads A and B once — arithmetic intensity rises from ~1 to ~50, exactly enough to saturate the Roofline compute ceiling.

Computing the Compute

H100 SXM5:

• 132 SMs × 4 Tensor Cores/SM = 528 Tensor Cores;
• 528 × 1024 = 540,672 FLOPS per clock cycle (BF16);
• At 1.83 GHz → 989 TFLOPS (540,672 × 1.83 GHz ≈ 989 TF).

3. Key Specs Across GPU Generations

GPU	Architecture	SMs	Clock	BF16 Compute	INT8 Compute	HBM Capacity	HBM Bandwidth	FP8 Compute
A100 80GB	Ampere	108	1.41 GHz	312 TF	624 TOPS	80 GB	2.0 TB/s	—
H100 SXM5	Hopper	132	1.83 GHz	989 TF	1979 TOPS	80 GB	3.35 TB/s	1979 TF
H200 SXM	Hopper	132	1.83 GHz	989 TF	1979 TOPS	141 GB	4.8 TB/s	1979 TF
H20	Hopper	78	1.5 GHz	148 TF	296 TOPS	96 GB	4.0 TB/s	296 TF
L40S	Ada	142	1.11 GHz	362 TF	733 TOPS	48 GB	0.866 TB/s	733 TF
B200 SXM	Blackwell	168	2.0 GHz	2.25 PF (FP4)	—	192 GB	8.0 TB/s	4.5 PF
B100	Blackwell	128	1.95 GHz	1.8 PF (FP4)	—	192 GB	8.0 TB/s	3.6 PF

A Practical GPU Selection Cheat Sheet for LLM Inference

• Large-model online serving (70B+): H200 (4.8 TB/s bandwidth + 141GB memory)
• Mid-size model online serving (7B-30B): H100 SXM5 or H20
• Offline batch processing: A100 80GB (still good value)
• Edge / small models: L40S (48GB, good value)
• Extreme compute / FP4: B200 (Blackwell flagship)
• Upgrading from A100: A100 → H200 is the best LLM inference upgrade path (bandwidth + memory jump together)

4. The CUDA Programming Model, Revisited

The Three Levels: Grid / Block / Thread

// CUDA kernel launch
my_kernel<<<numBlocks, threadsPerBlock>>>(arg1, arg2);
//                ↑               ↑
//              grid            block
//              (3D)            (3D)

• grid: the set of all blocks, scheduled across SMs;
• block: executes within a single SM; can synchronize and share shared memory;
• thread: a single execution unit with registers.

A CUDA C++ Kernel Template

__global__ void vector_add(const float* a, const float* b, float* c, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) c[i] = a[i] + b[i];
}
// Launch:
vector_add<<<ceil(N/256.0), 256>>>(d_a, d_b, d_c, N);

Shared Memory: SRAM Within a Block

__shared__ float tile[16][16];   // shared within the block
// multiple threads cooperate to load data from HBM into shared memory
tile[ty][tx] = a[row * K + tx];
__syncthreads();                   // wait until all threads in the block finish loading
// subsequent compute reuses tile without reading HBM again

The Essence of Shared-Memory Tiling

It is manual cache management — the programmer decides which data goes into SRAM, when to use it, and when to discard it. CPU caches are managed automatically by hardware; GPU shared memory is managed by software — more freedom for the programmer, and more responsibility (poor management makes things slower).

Warp Shuffle: Ultra-Fast Communication Among 32 Threads

// pass data among 32 threads without going through shared memory
float my_val = ...;
float other_val = __shfl_xor_sync(0xffffffff, my_val, 1);  // swap with the adjacent thread

Warp shuffle is 2-3× faster than shared memory and is commonly used in reductions (sum, max).

5. The Performance-Tuning Checklist

Writing the kernel is only half the battle — tuning it is the main arena of CUDA engineers. NVIDIA's recommended checklist:

1. Occupancy

occupancy = active warps / max warps per SM. Low occupancy → idle SM compute. Ways to raise it:

• Reduce per-thread resource usage (registers, shared memory);
• Increase the number of blocks (let each SM run multiple blocks);
• Tune launch bounds (__launch_bounds__ caps register usage).

But higher occupancy isn't always better — beyond ~80%, returns diminish, since latency is already hidden.

2. Coalesced Access

32 threads in a warp accessing consecutive memory → 1 HBM read; scattered access → 32 HBM reads.

// Good (coalesced):
a[threadIdx.x + blockIdx.x * blockDim.x]   // 32 threads access consecutive addresses

// Bad (uncoalesced):
a[threadIdx.y * N + threadIdx.x]            // 32 threads access non-consecutive addresses

Coalescing is the most important optimization for memory-bound operators — uncoalesced access is 5-10× slower than coalesced.

3. Bank Conflict

Shared memory is divided into 32 banks, each serving requests serially. If 32 threads access the same bank in one clock cycle → 32 serial accesses.

// Bad (bank conflict):
tile[threadIdx.x][0]   // all 32 threads hit bank 0

// Good (conflict avoided):
tile[threadIdx.x][threadIdx.x % 32]  // 32 threads hit different banks

4. Async Memcpy

The cp.async instruction (from H100) and the TMA unit make HBM→SRAM transfers asynchronous:

// Old way: synchronous load
tile[ty][tx] = a[row * K + col];   // wait for the HBM read
__syncthreads();

// New way: async load + pipelining
cp_async(&tile[ty][tx], &a[row * K + col]);
// compute other data while waiting (e.g. the matmul of the previous tile)

Async loads + double buffering can hide 100% of HBM latency.

5. Pipelining

Pipeline the load, compute, and write-back stages:

Naive:     load1 → compute1 → write1 → load2 → compute2 → write2 → ...
Pipelined: load1 ─────────────→
                    compute1 ─────────────→
                                  write1 ─────────────→
                                load2 ─────────────→
                                            compute2 ──→

Pipelining lets the SM compute other data while waiting on HBM — LLM inference kernels (like Marlin) all use 2-4-stage pipelines.

Tuning Priority Order

1. Fix the algorithm first (find the right algorithm/parallelization strategy) — 5-10× gain;
2. Then fix coalescing/bank conflicts (access patterns) — 3-5× gain;
3. Then fix occupancy/async/pipelining (latency hiding) — 1.5-2× gain;
4. Finally use Tensor Core/FP8 (hardware instructions) — 2-4× gain. Reversing the order wastes effort — tuning a wrong algorithm is futile.

6. The GPU Library Ecosystem

Don't write everything yourself — use mature libraries where possible:

Library	Purpose	Strength
cuBLAS / cuBLASLt	Matrix multiplication (GEMM)	NVIDIA's optimal GEMM implementation
cuDNN	CNN + attention + RNN	Standard DNN operator library
CUTLASS	Matmul templates (C++)	Customizable epilogue fusion
CUB / Thrust	Reduction/scan/sort	Standard parallel-algorithm library
NCCL	Multi-GPU communication	Optimal all-reduce/all-gather
TensorRT / TRT-LLM	Whole-graph optimization + deployment	NVIDIA's closed-source flagship
cuQuantum	Quantum simulation	Quantum-computing specific
nvJPEG / nvTIFF	Image codecs	Multimedia specific

Don't Reinvent the Wheel

99% of LLM inference doesn't need hand-written kernels — vLLM/TensorRT-LLM have already pushed matmul/attention/layernorm to the limit. Hand-written kernels are for the last 10% of extreme optimization — and they cost months of engineer time.

7. CUDA Profiling: Finding Bottlenecks with Nsight

Performance optimization is impossible without profiling:

• Nsight Systems: system-level tracing — find CPU-GPU synchronization, kernel launch overhead, and memory-transfer bottlenecks;
• Nsight Compute: single-kernel profiling — Roofline ceilings, occupancy, bank conflicts, and coalescing details;
• PyTorch Profiler: Python-level — find which ops are slow and which ops hog memory.

Typical workflow:

Nsight Systems → find the slowest kernel → analyze it alone in Nsight Compute → find the bottleneck (coalescing? occupancy?) → fix the kernel → re-measure

See Inference Benchmarking in Practice.

8. Distributed Inference and Multi-GPU

When a single GPU can't hold a large model (e.g. Llama-405B needs 800GB), multi-GPU is mandatory:

Strategy	Sharding Dimension	Communication	Best For
Tensor Parallel (TP)	Split each layer's weights into N parts	all-reduce per layer	Single machine, multi-GPU; latency-sensitive
Pipeline Parallel (PP)	Split layers into N stages	pipeline bubble	Cross-machine; throughput-heavy
Expert Parallel (EP)	Shard by expert	all-to-all	MoE models
Data Parallel (DP)	Data-parallel	all-reduce grads	Mostly training; rarely used in inference

TP is the most common for LLM inference — 8 GPUs in one machine split a 70B model into 8 shards, ~10GB per GPU. See Distributed Inference (TP/PP).

9. Trade-offs

• GPU model selection: online serving favors H200 (large bandwidth); training favors B200 (large compute); edge favors L40S (value);
• Tensor Core vs. ordinary CUDA Core: matmul must use Tensor Cores (mandatory); elementwise uses CUDA Cores;
• Precision choice: FP16/BF16 for training; INT8/FP8/INT4/FP4 for inference; prefer FP8 on H100+ (good accuracy, no calibration);
• Libraries vs. hand-written: use cuBLAS/cuDNN whenever possible; use CUTLASS or Triton for fusion; CUDA C++ only for the extreme;
• Compilation vs. manual tuning: torch.compile automates 80% of optimization; hand-tune the last 20% (find bottlenecks with Nsight).

Further Reading

• The GPU Memory Hierarchy and the Bandwidth Wall — the SRAM/L2/HBM hierarchy and bandwidth
• The Roofline Model and Compute Analysis — the unified analysis framework for compute and bandwidth
• Kernel Fusion and Custom Kernels — the Triton/CUDA programming models
• Weight-Only Quantization and Mixed Precision — using Tensor Core INT4/INT8/FP8
• Batching and Request Scheduling — enlarging batch to saturate compute and bandwidth
• Distributed Inference (TP/PP) — multi-GPU parallel strategies
• TensorRT and GPU Inference — NVIDIA's flagship inference engine
• Hardware Primer — GPU selection and configuration
• Benchmark Data & Tool Profiles — measured comparisons across GPUs

References

• NVIDIA H100 White Paper — Hopper architecture and Tensor Core details
• NVIDIA Blackwell Architecture White Paper — B200/B100 architecture and FP4
• NVIDIA CUDA C++ Programming Guide — official CUDA programming documentation
• NVIDIA CUTLASS Documentation — the GEMM template library
• NVIDIA Nsight Compute Documentation — kernel-level profiling
• Mark Harris. CUDA C++ Best Practices Guide — the authoritative CUDA tuning checklist
• Kirk & Hwu. Programming Massively Parallel Processors (4th Ed., 2022) — the authoritative CUDA textbook
• Patterson & Hennessy. Computer Organization and Design: RISC-V Edition — computer architecture fundamentals