What Is Warp Divergence? The Hidden GPU Throughput Killer

What is warp divergence? It's what happens when the 32 threads inside a CUDA warp disagree about which way to go at a branch. The GPU doesn't run both paths in parallel: it runs one path with the disagreeing threads switched off, then the other path with the roles reversed, and only then merges back to full width. It's the least glamorous line item in a kernel profile and one of the most reliable ways to leave GPU throughput on the table without realizing it.

We keep running into explanations of warp divergence that stop at "it's slow." This post gives you the actual mechanism, the minimum code that triggers it, the measured cost from a real profiling run, and where it hides inside the LLM serving kernels a lot of teams run every day: MoE routing, ragged attention, and sampling. If you rent GPUs by the hour, this is one of the few kernel-level concepts that changes how many hours you actually need.

If you'd rather see divergence measured than described, jump to Measuring the Slowdown on a Real Kernel.

What Is Warp Divergence? Warps, SIMT, and the Cost of a Branch

A modern NVIDIA GPU doesn't schedule individual threads. It schedules warps, groups of 32 threads that share a single instruction stream and a single program counter. That's the SIMT (single instruction, multiple threads) model, and it's the reason warp divergence exists at all: a warp can only be doing one thing at a time, so if its 32 threads want to do 32 different things, the hardware has to serialize them.

What a Warp Actually Is

Every thread in a warp executes the same instruction on every clock cycle it's active. There's no per-thread instruction pointer. When you write ordinary CUDA C++, the compiler and hardware handle warp assignment for you: threads 0-31 of a block form the first warp, 32-63 the second, and so on. As long as all 32 threads in a warp want to execute the same instruction, this is close to free. The GPU's execution units are built around exactly this assumption, and full occupancy on a warp means the SM is getting 32-wide work out of every issued instruction.

What Happens When Threads Inside a Warp Disagree

The trouble starts at a data-dependent conditional branch, something like if (threadIdx.x % 2 == 0), where the condition evaluates differently for different threads in the same warp. Per NVIDIA's own CUDA Programming Guide, "if threads of a warp diverge via a data-dependent conditional branch, the warp executes each branch path taken, disabling threads that are not on that path." The warp runs the if branch with an active mask that enables only the threads that took it and disables the rest, then runs the else branch with the mask flipped, and only reconverges once both paths finish.

This is why a two-way split inside a single warp doesn't just add a branch, it roughly doubles the instruction count that warp has to issue to get the same amount of useful work done: full-width execution twice, at half occupancy each time, instead of one pass at full width. The generalized version of that math, and where it stops mattering, is the next section.

One property worth holding onto, because it changes how you read a profiler later: divergence is warp-local. Per the CUDA guide, "different warps execute independently regardless of whether they are executing common or disjoint code paths." A warp that diverges doesn't stall the warps around it; it just does more work to finish its own instructions. That's good news for scheduling and bad news for silent cost, because a diverging warp's slowdown never shows up as a stall anyone notices. It shows up as lower throughput on a kernel that otherwise looks perfectly healthy.

A Minimal CUDA Code Example That Diverges

You don't need anything exotic to see this. Here's the textbook case, a kernel that branches on threadIdx.x:

__global__ void divergent_kernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= n) return;

    if (threadIdx.x % 2 == 0) {
        // Even lanes: one instruction stream
        data[idx] = data[idx] * 2.0f + 1.0f;
    } else {
        // Odd lanes: a different instruction stream
        data[idx] = sqrtf(data[idx]) + 3.0f;
    }
}

The Classic Case: A Data-Dependent If/Else Keyed on threadIdx

Inside every warp here, lanes 0, 2, 4... take the multiply-add path and lanes 1, 3, 5... take the square-root path. That's a 50/50 split inside every single warp in the launch, the worst-case pattern for a simple branch. The warp scheduler issues the multiply-add instructions with 16 lanes active and the rest masked off, then issues the square-root instructions with the mask flipped, then reconverges. Sixty-four total instruction issues to do 32 threads' worth of work in what looks, from the source code, like a single if.

Contrast that with branching on a value that's uniform across the warp, say if (blockIdx.x % 2 == 0). Every thread in a given warp shares the same blockIdx.x, so every thread in that warp takes the same path. No divergence, because the condition never disagrees within a warp, only between warps, and warps already execute independently of each other.

Predication vs. True Divergence: Why the Compiler Doesn't Always Serialize a Branch

