03 / Learning Path

Execution Model

Threads, warps, occupancy, and scheduler behavior.

Main Sections

Track your progress

Sub Topics

Topic 01

Thread - Smallest Unit with PC and Registers

Theory

A thread is the smallest unit of execution on the GPU. Each thread has its own program counter and private register state.

Threads are lightweight and execute the same kernel code on different data values.

Spec table

Property	Value
Private memory	Registers only
Max registers per thread	255
Has own program counter	Yes
Lives inside	A warp (group of 32)
Execution model	SIMT

Practical visual

Show 32 boxes labeled Thread 0 to Thread 31.
Inside each box show PC and Registers.
Label the full row as one warp.

Topic 02

Warp - 32 Threads and Real Scheduling Unit

Theory

GPU schedulers do not schedule individual threads; they schedule warps of 32 threads.

All active lanes in a warp execute the same instruction in lockstep under SIMT.

Spec table

Property	Value
Threads per warp	32
Execution model	SIMT
Scheduled by	Warp scheduler in SM
Instruction issue	Same instruction across lanes
Warps per SM (Hopper class)	Up to 64

Practical calculator

Warp occupancy calculator input: threads per block. Output: full warps plus wasted slots.

Threads per block	Full warps	Wasted threads
32	1	0
48	1	16
128	4	0
100	3	28

Topic 03

Thread Block (CTA) - Shared Memory and Sync Domain

Theory

A thread block (CTA) is a cooperative group scheduled on one SM for its lifetime.

Threads in the same block share shared memory and can synchronize with __syncthreads().

Spec table

Property	Value
Max threads per block	1024
Max warps per block	32
Shared memory	Shared across block
Can synchronize	Yes, within block
Placement	One SM, does not migrate

Practical table

Block size	Warps	Occupancy impact
32	1	Very low
128	4	Moderate
256	8	Good
512	16	High
1024	32	Max but less flexible

Topic 04

Grid - All Blocks in One Kernel Launch

Theory

A grid is the full set of blocks launched by one kernel call.

Blocks in a grid are independent and can be scheduled in any order across available SMs.

Spec table

Property	Value
Contains	All blocks for one kernel
Dimensions	1D, 2D, or 3D
Max blocks per dimension	65,535
Block communication	Not possible
Scheduled across	All available SMs

Practical visual

Grid
|- Block(0,0) -> SM0
|- Block(0,1) -> SM1
|- Block(1,0) -> SM0 (after reuse)
|- Block(1,1) -> SM2
`- Block(2,0) -> SM3

Order is not guaranteed; correctness must not depend on block execution order.

Topic 05

SIMT - One Instruction Across 32 Threads

Theory

SIMT issues one instruction to a warp and executes it across multiple threads with different data.

Compared with classic SIMD, GPU threads keep separate logical thread state and expose a scalar programming model.

SIMT vs SIMD

Property	SIMD (CPU)	SIMT (GPU)
Threads/lane count	4-16 lanes	32 threads
Own registers per lane/thread	No	Yes
Own program counter	No	Yes (logical)
Divergence handling	Limited	Supported with cost
Programming model	Explicit vectors	Scalar-like threads

Topic 06

Warp Divergence - Branch Cost and Predication

Theory

Divergence occurs when threads in one warp take different control-flow paths.

The warp executes each path with masking, which serializes branch paths and lowers efficiency.

Cost table

Scenario	Active threads	Efficiency
No divergence	32/32	100%
50/50 split	16 then 16	50%
1 thread on alternate path	31 then 1	~50%
All threads same branch	32/32	100%

Practical rule

BAD:
if (threadIdx.x % 2 == 0) { doA(); } else { doB(); }

BETTER:
result = condition ? A : B;  // predication-friendly form

Reduce branch entropy within each warp whenever possible.

Topic 07

Latency Hiding - Why Many Warps Matter

Theory

When one warp waits on memory, the scheduler switches to another eligible warp.

Latency hiding is achieved by maintaining enough ready warps to cover long memory delays.

Latency table

Memory type	Latency (cycles)	Warps needed to hide it
Register	1	0
Shared memory	30	1-2
L1 cache	30	1-2
L2 cache	200	6-8
VRAM (HBM)	600	18-20

Timeline visual

Cycle 1: Warp A requests VRAM (wait)
Cycle 2: Warp B executes
Cycle 3: Warp C executes
...
Cycle N: Warp A data returns and resumes

Throughput depends on scheduler having enough eligible work.

Topic 08

Occupancy - Active Warps vs Max Warps

Theory

Occupancy is active warps per SM divided by hardware maximum warps.

Higher occupancy often improves latency hiding, but peak occupancy is not always peak performance.

Limiting factors

Limiting factor	If too high	Effect
Registers per thread	Consumes register file	Fewer warps fit
Shared memory per block	Consumes shared memory	Fewer blocks fit
Threads per block	Poor packing	Lower active warp count
Balanced usage	Resource fit	Higher useful occupancy

Practical calculator

Inputs: threads per block, registers per thread, shared memory per block.
Output: occupancy percent and primary bottleneck resource.

Topic 09

Warpgroup (Hopper+) - 4 Warps for WGMMA

Theory

Hopper introduces warpgroup execution for wgmma instructions, combining 4 warps (128 threads).

This enables larger matrix operations and improved Tensor Core feeding compared with warp-only MMA patterns.

Comparison table

Unit	Threads	Instruction	Generation
Warp	32	mma	Ampere+
Warpgroup	128 (4 warps)	wgmma	Hopper+
Warpgroup benefit	Larger tile op	Higher Tensor Core feed	Hopper+

Topic 10

Thread Block Cluster (Hopper+) - Inter-Block Cooperation

Theory

Thread Block Clusters allow groups of blocks to cooperate within a GPC on Hopper-class GPUs.

Clustered blocks can use distributed shared memory and cluster-level synchronization primitives.

Updated hierarchy table

Level	Contains	Can sync?	Memory
Thread	Itself	-	Registers
Warp	32 threads	-	-
Block	Up to 1024 threads	Yes (__syncthreads)	Shared memory
Cluster	Up to 8 blocks	Yes	Distributed shared memory
Grid	All blocks	No	VRAM

Topic 11

Warp Execution States - Active, Eligible, Selected, Stalled

Theory

Warp schedulers evaluate warp readiness each cycle. State transitions determine which warps can issue instructions.

Understanding state distribution helps explain utilization and stall behavior in profilers.

States table

State	Meaning	Scheduler action
Active	Assigned resources on SM	Considered
Eligible	Operands and dependencies ready	Can be selected
Selected	Chosen this cycle	Issues instruction
Stalled	Waiting on dependency/data	Skipped for now

Practical flow

Memory request -> Stalled
Other eligible warps -> Selected
Data returns -> Eligible
Scheduler picks -> Selected

The scheduler continuously rotates through ready warps.

Topic 12

Scoreboard Stalls - Short vs Long

Theory

The scoreboard tracks register readiness after issued instructions, especially memory operations.

Long scoreboard stalls are often memory-bound symptoms tied to high-latency accesses.

Stall types table

Stall type	Source	Cycles	Mitigation
Short stall	Shared memory / L1	20-40	Increase locality and ready warps
Long stall	VRAM / HBM	400-700	Coalescing, prefetch, more warps
Execution dependency	Prior instruction result	4-8	Instruction scheduling/reordering
Sync stall	__syncthreads()	Varies	Reduce sync frequency

Practical profiling tip

Track Stall Long Scoreboard percent in Nsight Compute.
If high, prioritize access coalescing and occupancy-aware tuning.