GPU Architecture & the CUDA Execution Model

Topics We'll Cover

This tutorial introduces the hardware architecture that makes CUDA programming possible. No prior GPU knowledge required — we build from first principles.

• CPU vs GPU — Design Philosophy
• Inside an NVIDIA GPU
• The SIMT Execution Model
• Thread Hierarchy Overview
• How Hardware Maps to Software
• Compute Capability & GPU Generations
• Key Terminology Reference
• Summary & Next Steps

CPU vs GPU — Design Philosophy

CPUs and GPUs solve fundamentally different problems. Understanding why they're designed differently is the key to understanding when and how to use CUDA.

Die Layout Comparison

Most of a CPU die is dedicated to control logic and cache. Most of a GPU die is dedicated to arithmetic units.

The Analogy

Think of a CPU as a sports car — incredibly fast at getting one thing from A to B. A GPU is a bus fleet — slower per vehicle, but moves thousands of passengers simultaneously. When your problem involves moving thousands of passengers (pixels, matrix elements, neurons), the bus fleet wins.

Inside an NVIDIA GPU

An NVIDIA GPU is organised in a hierarchy: the full chip contains Graphics Processing Clusters (GPCs), each containing multiple Streaming Multiprocessors (SMs), each containing many CUDA cores.

GPU → GPC → SM → CUDA Cores

Streaming Multiprocessor (SM) — The Building Block

The SM is the fundamental compute unit. Each SM contains:

Inside a Single SM

The SIMT Execution Model

SIMT stands for Single Instruction, Multiple Threads. It's NVIDIA's execution model, closely related to SIMD (Single Instruction, Multiple Data) but with a crucial difference: each thread has its own program counter and register state.

Warps — The Fundamental Execution Unit

The GPU doesn't execute threads individually. Instead, it groups 32 consecutive threads into a warp. All 32 threads in a warp execute the same instruction at the same time.

Warp Divergence — The Performance Killer

When threads in a warp hit a conditional branch (like an if/else), some threads take one path and others take the other. The warp must execute both paths serially, disabling threads not on the active path. This is called warp divergence.

if (threadIdx.x < 16) {
    // Path A — threads 0-15 execute, threads 16-31 are IDLE
    doWorkA();
} else {
    // Path B — threads 16-31 execute, threads 0-15 are IDLE
    doWorkB();
}

Thread Hierarchy Overview

CUDA organises parallel work in a hierarchy of threads, warps, blocks, and grids. This hierarchy maps directly to the hardware.

Thread → Warp → Block → Grid

Concrete Example

Suppose you launch a kernel to process 1,000,000 elements with a block size of 256 threads:

Dimensionality

Grids and blocks can be 1D, 2D, or 3D. Use the dim3 type in CUDA to specify multi-dimensional configurations — useful for matrices (2D) or volumes (3D).

// 1D: processing a flat array
dim3 block(256);
dim3 grid((N + 255) / 256);
kernel<<>>(...);

// 2D: processing a matrix (rows × cols)
dim3 block(16, 16);           // 16×16 = 256 threads per block
dim3 grid((cols+15)/16, (rows+15)/16);
kernel<<>>(...);

// 3D: processing a volume
dim3 block(8, 8, 4);          // 8×8×4 = 256 threads per block
dim3 grid((X+7)/8, (Y+7)/8, (Z+3)/4);
kernel<<>>(...);

How Hardware Maps to Software

The CUDA programming model maps directly to hardware. Understanding this mapping is the key to writing performant code.

The Mapping

Blocks Are Scheduled onto SMs

Multiple blocks can be assigned to the same SM (if resources permit)

Occupancy

Occupancy is the ratio of active warps to the maximum number of warps an SM can support. Higher occupancy generally means better latency hiding — when one warp stalls on a memory access, the warp scheduler can switch to another active warp at zero cost.

What limits occupancy:

• Registers per thread — each SM has a finite register file. More registers per thread → fewer threads (warps) fit.
• Shared memory per block — each SM has limited shared memory. Larger shared memory use → fewer blocks fit per SM.
• Block size — if your block size doesn't divide evenly into the SM's warp capacity, some capacity is wasted.

Warp Scheduling — Zero-Cost Context Switching

Unlike CPU thread context switches (which are expensive), the GPU keeps the state of all active warps in registers simultaneously. When a warp stalls on a memory access, the warp scheduler can instantly switch to another ready warp — no save/restore overhead.

Compute Capability & GPU Generations

Each NVIDIA GPU has a compute capability version (e.g., 8.6) that determines which CUDA features it supports. The major number indicates the architecture generation; the minor number indicates incremental improvements.

Why Compute Capability Matters

• When compiling with nvcc, you target a specific compute capability with the -arch=sm_XX flag.
• Code compiled for a lower compute capability runs on higher GPUs (forward compatible), but not vice versa.
• Some features (like Tensor Cores or cooperative groups) require minimum compute capabilities.

# Target Volta and newer
nvcc -arch=sm_70 my_kernel.cu -o my_kernel

# Target Ampere (RTX 3090)
nvcc -arch=sm_86 my_kernel.cu -o my_kernel

# Target datacenter Blackwell (B200); RTX 5090 is sm_120, DGX Spark sm_121
nvcc -arch=sm_100 my_kernel.cu -o my_kernel

# Generate code for multiple architectures (fat binary)
nvcc -gencode arch=compute_70,code=sm_70 \
     -gencode arch=compute_86,code=sm_86 \
     -gencode arch=compute_100,code=sm_100 \
     my_kernel.cu -o my_kernel

Key Terminology Reference

A quick-reference glossary of the terms introduced in this tutorial.

Summary & Next Steps

What We Covered

• CPUs vs GPUs — latency-optimised vs throughput-optimised architectures with fundamentally different design philosophies.
• GPU architecture — the hierarchy from GPU chip → GPC → SM → CUDA cores, and the role of each component.
• SIMT execution — how 32 threads form a warp and execute in lockstep, and why warp divergence hurts performance.
• Thread hierarchy — the Thread → Warp → Block → Grid model, with concrete sizing examples.
• Hardware–software mapping — blocks run on SMs, the warp scheduler hides latency through zero-cost context switching, and occupancy determines utilisation.
• GPU generations — from Tesla to Blackwell, each generation's compute capability and key features.

Key Takeaways

Next Tutorial