The CUDA Programming Model

One cute dog, 786,432 pixels, and what actually happens between your image and the silicon.

August 2026 · a slow-motion prequel to Simon Boehm's CUDA matmul worklog, built on Stephen Jones' "How CUDA Programming Works" (GTC 2022)

There are two pieces of material on CUDA that I keep sending people. The first is Simon Boehm's worklog on optimizing a CUDA matmul kernel, which is excellent and also, for a beginner, brutal: by the third paragraph you are staring at a roofline analysis of SGEMM. The second is Stephen Jones' GTC talk "How CUDA Programming Works", which does something the blogs rarely do. It takes one picture, splits it into tiles, and follows those tiles down into the hardware until the whole machine makes sense.1 1The talk is also listed on NVIDIA On-Demand. Jones is one of the architects of CUDA, and his companion talk "How GPU Computing Works" (GTC 2021) covers the memory physics in more depth.

This post lives in the gap between the two. We are not going to write clever kernel code, and we are not going to optimize anything. We are going to take one image, give the GPU one honest job to do with it, and follow that job all the way down: through the die, the GPCs, the SMs, the blocks, the warps, the schedulers, and the two or three ideas (SIMT, divergence, occupancy) that every CUDA text mentions and almost none of them ground in a concrete example. By the end, the opening sections of Boehm's post should read like a natural next chapter rather than a cliff.

Everything below uses the numbers of one specific GPU, NVIDIA's H100 (SXM5), because it is the chip people actually rent today and because its structure is easy to hold in your head.2 2Jones' talk uses the A100 (108 SMs, 1.56 TB/s HBM2). The shapes are identical; only the counts change. H100 figures here come from NVIDIA's Hopper architecture whitepaper and the H100 datasheet. The ideas transfer unchanged to every CUDA-capable GPU back to 2007, including the one in your laptop.

1One dog, 786,432 small jobs

Here is the image we will use for the whole post: a puppy, 1024 pixels wide and 768 pixels tall.

The first job we will give the GPU is deliberately plain: convert the image to grayscale. For every pixel, read three color values, combine them with three multiplications and two additions, and write one number back.3 3The weights 0.299, 0.587, 0.114 are the Rec. 601 luma coefficients. Green dominates because human eyes are most sensitive to it.

gray = 0.299·R + 0.587·G + 0.114·B

The important property is not the arithmetic, it is the independence. Pixel (17, 512) does not care what happens to pixel (18, 512). There is no ordering, no shared running total, no communication. The problem is 786,432 separate five-operation jobs that happen to live in one array.

This is exactly the shape of work a GPU is built for. A CPU gives you a handful of very fast, very clever cores and asks you to feed each one a long sequential story. A GPU gives you tens of thousands of simple arithmetic lanes and asks you a different question: can you phrase your problem as an enormous number of small, independent jobs? For images the answer is almost always yes, which is why the GPU was born as a graphics chip in the first place. However, before we can hand the dog to those lanes, we need to know what the machine underneath actually looks like.

2The hardware, from the die down to one SM

Every CUDA tutorial says "the GPU has thousands of cores" and then draws a rectangle full of dots, which explains nothing. The machine has a real structure, three levels deep, and each level of the programming model corresponds to one of them. It is worth thirty seconds per level.

Seen from above, an H100 is eight large tiles called GPCs (graphics processing clusters) arranged around one shared 50 MB L2 cache, with stacks of HBM3 memory bolted onto both edges of the die. A thin strip of scheduling hardware, the GigaThread Engine, sits in front of everything; we will meet it properly in section 5.

Zoom into any one of the eight GPCs and you find the same thing repeated: 18 streaming multiprocessors, the SMs. The full GH100 die therefore carries 8 × 18 = 144 SMs; the H100 you actually rent ships with 132 of them enabled, because a few are always fused off for manufacturing yield.4 4Precisely: a GPC pairs its SMs into 9 TPCs (texture processing clusters) of 2 SMs each. The TPC level rarely matters when you are starting out, so the figure skips it.

Now zoom into one SM, because this is where the programming model stops being a metaphor. An SM is divided into four identical quadrants, and each quadrant contains three things that matter to us: one warp scheduler, a slice of the register file (16,384 registers of 32 bits each), and 32 FP32 arithmetic lanes, the things marketing calls "CUDA cores". Alongside them sit the tensor core and the load/store units. Beneath all four quadrants sits 256 KB of fast local SRAM, which is carved into L1 cache and "shared memory".

Two words in that figure deserve to be unpacked before we go anywhere else, because everything later leans on them: register and lane. Neither is complicated, but both are usually left undefined, and then the reader quietly carries the confusion for the rest of the article.

