CUDA Occupancy Calculator: What 100% Occupancy Really Means

The CUDA occupancy calculator answers one question: how many warps could be resident on a streaming multiprocessor at once, given your kernel's block size, register use, and shared memory footprint. It says nothing about how fast those warps actually run. That gap between "could be active" and "is fast" is where a lot of kernel-tuning time gets burned chasing a number that Nsight Compute will happily report at 100% on a kernel a differently configured version beats by a wide margin.

TL;DR: What Does the CUDA Occupancy Calculator Actually Measure?

• Definition: occupancy is active warps per SM over an SM's max, per Nsight Compute.
• Theoretical vs achieved: theoretical occupancy is the launch config's ceiling; achieved is what Nsight Compute measures live. A wide gap means imbalance, not a bad config.
• Register math, H100 SM: 64K registers and 2,048-thread cap mean 32 registers/thread hits 100% occupancy; 128 caps it at 25%.
• Block limits: Nsight Compute splits the ceiling into Registers, Shared Mem, and Warps to trace low occupancy to one resource.
• FlashAttention-4 added Blackwell tensor memory to stop spilling that capped Hopper tile size, trading occupancy for a bigger tile.
• Running ncu needs driver-level access most shared hosts block; Spheron's bare-metal GPUs pass it through.

What Is GPU Occupancy? Active Warps vs the Maximum Possible

A streaming multiprocessor doesn't run one thread block at a time and wait. It keeps several warps resident simultaneously and switches between them on every clock cycle a warp stalls, which is how the GPU hides memory and instruction latency without a CPU-style branch predictor or deep out-of-order window. Occupancy is the number that describes how many of those warp "slots" your kernel is actually using.

NVIDIA's own Nsight Compute Profiling Guide defines it plainly: occupancy is the ratio of active warps per multiprocessor to the maximum number of warps that multiprocessor can have active at once. If an SM can hold 64 warps resident and your kernel keeps 32 of them busy, that's 50% occupancy. Nothing in that ratio references instructions per second, memory bandwidth, or wall-clock time. It's a statement about how much of the SM's latency-hiding capacity a launch configuration is using, not a statement about speed.

This is also where occupancy gets confused with a related but distinct metric: warp execution efficiency, sometimes called warp divergence. Occupancy asks how many warps are resident. Divergence asks how many of the 32 threads inside each of those warps are doing useful work on a given instruction. Per NVIDIA's 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." A kernel can sit at 100% occupancy with every warp resident and still waste half its issued instructions on masked-off lanes from a bad branch. Our warp divergence deep dive covers that failure mode on its own; this piece is about the other one, the one where the warps are all there and all on the same path, and the kernel is still slower than it should be.

Theoretical Occupancy vs Achieved Occupancy

The occupancy calculator, wherever you read it from, actually reports two different numbers, and mixing them up is a common source of confused tuning sessions.

Theoretical occupancy is the ceiling. It's computed from your kernel's launch configuration, block size, registers per thread, shared memory per block, and the target device's hardware limits, before the kernel ever runs. It answers "given this configuration, what's the most warps that could ever be resident on one SM."

Achieved occupancy is what Nsight Compute measures while the kernel is actually executing on real hardware. Per the same Profiling Guide, a large gap between theoretical and achieved occupancy typically indicates workload imbalance: some thread blocks finish their work earlier than others, some blocks never get scheduled because the grid is too small to fill every SM, or the last "wave" of blocks straggles in with only a few SMs still busy while the rest sit idle waiting for the kernel to finish.

That distinction matters for where you spend tuning effort. A kernel with low theoretical occupancy has a launch-configuration problem: too many registers per thread, too much shared memory per block, or a block size that doesn't divide the SM's warp slots evenly. A kernel with high theoretical occupancy but low achieved occupancy has a scheduling or load-balancing problem instead, and no amount of register tuning will fix it. Nsight Compute reports both numbers side by side specifically so you don't have to guess which category you're in.

