VisualGPU

The problem › The front-end bill

2,048 threads cannot afford 2,048 fetchers

If the GPU's entire survival strategy rests on keeping 2,048threads resident on a single Streaming Multiprocessor (SM), we immediately hit a silicon-budget crisis. On a traditional CPU, every active thread requires a "front-end" — the fetch and decode logic that reads the instruction stream and prepares it for the arithmetic logic units. Replicating that fetch-and-decode machinery 2,048 times would instantly bankrupt the very transistor budget the GPU just saved by stripping out its complex control logic. How does the hardware manage 2,048 active threads without paying for 2,048 instruction fetchers?

The mechanism › Amortization

Fetch once, decode once, issue to 32 lanes

It pays off the debt through amortization. The SM does not give each thread its own front-end. Instead, the hardware groups threads together into bundles of 32, called warps.

1CUDA C++ Programming Guide Execution Model (SIMT architecture)

When a warp is ready to execute, the warp scheduler fetches an instruction exactly once, decodes it once, and then issues it across 32 parallel execution units, or lanes. Each lane executes that exact same instruction, but on its own private data, using its own registers. By amortizing a single fetch and decode operation over 32 threads, the budget balances perfectly.

Compute the number › balancing the front-end budget

resident threads per SM (max)	2,048
threads per warp	32
warps to track	2,048 / 32 = 64
hardware warp slots per SM	exactly 64
front-ends actually built	4 schedulers × 32 lanes each
Result	2,048 threads, tracked by 64 slots, fed by 4 front-ends — a 32:1 amortization

SourcesH100 whitepaper Table 4 · Fig. 7

The mechanism › Scalar in, vector out

SIMT is not SIMD

At this point, if you are familiar with CPU architectures, you might assume this is just Single Instruction, Multiple Data (SIMD) — the identical mechanism behind CPU vector extensions like AVX. But the GPU employs a fundamentally different abstraction: Single-Instruction, Multiple-Thread (SIMT).

In a traditional SIMD architecture, the programmer is painfully aware of the hardware's vector width, because the SIMD vector organization exposes that width directly to the software. You write explicit vector code against that width, manually packing elements into an AVX register and invoking a specific vector-add instruction. In SIMT, you write ordinary, scalar code for a single thread. You never declare a vector width in your program. The hardware takes your scalar code and dynamically groups the threads into warps at runtime. The programming model is completely scalar; only the physical execution is vectorized. This is why a warp is not just a vector register. The hardware handles the vectorization behind the scenes, allowing you to focus purely on the logic of one thread.

2CAQA 6e Ch. 4 (SIMT vs SIMD terminology)

The full story › Rigor check: SIMT is a relabeling

While "SIMT" is the official programming model terminology, the underlying hardware is fundamentally a multithreaded SIMD processor. The concept of using a single instruction stream to drive multiple execution lanes while masking inactive lanes is not unprecedented; it is a relabeling of data-level parallelism descending directly from classical vector architectures (Computer Architecture: A Quantitative Approach, Chapter 4).

Interactive

The SIMT illusion

Branch cutoff · if (i < …)

cycle

1/6

lanes active

32/32

warp efficiency

83%

Your code · one thread · no vectors anywhere

// scalar code — written for ONE thread
int i = threadIdx.x;
float x = data[i];
if (i < 16) {
  x = x * 2.0f;
} else {
  x = x + 1.0f;
}
out[i] = x;

Warp scheduler · fetch ×1 · decode ×1

int i = threadIdx.x;

Timeline

Fetched once, decoded once — issued to all 32 lanes. Each lane gets its own i.

1 step = 1 issue cycle for the warp · warp = 32 lanes (hardware chose the width — your code never did) · uniform cutoffs 0 and 32 skip a pass: no divergence, no penalty

The consequence › Lockstep

The price of disagreement

But this hardware abstraction demands a strict physical compromise. Because the 32 threads in a warp share a single instruction fetch and a single program counter, they are not truly independent. They must march in strict lockstep.

Predict first

A warp hits an equal-length if/else. 31 lanes take the if, 1 lane takes the else. What's the warp's throughput across the divergent region?

This creates a severe penalty for branching. If an if-else statement causes a warp to diverge — say, the hardware evaluates if (threadIdx.x < 16) and sends half the lanes one way and the rest another — the warp executes each branch path taken, disabling the threads that are not on that path. The hardware must serialize the execution. It runs the if path for the first 16 threads while the other 16 lanes sit completely dark and idle. Then it flips the active masks, idling the first group to run the elsepath. The paths are executed sequentially, and the warp's throughput falls in proportion to the number of distinct paths taken. The illusion of thread independence fractures the moment threads disagree on where to go, and the cost is paid in wasted cycles.

ChipDiagram · Level 1

The lane array — one crop of the SM

SM — unexplored

warp scheduler · 32 thread/clk[WP·?]

dispatch unit · 32 thread/clk[WP·?]

32 × FP32/INT32 execution lanes [DERIVED · 128 ÷ 4]

the rest of this box — three more of these blocks, the memory system, whatever feeds it — is 1.3's territory

uniform: all 32 lanes execute the same instruction

[WP] = verbatim H100 whitepaper · [DERIVED] = arithmetic shown · this level is a crop of Level 2 (chapter 1.3) · source of truth: context/H100-CHIP-SPEC.md

The full story › What I simplified

I simplified the strictness of the lockstep model. In the prose above, I state that threads in a warp share a single program counter and must always march in perfect lockstep, serializing entirely upon divergence. While this is the foundational mental model of SIMT, modern NVIDIA GPUs actually implement Independent Thread Scheduling. This allows threads within a warp to maintain their own program counters and execution states, enabling them to yield execution or diverge and reconverge at a sub-warp granularity.

This departure is honest for this chapter because the fundamental performance characteristic has not changed: divergent paths mapped to the same SIMD execution units still cannot execute simultaneously. The hardware must still mask off inactive lanes and serialize the execution of divergent branches. Teaching the lockstep model first establishes the exact cost of branch divergence before introducing the nuances of modern deadlock-prevention mechanisms.

Despite this branching penalty, the warp remains the fundamental scheduling unit of the GPU. But this brings us to a harsh mechanical reality. On a CPU, putting a thread to sleep and switching to another is a catastrophic event that costs microseconds of overhead. The GPU claims its warp schedulers can switch between warps on every single clock cycle. What makes that switch free?

One fetch drives 32 lanes: you write scalar code, the hardware picks the width. The bill arrives only when lanes disagree — every distinct path is a full serial pass with the rest of the warp masked dark.