A register is one 32-bit box that holds one number. That is the whole definition. "32 bits" just means the box is the size of one ordinary float or int. When your thread's code says int x or reads a pixel into p, those variables do not live in RAM somewhere far away; each one is parked in one of these boxes, physically inside the quadrant, closer to the arithmetic than any cache. The register file is simply the cabinet of all such boxes: 16,384 of them per quadrant, 65,536 per SM. Every thread gets its own private handful, and the compiler decides at build time how many boxes each thread of a given kernel needs, by counting how many values the thread must keep alive at once. Our grayscale thread needs roughly fifteen: its coordinates, its address arithmetic, the four channel bytes it loaded, the three partial products, the sum. Here is what that actually looks like for the one thread we keep zooming into.

A lane is one arithmetic machine. If you have wondered whether "128 FP32 lanes" means "128 floating-point calculations can happen per clock cycle", the answer is yes, exactly that. One lane takes numbers from registers, performs one FP32 operation per clock (a fused multiply-add, so technically two floating-point operations in one), and puts the result back into a register. It has no program counter and no cleverness; it is pure arithmetic, which is why the GPU can afford 128 of them per SM. And the arithmetic closes neatly: 132 SMs × 128 lanes = 16,896 lanes, each doing a multiply-add per clock at 1.98 GHz gives 16,896 × 2 × 1.98×10⁹ ≈ 67 TFLOPS, which is precisely the datasheet number from the table below. "CUDA core", "FP32 unit", and "lane" are three names for this same machine. What that word "clock" is doing in the middle of that sentence, and why it is allowed to set the arithmetic speed of a $30,000 chip, is the whole of the next section.

Here are the H100 numbers we will keep reaching for. You do not need to memorize them; you need to remember that every one of them is a hard, physical budget.5 5One detail for the curious: the 50 MB L2 is physically built as two halves joined by a crossbar, but every SM can reach all of it, so you may safely picture it as one shared pool sitting between the SMs and the HBM stacks.

3The clock: why arithmetic has a speed limit

We have now said "one operation per clock" twice without saying what a clock is. It is worth stopping on, because that single number is what converts a pile of silicon into a TFLOPS figure, and because, exactly like the memory story in section 6, it bottoms out in physics rather than engineering taste.

Start with the problem the clock exists to solve. Hand one lane two numbers to multiply. Inside that lane is a web of transistors wired so that, given the input voltages, the output voltages eventually arrange themselves into the right answer. The load-bearing word is eventually. The output does not appear the instant the inputs arrive; it goes through a period of nonsense first, while the signals work their way through, and only then settles on the answer. So the machine needs a rule for when it is safe to look. That rule is the clock: a square wave, ticking 1.98 billion times a second on an H100, broadcast to the entire chip. On each tick, every storage element grabs whatever value its input has settled to, and a new operation begins. Nobody looks in between.

So why does settling take time at all? Here we meet the same physics as the DRAM cell, in smaller form. A transistor is a switch, and what it switches is charge flowing onto a wire. Every wire on the chip has capacitance, because it is a conductor sitting near other conductors, so it behaves like a very small version of the bucket from section 6. Driving that wire from 0 to 1 means pushing charge into it through the transistor's own resistance, and voltage on a capacitor being charged through a resistance does not jump, it rises along a curve. The gate downstream sees nothing until that rising voltage crosses its threshold. That delay, a few tens of picoseconds, is the price of one logic gate.6 6Formally the wire charges as V(t) = V_dd(1 − e^−t/RC), the same exponential that governs the DRAM bucket. Smaller transistors mean smaller R and C, which is the real reason shrinking the process makes chips faster.

One multiply is not one gate, though. It is a chain of them, dozens deep, and the answer is only trustworthy once the slowest route through that chain has settled. Chip designers call that route the critical path, and it sets the rule the whole machine obeys: the clock period must be at least as long as the critical path. Run the tick faster than that and the machine does not run faster, it runs wrong.

Two sanity checks make the scale feel real. First, 1.98 GHz means a tick every 505 picoseconds, and in 505 ps light in vacuum travels about 15 centimetres. A signal in on-chip copper is slower than that, so within one tick a signal crosses a few centimetres at best, and the H100 die is about 2.9 cm on a side. The chip is, quite literally, as large as a signal can cross in about one beat of its own clock; this is why the clock is delivered as a carefully balanced tree and why the machine is built from 132 largely independent SMs rather than one enormous unit. Second, if a single gate costs tens of picoseconds, a 505 ps budget buys a chain only a few dozen gates deep, which is precisely why arithmetic units are pipelined: a multiply-add is chopped into a handful of shorter stages, each finishing comfortably within one tick. Like an assembly line, several operations are in flight at different stages, and one finishes every tick. That is what "one operation per clock" really means: not that a multiply takes 505 ps end to end, but that a completed multiply falls off the end of each lane on every tick.7 7An FP32 fused multiply-add on recent NVIDIA hardware takes roughly four clocks from input to output, but each lane accepts a fresh pair of operands every clock. Latency four, throughput one: the assembly line, exactly.

Now the TFLOPS number assembles itself, and every term in it is something we have already met.

