A warp is the unit of execution, and every kernel bug starts there

TL;DR: The hardware issues each instruction to a warp of lanes, so a kernel's cost is the number of warp instructions issued plus the stalls between them. Thread count is the wrong unit. Three habits follow: count the tail warp, price a divergent branch as the sum of its paths rather than the longer one, and never place a block barrier where some lanes might not arrive.

Where you are. First lesson of the kernels course. Foundations taught you to name the binding resource before proposing a fix. This module moves that habit inside a single kernel: which unit the machine schedules, what that unit costs, and where the code you wrote per thread stops matching the instructions the machine runs.

You write threads, the machine runs warps

A kernel is written as if one thread executes it. The compiler and the hardware do something different: they group lanes into warps, and each instruction is issued once per warp with a mask that says which lanes take effect. On NVIDIA hardware a warp is 32 lanes; the box at the end gives the current figures for other vendors. Everything in this lesson follows from that one fact.

Two consequences arrive immediately. A warp with one active lane occupies the same issue slot as a warp with all 32, so idle lanes are paid for. And a warp cannot issue its next instruction until the previous one's dependencies resolve, so when a warp waits on memory, the scheduler needs a different warp with something ready. The second consequence is the subject of the next lesson. This one is about the first.

Follow one warp through the simplest kernel there is:

__global__ void scale(float* y, const float* x, float a, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) y[i] = a * x[i];
}

Per warp, the machine issues a handful of integer instructions to form i, one compare, then a predicated load, a multiply and a store. Call it six warp instructions for 32 elements. That count, not the thread count, is what the SM spends. When you read a kernel from now on, read each line as "what does this do for 32 lanes at once", and the cost model falls out.

Count the tail

The grid rarely divides the work evenly. Take a hypothetical launch: n = 1,000,000 elements, 256 threads per block.

blocks          = ceil(1,000,000 / 256) = 3,907
threads launched = 3,907 × 256 = 1,000,192
warps launched   = 3,907 × 8 = 31,256
warps with work  = ceil(1,000,000 / 32) = 31,250
warps with none  = 6
waste            = 6 / 31,256 = 0.02%
sanity: 192 surplus threads is 6 warps exactly, so the arithmetic closes

Those six warps still issue the index arithmetic and the compare before every lane predicates off. At this size it is noise. Now change one number: 100 threads per block instead of 256.

warps per block  = ceil(100 / 32) = 4      (the fourth has 4 active lanes)
lane slots/block = 4 × 32 = 128
idle lane slots  = 128 − 100 = 28 per block, on every block
waste            = 28 / 128 = 21.9%
sanity: 96 threads would need 3 warps and waste nothing

The tail that matters is not the end of the grid. It is a block size that is not a multiple of the warp width, because that waste repeats in every block rather than once. The rule is short: block sizes are multiples of 32, always, and the reason is the mask.

Price a branch as the sum of its paths

Divergence is the second place the per-thread view lies. Suppose the body of a kernel is:

if (x[i] > 0.0f) { /* path A: 10 instructions */ }
else             { /* path B:  6 instructions */ }

If every lane in a warp takes the same path, the warp issues 10 or 6 instructions. If lanes disagree, the warp issues path A with one mask, then path B with the complementary mask: 16 instructions plus the bookkeeping to reconverge. The cost is bounded below by the longer path and above by their sum, and it never drops below the longer path however few lanes take it.

The question to ask of any condition is whether it is warp-uniform: do the 32 lanes of a warp always agree? Some conditions are uniform by construction and cost nothing.

The third row is the one people write without noticing. The fix is a remap: give each lane two consecutive elements and handle the even and odd cases with arithmetic, or split the work into two launches whose conditions are uniform. The fourth row is harder because the data decides. Sorting or partitioning the input so that warps see homogeneous runs is sometimes worth a pass of its own; predication of a short path is often cheaper than avoiding it.

The profiler reports this directly as the average number of active lanes per executed instruction. A kernel at 16 out of 32 is issuing half its instructions to masked lanes, and no amount of occupancy tuning will get that half back.

The barrier rule, and the bug that hides

The third habit concerns synchronisation. A block barrier must be reached by every thread in the block, or the behaviour is undefined. The programming guide permits it inside a conditional only when the condition evaluates identically across the whole block. The bug that violates it looks like careful code:

// wrong: lanes past n return before the barrier
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i >= n) return;
tile[threadIdx.x] = x[i];
__syncthreads();
// ... reduce tile ...

In the tail block, some lanes return and never arrive at the barrier. On one driver this hangs, on another it proceeds with a partial tile, and the kernel that passed every test at sizes that divided evenly fails at the first size that did not. The fix keeps every thread alive and guards the work rather than the barrier:

// right: every thread reaches the barrier; the work is predicated
int i = blockIdx.x * blockDim.x + threadIdx.x;
tile[threadIdx.x] = (i < n) ? x[i] : 0.0f;
__syncthreads();

The related bug lives inside the warp. Older code assumed that the 32 lanes of a warp execute in lockstep, and used that assumption to skip synchronisation in the last steps of a reduction. Current hardware schedules lanes independently between explicit synchronisation points, so that code is wrong even though it may still pass. The rule is to use the explicit warp-level synchronisation primitives and the synchronising variants of the shuffle intrinsics, and to treat any inter-lane communication without them as a race. This class of bug is the reason the lesson title says every kernel bug starts at the warp: the per-thread view has no place to put a lockstep assumption, so nobody notices they made one.

What to do with this in the room

An interviewer who hands you a kernel wants to hear the warp view before the fix. Read the block size and say whether it is a multiple of 32. Read each condition and say whether it is warp-uniform. Find every barrier and say whether every thread reaches it. Only then talk about memory. A candidate who starts with coalescing on a kernel whose block size is 100 has optimised the wrong 22 percent.

The decision is small and firm: block sizes are multiples of 32, conditions are made warp-uniform where the remap is cheap, and barriers are unconditional. The condition that reverses the second rule is a data-dependent branch whose short path is a few instructions; there, predication costs less than any remap, and you leave it.

Do this before moving on

Take the scale kernel above and write, for each line, the number of warp instructions it issues and whether its condition is warp-uniform. Then redo the tail arithmetic for n = 1,000,000 at 256 threads per block, and again at 100 threads per block. You should reach 1,000,192 launched threads, 31,256 warps, 6 warps with no work and 0.02 percent waste for the first, and 21.9 percent idle lane slots in every block for the second. Finally, rewrite the i % 2 condition as a remap where each lane handles elements 2i and 2i + 1, and say why the branch disappeared.

Go deeper

• GPU Execution Model takes the same hierarchy from the hardware side, with the residency bound worked and the independent-scheduling caveat spelled out.
• CUDA Programming Model has the launch and synchronisation rules that the barrier example above depends on, including what a stream orders.
• Occupancy and Register Pressure is the next question once the warps are counted: how many the SM can hold, and the calculator that says which resource decides.
• Warp divergence and why it costs you is the branch-pricing argument as an interview answer, with the follow-ups that usually come after it.
• Write a CUDA vector add and explain the launch is the grid arithmetic from this lesson delivered as a screening question, copies included.
• Memory-Bound vs Compute-Bound Kernels is where the warp count meets the roofline, and it decides whether any of this lesson's savings were the binding term.

Key takeaways

• The SM issues instructions per warp; a kernel costs warp instructions and stalls, not threads.
• A block size that is not a multiple of 32 wastes lane slots in every block, not just at the grid tail.
• A divergent branch costs the sum of its paths; ask of every condition whether the 32 lanes agree.
• Guard the work, never the barrier, and use explicit warp synchronisation for any inter-lane communication.

Check yourself

Answer before you look. Recalling it is what makes it stick; recognising it does not.

1. 1A kernel launches blocks of 200 threads. What fraction of lane slots is idle in every block, before any tail effect at the end of the grid?

   About 11 percent: 7 warps hold 224 lane slots for 200 threadsZero; 200 threads all do workAbout 22 percent, the same as a block of 100It depends on the grid size

2. 2A branch has a 10-instruction path and a 6-instruction path. Within one warp, 31 lanes take the short path and one lane takes the long path. How many instructions does the warp issue for the branch body?

   10, the longer pathAbout 6.1, the lane-weighted average16, both paths with complementary masks6, because almost every lane took the short path

3. 3Why is an early return before a block barrier a bug even when the kernel produces correct output in testing?

   Answer it, then reveal