FDEInterviews logo
ML Infrastructure & GPUs / 05
mediumNVIDIAGoogleCoreWeave

What is occupancy, and how do you balance it against register and shared-memory usage when choosing block size?

Everyone says 'maximize occupancy.' The candidates who get hired at NVIDIA know when 25% occupancy beats 75%, and can explain the resource math that decides it.

Updated Sep 2026 · Grounded in real Forward Deployed Engineer interview loops and written to a senior-engineer editorial bar.

TL;DR: Occupancy is resident warps over the hardware max (64 warps / 2,048 threads per SM), capped by whichever of registers, shared memory, or block limits binds first. It is a means to latency hiding, not the goal: a register-heavy GEMM at 25% occupancy often beats the same kernel forced to 75%, so you profile stall reasons before touching the knob.

How to approach it

Define occupancy precisely, list the three resources that cap it, then make the senior move: occupancy is a means (latency hiding), not the goal. Showing you know when lower occupancy wins is the difference between a textbook answer and a practitioner's answer.

A strong answer

Occupancy is the ratio of active warps per SM to the hardware maximum (64 warps, meaning 2,048 threads, on recent datacenter parts). It matters because warps are how the SM hides latency: when one warp stalls on a 400-cycle HBM load, the scheduler issues from another. Too few resident warps and the SM sits idle during stalls.

Three resources cap how many blocks fit on an SM, and the worst one wins:

  • Registers: an SM has 64K 32-bit registers. At 32 registers/thread you can host 2,048 threads (100%); at 128 registers/thread, only 512 (25%).
  • Shared memory: ~100-228 KB per SM depending on configuration. A block asking for 48 KB means at most 2-4 blocks resident.
  • Threads/blocks per SM: hard limits (2,048 threads, 32 blocks per SM), which is why tiny 32-thread blocks cannot reach full occupancy. 32 blocks × 32 threads = 1,024 threads, a 50% cap.
rendering diagram…

"The worst one wins" is a min() you should be able to evaluate live, so run one kernel through all three caps (arithmetic verified). Say the kernel compiles to 96 registers per thread, launches 256-thread blocks, and each block uses 32 KB of shared memory. Registers: 65,536 / 96 = 682 threads, which fits only 2 whole blocks of 256, so 512 threads. Shared memory: 228 / 32 = 7 blocks. Thread limit: 2,048 / 256 = 8 blocks. The minimum is the register cap at 2 blocks: 512 resident threads, 25% occupancy, and the other two caps are irrelevant slack. That final sentence is the actionable one: this kernel's occupancy responds only to register pressure, so shrinking shared memory or resizing blocks accomplishes nothing, and the profiler-guided question becomes whether 16 warps hide this kernel's particular stalls or whether trimming registers (without spilling) pays. Being able to name the binding resource, not just the percentage, is what turns the calculator output into a tuning decision.

Block size guidance: always a multiple of 32 (partial warps waste lanes), and in practice 128-256 is the right default. Enough warps per block to sync efficiently, small enough to give the scheduler several blocks per SM so an SM is not left half-empty by a finishing block's tail. 1,024-thread blocks are usually a mistake: one resource-hungry mega-block quantizes badly.

Now the part interviewers actually score: more occupancy is not always faster. Occupancy buys latency hiding; once you have enough warps to cover your stall pattern, extra warps buy nothing and may hurt, because more resident threads share the same L1/L2 and increase cache thrash. Meanwhile registers are the fastest storage on the chip. A kernel that keeps a 64-element accumulator tile in registers at 25% occupancy routinely beats the same kernel squeezed to 64 registers/thread at 75% occupancy, because instruction-level parallelism within each thread also hides latency (the classic Volkov result, and the design point of CUTLASS-style matmuls: high register tiles, modest occupancy). Forcing registers down with -maxrregcount or __launch_bounds__ can backfire by spilling to local memory, which is HBM.

So the actual workflow: profile first. Nsight Compute tells you achieved occupancy and, more usefully, why warps stall (memory latency vs dependency vs issue). If the kernel is memory-latency-bound with low occupancy, cut register/shared usage or shrink blocks. If it is already saturating bandwidth or compute, occupancy is the wrong knob entirely. Use the occupancy calculator (or cudaOccupancyMaxPotentialBlockSize) for a starting point, then sweep block sizes, since it is cheap to test 128/256/512 empirically.

What interviewers probe next

  • "Your kernel is at 30% occupancy. Is that a problem?" Only if the profiler shows latency-bound stalls; a register-heavy compute-bound GEMM at 30% can be at peak FLOPs. Look at stall reasons before touching anything.
  • "What happens if I force fewer registers per thread?" Occupancy rises, but past the compiler's comfort the surplus values spill to local memory; you trade fast registers for HBM traffic, often a net loss.
  • "Why are blocks of 32 threads bad if 32 is the magic number?" Per-SM block-count limits cap total threads, and per-block scheduling overhead grows; a multiple of 32 is not the same as exactly 32.
  • "How does shared memory tiling interact with this?" Bigger tiles mean more reuse per HBM byte but fewer resident blocks; you are trading arithmetic intensity against latency hiding, and the optimum is found by sweeping, not by formula.

Common mistakes

  • Declaring "maximize occupancy" as the strategy. Interviewers at NVIDIA specifically hold the low-occupancy counterexample in reserve for this.
  • Not knowing the actual resource numbers (64K registers, ~2K threads/SM). The question is arithmetic, and hand-waving it shows.
  • Ignoring spills: recommending -maxrregcount without mentioning local-memory spill is a red flag.
  • Choosing block size by superstition ("we always use 256") with no mention of profiling or sweeping. 256 is a fine default; presenting it as a law is the tell.

Key takeaways

  • Occupancy = min over the three resource caps; know the 64K-register, 2,048-thread numbers cold.
  • It buys latency hiding only; past "enough warps," more occupancy can hurt via cache pressure.
  • Register-rich, low-occupancy GEMMs (CUTLASS, Volkov) are a feature, not a bug. Profile stall reasons before tuning.
That one was free — and so are 10 answers per topic without an account. Signing in doubles that to 20, opens the Plus lessons in the courses, and remembers which topics you keep getting wrong.no card · Google sign-in · nothing to cancel
HOW DID IT GO?
0
READING SIGNED OUT

Signing in doubles your free answers, from 10 to 20 per topic, and the site starts remembering you: mastery per topic, bookmarks, and a next-focus recommendation. Free, no card.

Sign in free
UP NEXT ON YOUR JOURNEY
FEDITOR'S NOTE

NVIDIA interviewers hold the low-occupancy counterexample in reserve specifically for candidates who declare 'maximize occupancy' as the strategy; a register-heavy GEMM at 25% can be at peak FLOPs because instruction-level parallelism also hides latency (the Volkov result, and the CUTLASS design point). Recommending -maxrregcount without mentioning that it can spill to local memory in HBM is a red flag. Presenting block size 256 as a law rather than a profiled default is the tell that you choose by superstition.

DISCUSSION · 0

No comments yet — be the first to share your approach.