That last box is the punchline, and it is the reason the GPU exists in the shape it does. The clock cannot be raised much: settling time is set by resistance and capacitance, and pushing the tick faster means raising the voltage, which raises power as the square, and an H100 is already a 700-watt space heater. Clock speeds have accordingly barely moved in fifteen years, hovering around 1 to 2 GHz for GPUs and 3 to 5 GHz for CPUs. If you cannot make each lane finish work more often, the only remaining way to do more arithmetic per second is to build more lanes and keep them all busy. That is precisely what an H100 is: 16,896 unhurried lanes rather than a handful of frantic ones. It is also, once again, the thesis of this whole post arriving from a new direction, the one Jones states outright: the programming model looks the way it does because the machine looks the way it does, and the machine looks the way it does because of physics.

4Carving the image: grids, blocks, threads

Now we connect the two worlds. CUDA asks you to describe your work as a hierarchy with exactly three levels, and the reason it has three levels is that the hardware does.

You cut the problem into blocks, and you decide how many threads each block contains. All the blocks together form the grid. For our dog we will pick blocks of 16 × 16 threads, so one block owns one 16 × 16 tile of pixels, and the grid is 64 × 48 = 3,072 blocks. Each thread does the whole job for exactly one pixel: read it, compute the gray value, write it back.

In code, the whole arrangement looks like this. Do not worry about the syntax; worry about the first two lines of the function body.

__global__ void rgb_to_gray(const uchar4* rgb, unsigned char* gray,
                            int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;  // which column am I?
    int y = blockIdx.y * blockDim.y + threadIdx.y;  // which row am I?
    if (x >= width || y >= height) return;           // guard the edges

    uchar4 p = rgb[y * width + x];                   // read my pixel
    gray[y * width + x] =
        (unsigned char)(0.299f*p.x + 0.587f*p.y + 0.114f*p.z);
}

dim3 threadsPerBlock(16, 16);   // 256 threads = 8 warps
dim3 numBlocks(64, 48);          // 1024/16 × 768/16 = 3,072 blocks
rgb_to_gray<<<numBlocks, threadsPerBlock>>>(d_rgb, d_gray, 1024, 768);

The kernel is written as if it were a single-threaded C++ program for one pixel. There is no loop over the image anywhere. The launch line at the bottom says "run this function 786,432 times", and the two index variables are how each of those 786,432 executions finds out which pixel it owns: the hardware fills in blockIdx and threadIdx for every thread automatically.8 8When the image size is not a clean multiple of 16 you launch numBlocks((width+15)/16, (height+15)/16), and the edge guard in the kernel sends the surplus threads home. This is why the guard is idiomatic even when it looks unnecessary. Stephen Jones makes a point of this in the talk, and it is worth repeating: this indexing line is the single most load-bearing line in all of CUDA. Everything after it is ordinary C++.

The three software levels now map one-to-one onto the three hardware levels from section 2, and this mapping is the entire CUDA programming model in one picture.

5How blocks land on SMs

At launch, our 3,072 blocks form a queue of work, and the hardware work distributor (NVIDIA calls it the GigaThread Engine, the thin bar from Figure 3) starts dealing them out to SMs. Three rules govern the deal, and each one shapes how you must write CUDA programs.

Rule one: a block never spans two SMs. Once block (36, 20) is placed, all 256 of its threads live on that one SM together until they finish. This is what makes a block more than a bookkeeping unit: threads within a block can share data through the SM's shared memory and can synchronize with each other, precisely because they are guaranteed to be alive on the same physical hardware at the same time. Jones calls the block CUDA's "quantum of parallelism", the largest group of threads you may assume run together.

Rule two: you get no promises about order or placement. Block (0, 0) may run first or last; blocks (5, 3) and (5, 4) will usually land on different SMs; nothing may depend on any of it. This sounds like a limitation, and it is, deliberately. Because blocks are independent, the same 3,072-block launch runs correctly on a laptop GPU with 30 SMs and on an H100 with 132: the distributor just deals the same deck onto however many SMs exist. Independence is the price of scalability, and CUDA makes you pay it up front.

Rule three: blocks are spread out wide, not packed SM by SM. Each SM has a limited set of wires into the memory system, so no single SM can come close to using the chip's 3.35 TB/s on its own. Spreading blocks across all 132 SMs puts as many independent memory streams in flight as possible.

How many blocks fit at once? An SM accepts blocks until some resource runs out. Our grayscale blocks have 256 threads each and use almost nothing else, so the binding limit is the SM's 2,048-thread ceiling: 8 blocks per SM. Across 132 SMs that is 1,056 blocks resident at the same time, which is 270,336 threads in flight, on one chip.9 9Section 11 returns to "some resource runs out" properly; it is the occupancy game, and it is the second-biggest performance lever after memory layout. Our grid of 3,072 blocks therefore runs as roughly three waves: the distributor fills every SM, blocks finish and drain, new blocks drop into the gaps, and after the third wave the dog is gray.

