PidokuInfra
10Occupancy and Warp Divergence
On this page

Occupancy and Warp Divergence

Intermediate Advanced 1h 15m Difficulty 4/5 Topic 10 of 12

Prerequisites 02, 08

1. What is it?#

Two related properties of how well a kernel uses the SM.

Occupancy = active warps / maximum possible warps. Determines how well memory latency is hidden.

Warp divergence = threads in a warp taking different branches, forcing serial execution of both paths.

Occupancy 100%: 64 warps resident, always one ready → latency hidden
Occupancy 25%:  16 warps resident, often all stalled → latency exposed

No divergence:  all 32 lanes do useful work
Full divergence: 32 branches × 1 useful lane each = 1/32 efficiency

2. Why do they exist?#

Occupancy exists because GPUs hide latency with parallelism rather than avoiding it with caches (Section I.09). More resident warps = more chances that one has its data ready.

Divergence exists because a warp is physically one instruction stream with 32 lanes (Section VI.02). Different lanes can’t execute different instructions; they can only be masked off.

3. Simple analogy#

Occupancy: a call centre. If each agent spends 90% of their time on hold waiting for a database, you need 10 agents to keep one conversation active at all times. More agents (warps) means fewer moments where everyone is on hold.

Divergence: a squad marching in lockstep. Order: “those with red badges, turn left; others, turn right.” The squad must do both movements sequentially, with half standing still each time. Two movements’ worth of time for one movement’s worth of progress.

4. Tiny example#

Divergence, measured:

CUDA C++
// NO DIVERGENCE: the branch is uniform within each warp
__global__ void uniform_branch(float* x, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        if ((i / 32) % 2 == 0) x[i] = expensive_a(x[i]);   // whole warps take the same path
        else                   x[i] = expensive_b(x[i]);
    }
}

// FULL DIVERGENCE: alternate lanes take different paths
__global__ void divergent_branch(float* x, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        if (i % 2 == 0) x[i] = expensive_a(x[i]);         // lanes within a warp differ
        else            x[i] = expensive_b(x[i]);
    }
}

The divergent version takes ~2x as long, because each warp executes both branch bodies with half its lanes masked.

Measure with:

Shell
ncu --metrics smsp__thread_inst_executed_per_inst_executed.ratio ./bench
# 32.0 = no divergence;  16.0 = 2-way;  1.0 = fully divergent

5. Technical explanation#

Occupancy arithmetic#

Limits per SM (Hopper):
  max warps          64
  max blocks         32
  registers          65,536 per partition × 4 partitions
  shared memory      up to 228 KB

occupancy = min(
    64,
    32 × warps_per_block,
    (65536 / (regs_per_thread × 32)) × 4    [approximately],
    (228KB / shared_per_block) × warps_per_block
) / 64

Get the inputs from nvcc --ptxas-options=-v:

ptxas info: Used 64 registers, 16384 bytes smem, 380 bytes cmem[0]

Why occupancy matters — and when it doesn’t#

HIGH occupancy helps when:
  - the kernel is memory-latency-bound (many dependent loads)
  - there is little instruction-level parallelism per thread

HIGH occupancy does NOT help when:
  - the kernel is bandwidth-saturated (more warps just means more waiting)
  - the kernel has high ILP (each thread has independent work in flight)
  - the kernel is compute-bound

Volkov’s classic result: a matmul kernel at 33% occupancy, using many registers per thread to hold a large accumulator tile, beat a 100%-occupancy version. More work per thread meant more independent instructions in flight, which hides latency just as effectively as more warps.

So: occupancy is a means, not a goal. Target 40-60% and optimize for actual throughput. Modern GEMM and attention kernels typically run at 25-50% occupancy by design.

Reducing register pressure (if you need to)#

CUDA C++
__global__ void __launch_bounds__(256, 4) my_kernel(...) { ... }
//                                  ^     ^
//                    max threads/block   min blocks/SM

This tells the compiler to limit registers so at least 4 blocks fit. But if the kernel genuinely needs more registers, it will spill to local memory (which is global memory), usually making things worse. Check lmem in the ptxas output.

Divergence patterns to avoid#

BAD:  if (threadIdx.x % 2)                  alternate lanes diverge
BAD:  if (data[i] > threshold)              data-dependent, unpredictable
BAD:  for (j = 0; j < lengths[i]; j++)      variable trip count per lane
BAD:  switch (type[i])                      many-way divergence

OK:   if (threadIdx.x < 32)                 warp-aligned
OK:   if (blockIdx.x % 2)                   block-uniform
OK:   if (i < n)                            only the last warp diverges

Fixing divergence#

