A kernel launch creates a grid of thousands of threads, grouped into blocks that the hardware places on SMs independently, and each block is cut into warps of 32 threads that execute one instruction at a time, together. The warp exists for the hardware (one instruction fetch feeds 32 lanes); the block exists for you (threads that must cooperate have to live on the same SM). That last fact explains the two things this chapter is really about. Branches cost nothing when a warp agrees and cost both paths when it doesn’t, so the question is never “is there an if?” but “do the 32 threads of a warp disagree?”. And occupancy, how many warps an SM can hold at once, is set by whichever resource runs out first: thread slots, block slots, registers or shared memory. It matters only as far as it keeps enough work in flight to hide latency.
A CUDA program is an ordinary C++ program running on the CPU (the host) that occasionally hands a function to the GPU (the device) to run in parallel. The two have separate memories: your 12 GB of VRAM is not addressable by the CPU the way RAM is, and the CPU’s memory is not directly visible to kernels. So every CUDA program has the same skeleton, which you already ran in lab 0:
cudaMalloc.cudaMemcpy(..., cudaMemcpyHostToDevice).kernel<<<blocks, threads>>>(args).cudaMemcpy(..., cudaMemcpyDeviceToHost).cudaFree.Click through the host side of lab 0’s test to see each call and the trap hiding in it:
The most important line is the one that looks least important. vecadd<<<...>>>(...) does not run the kernel; it enqueues it and returns immediately, typically in a few microseconds, while the GPU starts working. The CPU is free to prepare the next launch. That is how a well-written program keeps the GPU fed, and it is why the two kinds of errors in check.h exist: a bad launch configuration is reported at once by cudaGetLastError(), but an out-of-bounds access inside the kernel only surfaces at the next synchronization point, possibly many lines later.
cudaMemcpy back to the host is what waits. The transfer lanes are also the honest cost accounting: the PCIe link is an order of magnitude slower than VRAM.On the CPU. The kernel reads and writes 12 bytes per element from VRAM at ~600 GB/s, but getting those bytes onto the card and back costs the same 12 bytes over PCIe at, at best, a tenth of that speed. The GPU’s entire advantage is eaten by the trip. The lesson is architectural, and it is how every serving engine is built: data lives on the GPU. Weights are uploaded once at startup; activations and the KV cache never leave VRAM; only token IDs and logits cross the bus. A kernel that is fast in isolation but requires a round trip to the host is not fast.
Now the device side. Click each line:
Three ideas to keep from it.
The loop became the grid. The CPU version is for (i = 0; i < n; ++i) c[i] = a[i] + b[i];. The kernel has no loop: each thread is one iteration, identified by its coordinates. This style is called SPMD, single program, multiple data: every thread runs the same function, and the only thing that differs is its ID and therefore its data.
Function qualifiers say where code runs. __global__ marks a kernel: runs on the device, launched from the host, returns void. __device__ marks a helper callable only from device code (a function a kernel calls; it does not launch anything). __host__ is the default for ordinary functions, and __host__ __device__ compiles one function for both sides, which is how lab 1’s path_a serves as both the kernel’s math and the CPU reference.
The grid is rounded up, so the guard is mandatory. Blocks all have the same size. For n = 1,000,003 and 256-thread blocks you need ⌈1,000,003 / 256⌉ = 3,907 blocks = 1,000,192 threads, and the last 189 threads have no element. Without if (i < n) they would read and write past the end of the arrays. The integer idiom for the rounding is (n + B - 1) / B.
blockIdx.x * blockDim.x is computed in 32-bit unsigned arithmetic. It wraps silently once a grid covers more than 232 elements, which is 16 GB of floats, so it is rare on your card but common on an 80 GB H100. Lab 0 casts to size_t before multiplying. Lab 1 uses int where sizes are known to be small; decide deliberately, every time.
Section 2 used three built-in variables as if their meaning were obvious. It isn’t, and most confusion about CUDA comes from mixing up two different hierarchies that share vocabulary: the one you write and the one that runs. This section builds both from the bottom, and for each level answers the question that matters: why does it exist?
Start from chapter 1. A GPU spends its area on arithmetic, not on control logic. But every instruction still has to be fetched, decoded and scheduled, and that machinery is expensive. The trick is to pay for it once and use it 32 times: the hardware fetches one instruction and executes it on 32 lanes at once, each lane with its own registers. The group of 32 threads that shares that one instruction stream is a warp.
Threads often need to work together: load a tile of data once and let everyone reuse it, or add up partial results. That needs two things: fast memory they all share, and a way to wait for each other. Both are physical parts of an SM: shared memory is on-chip storage inside one SM, and the barrier __syncthreads() is implemented by that SM. So threads can only cooperate if they are guaranteed to run on the same SM, at the same time.
That guarantee is exactly what a block is. You choose its size (up to 1,024 threads); the hardware places the whole block on one SM, keeps it there until it finishes, and cuts it into warps. The block is the unit of cooperation and of placement; the warp is the unit of execution. A warp never spans two blocks.
A launch creates a grid of blocks, usually far more than the GPU can hold at once. The price of being a block is independence: blocks cannot wait for each other and get no ordering guarantee. Block 7 may finish before block 0 starts. In return the hardware can run them in any order, on any SM, as many at a time as fit. That is transparent scalability: your 48-SM card and a 170-SM RTX 5090 run the same binary, and the bigger one simply drains the queue of blocks faster.
Every thread runs the same code, so the only thing that tells it which data is its own is its ID. Some IDs are per thread, some are shared by a whole block. Get this table into your fingers:
| Name | What it is | Same for |
|---|---|---|
threadIdx.x/y/z | position of the thread inside its block | no one else in the block; repeats in every block |
blockIdx.x/y/z | position of the block in the grid | every thread of the block (so every warp in it) |
blockDim, gridDim | shape of the launch | everyone |
tid = x + y·blockDim.x + z·blockDim.x·blockDim.y | the thread’s linear position in its block (you compute it) | no one else in the block |
warp = tid / 32 | which warp of the block | the 32 threads of that warp |
lane = tid % 32 | position inside the warp, 0–31 | no one else in the warp |
blockIdx.x * blockDim.x + threadIdx.x | position in the whole grid | no one else anywhere |
The row people trip on is the first one. Thread 5 of block 0 and thread 5 of block 1 both have threadIdx.x == 5; that is why every kernel adds blockIdx.x * blockDim.x to get a globally unique index. And lane equals threadIdx.x % 32 only when blockDim.x is a multiple of 32. In general a warp is cut from the linear tid, which section 4 shows in two dimensions.
threadIdx.x and global i?They share blockIdx.x = 10, blockDim.x = 256, gridDim.x, and of course warp number 3. They differ in threadIdx.x = 96…127 (warp 3 starts at 3×32), in lane = 0…31, and in i = 10×256 + threadIdx.x = 2,656…2,687: 32 consecutive elements. Everything in the first group is uniform across the warp, everything in the second varying. Section 6 turns that distinction into the most useful tool in this chapter.
If the warp is what the hardware runs and the block is only a grouping, why not make every block a single warp? Three reasons:
And the opposite, one enormous block? A block must fit on one SM, so it is capped at 1,024 threads. More importantly, one block runs on one SM: the other 47 would sit idle. The block is the unit that gets spread across SMs, so you want many of them.
Blocks of 128–512 threads (always a multiple of 32), and many more blocks than SMs, thousands if the data allows. Every SM then holds several blocks at once, and the few left over at the end (the “tail”) are a small fraction of the run.
It can deadlock. If the grid has more blocks than fit on the GPU at once, the resident blocks spin forever waiting for blocks that can never be scheduled, because the spinners are occupying the SMs those blocks need. Even when it happens to work on your card, it can hang on a smaller GPU or under a profiler. The standard answer is to split the kernel in two: the end of a kernel launch is the global barrier, and it costs only a few microseconds. (Cooperative launches offer a real grid-wide sync, but only for grids small enough to be entirely resident.)
A __syncthreads() must be reached by every thread of the block. Placing one inside an if that only some threads take, or one in the then branch and another in the else (they are different barriers), is undefined behaviour: wrong results or a hang. Keep barriers in code that all threads of the block execute.
If you want to see all of this move before reading on, jump to the film at the end of section 10: five minutes following one thread through a launch, from the CPU call to its last store.
Grids and blocks can have up to three dimensions, described by the type dim3 (x, y, z, with unused dimensions set to 1). You pick the dimensionality that mirrors the data: 1D for vectors, 2D for images and matrices. The limits on your card, which have not changed in years:
| Limit | Value |
|---|---|
| Threads per block (x · y · z) | ≤ 1,024 in total; x, y ≤ 1,024; z ≤ 64 |
| Blocks per grid | x ≤ 231−1; y, z ≤ 65,535 |
Take an image 1,000 pixels wide and 700 tall, one thread per pixel, with 32×8 blocks (256 threads):
dim3 block(32, 8); // x = 32 columns, y = 8 rows
dim3 grid((1000 + 31) / 32, (700 + 7) / 8); // 32 x 88 blocks
blur<<<grid, block>>>(in, out, 1000, 700);
// inside the kernel
int col = blockIdx.x * blockDim.x + threadIdx.x; // x is the column...
int row = blockIdx.y * blockDim.y + threadIdx.y; // ...y is the row
if (row < height && col < width) // BOTH guards
out[row * width + col] = ...; // row-major: row, then columnArrays are indexed [row][col], “y then x”, but launches are written dim3(x, y), “x then y”. Put the column in x. The reason is not taste: threads with consecutive threadIdx.x end up in the same warp, as the explorer below shows, and you want a warp to touch consecutive addresses, which in a row-major array means consecutive columns. Chapter 3 turns this into the coalescing rule.
Warps are cut from a single line of threads, so before cutting, the hardware flattens the block: x varies fastest, then y, then z, exactly like a row-major array.
tid = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y;
warp = tid / 32; // cut every 32
lane = tid % 32;With the 32×8 blocks above, one row of the block is exactly 32 threads, so each warp is one row: threadIdx.y (and therefore row) is the same for the whole warp, and threadIdx.x runs 0–31. With 16×16 blocks a row is only 16 threads, so each warp covers two rows of 16. Pick shapes below and click threads:
The explorer’s last line anticipates chapter 3: memory is fetched in 32-byte sectors, and a warp’s store costs one trip per sector it touches. A warp writing 32 consecutive floats of one row touches 4 sectors, the minimum. Notice that 16×16 and even 8×32 still hit the minimum (each row piece is a whole number of sectors), and that 4×64 does not.
row from x, col from y, and swap the launch too (same dim3 block(32, 8), but the grid sized so that x covers the 700 rows and y the 1,000 columns). Is it still correct? Is it as fast?Correct, yes: x and y are just two axes. As long as the kernel and the launch are swapped consistently, every (row, col) is still covered exactly once and both guards still protect the edges. (Swap only one of the two and you skip pixels or write out of bounds.)
As fast, no. Now consecutive lanes have consecutive rows, so a warp’s 32 addresses are a whole image row apart: 4 KB each. That is the explorer’s “x → row” setting with 32×8: watch the sector count go from 4 to 32 for the same 128 useful bytes, 8× the memory traffic. The rule: threadIdx.x walks the dimension that is contiguous in memory. For row-major arrays that is the column. Chapter 3 calls this coalescing.
Lab 0 printed a line you were told to ignore: ptxas info : Used 12 registers. Here is where it comes from.
.cu file. Host code goes to your normal C++ compiler. Device code goes to PTX, a stable virtual instruction set, and then ptxas assembles it into SASS, the real machine code of one GPU family. -arch=sm_120 is shorthand for “SASS for sm_120, plus the PTX”, so a future GPU can still JIT-compile it.Register allocation happens in ptxas, which is why it is the one reporting. “Used 12 registers” means each vecadd thread needs 12 32-bit registers: the three pointers (two registers each, since pointers are 64-bit), n, the index, the loaded values. Twelve is tiny. In section 9 you will see why the number matters: registers are the resource most likely to limit how many threads an SM can hold.
You can read the SASS yourself with cuobjdump -sass ./divergence. You don’t need to write it, but learning to skim it is a real skill: it is the only ground truth about what the compiler did. Here is the heart of lab 1’s branch, abridged:
ISETP.GE.AND P0, PT, R0, UR5, PT // P0 = (i >= n), one bit per thread
@P0 EXIT // the bounds guard: those threads leave
...
BSSY.RECONVERGENT B0, ... // note where the two paths will meet again
@P0 BRA path_b // threads with P0 jump; the others fall through
FFMA R5, R5, R2, 0.001 // path A: x = x * 0.999 + 0.001 (one fused op)
...
FMUL R5, R5, 0.999 // path B: the same map as two instructions
FADD R5, R5, 0.001
...
BSYNC.RECONVERGENT B0 // wait until both paths have finishedNotice the @P0 prefix: predication. Every SASS instruction can be guarded by a per-thread predicate bit, and threads whose bit is false simply don’t execute it. That is the mechanism behind everything in section 7.
SASS for sm_120 runs on compute capability 12.x GPUs and nothing else. PTX runs anywhere newer, via the driver’s JIT compiler (with a one-time delay at startup). Libraries like PyTorch ship SASS for several architectures plus PTX for the newest, which is why their binaries are huge. You will later meet sm_120a: the a unlocks architecture-specific instructions (some tensor-core formats) at the price of that forward compatibility.
An SM on your card is split into four quadrants (also called sub-partitions or processing blocks). Each quadrant has its own warp scheduler, a quarter of the 256 KB register file, 32 of the SM’s 128 FP32/INT32 lanes, and one tensor core. The L1 cache / shared memory (128 KB) is shared by the whole SM.
When a block arrives on an SM, its warps (cut as in section 4) are spread over the four quadrants, and each quadrant’s scheduler owns its warps for their whole life. A block whose size is not a multiple of 32 wastes the tail of its last warp: a 48-thread block occupies two warp slots, one half empty.
The warp is the unit the hardware actually schedules. A scheduler fetches one instruction and issues it to all 32 threads of a warp at once, each thread using its own registers: same instruction, different data. That is SIMD hardware, and NVIDIA calls the programming model on top of it SIMT, single instruction, multiple threads, because you write the code of one thread and the hardware groups 32 of them. The grouping is invisible for correctness and decisive for performance.
The most useful mental tool for everything that follows: for any value in a kernel, ask whether it is the same for all 32 threads of a warp.
n, blockIdx.x, blockDim.x, pointers passed as arguments, loop counters of loops with uniform bounds, threadIdx.x / 32 (the warp index).threadIdx.x, i, anything loaded from a[i], anything computed from a varying value.The hardware cares about this distinction so much that it has a separate uniform datapath for it: in the SASS above, the UR registers and U-prefixed instructions (UISETP, BRA.U) hold and process values the compiler proved are the same for the whole warp, computed once instead of 32 times. Section 7 applies the distinction to branches.
A warp has one instruction stream. So what happens at an if whose condition differs between threads? The hardware evaluates the condition per thread into a predicate, then:
then path with the else-threads masked off, then the else path with the then-threads masked off, then reconverges. Every thread waits through the path it did not take.So the cost model is simple: a warp pays for every path that at least one of its threads takes. Step through it:
The consequence, which is the thing to take away if you have ever struggled to “avoid branches” in shaders: branches are not the problem; disagreement within a warp is. Check the condition’s uniformity:
| Condition | Uniform within a warp? | Cost |
|---|---|---|
if (n > 1024), if (mode == 2) | yes: kernel arguments | free |
if ((threadIdx.x / 32) % 2 == 0) | yes: depends only on the warp index | free: different warps take different paths, each takes one |
if (threadIdx.x % 2 == 0) | no: neighbours disagree | both paths, every warp |
if (threadIdx.x % 32 != 0) | no: 31 against 1 | both paths, every warp: one dissenter is enough |
if (i < n) | no, but only in one warp | negligible: one warp of the grid has idle lanes |
if (x[i] > 0) on data | depends on the data | sorted data: rarely diverges; random data: usually diverges |
divergence, path A costs 1 instruction per step and path B costs 2. If all-A takes time t, predict all-B, warp-split, lane-split and one-lane.all-B ≈ 2t (twice the instructions). warp-split ≈ 1.5t: half the warps run A, half run B, each warp runs only its own path, so the average is (t + 2t) / 2. lane-split ≈ 3t: every warp diverges and pays A plus B. one-lane ≈ 3t as well: one lane taking B forces the whole warp through B while the other 31 wait, so it is exactly as bad as the 50/50 split. That last one is the one people get wrong. The cost of divergence depends on whether a warp disagrees, not on how many threads disagree.
The most common divergent branch in real code is the guard if (i < n). For n = 1,003 and 256-thread blocks, 4 blocks (32 warps) are launched; 31 warps are entirely inside the array, and one, covering i = 992…1,023, has 11 active lanes and 21 idle. It is a one-sided branch (there is no else), so the warp doesn’t even pay a second path: it just runs the body with 21 lanes masked, as the stepper’s last case shows. In a 2D kernel the same logic applies at the edges: in the 1000×700 example, the last column of blocks contains 704 partially-idle warps out of 22,528 (3%), and at the bottom, because each warp is one full row, whole warps fall outside the image and exit immediately without costing anything.
A loop whose trip count varies per thread is a branch at the end of every iteration. A warp runs until its slowest thread finishes: if one thread needs 1,000 iterations and the rest need 10, the warp runs 1,000, mostly with a single active lane. This is the real-world killer: ragged work per thread (variable-length sequences, sparse rows of different lengths, early-exit searches). The fixes are about work assignment, not syntax: sort or bucket the work so that neighbours get similar amounts, or split long items across several threads.
For a short conditional like y = (x > 0) ? x : 0.1f * x, the compiler usually doesn’t branch at all. It computes both sides and selects the result per thread (FSEL), or predicates a few instructions with @P. Both “paths” always execute, but they are one or two instructions each, which is cheaper than a jump. So for short divergent branches, divergence is simply not a concept you need to worry about. It becomes one when the paths are long, as in lab 1, or when they contain memory accesses.
Since the Volta generation (2017), each thread has its own program counter: independent thread scheduling. The two sides of a divergent branch may now interleave (A for a while, then B, then A again) instead of strictly running one after the other, and threads of a warp are no longer guaranteed to be in lockstep at any given point. What this changes:
__syncwarp(), or warp primitives like __shfl_sync that take a mask. You will use these in chapter 5.What it does not change: at any instant a warp still executes one instruction for one subset of its threads. A divergent warp still pays for both paths. The cost model above is unchanged.
1. Make conditions depend on uniform values whenever the algorithm allows: branch on the warp or block index, not the thread index. 2. If the condition depends on data, arrange the data so that neighbours agree: sort, bin, or launch separate kernels per case. 3. Keep divergent paths short, and let the compiler predicate them. 4. Hoist uniform work out of divergent regions. 5. Ignore divergence from bounds checks; it is noise.
Each quadrant can hold up to 12 warps (48 per SM) but issues one instruction per cycle. Why keep so many? Because most of them are waiting. A load from DRAM takes several hundred nanoseconds, so several hundred cycles; even an FMA takes a few cycles before its result can be used by the next dependent instruction. Every cycle, the scheduler looks at its warps, picks one whose next instruction has its inputs ready, and issues it. Switching between warps costs nothing, because every resident warp’s registers stay in the register file permanently: there is no context to save, unlike a CPU thread switch.
This is chapter 1’s Little’s law seen from inside one SM. To keep an issue slot busy through a latency of L cycles, you need about L cycles’ worth of independent instructions available, and there are exactly two places they can come from:
Either works. That is why the next section ends with a warning.
Occupancy is the number of warps resident on an SM divided by the maximum (48 on your card). A block is admitted to an SM only if all of its resources are available at once. It needs thread slots for all its threads, a block slot, registers for all its threads, and its shared memory. So four independent limits apply, and the tightest one wins:
| Resource per SM (CC 12.0) | Capacity | Consumed per block |
|---|---|---|
| Thread slots | 1,536 (48 warps) | threads per block |
| Block slots | 24 | 1 |
| Registers | 65,536 (32-bit) | registers/thread × threads, allocated per warp in units of 256 |
| Shared memory | up to 100 KB | what the kernel declares, + 1 KB reserved by CUDA per block |
Worked through for your card, block size alone first:
| Threads/block | Blocks that fit | Warps | Occupancy | Why |
|---|---|---|---|---|
| 32 | 24 | 24 | 50% | block slots run out first |
| 64, 96, 128, 192, 256, 384, 512, 768 | 1536 / size | 48 | 100% | divides 1,536 evenly |
| e.g. 448, 640 | 3, 2 | 42, 40 | 88%, 83% | doesn’t divide 1,536: the leftover slots stay empty |
| 1,024 | 1 | 32 | 67% | a second block doesn’t fit in the remaining 512 slots |
That last row surprises people coming from datacenter GPUs, where an SM holds 2,048 threads and 1,024-thread blocks are perfect. 1,536 is not a power of two: on your card, 1,024-thread blocks are a bad default, and 256 or 512 are good ones.
Registers are where the real cliffs are. For full occupancy each of 1,536 threads can have at most 65,536 / 1,536 ≈ 42 registers, and because allocation is per warp in units of 256 registers (8 per thread), the practical limit is 40. At 41 registers each warp’s allocation rounds up to 48 per thread and occupancy drops to 83%. At 64 it is 67%, at 80 it is 50%, at 128 it is 33%. One extra variable can cost you a sixth of your warps. Try it:
tile kernel keeps 64 accumulators per thread and compiles to 80 registers. What is its occupancy with 256-thread blocks, and what happens if you force it to 100% with __launch_bounds__(256, 6)?At 80 registers: 80 × 32 = 2,560 registers per warp, so each quadrant (16,384 registers) holds 6 warps: 24 warps, 3 blocks, 50%. __launch_bounds__(256, 6) tells the compiler “I want 6 blocks of 256 per SM”, which caps it at 40 registers. It obeys, but 64 accumulators don’t fit in 40 registers, so it spills the rest to local memory (per-thread memory that lives in the cache hierarchy and ultimately DRAM). The build prints it: Used 40 registers, 428 bytes spill stores, 504 bytes spill loads. You get 100% occupancy and a kernel doing hundreds of extra memory operations per thread. Predict which is faster, then measure; most of the time the spilling version loses.
High occupancy helps only when the kernel needs more warps to hide latency. Section 8 says there is another source of in-flight work: independent instructions within a thread. Some of the fastest kernels you will write in Part II (matrix multiply with register tiling, FlashAttention) deliberately run at 25–50% occupancy, using many registers per thread to hold data close and keep many independent operations in flight. Vasily Volkov’s talk Better Performance at Lower Occupancy made this point famous in 2010, and it has only become more true.
Back to Little’s law. Chapter 1 estimated that even 100% occupancy with one 4-byte load per thread was barely enough to cover DRAM latency. So vecadd should start losing bandwidth as soon as occupancy falls noticeably below 100%, and collapse toward the bottom of the range. vecadd4 keeps 4× more bytes in flight per thread, so it should hold its bandwidth down to a much lower occupancy, roughly a quarter of what vecadd needs, before it starts to drop. If your measurements show that, you have just seen TLP and ILP trade against each other on real hardware.
Here is everything in this chapter, plus a preview of chapter 3, written as the program the hardware runs. It is a model, not a spec: real schedulers are smarter and the numbers differ between generations. But every performance topic in Part I is one line of it. Read the four tabs in order.
# HOST (your C++): enqueue and return, the CPU does not wait
kernel<<<grid, block, shared_bytes>>>(args)
# GPU-WIDE BLOCK DISTRIBUTOR (hardware)
waiting = all blockIdx of the grid # no order promised
while waiting:
for sm in the 48 SMs:
while waiting and sm.fits(block): # section 9: thread slots, block slots,
b = waiting.pop() # registers, shared memory
sm.reserve(b.registers, b.shared)
# cut the block into warps: flatten x-first, then every 32
for w in range(ceil(threads_per_block / 32)):
warp = threads with tid 32*w .. 32*w + 31
warp.mask = the lanes that exist # partial last warp
warp.pc = first instruction of the kernel
sm.scheduler[w % 4].add(warp) # spread over 4 quadrants
when a block's warps have all exited:
sm.release(b) # room for the next waiting block# EVERY CLOCK CYCLE, in every SM, for each of its 4 schedulers:
ready = [w for w in scheduler.warps # up to 12 warps each
if inputs_of(w.next_instruction) are ready # "scoreboard": loads back?
and not w.waiting_at_barrier]
if ready:
w = pick_one(ready) # one warp per scheduler per cycle
execute(w, w.next_instruction) # tab 3
else:
stall # nothing ready: the issue slot is wasted
# (section 8: more warps or more ILP)
# Switching warps costs nothing: every resident warp keeps its registers
# in the register file the whole time. There is no context to save.execute(w, instr):
if instr is arithmetic: # FFMA, IADD, ...
for lane in 0..31, all at once:
if w.mask[lane] and instr.predicate(lane): # the @P guard
reg[lane][dst] = op(reg[lane][a], reg[lane][b])
# the result is usable a few cycles later; dependent instructions wait
if instr is a branch on cond: # section 7
yes = w.mask & cond # per-lane bits
no = w.mask & ~cond
if yes == w.mask or no == w.mask: # UNIFORM: everyone agrees
jump (or fall through) with the whole warp # costs one instruction
else: # DIVERGENT
run the "yes" path with mask = yes
run the "no" path with mask = no # the warp pays for both
at the reconvergence point: mask = yes | no if instr is a global load or store: # chapter 3
addrs = [address(lane) for each active lane]
sectors = { a // 32 for a in addrs } # 32-byte pieces touched
fetch each sector from L1, else L2, else DRAM # fewer = faster
mark dst "pending" # this warp is not ready now;
# ... hundreds of cycles later, when every sector has arrived:
mark dst "ready" # the scheduler may pick it again
if instr is a shared-memory access: # chapter 3
bank(lane) = (address(lane) // 4) % 32
rounds = most DIFFERENT addresses in any one bank # same address: broadcast
takes `rounds` passes # 1 = conflict-free
if instr is __syncthreads(): # section 3
w.waiting_at_barrier = True
when every warp of the block has arrived:
release them all, with their shared-memory writes visibleNow map each line to what you learned:
| Line | Topic | Where |
|---|---|---|
sm.fits(block) | occupancy: which resource runs out first | section 9 |
| flatten x-first, cut every 32 | warp shape, partial warps, why x walks columns | section 4 |
pick_one(ready) / stall | latency hiding: TLP and ILP | section 8, chapter 1 |
yes / no masks | divergence costs both paths | section 7 |
sectors = {…} | coalescing | chapter 3 |
rounds = … | bank conflicts | chapter 3 |
__syncthreads() | cooperation inside a block, tiling | section 3, chapter 3 |
Since Volta each thread has its own program counter, so the two sides of a divergent branch may interleave instead of running strictly one after the other (section 7); the cost model above is unchanged. Real schedulers also have policies (which ready warp to prefer) and pipelines of different widths for different instruction types. None of that changes how you reason about a kernel; the lines above do.
Ten sections took a kernel launch apart. This film puts it back together by following one thread from birth to exit: element 2,690 of mix, a kernel that adds two arrays and then treats odd and even elements differently. The thread tells the story itself, and almost nothing in it is new. Watch for three moments the chapter promised. The thread spends nearly its whole life waiting for one memory load, and its SM doesn’t mind, because the scheduler always has another warp ready (section 8). When odd and even lanes disagree, the warp runs both paths, one after the other (section 7). And with room for six blocks per SM, 1,024 blocks pass through the chip in three and a half waves, the last one ragged (section 9).
The whole launch takes about ten microseconds. The clock in the corner shows where the film is in that real time, and how many times slower than life you are watching it.
Real: 48 SMs, 6 blocks of 256 threads per SM for this kernel (the 1,536-thread limit of section 9), 288 resident blocks and 3.6 waves, 4 schedulers with up to 12 warps each, 32-byte sectors, a coalesced warp access becoming 4 sectors, both paths of a divergent branch being issued, fire-and-forget stores, a few microseconds of launch latency. Staged: the exact nanosecond timings (a loaded memory round trip of ~800 ns and block lifetimes of ~1.5 µs are plausible for a memory-bound kernel, not measured), which DRAM chip holds a[2690], the on-chip route of a request, and the order in which the scheduler picks warps. The same story is also told as eight still plates and in a first, flat cut of this film.
The film draws the GPU as a map. Figure 7 tells the same story on the hardware itself: the CPU on its motherboard, the PCIe slot and a riser cable, and an RTX 5070 with its cooler off. Watch block 10 land in one of SM 10’s six slots, its two loads leave the die for a GDDR7 chip and come back, and 1,024 blocks pass through 288 slots in waves. Point at any SM to see which blocks it holds.
Two programs, four experiments, each one a prediction from this chapter. Write your predictions down before running.
cd mlsys/labs/lab01-warps-and-occupancy
make # read the ptxas lines: 10, 12, 22, 80 and 40 (+spills) registers
make test # every kernel against a CPU reference
make bench # A: divergence, then B, C, D: occupancy| Exp. | What it does | Prediction to write down |
|---|---|---|
| A | Five branch patterns over two arithmetic paths (section 7) | time of each, relative to all-A |
| B | vecadd with 32 … 1,024 threads per block | which block sizes lose bandwidth, and why |
| C | vecadd and vecadd4 at occupancy 100% → 17% | where each starts to lose bandwidth |
| D | 80-register tile kernel, free vs. capped at 40 | which one is faster |
Each experiment prints the theoretical occupancy from cudaOccupancyMaxActiveBlocksPerMultiprocessor, the runtime’s own calculator, next to the time. If it disagrees with the widget above, the runtime is right; tell me and I’ll fix the widget. Paste the four PASTE: lines back.
In divergence.cu, replace path B with the same FMA as path A but different constants. Rebuild, run, and look at the SASS with cuobjdump -sass divergence | grep -E "BRA|FFMA|FSEL". Did the compiler keep the branch, or merge both paths into one instruction stream with a per-thread constant? Divergence is a property of the machine code, not of the source.
The warp exists for the hardware: fetching and scheduling an instruction is expensive, so the GPU does it once for 32 threads that then execute it together. The block exists for the programmer: threads can only share memory and synchronize if they are guaranteed to be on the same SM at the same time, and a block is exactly that guarantee. The block is the unit of cooperation and placement; the warp is the unit of execution.
out[row][col] (x → column) coalesce?The block flattens x-first, so tid = x + 16y; warp 1 is tid 32–63, which is rows y = 2 and 3 of the block, all 16 columns each. Its store touches two runs of 16 consecutive floats, 64 bytes each: 2 sectors per row, 4 in total, the same as a full 32-wide row. It coalesces fine at the sector level; a 4-wide block would not.
Launches are asynchronous: they enqueue work so the CPU can keep feeding the GPU. Configuration errors are reported immediately by cudaGetLastError(), but faults during execution (like an out-of-bounds access) only appear at the next synchronizing call, possibly far from the launch that caused them. When debugging, synchronize right after the launch.
(5,000 + 255) / 256 = 20 blocks = 5,120 threads, 120 idle. Threads 5,000…5,119: warp 156 (threads 4,992…5,023) is partially idle, and warps 157…159 are entirely idle and exit at once. So only one warp has a mix of active and idle lanes.
Blocks are scheduled independently, in any order, as resources free up; a grid can have more blocks than fit at once, so a block waiting for another that isn’t resident yet could wait forever. Independence is what buys transparent scalability. For a global barrier, end the kernel and launch another: the kernel boundary is the grid-wide sync.
-arch=sm_120 put in your binary?PTX is NVIDIA’s virtual, forward-compatible instruction set; SASS is the real machine code for one GPU family. -arch=sm_120 embeds SASS for sm_120 plus the PTX, so the program runs natively on compute capability 12.x and can be JIT-compiled by the driver on newer GPUs.
if (lane == 0) { long work }.A warp pays for every path that at least one of its threads takes. With lane == 0, every warp executes the long work (one lane active, 31 masked) in addition to whatever the others do, so it costs as much as if all 32 threads had done it.
blockIdx.x < 10, threadIdx.x < 16, threadIdx.x < 64 (block of 256), data[i] > 0?blockIdx.x < 10: never, it is uniform across the block. threadIdx.x < 16: yes, it splits warp 0 in half. threadIdx.x < 64: no, the boundary falls exactly between warps 1 and 2. data[i] > 0: depends on the data; neighbouring values that agree mean no divergence.
Threads have individual program counters, so divergent paths can interleave and threads of a warp are not guaranteed to be in lockstep. Code that relied on implicit warp lockstep must use __syncwarp() or *_sync primitives. It did not make divergence cheaper: a warp still executes one path at a time.
Threads allow 6 blocks. Registers: 48 × 32 = 1,536 per warp; 16,384 / 1,536 = 10 warps per quadrant = 40 warps, so 5 blocks. Shared memory: 100 KB / (20 + 1) KB = 4 blocks. Block slots allow 24. The minimum is 4 blocks = 32 warps = 67%, limited by shared memory.
Occupancy only matters as a source of independent work to hide latency. A kernel whose threads each have many independent operations in flight (instruction-level parallelism, typically using many registers to hold data) can saturate the hardware with fewer warps, and the registers that lowered its occupancy are what let it reuse data instead of re-reading memory.
-arch=sm_XX shorthand.__syncwarp.