6What the memory system wants: a short physics detour

The next two sections keep using phrases like "128 consecutive bytes", "one memory transaction", and "the natural stride of the memory system", and none of them mean anything unless you know how a memory chip physically hands over bytes. Jones spends the first third of his talk on exactly this, and he is right to: it is the floor the whole building stands on. The good news is that it comes down to three small facts, all of them plain physics, none of them requiring a computer-science background. We will take them one at a time.

Fact one: a bit is a bucket of charge. One bit of your image lives in a DRAM cell, and a DRAM cell is a microscopic capacitor: a bucket that either holds electrons or does not. Full bucket means 1, empty bucket means 0. A tiny switch (a transistor) connects the bucket to a shared wire called the bit line, and at the end of that wire sits a sense amplifier, whose job is to listen. The charge in one bucket is so small that it cannot announce itself; the amplifier senses the faint nudge on the wire and boosts it into a clean, readable 0 or 1. Two consequences follow directly. First, reading drains the bucket: the act of measuring the charge destroys it, so whatever was read must be written back afterwards. Second, filling and draining buckets of electrons takes real, physical time. Every performance rule in this post traces back to these two sentences.

Fact two: the buckets live in huge grids, and you can only read a whole row at a time. A memory chip packs millions of these cells into a 2D matrix of rows and columns (that is where the name "random access memory" comes from: pick any row, then any column). Here is the crucial mechanical detail: the chip cannot read one cell by itself. To read anything at all, it must activate one entire row, draining every bucket in that row into the strip of sense amplifiers at the edge of the grid. The activated row is called the open page, and it is about 1 KB wide. Once a page is open, reading along it is fast and cheap: the data is already sitting in the amplifiers, which hold strong, clean voltages you can read again and again. But switching to a different row is expensive: first the current row's contents must be written back into their buckets (remember, reading drained them), and only then can the next row be drained in. A row switch costs roughly three times as much as a read along the open page.

Fact three: because of facts one and two, memory is served in trays. Since opening a row is the expensive step, the memory system never sells you a single byte; it always hands over a sizeable run of consecutive bytes from the open page in one go (the chip feature behind this is literally called burst mode). Think of a bakery whose oven only bakes full trays: you may ask for one bun, but a whole tray comes out regardless. On NVIDIA GPUs the tray is 128 bytes, and the street of addresses is pre-cut into fixed trays: bytes 0–127 are one tray, 128–255 the next, and so on. "Aligned chunks" means exactly this pre-cut: a request is served by whichever tray its address falls into. And "the natural stride of the memory system", the phrase that bothered you, means nothing more than the tray size: 128 bytes is the step at which this machine hands out data at full speed. Use all 128 bytes of every tray that arrives and you get the full 3.35 TB/s. Use 4 bytes of each tray and 97% of everything shipped is waste; that is the cliff in the measurement above.

Hold that number: one tray is 128 bytes, and one of our pixels is 4 bytes, so one tray holds exactly 32 pixels sitting side by side. The number 32 is about to stop being arbitrary.

7Warps: the unit the hardware actually runs

So far the story has been grids, blocks, and threads, which is the vocabulary you program in. The hardware, however, does not execute threads one at a time. Inside the SM, the 256 threads of each block are cut into groups of 32 called warps, and the warp is the unit the machine actually fetches instructions for, schedules, and runs. Threads are numbered within a block x-fastest, so in our 16 × 16 block, warp 0 is rows 0 and 1 of the tile (two runs of 16 pixels), warp 1 is rows 2 and 3, and so on: 8 warps per block, 24,576 warps for the whole dog.

Why 32, and why gang threads together at all? You already know the answer; it was the last line of the physics detour. The memory system serves 128-byte trays, and one of our pixels is 4 bytes. When the 32 threads of a warp sit on 32 pixels lying side by side in a row, their combined request is exactly 32 × 4 = 128 consecutive bytes: one tray, one memory transaction, nothing wasted and nothing extra shipped. The warp width is not an arbitrary batch size. It is the width at which "one thread per pixel" and "one full tray" become the same request. This matching of threads to trays is called coalescing, and it is the property Boehm's Kernel 2 is about.

This also explains a rule of thumb Jones states outright in the talk: give a block at least 128 threads, which is 4 warps. The reason is now easy to say in tray language. An SM has four quadrants, each with its own scheduler, and you want all four issuing loads at once; four warps requesting their trays together puts four full trays in flight simultaneously, which keeps the memory system fed at the pace it was built for. A block with fewer than four warps can leave quadrants sitting idle whenever it is the only block that has landed on the SM. Our 256-thread block clears the bar comfortably: eight warps, two for every quadrant.

8The four warp schedulers

We can finally answer a question that confuses almost everyone at first contact: the SM diagram shows four warp schedulers, so what are they for, and four of what, exactly?