Not every if in your source becomes a diverging branch in the compiled kernel. For short, simple branch bodies, NVIDIA's compiler will often use predication instead: it compiles both paths into a single instruction stream and uses a per-thread predicate bit to suppress the store or the side effect for threads that shouldn't take that path. Predication avoids the reconverge-and-serialize cost because there's only one instruction stream to begin with, just with some lanes' results discarded. The compiler generally reserves this for branches short enough that computing both unconditionally is cheaper than the branch-and-merge overhead. Longer branch bodies, branches containing function calls, loops with data-dependent trip counts, or backward branches (loops) are the cases where the compiler falls back to true divergent execution with an active mask, which is what makes the if/else in the example above expensive rather than free. This is also why divergence bugs are notoriously hard to spot by reading source: the same branch shape can compile to two different execution strategies depending on what's inside it.

Measuring the Slowdown on a Real Kernel

The honest answer to "how slow is warp divergence" is a formula, not an adjective: for a warp that splits into k distinct execution paths, roughly 32/k threads stay active per issued instruction. A 2-way split (like the example above) leaves about 16 active lanes per instruction phase, an efficiency of roughly 50%. A 4-way split drops that to about 8 active lanes, 25% efficiency. A routing pattern with poor locality, where each thread in a warp picks from a wide set of destinations, such as a wide MoE expert table, compounds this active-mask cost with the memory-access divergence of each path pulling from a different address, which is why divergence-heavy kernels tend to underperform their compute-bound theoretical peak by far more than the 32/k instruction-count math alone would suggest.

Reading Warp Execution Efficiency in Nsight Compute

You don't have to take the formula on faith. NVIDIA Nsight Compute (ncu) reports a metric called warp execution efficiency, the average percentage of threads active per issued instruction, directly from hardware performance counters. Profile a kernel with a clean, non-divergent hot loop and you'll typically see it sitting close to 100%. Profile the divergent_kernel example above and you should see it drop to roughly 50%, matching the 32/2 math for a 2-way split.

ncu needs privileged access to hardware performance counters, which is exactly the access most serverless and shared GPU platforms restrict for tenant isolation; it's a documented, structural limitation covered in more depth in our Nsight Compute and PyTorch Profiler production guide, which walks through the full profiling workflow the metric above comes from.

What the Published Numbers Actually Say

We wanted to know whether newer GPUs changed any of this. They haven't. Independent Thread Scheduling, the per-thread progress feature Volta introduced, changed how threads within a diverged warp can interleave and make independent forward progress (useful for producer-consumer patterns within a warp), but it never claimed to change the cost of executing a diverging path. What actually changed generation to generation is the compiler's branch-reconvergence machinery, when and how the hardware decides two diverged paths can merge back together, not the underlying execution cost model. If your mental model is "newer architectures handle branches better," the data doesn't support it: the tax is architectural, not a bug later generations patched.

What Is Warp Divergence Costing You in LLM Serving Code?

Warp divergence reads like a CUDA-101 problem, but it lands squarely in production LLM serving, because three of the most common patterns in modern inference kernels are, structurally, exactly the branch pattern above.

MoE Expert Routing: One Token, One Branch, One Diverging Warp

This is the pattern we see break kernel throughput most often in production MoE decode paths. Mixture-of-experts models route each token to one of several experts based on a learned gate, and a naive implementation writes that as a per-token conditional: "if this token routed to expert 3, run expert 3's weights; if expert 7, run expert 7's." Pack a batch of tokens into a warp and route them independently, and you've reconstructed the divergent kernel pattern with the number of experts standing in for the number of branch paths. A batch where 32 tokens in a warp spread across, say, 8 experts is roughly an 8-way split, the 32/k formula putting expected efficiency around 12.5% before you even account for the memory-access divergence of each path pulling a different expert's weights from HBM.

Ragged Attention Masks and Variable Sequence Lengths at Batch Time

Batched inference rarely has uniform sequence lengths. When a warp processes attention for a batch of requests with different context lengths, the threads mapped to shorter sequences hit their sequence boundary earlier and take a different path (an early-exit or a masked no-op) than threads still processing a longer sequence's remaining positions. That's a divergent branch keyed on sequence length rather than thread index, but the cost model is identical: threads past their sequence's end sit masked off while the warp keeps issuing instructions for the threads still working, so a batch with high length variance pays for its longest member on every warp that mixes lengths.

Sampling and Verification Branches: Top-K and Speculative Decode Accept/Reject