1. WARP-ALIGN the condition
   if (threadIdx.x / 32 == 0)   instead of   if (threadIdx.x % 2 == 0)

2. SORT/PARTITION the data so similar work lands in the same warp
   This is what MoE implementations do: sort tokens by expert, then each
   warp handles tokens for one expert.

3. PREDICATE instead of branch (for short bodies the compiler does this)
   x = cond ? a : b;    →  a select instruction, no divergence

4. USE WARP PRIMITIVES to redistribute work
   __ballot_sync to find which lanes have work, then compact

5. ACCEPT IT if the branch bodies are cheap
   Divergence on two 3-instruction bodies costs 3 instructions. Fine.

Point 2 is the important one for inference. MoE routing sends different tokens to different experts — maximally divergent if handled naively. Production MoE kernels sort tokens by expert id first, turning divergence into a grouped GEMM (Section XIII.03).

6. Under the hood#

Nsight Compute’s “Warp State Statistics” tells you why warps stalled:

Stall Reason                    Meaning                          Fix
─────────────────────────────────────────────────────────────────────────
Long Scoreboard                 waiting on global memory         more occupancy, better access
Short Scoreboard                waiting on shared memory         reduce bank conflicts
Barrier                         at __syncthreads()               fewer barriers, more work between
MIO Throttle                    memory pipe saturated            reduce memory instructions
LG Throttle                     local/global pipe saturated      vectorize
Math Pipe Throttle              ALU saturated                    you're compute bound (good!)
Not Selected                    ready but another warp chosen    fine, means you have parallelism
No Instruction                  instruction cache miss           unusual; huge unrolled loops

“Long Scoreboard” dominating with low occupancy is the classic “add more warps” signal. “Math Pipe Throttle” dominating means you’re compute-bound and doing well.

7. Performance implications#

Occupancy 12% → 50% on a latency-bound kernel:      2-4x faster
Occupancy 50% → 100% on the same kernel:            1.1-1.3x
Occupancy increase on a bandwidth-saturated kernel:  ~0x
Full divergence eliminated:                          up to 32x on that region
Register spilling eliminated:                        2-5x

8. Production implications#

  • Read occupancy in profiles, don’t chase it. 25-50% is normal for good kernels.
  • Check for register spilling in any custom kernel: lmem should be 0.
  • Divergence in inference shows up in: MoE routing, variable-length sequence handling, sampling with per-request parameters, and early-exit schemes. Each has a sort-or-group fix.
  • When evaluating a new kernel, look at warp stall reasons before anything else. They tell you what to fix.

9. Common mistakes#

Maximizing occupancy as a goal. It’s a means.

Using __launch_bounds__ to force occupancy and causing spills.

Ignoring divergence in data-dependent code.

Assuming the last-warp divergence from if (i < n) matters. It affects one warp out of thousands.

Not checking stall reasons. They’re the diagnosis; occupancy is a symptom.

10. Hands-on exercise#

A. Measure divergence. Implement both kernels from section 4 with a genuinely expensive body. Measure the ratio and thread_inst_executed_per_inst_executed. Confirm ~2x.

B. Occupancy sweep. Take a memory-latency-bound kernel. Artificially vary its register usage (via __launch_bounds__ or extra local variables) and plot occupancy vs achieved bandwidth. Where does the curve flatten?

C. Volkov’s result. Write a kernel where each thread computes 1 output element and one where each computes 8 (using more registers). Compare occupancy and throughput. Does the lower- occupancy version win?

D. Stall reasons. Profile three different kernels (a memory-bound elementwise, a GEMM, and a reduction) and compare their warp stall distributions. Explain each.

E. MoE divergence. Simulate MoE routing: 1024 tokens, 8 experts, random assignment. Implement (i) a naive kernel where each thread branches on its token’s expert, and (ii) a version that sorts tokens by expert first. Measure both.

11. Interview questions#

  1. Define occupancy. What limits it?
  2. Is higher occupancy always better? Explain with an example.
  3. What is warp divergence and what’s the worst-case cost?
  4. Give three ways to reduce divergence.
  5. Why is MoE routing divergence-prone, and how do production kernels handle it?
  6. What is register spilling and how do you detect it?
  7. You profile a kernel and see “Long Scoreboard” as the dominant stall reason at 15% occupancy. What do you do?

12. Further reading#

  • [ESTABLISHED] Volkov, “Better Performance at Lower Occupancy” (GTC 2010)
  • [REFERENCE] Nsight Compute warp state documentation
  • [REFERENCE] CUDA Best Practices Guide, “Execution Configuration Optimizations”
  • Next: 11 — Tensor cores
esc

↑↓ navigate↵ openesc close