Let us answer it with the dog on the machine, concretely. From section 5 we know each SM is holding 8 of our blocks, and each block is 8 warps, so the SM is hosting 64 warps. Those 64 warps are dealt out to the four quadrants round-robin as each block lands: warp 0 of a block goes to quadrant 0, warp 1 to quadrant 1, warp 2 to quadrant 2, warp 3 to quadrant 3, warp 4 back to quadrant 0, and so on, and a warp stays with its quadrant for its whole life. So each of our blocks contributes exactly 2 warps to each quadrant, and each warp scheduler ends up juggling 16 warps, drawn from all 8 resident blocks. One picture to correct before it settles: it is tempting to imagine "each quadrant owns two blocks", and the arithmetic even matches, but that is not the mapping. Blocks are never assigned to quadrants; warps are. A block's 8 warps are spread across all four.

Now separate two verbs that the phrase "the SM can process 2,048 threads" quietly mixes together, because they are different by a factor of sixteen. The SM hosts up to 2,048 threads: all of them alive, all of their registers allocated, all listed with the schedulers. But the SM executes at most 4 warps' instructions per clock, one per scheduler: 128 threads' worth. Every clock cycle, each scheduler looks across its 16 resident warps, asks which are eligible (next instruction known, operands sitting ready in registers, not waiting for anything), picks one, and issues that warp's next instruction to its 32 lanes. The other 15 warps of the quadrant simply remain as they are: fully formed, registers intact, going nowhere.

Where do the 15 waiting warps of a quadrant physically be, while they wait? This question sounds naive and is actually the crux, so here is the plain answer: nowhere special, because a warp is not a thing that moves. A resident warp is its 32 threads' registers, sitting in the quadrant's register file (the pockets of Figure 6, parked on their shelves), plus a program counter remembering which line each thread is on, plus one row in the scheduler's scoreboard. It is not stored in L1, it is not stored in shared memory, and nothing about it goes back to HBM. That is exactly why "switching warps" costs zero clocks: on a CPU, switching threads means saving one thread's registers out to memory and loading another's back in, thousands of cycles; here, both warps' registers are already on their shelves, and the scheduler merely points at a different shelf. Nothing is copied. Nothing moves.

One more confusion is worth dismantling while we are here, because the timeline figure below invites it: why would a warp that has computed go back to waiting at all? Once it is done, it is done, no? The answer is that the wait sits in the middle of each thread's work, not after it. Follow one warp of our kernel through its own five lines. First comes address math: a few clocks of integer arithmetic to compute x, y, and the memory address. The lanes do this immediately; nothing to wait for. Then comes the load: the warp asks for its 32 pixels, a request departs, and one 128-byte tray must now physically travel from the HBM chips, through L2 and L1, into the warp's registers. That journey takes several hundred clocks, and during it the warp cannot run its next line, because the R, G, B values that line needs have not arrived. The warp is stalled mid-recipe, a cook who has sent out for flour. So yes, something very much is being transferred during the yellow "waiting" stretch: the warp's own data is on the road. When the tray lands in its registers the warp becomes eligible again, a scheduler picks it, it does its three multiply-adds, writes the gray byte back, and only then is it actually finished.

Seen this way, "most threads are waiting most of the time" stops being alarming, because it was never the threads we needed to keep busy: it is the lanes. A waiting warp costs nothing. It occupies no lane, burns no cycle, and holds only shelf space it was going to hold anyway. The design goal is not zero waiting (physics has voted on that already); the goal is that at every clock, each scheduler can find some warp whose tray has arrived. While warp 3's pixels are on the road, warps 5, 9, and 14 take their turns issuing their own loads, so trays stream back continuously, and by the time the scheduler's attention returns to warp 3, its delivery is sitting in its registers.

And this finally answers "what can I even do as a programmer, if there are only four schedulers?" You have two levers, and you have already met both. First, keep every scheduler's shelf unit full: a scheduler choosing among 16 warps almost always finds an eligible one, while a scheduler with only 4 warps to pick from will hit clocks where nobody's data has arrived and the lanes sit idle. That is occupancy, and it is set by your block size and your kernel's resource appetite, which is section 11. Second, make every tray count, the coalescing rule of section 7: full trays mean fewer loads for the same pixels, so there is simply less waiting to hide. For our grayscale kernel, which does only five arithmetic operations per 4-byte load, the loads set the pace entirely, and these two levers are essentially the whole game.

9SIMT, or why threads feel independent

There is an apparent contradiction sitting in the earlier sections. Section 4 said each thread is written like an independent single-threaded C++ program with its own variables. Section 8 said the hardware issues one instruction to 32 threads at once. Both are true, and the name for the trick that reconciles them is SIMT: single instruction, multiple threads.

