Foundations — SIMT (single instruction multiple thread)

This page builds every word and symbol the parent SIMT note leans on. If a term appears there without explanation, you will find its picture here. Read top to bottom — each block only uses things defined above it.

1. Thread — the smallest worker

The parent note keeps saying "32 threads do the same step". Before that sentence means anything, you must picture a single worker first.

Look at figure 1: each blue box is one thread. Every thread has its own scratch space (its registers) and its own answer, but they are all reading from the same recipe (the same program). That shared recipe is the whole trick — hold onto it.

2. threadIdx.x — "which worker am I?"

Every thread needs a name so it knows which data to touch. In CUDA that name is threadIdx.x.

Why does the topic need this? Because all threads run the same code. The only way thread #7 does different work from thread #3 is by using its badge number to pick a different array slot:

int val = data[threadIdx.x];   // thread 7 reads data[7]

Same recipe line, different plate — that is SIMT in one line of code.

3. Lockstep — "everyone on step 4 together"

The parent note's key phrase "execute in lockstep" is why the GPU saves hardware. If all 32 threads are guaranteed to be on the same instruction, you only need one instruction-fetch unit feeding all 32 workers, instead of 32 separate fetch units.

Figure 2 contrasts the two possible worlds: on the left every worker has its own drummer (expensive, like a CPU with many cores); on the right one drummer beats for all 32 (cheap — this is SIMT). The cost you pay for the cheap version comes later, when workers want to be on different steps (that is divergence, below).

4. Warp — the bundle of 32

Why 32 and not 8 or 100? Three plain reasons the parent lists:

• One fetch feeds many: one instruction decode drives 32 arithmetic lanes — good ratio of overhead to work.
• 32 is a power of two ($32=2^5$), so "which lane am I inside my warp?" is a cheap bit operation: threadIdx.x & 31. (The & is bitwise AND, and masking with 31 keeps the low 5 bits — that is just the remainder after dividing by 32.)
• Memory hardware likes 32-wide chunks (you will meet this as coalescing).

5. Active vs masked — the on/off switch per lane

Because 32 threads share one instruction, what happens when one thread should sit out? The hardware cannot skip that worker's slot in the bundle — the slot still ticks. Instead it masks the lane: the lane runs the instruction but throws the result away.

Figure 3 shows the mask as a strip of 32 boxes: green = kept, gray = discarded. This is the mechanism behind the parent's phrase "masked off if diverged." No lane is ever truly removed — it is only silenced.

6. Divergence — the two-recipe problem

Picture it with figure 4:

• Phase A: the "if" lanes are active (green), the "else" lanes masked (gray). The warp runs branch A.
• Phase B: the masks flip. The "else" lanes are active, the "if" lanes masked. The warp runs branch B.

So the warp spends time A + time B, even though each individual thread only needed one of them. This is exactly the parent's formula:

$T_{\text{diverged}}=T_A+T_B\text{vs}{T}_{\text{converged}}=\max (T_A,T_B)$

Why max when converged? If all threads take the same branch, no masking happens — the warp runs that one branch once, taking whichever time that branch costs. The max in the general case just means "the longest single path", since a converged warp only ever walks one path.

7. Latency and throughput — waiting vs flowing

Two words the parent uses constantly. Keep them separate:

8. Occupancy — how many warps are "in the oven"

Why the topic cares: while one warp waits out its 400-cycle memory latency, the scheduler runs other warps. The more warps you have ready ($W$), the more waiting you can hide. That is the parent's rule:

$W\ge L\cdot T$

Read in words: the number of ready warps must be at least the latency times the throughput you demand. To hide $L=400$ cycles at $T=1$ instruction/cycle you need $W\ge 400$ warps' worth of independent work. Since a single SM holds far fewer, GPUs use many SMs and demand high occupancy — this is the whole reason occupancy matters.

9. Streaming Multiprocessor (SM) — the workshop

Everything above — warps, masks, occupancy — happens within one SM. When the parent says "warps per SM", the SM is this workshop. See Streaming Multiprocessor for its full anatomy.

How the pieces feed the topic

Each arrow means "you must understand the source before the target makes sense." The whole graph funnels into the SIMT model — the parent topic.

Where each foundation is used next

• Thread + warp → CUDA Thread Hierarchy (thread → warp → block → grid)
• Masks + divergence → Warp Divergence
• 32-wide access → GPU Memory Coalescing
• Latency vs throughput, occupancy → GPU Occupancy
• Lockstep vs vector lanes → SIMD vs SIMT Comparison
• Why many warps, not deep pipelines → GPU vs CPU Architecture and Instruction-Level Parallelism

Equipment checklist

Test yourself — cover the right side and answer before revealing.