Intuition The one core ideaA GPU runs threads not one-by-one but in fixed bundles of 32 , and every thread in a bundle is forced to run the same instruction at the same moment . Everything in this topic — warps, scheduling, divergence, occupancy — is just the consequence of that one design choice: bundle the threads, share the control, hide the waiting.
This page assumes nothing . Before you read Warps and warp scheduling proper, we build every word, symbol, and picture it leans on. Read top to bottom; each block earns the next.
Before warps, we need the atom they are made of.
A thread is one worker running one copy of your program , working on its own little piece of the data. If you add two arrays of a million numbers, you launch a million threads and thread number i computes C[i] = A[i] + B[i].
Picture: one tiny worker holding one index number i, standing over one slot of memory.
Threads never live alone. They come in a nested set of containers, which we draw next.
A grid is the whole army of threads for one launch.
A block is a team inside the grid — threads in the same block can talk via fast shared memory.
A warp is the smallest squad the hardware actually moves : 32 threads glued together.
This nesting is exactly the subject of Thread-Blocks-and-Grids ; here we only need to see the boxes.
==threadIdx== is the coordinate label each thread carries so it knows which piece of data is mine . It has three parts written threadIdx.x, threadIdx.y, threadIdx.z.
Picture: in figure s01, each cell in a block is a thread; its (x, y, z) is just its row/column/layer address inside that block.
Intuition Why three numbers and not one?Real data is often 2D (an image) or 3D (a volume). Letting a thread say "I am at column x, row y" makes it natural to map thread → pixel. But the hardware underneath is a plain straight line of workers, so we must flatten those coordinates into a single number. That flattening is the next symbol.
==blockDim== tells you how many threads sit along each axis of a block: blockDim.x, blockDim.y, blockDim.z. If blockDim = (16, 8, 1) the block is 16 wide, 8 tall, 1 deep → 16 × 8 = 128
Picture: the width, height, and depth of the grid of cells in figure s01.
The hardware lines all threads up single-file. To find a thread's spot in that line we read the block row by row (x fastest, then y, then z). This is called row-major order.
Worked example Reading the formula on a17 × 8 (16, 7, 0), i.e. last column, last row, blockDim = (17, 8, 1).
lin = 0 × ( 17 × 8 ) + 7 × 17 + 16 = 0 + 119 + 16 = 135
The picture shows why row-major matters: it decides which threads land in the same warp . A single warp is a slice of 32 consecutive line positions — so a warp can span the boundary between two rows.
Now we can define the warp arithmetically.
Two new symbols appeared. Meet them:
⌊ x ⌋ ⌊ x ⌋ round down to the nearest whole number . ⌊ 135/32 ⌋ = ⌊ 4.21 ⌋ = 4
Picture: stand on a number line, then step left to the nearest integer tick.
⌈ x ⌉ ⌈ x ⌉ round up to the nearest whole number . ⌈ 136/32 ⌉ = ⌈ 4.25 ⌉ = 5 5 warps.
Picture: step right on the number line to the nearest integer tick.
Intuition Why we need ceiling hereThe hardware can only hand out whole warps — you cannot ask for "4.25 warps". So the count of warps is always rounded up : any leftover threads still cost a full warp. That is where waste comes from, and why the parent note begs you to keep block sizes a multiple of 32.
Common mistake Floor vs ceiling mix-upUse floor to ask "which warp is a given thread in?" (index → bucket).
Use ceiling to ask "how many warps do these threads need?" (count → containers).
Lockstep means every thread in a warp executes the exact same instruction on the same clock tick . Not "similar work" — the identical instruction, just on different data.
Picture: 32 rowers pulling on the same drumbeat; one instruction fetch drives all 32 oars.
SIMT = Single Instruction, Multiple Thread . One instruction stream, 32 threads riding it, each with its own data and its own registers. Compare and contrast this with plain SIMD in SIMT-vs-SIMD .
But branches exist — sometimes some threads should skip an instruction. The hardware handles this with a mask.
The active mask is a 32-bit on/off switch , one bit per thread. A 1 means "this thread does the instruction"; a 0 means "sit this one out, produce nothing."
Picture: a strip of 32 light-switches over the warp; the fetched instruction fires, but only lit threads act.
Intuition Why the mask makes divergence expensiveA warp cannot run two different instructions at once. If 20 threads want the if branch and 12 want the else, the warp runs the if code with 12 switches off, then runs the else code with 20 switches off — the two paths happen one after the other . That serial replay is exactly warp divergence .
A cycle is one tick of the GPU's clock — the smallest unit of time the hardware measures. "400 cycles" just means "400 ticks of waiting."
L L number of cycles you must wait for a result — especially reading from far-away global memory, which is slow (hundreds of cycles). See GPU-Memory-Hierarchy for where that distance comes from.
Picture: a warp fires a "fetch from memory" arrow; the answer does not come back for L
A warp is stalled when it cannot make progress — usually waiting out that latency L __syncthreads().
Intuition The whole point of having many warpsOne stalled warp is a wasted engine — unless another ready warp can jump in and fill the empty ticks. That swap is latency hiding : while warp 0 waits L
Occupancy = how full the SM is with warps , as a fraction:
Occupancy = Max Warps per SM Active Warps per SM necessary but not sufficient for speed — explored in Occupancy-vs-Performance .
Two symbols support it:
SM (Streaming Multiprocessor): one physical engine on the GPU that holds many warps and contains the schedulers. Think of it as one factory floor; the whole GPU has many.Warp scheduler: the traffic controller inside an SM that each cycle picks a ready warp to issue.
Why occupancy can't be pushed to 1 for free: each warp claims registers (Register-Pressure ) and shared memory. More warps → less per warp. That trade-off is the bridge from these foundations into the real tuning story.
Thread one worker one index
threadIdx and blockDim coordinates
Linear index row-major flatten
Floor and ceiling rounding
Warp 32 consecutive threads
Active mask on-off switches
Warp divergence serial replay
Latency hiding swap warps
Warps and warp scheduling
Self-test: cover the right side and answer each aloud before you read on.
What is a thread, in one sentence? One worker running one copy of the program on its own data slice.
What do threadIdx and blockDim describe? A thread's coordinate label, and how many threads sit along each axis of the block.
Flatten thread (3,2,0) in a (16,4,1) block. 0 + 2 × 16 + 3 = 35
What does ⌊ x ⌋ ⌈ x ⌉ Floor rounds down; ceiling rounds up.
Which one gives a thread's warp ID, and which gives the warp count? Floor gives the warp ID; ceiling gives the number of warps.
How many warps for a block of 136 threads? ⌈ 136/32 ⌉ = 5
What does "lockstep / SIMT" mean? All 32 threads run the same instruction on the same clock tick, each with its own data.
What is the active mask? A 32-bit on/off switch deciding which threads act on the current instruction.
Why is divergence slow? Both branches run one after the other with opposite masks, since a warp runs only one instruction at a time.
What is a stall, and how is it hidden? A warp unable to progress; another ready warp is swapped in for free to fill the empty cycles.
Give the formula for warps needed to hide latency L W ⌈ L / W ⌉
Define occupancy. Active warps per SM divided by max warps per SM.