The line add<<<blocks, threads>>>() looks like a function call with some weird symbols bolted on. It isn't. Understanding that one difference is the gap between thinking you write GPU code and actually writing it.
When you call a normal function, you wake up one path of logic and your processor walks down it, top to bottom, alone. When you launch a kernel, you are not calling anything. You are stamping out thousands of copies of the same instructions and handing one to every worker at once.
↳ So the right way to picture a kernel is not as a function but as a mold. Let's look at how this changes the code.
The code inside the mold is the instruction manual for a single thread. Nowhere in that body do you say how many times it runs. The count is decided entirely by the strange syntax outside, the triple angle brackets.
Write add<<<1, 1>>>() and it runs once. Write add<<<2, 64>>>() and the exact same code runs 128 times. Write add<<<4096, 512>>>() and it runs two million times, virtually all at once. One line, one mold, and the hardware pours millions of threads into it.
Now the objection that should immediately bother you. If all those millions of threads run the identical code, isn't that pointless? Wouldn't every thread read the same data, do the same math, and write the same answer—a giant pile of redundant work?
It would, if the body were empty. But it isn't, because every thread knows who it is. You never write a fixed data address in the code. Instead you write a formula, and each thread plugs its own identity into that formula to land on its own slice of data.
↳ Its position is its thread number inside its block, plus its block number times the size of a block. So the same single line of code sends every thread to a different address.
Walk it through concretely. You launch two blocks of 64 threads, so 128 threads total covering a 128-element array:
5 + (1 * 64) = 69, so it works on element 69.5 + (0 * 64) = 5, so it works on element 5.Same code, same formula, completely different targets. No overlap, no redundancy. The identity math is what turns one blueprint into millions of non-colliding workers.
Now the question that reveals the real engineering. Your hardware is a fixed size. Say you have an old chip with only two of those processing units. And you launch four thousand blocks. A normal program would need a loop to grind through that many. How does the GPU run four thousand blocks on two units without a loop and without crashing?
↳ The answer is that the GPU never looks at the whole grid at once. It only ever looks at one block at a time.
You launch four thousand blocks and they all go into a queue. The two processing units each grab one block, run all its threads to completion, and the instant a unit finishes its block, it reaches into the queue and pulls the next waiting one. Block 0 and block 1 run first, then block 2 and block 3, and so on, the hardware chewing through the queue at its own pace. You never wrote a loop. The hardware serialized the work itself, according to how many units it physically has.
And this is the quiet genius of the block. It is a booking ticket. The hardware looks at a block, checks if it has room, admits it if so, queues it if not. Which means the same code you wrote runs on that two-unit chip slowly but correctly, and five years later runs on a hundred-unit chip blazingly fast, grabbing a hundred blocks at once, without you changing a single line. The block is what decouples your code from the size of the machine.
Now the cost that almost everyone trips on, and it ties straight back to how the hardware actually executes. Remember that the real unit of execution is the warp, a fixed bundle of 32 threads that move together. The hardware cannot mix threads from different blocks into one warp, and it always reserves resources for a full warp of 32 even if you use fewer.
Watch what happens when you launch tiny blocks. Say two threads per block. The hardware still reserves a whole warp of 32 lanes for them, runs the fetch and decode for all 32, burns the power for all 32, and then only 2 lanes do any useful work while the other 30 sit switched off.
↳ That is roughly 6 percent efficiency. You paid for a full warp and used a sixteenth of it.
And add<<<1, 1>>>(), a single thread, is the worst case of all. One slow GPU lane running alone, no neighbors to hide memory delay, every other unit idle, plus the cost of shipping data across to the GPU and back. You took a machine built to run a million things at once and asked it to run one. It loses badly to a single CPU core, which runs at a much higher clock and is actually built for solo work.
This is exactly why real GPU code never uses block sizes like 2 or 37. A block of 37 fills one warp of 32 and then spills 5 threads into a second warp, wasting most of that second warp's 32 lanes.
↳ Block sizes are always multiples of 32, and usually 128, 256, or 512, so the warps pack full and no lanes go idle. The magic number 32 from the hardware reaches all the way up into how you choose to launch your code.
↳ The cleanest way to remember the warp penalty:
Blocks must pack in 32s. 37 threads = 64 slots paid for.
So the whole picture. A kernel is a mold, not a function. The angle brackets decide how many workers pour into it, and the body is just one worker's instructions. Redundancy is avoided because every worker computes its own identity and lands on its own data.
The work scales across any size of chip because blocks queue up and get pulled as units finish, no loop, no recompile. And the price of all this is that you must respect the warp, launching threads in full bundles of 32 or watching the hardware idle on your behalf.
Write a kernel thinking it's a function and the hardware confuses you.
Write it knowing it's a blueprint stamped across an army, and it finally makes sense.