Before any formality: if your instinct says this is simple, one instruction such as gray = 0.299R + 0.587G + 0.114B, executed by 32 threads at the same moment, each on its own pixel, then your instinct is exactly right, and that really is the whole idea. Stated one notch more precisely: you write the scalar program from section 4, one pixel's worth of work, and every thread carries its own state, its own registers, its own threadIdx, its own program counter. You never write a vector operation. The vector exists only in the hardware: when a warp's 32 threads all sit at the same instruction, which is the overwhelmingly common case, the scheduler fetches that instruction once and runs it across all 32 lanes in the same clocks, each lane on its own registers.11 11Since the Volta generation (2017), each thread really does keep its own program counter, so diverged threads can make independent progress and even synchronize; before that, a warp had a single shared PC. The performance picture in the next section is unchanged.

The fine print is the word "when": the lanes execute together when the 32 threads agree on what instruction comes next. Our grayscale kernel agrees always, because every pixel takes the identical path through the code. The moment we write a branch that depends on the data, agreement can break, and we finally get to see SIMT's cost model.

10Warp divergence, measured on the dog

Let us give the GPU a second job, chosen to be maximally unkind. We brighten every pixel whose coordinates sum to an even number and darken every pixel whose coordinates sum to an odd one, a checkerboard at single-pixel scale:

if ((x + y) % 2 == 0) out[i] = brighten(p);   // half of every warp…
else                  out[i] = darken(p);     // …the other half

Think about what any one warp now has to do. Its 32 threads sit on 32 neighbouring pixels, so by construction 16 of them want the brighten path and 16 want the darken path. However, the warp has one instruction stream. The hardware resolves this the only way it can: it runs both paths, one after the other. First it executes the brighten instructions with the 16 odd-pixel lanes switched off, then the darken instructions with the other 16 lanes switched off. The lanes that sit out an instruction are masked: they burn the clock cycle and produce nothing. This is warp divergence.

Note what warp divergence is not. It is not an error, the result is still correct; and it is not global, the cost is confined to the diverged region of the code, and the warp reconverges when both sides rejoin. It also has nothing to do with other warps: each warp diverges or not on its own. The checkerboard is the worst case precisely because its condition, threadIdx-parity, is engineered to split every single warp down the middle.

Most real branches are not like that, and this is the part the tutorials rarely show. Let us run a job you have used this week: portrait mode. Keep the dog sharp, blur the background. A small segmentation mask tells each thread which kind of pixel it owns:

if (isDog[i]) out[i] = p;              // keep the puppy sharp
else          out[i] = blurred[i];     // soften the background

The condition still varies per pixel, so some warps must diverge. But which ones? A warp covers 32 neighbouring pixels, and isDog is spatially coherent: deep inside the puppy all 32 lanes say "dog" and take one path together, out in the background all 32 say "background" and take the other path together. Neither case costs anything. The only warps that run both paths are the ones whose 32 pixels straddle the silhouette. We can count them, because we have the actual mask:

This is the picture to keep. Divergence is charged per warp, so its price is set by how your branch condition varies across 32-pixel neighbourhoods, not by how often it is true. The checkerboard flips at pixel scale and poisons every warp. The silhouette varies at object scale and touches 2.7% of them, and even those pay at most double, so the whole feature costs the branchy part of this kernel under 3% extra. And a condition that is constant across each warp, say if (y < 384) to process only the top half of the image, diverges nowhere at all: every warp's 32 threads agree, and each warp simply takes its one path.12 12For very short branch bodies the compiler often skips branching entirely and issues predicated instructions, both sides executed with lane masks. The cost model stays the same: you pay for both sides when a warp is mixed.

11So how do you pick a block size?

We chose 16 × 16 back in section 4 without justification. Time to pay that debt, because this is the question everyone actually has: the tutorials all say "we launch blocks of 256 threads" as if it were obvious, and it is not. The choice is constrained from five directions, and once you can name them, the standard advice stops being folklore.

First, the hard limits. A block may have at most 1,024 threads (whatever its shape), and an SM hosts at most 2,048 resident threads, 64 warps, and 32 resident blocks. These are architectural constants you cannot negotiate with. There is, by the way, no matching hard minimum: a block of a single thread is perfectly legal, it is just a bad idea, because the hardware will still dedicate a whole 32-lane warp to it. The practical floor is one warp, and the comfortable floor, from section 7, is four.

Second, warps are the allocation unit. The hardware runs whole warps, so a block of 250 threads is silently rounded up to 8 warps of hardware with 6 lanes permanently dead. A block of 50 threads occupies 2 warps and wastes 14 of their 64 lanes, a 22% tax paid on every single cycle. This is why block sizes are always multiples of 32. Not convention: arithmetic.

Third, blocks below 128 threads starve the SM's front end. The SM has four warp schedulers, and it wants at least four warps in flight from the moment a block lands, with their combined memory requests wide enough to fill the 128-byte trays of section 6. Jones is blunt about this in the talk: never fewer than 128 threads per block. Tiny blocks also collide with the 32-resident-block cap: blocks of 32 threads max out at 32 × 32 = 1,024 threads, half the SM's capacity, before any other limit even comes into play.

