What is a Warp?

A warp is a group of 32 threads that the GPU picks up, schedules, and runs as one unit. All 32 lanes share a single instruction stream, and each lane keeps its own registers and its own result.

So: one instruction, 32 pieces of data.

The easiest way to picture it is a squad of 32 workers standing in a line. The leader shouts one order and all 32 do it at the same time. Nobody gets a private instruction. If worker 7 needs to do something different, worker 7 sits out while the rest finish, and then the rest sit out while worker 7 finishes. That waiting is the price of this design, and it is the reason warp behavior decides how fast your kernel runs.

You will not find warps anywhere in your C++ code. You write plain scalar code, and the compiler plus the hardware decide the rest. The CUDA Programming Guide compares warp behavior to a CPU cache line: ignore it and your program is still correct, ignore it while tuning and you leave a lot of performance on the table.

The words you need

• A thread is the smallest unit you program.
• A lane is a thread's position inside its warp, numbered 0 to 31.
• A warp is 32 lanes that share one instruction stream.
• A thread block is a group of warps that runs on one SM.

Your block size decides the split. A 256-thread block always becomes 8 warps, because warps are built from consecutive thread IDs, and thread 0 always lands in warp 0.

This whole scheme has a name: SIMT (Single Instruction, Multiple Thread). It looks like SIMD from the outside, but you never write the vector width, which is why the same CUDA code runs on every GPU generation without edits.

Why 32?

Because 32 is the number that makes one lane's 4-byte float line up with the rest of the machine:

Hardware piece	Size	What 32 lanes do to it
Memory sector	32 bytes	8 lanes fill one sector
Cache line	128 bytes	32 lanes fill exactly one line
Shared memory banks	32 banks of 4 bytes	one lane per bank, no conflict
Registers	32 bits	one register slot per lane

None of that is a coincidence. When a warp reads 32 floats that sit next to each other in memory, the hardware moves all of them in one trip.

The cost of the alignment shows up when a block does not divide evenly. A 100-thread block still gets 4 warps, and the last one runs with 4 active lanes and 28 doing nothing the entire kernel. Round your block size up to a multiple of 32 unless you have measured a reason not to.

How the GPU hides waiting

Memory is slow. A load from global memory can take hundreds of cycles, and the warp that asked for it has nothing to do until the data arrives.

So the warp scheduler does not wait. It sets that warp aside and issues an instruction from a different warp this cycle. Every resident warp keeps its state on the chip, so jumping between them costs nothing at all. That is how a GPU stays busy while most of its warps are waiting on memory.

More resident warps usually means more cover for that waiting, which is what occupancy measures. It is not a score to maximize. Kernels full of independent work can hide latency with few warps, and big GEMM kernels often run great at single-digit occupancy. Squeezing in more warps can cost you per-warp registers and make things slower.

When the lanes disagree

Divergence is what happens when one instruction has to serve lanes that want different things. Take a plain if:

if (data[idx] > 0.5f) {
    data[idx] = data[idx] * 4.0f;   // some lanes want this
} else {
    data[idx] = data[idx] + 2.0f;   // the others want that
}

The warp cannot do both at once. It runs one path with the lanes that want it, then the other path with the remaining lanes, and then everyone meets again. The inactive lanes are masked off: they are not running, but their registers are held for them.

Both paths still cost time. If a warp splits in two, it pays for both sides, and while it runs the first side the other lanes sit out. Divergence never breaks correctness, it just burns cycles.

What actually helps:

• Sort or bucket your data so a warp sees one case.
• Write branchless code (x = pred ? a : b;) and let the compiler use predication.
• Push rare cases out of the hot loop.
• Measure the cost with the warp state statistics in Nsight Compute instead of guessing.

Volta removed the lockstep promise

Old GPUs gave every warp one program counter and one active mask. Lanes ran in lockstep, and clever code leaned on that: a reduction with no synchronization at all was fine, because every lane was guaranteed to be at the same line as its neighbors.

Compute capability 7.0 (Volta, 2017) changed it. Each lane now has its own program counter and call stack, and lanes can drift apart and rejoin at a finer grain than before. That flexibility is good, and it also means that old assumption is now a bug.

The fix is one line: __syncwarp() where your code needs the lanes to be in the same place. Use the _sync variants of the warp intrinsics too, and pass a mask that names exactly the lanes taking part. Getting the mask wrong is not a performance issue, it is undefined behavior.

Warps are memory units too

Every lane can read whatever address it likes. The hardware just prefers that neighbors read neighbors.

Thirty-two floats side by side is 128 contiguous bytes, one trip, every byte useful. Spread the same lanes 16 bytes apart and the request spans 16 sectors, and only a quarter of the bytes that arrive are ever used. On a T4 that pattern drops throughput from 206 GB/s to about 15 GB/s.

That is why layout beats micro-tuning so often. A struct-of-arrays where neighboring lanes touch neighboring fields usually buys more than any instruction-level trick. Shared memory plays by the same rules, with bank conflicts as the failure mode.

Talking to your warp

Once you accept that 32 lanes move together, they become the cheapest place to cooperate, because nothing else in the memory system is involved.

// Add up one value per lane. All 32 lanes must call this.
__inline__ __device__ float warp_sum(float v) {
    #pragma unroll
    for (int offset = 16; offset > 0; offset >>= 1) {
        v += __shfl_down_sync(0xFFFFFFFFu, v, offset);
    }
    return v;  // lane 0 now holds the total
}

Five steps, because 32 is 2 to the 5th. No shared memory, no barrier, no trip to DRAM. The other functions you will reach for are __ballot_sync (which lanes said yes), __any_sync and __all_sync (did any or all say yes), and __shfl_xor_sync (swap with a partner lane). Cooperative Groups wraps the same idea in a type: tiled_partition<32>() hands you one warp as a group object.

Warps in a real kernel

My tiled matrix multiplication walkthrough is where all of this lands. The cooperative load works only because neighboring lanes read neighboring addresses. The compute phase is a dot product per lane. In production GEMM, the warp is the unit that owns Tensor Core fragments: one mma instruction writes 128 accumulator values across 32 lanes, four registers each.

On Hopper and Blackwell, kernels increasingly hand different jobs to different warps, with some warps feeding data and others doing the math. The warpgroup is that idea formalized: four warps, 128 threads, acting as one bigger unit.

What to remember

• A warp is 32 lanes, one instruction stream, one active mask, 32 private register banks.
• 32 lines up with sectors, cache lines, shared memory banks, and register slots, which is why one warp read of 32 floats is so cheap.
• Divergence costs speed, not correctness. Make the common path uniform and let predication handle small branches.
• After Volta, lockstep is not promised. Reach for __syncwarp() when your code needs the lanes together.
• The cheapest coordination in CUDA lives inside a warp. Start reductions with shuffles before you reach for shared memory.

Keep reading

• Warps and SIMT and the SIMT Execution Model, CUDA Programming Guide
• Independent Thread Scheduling, CUDA Programming Guide
• Warp Shuffle Functions and Using CUDA Warp-Level Primitives
• Nsight Compute Profiling Guide, for sectors and warp state statistics
• NVIDIA Tesla: A Unified Graphics and Computing Architecture, where the name comes from