What Actually Limits Occupancy: Registers, Shared Memory, Block Size

Three resources cap how many thread blocks, and therefore how many warps, an SM can hold resident at once, and every one of them is finite per SM regardless of how big your grid is:

• Registers. Every SM has a fixed register file, allocated in full to whatever threads are currently resident. A kernel that uses more registers per thread leaves room for fewer resident threads.
• Shared memory. Each SM has a fixed pool of on-chip shared memory, split among the resident thread blocks that request it. A kernel with a large __shared__ allocation per block leaves room for fewer resident blocks.
• Block size and hardware block/warp limits. Every architecture also caps the raw number of resident blocks and warps per SM independent of registers or shared memory, so a kernel with tiny blocks can hit that ceiling before it exhausts either resource.

NVIDIA's own Hopper Tuning Guide gives the concrete numbers for the H100's SM: a maximum of 64 concurrent warps per SM, the same ceiling as the prior Ampere generation, which works out to 2,048 resident threads at 32 threads per warp, a register file of 64K 32-bit registers per SM, and a maximum of 255 32-bit registers per thread. Those numbers are what every occupancy calculation for an H100 kernel is actually working from, whichever tool reports the final percentage.

Why More Registers Per Thread Can Lower Occupancy but Raise Speed

Here's the part that trips people up: cutting a kernel's register use to raise occupancy is not automatically a win, and the reverse move, deliberately using more registers per thread, is a standard optimization technique in hand-tuned GEMM and FFT kernels.

The mechanism is straightforward once you separate two different ways a GPU hides latency. One is thread-level parallelism: keep many warps resident, and when one stalls on a memory fetch, switch to another that's ready to issue. That's what occupancy measures. The other is instruction-level parallelism inside a single thread: give a thread enough independent work and enough registers to hold several partial results at once, and it can keep issuing useful instructions between its own memory fetches without ever needing a second warp to cover for it. A kernel that leans on the second mechanism can run fewer resident warps and still keep the SM's execution units fed, because each warp is doing more useful work per instruction it issues instead of waiting more often.

NVIDIA's CUDA C++ Best Practices Guide frames the underlying scheduling behavior this way: "if the GPU must wait on one warp of threads, it simply begins executing work on another." That's the occupancy mechanism working as designed. What it doesn't say, and what NVIDIA's Nsight Compute guide says explicitly, is that this is the only way to keep an SM busy. Register-heavy, ILP-driven kernels are the other way, and they trade occupancy for it on purpose.

This Isn't a New Trade, Just a Newly Visible One

Hand-tuned dense linear algebra and FFT kernels have leaned on this register-for-occupancy trade for years, well before attention kernels made it visible to a wider audience. The mechanism is the one described just above: restructure a kernel so each thread carries more independent work and more live registers, and it needs fewer resident warps per block, because its own instruction-level parallelism keeps the SM's pipelines fed instead of relying on a second warp to cover every stall. It's the same latency-hiding job the CUDA C++ Best Practices Guide describes the scheduler doing automatically, just taken on deliberately at the kernel-authoring level instead. That's also the conceptual ancestor of the register-vs-tile-size tradeoff every modern attention kernel author still makes, including the one covered two sections down.

A Worked Register Math Example on a Modern SM

Put real numbers behind it using the H100 SM specs above: a 64K 32-bit register file and a 2,048-thread (64-warp) maximum per SM. Register pressure alone determines how many threads can be resident, before shared memory or block-size limits are even considered:

The math is just division: 65,536 registers divided by registers-per-thread gives the register-limited thread count, and that count divided by 2,048 gives occupancy relative to the SM's maximum. In practice the compiler allocates registers in fixed-size blocks and the hardware rounds down to whole warps and whole blocks, so real kernels land slightly under these clean numbers, but the direction and the order of magnitude are exactly right: doubling registers per thread roughly halves theoretical occupancy on the same SM. Whether that's a bad trade depends entirely on what those extra registers buy the kernel, which is the question the rest of this post is about.