Fourth, big blocks are coarse Tetris pieces. Resident blocks must fit whole: the SM packs blocks until threads, registers, or shared memory runs out, and whichever runs out first wins. Our grayscale block asks for 256 threads, a few registers per thread, and no shared memory, so 8 of them fill the 2,048-thread budget exactly. Now imagine 1,024-thread blocks: only 2 fit, and if the kernel's register appetite leaves room for just 1.9 of them, one block is all you get, and occupancy falls off a cliff. Smaller pieces pack tighter and degrade gracefully. This packing ratio, resident warps as a fraction of the 64 possible, is called occupancy, and after memory access patterns it is the biggest lever you have: it decides how much latency-hiding budget the schedulers of section 8 get to work with.

Two of the three bars in that figure have gone unexplained until now: we have played the whole game with threads and waved at "registers and shared memory". Time to fix that, because "register appetite" and "shared-memory appetite" sound like personality traits and are actually plain numbers you can read off your build.

Both of the kernels we are about to weigh are blurs, so let us pin down what a blur actually does before using it to explain anything. Grayscale was a per-pixel job: one thread, one pixel, done. A blur is the first job in this post that is not. To blur, a thread replaces its pixel with the average of its neighbourhood, so with a radius of 2 it must read the 5 × 5 = 25 pixels around itself, add them up, and divide. That is the entire algorithm, and here it is running on the pixel we have followed since section 4.

Registers per thread: the compiler decides, at build time. When your kernel is compiled, the compiler counts how many values each thread must keep alive at once, exactly the way we counted the grayscale thread's pocket in Figure 6, and assigns that many registers to every thread of the kernel (there is a cap of 255 per thread). Simple straight-line code stays lean. You never set the number directly, but you can read it (nvcc --ptxas-options=-v prints it), and the Tetris charges you for it: a block's slice of the register file is 256 threads × N. The cliff is startlingly sharp. At 32 registers per thread a block costs 8,192, and exactly 8 blocks fill the 65,536-register file alongside the full 2,048 threads. At 33 registers per thread, one more, a block costs 8,448, only 7 fit, and the SM silently loses a whole block's worth of latency-hiding warps.

So what makes a thread need many registers at once? This is where unrolling comes in, and it is worth spelling out because the word gets used as if everyone knows it. Our blur thread has 25 pixels to fetch. Written the natural way, that is a loop, and a loop does one trip at a time: fetch a neighbour, wait for it to arrive, add it to the running total, go back and do the next one. The thread only ever holds a couple of values, so it is cheap in registers, but it also spends 25 waits end to end, one per trip, and section 8 taught us what a wait costs. Unrolling means the compiler rewrites those 25 trips as 25 separate lines of code with no loop at all. Now it can issue all 25 requests back to back without waiting in between, so the 25 journeys overlap and the thread waits roughly once instead of 25 times. That is a real speedup, and it has a price: every request that is in flight needs its own register to land in when it arrives. Twenty-five values in the air means twenty-five landing spots, plus the weights and partial sums, and the register count climbs from a handful to dozens.

Shared memory per block: you decide, explicitly. Shared memory is the block-owned scratchpad carved out of the SM's 256 KB of fast SRAM, and a kernel uses exactly as much as its __shared__ declarations ask for. Our grayscale kernel declares none, which is why its bar sat at zero. The reason a blur wants some is the second fact from Figure 26: neighbouring threads read almost the same pixels, so if every thread fetches its own window from HBM, the block drags each pixel across the bus dozens of times over. The fix is to fetch the block's whole neighbourhood into shared memory once, then let all 256 threads read from that scratchpad. And here the apron appears: to write a tile of output, a block must read a slightly larger tile of input, because the threads at the tile's edge need neighbours that lie outside it. The apron is that surrounding ring, and its width is the blur radius.

So let us make the Tetris concrete with three kernels for the dog that you could genuinely write, all launched with the same 256-thread blocks:

• Kernel A, our grayscale: one load per thread, about 15 registers, no shared memory. Threads run out first: 8 blocks, 64 warps, 100% occupancy.
• Kernel B, the radius-2 blur, fully unrolled: 25 loads per thread, all in flight at once, so the compiler lands on 64 registers per thread. One block now costs 256 × 64 = 16,384 registers, and 65,536 ÷ 16,384 = 4. The register file is the wall: 4 blocks, 1,024 threads, 32 warps, 50% occupancy, while half the thread budget sits unused.
• Kernel C, the radius-16 blur, tiled through shared memory: registers stay modest because the loop is left rolled, but each block caches the 48 KB apron of Figure 28. Four blocks ask for 192 KB of the 228 available; a fifth would need 240. Shared memory is the wall: again 4 blocks, again 50% occupancy, with threads and registers both half idle.