Decoding-time kernels have their own branch surface. Top-k or nucleus sampling involves per-token conditional logic to decide which candidates survive a threshold, and speculative decoding's verification step is a per-token accept/reject branch: does the draft model's proposed token match what the target model would have generated. Both are data-dependent per-thread outcomes packed into a warp, and both diverge for the same structural reason as the toy if/else above: different threads in the same warp landing on different sides of a data-dependent condition.

How Serving Frameworks Avoid It

None of the fixes here are "don't branch." They're all versions of the same move: restructure the work before it reaches the warp, so that by the time a warp executes, every thread in it agrees on the path.

Token Permutation and Grouped GEMM Instead of Per-Token Branching

The standard fix for MoE routing divergence is to sort or permute tokens by assigned expert before the compute kernel runs, so that tokens routed to the same expert end up contiguous and get processed by the same warps with the same weights, no per-token branch required inside the hot kernel. The heavy lifting then becomes a grouped GEMM (one matrix multiply per expert group, on contiguous data) instead of a conditional dispatch per token. A 2026 engineering writeup from Cursor on redesigning an MoE decode kernel took this further at the warp level rather than just the batch level: "each warp in warp decode is independent and gets a single, stable assignment for its entire lifetime: produce one output scalar," removing not just per-token branching inside a warp but cross-warp coordination entirely. That gain didn't come from a faster GPU. It came from a kernel where no warp is ever asked to disagree with itself.

Block-Sparse and Tile-Level Kernels That Mask Instead of Branch

Ragged attention gets a related fix: rather than let each thread branch on whether its position is past the sequence boundary, block-sparse and tile-level attention kernels compute masking as a data operation, multiplying by zero or skipping whole tiles that fall entirely outside the valid range, rather than as a per-thread conditional. A tile either participates or it doesn't; there's no data-dependent branch inside a warp deciding thread by thread. This is the same idea as compiler predication from the code example earlier, applied deliberately at the kernel-design level instead of hoping the compiler finds it.

Why Triton and Tile DSLs Push Divergence Decisions Out of Your Hands

The broader industry trend is to stop asking kernel authors to reason about warps at all. Our OpenAI Triton kernel development guide covers this directly: Triton abstracts away thread indices, shared memory management, and warp-level synchronization, so you write tile-level operations in Python and the compiler decides how to map them onto warps, including how to handle boundary conditions with masks rather than branches. CUDA 13's tile programming APIs, covered in our CUDA 13 tile programming guide, move in the same direction from the C++ side: express computation on tiles, not threads, and the compiler owns the divergence-avoidance decisions that used to sit in a kernel author's if statements. Neither approach eliminates divergence as a physical phenomenon. They eliminate the chance a kernel author introduces it by accident, which is where most of it comes from in practice. Serving stacks built on top of these primitives, benchmarked head to head in our vLLM vs TensorRT-LLM vs SGLang comparison, inherit whichever of these strategies their kernel authors chose, which is one reason the same model can throughput differently across engines even on identical hardware.

Reproducing These Numbers Yourself

Everything in the "Measuring the Slowdown" section is reproducible in under an hour: compile the divergent kernel example, profile it with ncu --metrics smsp__thread_inst_executed_per_inst_executed.ratio, and compare against a non-branching control kernel. The only requirement is hardware performance counter access, which is where the GPU you rent matters. Spheron offers H100 instances as either a fully provisioned VM or a bare-metal instance with per-minute billing and no minimum rental period, and bare-metal allocation is what passes NVIDIA driver capabilities through to the container so ncu doesn't fail with ERR_NVGPUCTRPERM, the counter-permission error that blocks profiling on most shared, serverless GPU hosts. As of this post's publish date, H100 SXM5 on-demand on Spheron starts at $5.07/hr against $6.98/hr on Azure and $7.00/hr on AWS, which puts a short profiling session at well under a dollar. None of this fixes a divergent kernel, that's still a routing, masking, or tiling decision you make in the kernel code, but it does mean the cost of measuring the problem isn't the thing stopping you from finding it.

Pricing fluctuates based on GPU availability. The prices above are based on 04 Sep 2026 and may have changed. Check current GPU pricing → for live rates.

Warp divergence isn't a bug you fix once. It's a property of the SIMT execution model that every branchy kernel pays for, on every architecture, generation after generation. The teams that stay ahead of it aren't the ones with newer GPUs, they're the ones who've restructured routing, masking, and sampling so their warps never have to disagree in the first place.

Reproducing the warp execution efficiency numbers above takes a short bare-metal session with full Nsight Compute access, not a monthly commitment.

Spheron H100 instances →