The CUDA Occupancy Calculator: Reading It Correctly

The CUDA occupancy calculator, in whichever form you're using it, answers a single question: at your kernel's actual block size, register count per thread, and shared memory use per block, what fraction of an SM's maximum warp capacity could be resident at once. It's a launch-configuration diagnostic. It is not a benchmark, and it has no way to know how memory-bound, compute-bound, or divergence-heavy your kernel's inner loop actually is.

The Old Spreadsheet Is Deprecated, Here's What Replaced It

If you're picturing a standalone Excel file where you type in registers-per-thread and block size and it spits out a percentage, that tool is deprecated. NVIDIA's own archived documentation for the CUDA Occupancy Calculator states that the standalone spreadsheet is no longer maintained, and that the same occupancy analysis now lives inside Nsight Compute. The practical difference matters: the spreadsheet took manual input and produced a theoretical estimate. Nsight Compute reads your kernel's actual compiled register count and shared memory allocation directly, and reports both theoretical occupancy from that real configuration and achieved occupancy measured from an actual run, side by side, on the actual GPU you're targeting rather than a generic model of it.

Block Limit Registers vs Block Limit Shared Mem vs Block Limit Warps

This is the part of the occupancy calculator most people skip past to get to the single percentage at the top, and it's the most useful part. Per the Nsight Compute Profiling Guide, the Occupancy section reports separate block limits for SM, Registers, Shared Memory, and Warps, so a low-occupancy kernel can be traced to one specific resource constraint instead of read as a single, unexplained number.

Each limit answers a different question about the same kernel:

• Block Limit Registers is how many thread blocks fit before the SM's register file runs out, given the compiled register count per thread.
• Block Limit Shared Mem is how many blocks fit before the SM's shared memory pool runs out, given the __shared__ allocation per block.
• Block Limit Warps is the raw hardware ceiling on resident warps per SM, independent of registers or shared memory.
• Block Limit SM is the architecture's cap on resident blocks per SM regardless of how small each block's resource footprint is.

Theoretical occupancy is whichever of these hits its wall first, so the fix is specific to whichever limit is binding: too many registers means shrinking the kernel's live-value footprint or splitting work across more, smaller operations; too much shared memory means a smaller tile size or double-buffering scheme; hitting the warp or block ceiling directly with small, resource-light blocks usually means the block size itself needs to grow. Reading which column is red before touching any code is the whole point of the section. For the full profiling workflow this section comes from, including how to capture the trace ncu needs on a remote GPU, see our Nsight Compute and PyTorch Profiler production guide.