And here is why this Tetris is worth playing rather than merely admiring: the walls can be moved. Take kernel C. Its 48 KB came from a choice, not a law: we chose to cache each pixel as three floats, 12 bytes, because floats are convenient to average. Cache the tile in its original 4-byte form instead and convert to float on read, and the same 4,096 pixels take 4,096 × 4 = 16 KB. Now eight blocks want 128 KB, comfortably under 228, so shared memory stops being the wall, the thread ceiling binds again, 8 blocks land, and the schedulers get their 64 warps back. Not one arithmetic operation per pixel changed; the block simply got cheaper to host, so twice as many warps are available to hide the memory waits. This is the move Jones works through live in the talk (his kernel drops from 48 to 32 KB per block, one more block fits, and throughput rises by a third), and it is why he calls occupancy the second biggest lever after memory access patterns.

Fifth, shape follows memory. For row-major image data you want each warp's 32 threads to sit along a row, so the x-dimension of the block should cover whole warps. Our 16 × 16 choice makes each warp a pair of 16-pixel row segments (two 64-byte transactions); 32 × 8 would make each warp one full 32-pixel run (one 128-byte transaction). Both coalesce, and on a kernel this simple they benchmark within noise of each other, but 32 × 8 is the shape a CUDA veteran reaches for first, precisely because it makes warp and row line up perfectly.

Put together, the constraints squeeze from both ends: hard caps and coarse packing push down from 1,024, warp arithmetic and scheduler appetite push up from 32, and what survives is the familiar recipe. Use a multiple of 32, at least 128, usually 128 or 256 threads; shape it so the x-dimension spans whole warps; then stop guessing and measure.13 13The measuring tools: Nsight Compute reports achieved occupancy per kernel, and cudaOccupancyMaxPotentialBlockSize() computes a good block size for a kernel's actual register and shared-memory appetite at runtime. Block size is a knob with maybe 2× of performance in it, set by five mechanical constraints, and every one of those constraints is something you have now seen in the silicon: warps because of memory transactions, 128 because of four schedulers, packing because blocks are indivisible residents of one SM.

12What we left out, and where to go

The dog is gray, blurred, and once, regrettably, checkerboarded. Here is the mental model the whole post compresses into, five sentences long:

• The GPU is 132 copies of one machine, the SM, and the SM is four schedulers, each steering 32 lanes and a pool of waiting warps.
• Arithmetic speed is lanes × 2 × clock rate, and the clock is pinned near 2 GHz by settling time and heat, which is why the machine is wide rather than fast.
• You describe work as a static grid of independent blocks; the hardware deals blocks onto SMs in waves, in no promised order.
• Within a block, threads run as warps of 32; a warp reading 32 neighbouring values is one memory transaction, and that identity is the whole reason the number 32 exists.
• SMs hold 16× more threads than lanes so schedulers can hide memory latency by always finding an eligible warp; your block size decides how full that pool gets.
• A warp with a disagreeing branch runs both sides; the price of an if is the spatial frequency of its condition, not its truth rate.

What did we skip? Deliberately, a lot: how the image gets to the GPU in the first place (cudaMemcpy, and why the PCIe trip often costs more than the kernel), shared memory and __syncthreads(), which turn a block from 256 loners into a cooperating team, the L1/L2 cache hierarchy, tensor cores, and streams, the mechanism for overlapping copies with compute. Each of these earns its keep only after the model above feels boring.

Where to go, in order:

1. Watch "How CUDA Programming Works" (Stephen Jones, GTC 2022), ideally today, while this post is fresh. You will recognize every picture, and he goes further than we did on memory physics, occupancy Tetris, and concurrency. His 2021 talk "How GPU Computing Works" is the natural double feature.
2. Write the grayscale kernel yourself. Then write the blur: each thread reads a small window around its pixel and writes one value. You will feel the difference between the window (slides) and the block (does not), and you will meet the edge guard in earnest.
3. Read Boehm's matmul worklog, now. His Kernel 1 is our dog with matrices: one thread per output element, indexed by the same two lines. His Kernel 2 is sections 6 and 7 of this post applied ruthlessly. His Kernel 3 introduces shared memory, the first thing we skipped. The cliff should now be a staircase.
4. Keep the CUDA C++ Programming Guide open as the reference for every constant this post quoted.

There is now a part two. How a Matmul Kernel Gets Fast takes everything above and puts it to work, rewriting one matrix-multiply kernel ten times, from 1.3% of the hardware to 93.7%, with every byte and every multiply accounted for. Coalescing, shared-memory tiling, the outer-product trick that lets one thread do 64 multiply-adds from 16 loads, occupancy, vectorized loads, autotuning and warptiling. It is the natural next step from here, and it assumes exactly this post and nothing more.

Jones ends his talk by saying that CUDA is the way it is because of the laws of physics. After an hour with one dog and one die, the claim should feel less like a slogan: every rule we met, 32-wide warps, 128-thread blocks, independent blocks, oversubscribed SMs, traced back to the cost of moving bytes through real capacitors and real wires. The programming model is not an API someone invented. It is the shape of the machine, with names.