Execution / Warp Behavior

GPU Execution Model

This page explains how warps, divergence, and scheduling behavior affect real GPU utilization. The goal is to connect abstract CUDA concepts to practical outcomes like underutilized lanes, hidden latency, and inconsistent kernel efficiency.

Back to GPU

Precision Tool

Warp Divergence Visualizer

Enter a CUDA branch condition and visualize how a warp serializes active and waiting lanes across execution passes.

KERNEL CONDITIONAL

QUICK EXAMPLES

Valid expression

EXECUTION METRICS

Efficiency Ratio

33.3%

Active Threads: 11/32

Waiting Threads: 21/32

Passes: 3

Overhead: 66.7% serialization

Warp View

Warp Divergence

Pass 1 of 2
00
01
02
03
04
05
06
07
08
09
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
Pass 1 of 2
Speed:

PASS 1

Executing the if-branch. Non-matching threads are waiting.

11 threads active, 21 threads waiting

PASS 2

Executing the else-branch. Previously active threads become idle.

21 threads active, 11 threads waiting

DIVERGENCE EXPLAINER

WHAT HAPPENED

Only 11 of 32 threads take the if-branch. 21 threads sit idle during pass 1, then 11 sit idle during pass 2. Efficiency: 33.3%.

HOW TO FIX IT

Consider restructuring data so threads in the same warp process elements of the same type - eliminating the modulo branch. Warp divergence is unavoidable sometimes, but minimizing it is key to high GPU occupancy and throughput.

DETAILED METRICS

Efficiency Ratio: 33.3%

Active Threads: 11 / 32

Wasted Slots: 11

Execution Passes: 3

Serialization Overhead: 66.7%

Threads Taking If-Branch: 11

Threads Taking Else-Branch: 21

Branch Imbalance: 31.3%

Recommended Fix: Group similar data per warp to avoid modulo-driven divergence.

High divergence

Condition: threadIdx.x % 3 == 0

Key concept

A warp does best when many lanes stay active together. The more lanes peel off into separate execution paths, the lower your effective utilization becomes.

Common mistake

Developers often focus only on core count or clock speed while ignoring control-flow structure. Divergence can quietly dominate performance even when hardware looks strong on paper.

How a warp actually executes

A GPU does not schedule individual threads. On NVIDIA hardware it schedules warps — groups of 32 threads that issue the same instruction in the same cycle. This is the SIMT (single instruction, multiple thread) model, and it is the single most important thing to understand about GPU performance, because almost every surprising result traces back to it. Thirty-two threads share one instruction pointer, so they move through your kernel together whether or not that is what your code implies. AMD hardware uses the same idea with a 64-thread wavefront (32 on RDNA), so the reasoning transfers.

When threads inside one warp hit a branch and disagree — some take the if, others the else — the hardware cannot run both at once. It executes one path with the disagreeing lanes masked off, then the other path with the mask inverted. Both paths cost full time, and the masked lanes contribute nothing. That is warp divergence. A two-way branch can halve throughput; a 32-way switch where every lane picks a different case can, in the worst case, cost 32× the time of a uniform warp while doing the same amount of useful work. Since Volta, independent thread scheduling lets warps reconverge more flexibly, but the cost of executing both sides remains.

The critical detail people miss: divergence only costs you within a warp. If threads 0–31 all take the true branch and threads 32–63 all take the false branch, those are two different warps and there is no penalty at all. This is why the fix for divergence is almost never removing the branch — it is reorganizing the data so that threads that will take the same path land in the same warp. Sorting or bucketing work by branch condition before launching the kernel routinely recovers most of the loss.

Occupancy and latency hiding

A GPU tolerates slow memory not by making it faster but by having enough other work queued to run while it waits. A read from global memory costs hundreds of cycles. Rather than stall, the streaming multiprocessor swaps to another resident warp and issues its instruction instead. Occupancy — the ratio of resident warps to the hardware maximum — measures how much of that alternative work is available.

Occupancy is capped by whichever per-SM resource runs out first: registers per thread, shared memory per block, or the block and warp limits themselves. A kernel using many registers per thread lets fewer warps stay resident, which reduces the pool available to hide latency. This is why register pressure shows up as a memory-latency problem even though nothing about the memory changed.

The common misconception is that higher occupancy is always better. It is not. Occupancy only needs to be high enough to cover the latency actually present; past that point, pushing it higher usually means cutting registers per thread, which forces spills to local memory and makes the kernel slower. Treat occupancy as a floor to clear, not a number to maximize. Measure with the occupancy estimator, then confirm against real kernel timings rather than trusting the ratio alone.

Frequently asked questions

What is a warp?

A group of 32 threads that execute the same instruction together on NVIDIA GPUs. It is the smallest unit the hardware actually schedules. AMD calls the equivalent a wavefront, sized 64 or 32 depending on architecture.

Does every if statement cause divergence?

No. A branch is free when all 32 threads in the warp evaluate it the same way — the warp simply takes one path. Divergence only occurs when threads inside the same warp disagree.

How do I fix warp divergence?

Reorganize data so threads taking the same path share a warp — sort or bucket by the branch condition before launching. Replacing short branches with predicated arithmetic also helps, since the hardware can compute both results and select without serializing.

Is 100% occupancy the goal?

No. You need enough resident warps to hide memory latency; beyond that, raising occupancy usually means fewer registers per thread and spills to local memory, which costs more than the extra warps gain.