GPU Architecture · Course Home
An eleven-lesson path through how an NVIDIA GPU runs your code — from the silicon, to the CUDA model, to the driver and tools that connect them, and into making it fast.
Every topic is placed on one of three maps: the physical hardware (SMs, warps, memory tiers), the logical CUDA model (kernel, grid, block, thread), and the operational layer that connects them (compiler, driver, runtime, tools). Read the lessons in order — each builds on the last — and keep the glossary open beside them. This page is home; every lesson links back here, and lessons link to each other with previous/next.
Foundations & compute
Reframes CPU intuition: a GPU trades single-task latency for aggregate throughput, hiding memory stalls by keeping thousands of threads in flight. Introduces the two maps the whole course rests on — physical hardware vs. the logical CUDA model.
The physical compute hierarchy — Streaming Multiprocessors, sub-partitions, CUDA-core lanes, and warp schedulers — and how a 32-thread warp executes one instruction in lockstep (SIMT). Explains why simple cores let a chip hold thousands of them, and what warp divergence costs.
The CUDA model — kernel, grid, block, thread — and how blocks map onto SMs and warps onto schedulers. Covers the global-index formula and why independent blocks are what let one compiled program scale across any GPU.
Memory & cooperation
Registers → shared/L1 → L2 → DRAM, and the crucial split between latency (hideable) and bandwidth (a hard ceiling). Introduces arithmetic intensity, the roofline, and tiling as the way past the bandwidth wall.
The five logical spaces — registers, local, shared, global, constant — and how each maps onto the physical tiers. Clears up the traps: local memory lives in DRAM, and shared memory's banks can serialize on conflicts.
__syncthreads(): the block barrierThe barrier that lets a block's threads cooperate through shared memory safely, and the data race it prevents. Establishes the two-sided tiling pattern every optimized kernel is built on.
Execution & operations
One kernel launch traced end to end — nvcc's embedded device code, lazy context init, how <<<>>> becomes cudaLaunchKernel, and how work crosses into the GPU via pushbuffer and doorbell. Introduces the third axis: the operational/systems layer between your code and the silicon.
How bytes move (DMA, pinned memory, copy engines), how the host learns a kernel finished (streams, events, spin-vs-block), and why CUDA errors are asynchronous and "sticky." The async host↔device contract — and why a crash can surface on an innocent later line.
How one GPU serves many streams, threads, and processes — stream concurrency, per-process time-slicing, MPS, MIG, and CUDA graphs. The core rule: real overlap depends on spare capacity, not intent.
What debug metadata travels in your binary, and precisely which data source each tool reads — nvidia-smi, Nsight Systems, Nsight Compute, compute-sanitizer, cuda-gdb. Ends the arc on troubleshooting: match the symptom to the right instrument.
Optimization
The second kind of arithmetic unit beside the CUDA cores — it multiplies whole 16×16 matrix tiles in one shot (D = A×B + C), which is why anything shaped like a matrix multiply (neural nets, attention, dense linear algebra) rides it. Opens the optimization arc, with a runnable WMMA kernel and the core lesson: a Tensor Core is only as fast as you can feed it.
The canonical nomenclature for the course, one section per lesson, colour-coded by axis (physical / logical / operational). The single source of truth for every term the lessons use.
Curated, verified external sources — NVIDIA docs and whitepapers, the PMPP textbook, microbenchmark papers, and communities — plus links to these in-workspace cards. Where the lessons draw their facts.
Every CUDA tool the course touched plus the wider ecosystem (compile, inspect, profile, debug, monitor), each with its data source, an example, and the best article. Closes with end-to-end scenario tutorials and a deprecated-tools table.
Pascal → Volta → Turing → Ampere → Hopper → Blackwell → Rubin: what changed each generation and the one capability that makes each worth paying more for, with concrete products (V100, A100, H100, B200…). Maps every datacenter part to its GeForce sibling, and shows why a pricier GPU is often the cheaper one per token.
Which of the ten lessons still applies on AMD hardware, what is merely renamed (SM→CU/WGP, warp→wavefront, PTX→nothing), and the five places the architectures genuinely diverge. Includes a full vocabulary map, a lesson-by-lesson audit, the tool equivalents, and the porting hazards that cause silent bugs.
The hard numbers behind occupancy — register-file size, per-thread and per-block budgets, and the two-budget rule. A quick-lookup card for the occupancy math introduced in Lesson 2.
A step-by-step to a working CUDA toolchain on Windows 11 via WSL2 — no Visual Studio required. The setup that gets you from zero to compiling the code examples below.
Browse the folder: code/ — each program below opens syntax-highlighted, with a link to the raw, compilable .cu source.
Queries your GPU and prints its SM count, cores, register file, cache sizes, and memory bandwidth — confirming the lesson numbers on your actual hardware — then runs a first vector-add kernel. Your hands-on entry point for Lessons 2–4.
Reverses each block's segment through shared memory, which only works because of __syncthreads(); a -DSKIP_SYNC build plus compute-sanitizer racecheck expose the race when the barrier is removed. Makes Lesson 6's barrier necessity concrete and observable.
Adds two 2D matrices two ways — a flat 1D grid and a 2D grid — producing identical results. Demonstrates that grid dimensionality is a programmer convenience laid over linear memory, not a hardware requirement (Lesson 3).
Computes one matrix multiply as scalar FP32 on the CUDA cores and as 16×16 tiles on the Tensor Cores via WMMA, checks both against a CPU reference, and times them so the speedup is a measured number. The hands-on companion to Lesson 11.
Where the journey goes next. Lessons 1–10 cover how the GPU runs your code across all three maps; Lesson 11 opens the optimization arc with the Tensor Core. Still ahead — a shared-memory-tiled matmul (WMMA fed from on-chip tiles, and a plain FP32 version measured against your 320 GB/s), profiled under Nsight Compute so every concept here becomes a number you can move. The tools card and resources point the way.