When Low Occupancy Is Fine (and When It Isn't)

None of the above means occupancy doesn't matter. It means it's conditional on what's actually bottlenecking the kernel, and that split is the real answer to "is my occupancy good enough."

Compute-Bound Kernels: FlashAttention's Register-vs-Tile-Size Tradeoff

Compute-bound kernels, ones already spending most of their cycles on tensor-core matrix multiplies rather than waiting on memory, have room to trade occupancy for a bigger, more efficient unit of work per warp. FlashAttention's kernel authors make exactly this call. On Blackwell, FlashAttention-4's design adopted UMMA, a matrix multiply issued by a single thread, along with dedicated tensor memory (TMEM) for holding accumulators, specifically to avoid the register spilling that had been constraining tile size on Hopper's warpgroup MMA. Spilling registers to local memory is one of the costs that makes low occupancy actively harmful even in a compute-bound kernel: a spilled register turns into a slow DRAM round trip on every access, not just an extra register file entry. Moving accumulators into TMEM freed up the register budget to run larger tiles without the spilling penalty, a direct, shipping example of choosing a lower-occupancy, register-heavy configuration because the bigger tile it enables does more useful compute per warp launched. Our FlashAttention explainer covers what that kernel-level tiling buys in inference cost terms; the register and occupancy mechanics behind it are usually left out of that conversation, and they're the reason the newest kernel generation needed new hardware (TMEM) to keep making the same trade.

The same tradeoff shows up wherever kernel authors write tile-level code by hand. Our OpenAI Triton kernel development guide walks through the same register-spill failure mode from the other direction: a Triton kernel with num_warps or BLOCK_SIZE set too high fills the register file and spills to local memory, invisibly, with the kernel still producing correct output at a lower speed. The difference with Triton and with CUDA 13's tile programming APIs is who makes the call: increasingly, it's the compiler deciding how to map a tile-level operation onto warps and registers, rather than a kernel author hand-picking a register count and checking the occupancy calculator afterward.

Memory-Bound Kernels: Where Low Occupancy Actually Hurts

Memory-bound kernels don't get the same latitude. When a kernel's bottleneck is HBM bandwidth or memory latency rather than compute throughput, occupancy is doing real work: more resident warps means more outstanding memory requests in flight at once, which is what actually saturates the memory bus and hides the hundreds of cycles a DRAM access takes. Cut occupancy on a memory-bound kernel and there usually aren't enough independent memory operations in flight to keep the bus busy, so throughput drops in direct proportion to the warps you removed. This is exactly the case the Nsight Compute Profiling Guide is describing when it states that low occupancy always reduces the ability to hide latencies, resulting in overall performance degradation, even while noting in the same breath that higher occupancy doesn't guarantee higher performance either. Both halves of that sentence are true at once because they're describing different kernel types.

Reading a single occupancy percentage in isolation, without knowing whether the kernel in question is compute-bound or memory-bound, is the mistake this entire post is arguing against. It's the same trap covered from the utilization side in our GPU spec sheet vs real-world performance piece: a single averaged number, whether it's tensor-core utilization or occupancy, can describe two completely different underlying conditions, and the only way to tell them apart is to look at the roofline position of the kernel alongside it.

What This Means for Choosing a GPU for Your Kernel

Occupancy math isn't portable across GPU generations, because the inputs to it, register file size, shared memory per SM, and max warps per SM, change with the architecture. A kernel tuned to 100% theoretical occupancy on one SM design can land at a different percentage on the next generation's SM without a single line of code changing, simply because the register file or shared memory budget per SM moved. That's a real planning cost if you're choosing hardware for a custom kernel rather than running someone else's pre-tuned inference stack: the occupancy number your kernel gets on paper is a property of the specific SM you compile and profile against, not a fixed characteristic of the kernel itself.

The only way to know where a specific kernel actually lands, on the specific GPU you're planning to rent, is to compile it, run it, and read Nsight Compute's achieved occupancy and block-limit breakdown on that hardware. Spheron's bare-metal GPU instances pass those driver capabilities through directly, and with per-minute billing and no minimum rental period, a short profiling session, compile the kernel, run ncu, read the Occupancy section, costs a fraction of an hour's rate rather than a monthly commitment. As of this writing, Spheron lists H100 on-demand at $2.64/hr and spot at $2.04/hr; that's a hardware-access argument for getting the profiler to run at all, not a claim that Spheron tunes the kernel for you, and a team that only cares about served-inference throughput rather than kernel-level tuning is better served by a pricing or serving guide than by occupancy math.

Pricing fluctuates based on GPU availability. Spheron rates above are live as of 08 Sep 2026; other providers reflect their most recent published rates and may have changed. Check current GPU pricing → for live rates.

Occupancy is a ceiling on how much latency-hiding parallelism your launch configuration allows, nothing more. Reading it as a speed score is how a correctly-tuned register-heavy kernel gets "optimized" back down to a lower, worse configuration by someone chasing 100% on a spreadsheet number that was deprecated years ago.

Profiling the real occupancy and register-spill behavior of a custom kernel takes full ncu access, not a guess from a spreadsheet.

On-demand H100 on Spheron →