GPU Occupancy Calculator, in depth: the exact arithmetic, a CI-ready calculator, the CUDA occupancy API and reading the result

The inputs and where they come from

Three inputs come from the kernel and one from the device. From the kernel: threads per block, registers per thread, and shared memory per block (static plus dynamic). From the device: its compute capability, which fixes the per-SM limits. Get the kernel numbers from the compiler rather than guessing them. Compiling with nvcc -Xptxas -v prints, per kernel, a line such as Used 72 registers, 0 bytes smem along with any spill stores and loads. That figure is static shared memory only; a tile declared extern __shared__ is dynamic, sized at launch, and you must add it yourself. cuobjdump --dump-resource-usage reports the same for a built binary, and at runtime cudaFuncGetAttributes returns numRegs and sharedSizeBytes for the kernel actually loaded.

The per-SM limits below are from NVIDIA's compute capability table for three common data centre and workstation targets. Other architectures differ, so put the table in code and extend it from the programming guide rather than from memory.

The one-kilobyte gap between the per-SM and per-block maxima is not a typo: the driver reserves shared memory for each resident block on these architectures. Query it with cudaDevAttrReservedSharedMemoryPerBlock instead of hard-coding it.

The arithmetic, including the granularities

Each limit gives a number of blocks; resident blocks is the minimum. Two of the limits are simple division. Warp slots allow max_warps // warps_per_block blocks, where warps per block is threads divided by 32, rounded up, so a 100-thread block costs four full warps. Block slots allow max_blocks blocks regardless of size, which is why tiny blocks of 32 or 64 threads cannot fill an SM.

Registers are where hand arithmetic usually goes wrong. Registers are allocated per warp, in units of 256, so a warp's cost is ceil(regs * 32 / 256) * 256. In effect the per-thread count rounds up to a multiple of eight: 65 registers cost as much as 72. The register-limited warp count is then rounded down to a multiple of four, which reflects each SM being split into four scheduler partitions that each own a quarter of the register file. Only then do you divide by warps per block. NVIDIA's spreadsheet calculator and its cuda_occupancy.h header model these granularities; a formula that skips them can be optimistic by a block.

Shared memory per block is the kernel's static plus dynamic bytes, plus the per-block reservation, rounded up to the allocation unit. Blocks allowed is the SM's configured carveout divided by that. The carveout matters: shared memory and L1 come from one pool, and the driver picks a split per kernel. You can state a preference with cudaFuncAttributePreferredSharedMemoryCarveout, and any kernel wanting more than 48 KB of dynamic shared memory per block must opt in with cudaFuncSetAttribute and cudaFuncAttributeMaxDynamicSharedMemorySize or the launch fails.

A calculator you can run in CI

The calculator below is small enough to read in one sitting and is suitable for a CI check. It encodes the table above and the granularity rules, returns the binding limit by name, and is meant to be cross-checked against the runtime API on real hardware.

from dataclasses import dataclass
from math import ceil

@dataclass(frozen=True)
class Arch:
    name: str
    max_warps: int            # resident warps per SM
    max_blocks: int           # resident blocks per SM
    smem_sm_max: int          # largest shared memory carveout, bytes
    regs_per_sm: int = 65536
    reg_unit: int = 256       # registers allocated per warp in units of 256
    warp_gran: int = 4        # register-limited warps round down to a multiple of 4
    smem_unit: int = 128      # shared memory allocation unit, bytes
    smem_reserved: int = 1024 # per-block reservation; query it on the device

A100 = Arch("sm_80", 64, 32, 164 * 1024)
SM86 = Arch("sm_86", 48, 16, 100 * 1024)
H100 = Arch("sm_90", 64, 32, 228 * 1024)

def round_up(x, unit):
    return ceil(x / unit) * unit

def occupancy(arch, threads, regs, smem_bytes, carveout=None):
    wpb = ceil(threads / 32)
    limits = {"warps": arch.max_warps // wpb, "blocks": arch.max_blocks}
    regs_per_warp = round_up(regs * 32, arch.reg_unit)
    warps = arch.regs_per_sm // regs_per_warp
    limits["registers"] = (warps - warps % arch.warp_gran) // wpb
    per_block = round_up(smem_bytes + arch.smem_reserved, arch.smem_unit)
    limits["shared"] = (carveout or arch.smem_sm_max) // per_block
    blocks = min(limits.values())
    binding = [k for k, v in limits.items() if v == blocks]
    return blocks, blocks * wpb / arch.max_warps, binding

A zero result is a real outcome, not a rounding curiosity. A 1,024-thread block at 72 registers per thread needs 73,728 registers, more than the SM has, so the launch fails with a too-many-resources error. Make the CI check fail on zero, and warn when the binding limit changes between commits.

Worked example: a tiled kernel on sm_90

Take a tiled kernel on an H100: 256 threads per block (eight warps), 40 KB of shared memory per block, and ptxas reports 72 registers per thread. Work each limit.

• Warp slots: 64 / 8 = 8 blocks.
• Block slots: 32 blocks.
• Registers: 72 x 32 = 2,304 per warp, already a multiple of 256. 65,536 / 2,304 = 28.4, so 28 warps, which is already a multiple of four. 28 / 8 = 3 blocks.
• Shared memory: 40 KB + 1 KB reserved = 41 KB per block. With the full 228 KB carveout, 228 / 41 = 5.56, so 5 blocks.

Resident blocks are min(8, 32, 3, 5) = 3: 24 warps of 64, or 37.5 percent theoretical occupancy, bound by registers. If you cap the kernel at 64 registers, each warp costs 2,048, the SM holds 32 warps and four blocks: 50 percent. Cap further, at 40 or even 32, and nothing improves past five blocks, because shared memory now binds at 62.5 percent. That is the key reading skill: after each change, ask which limit is binding now, because pulling a lever that is not binding costs performance and buys nothing.

The block size sweep tells the same story from another angle. At 72 registers and 40 KB, blocks of 64 or 128 threads are capped by shared memory at five blocks (15.6 and 31.3 percent), 192, 256 and 384 threads all land on 24 warps, and 512 threads fits a single block for 25 percent. Larger blocks are not automatically better.

The authoritative answer: the CUDA occupancy API

On a real device, ask the runtime instead of trusting your table. The occupancy functions use the loaded kernel's actual resource usage and the device's actual limits.

// Authoritative theoretical occupancy for this kernel on this device.
int blocks_per_sm = 0;
size_t dyn_smem = 40 * 1024;
cudaFuncSetAttribute(tiled_kernel,
                     cudaFuncAttributeMaxDynamicSharedMemorySize, (int)dyn_smem);
cudaOccupancyMaxActiveBlocksPerMultiprocessor(&blocks_per_sm, tiled_kernel,
                                              256, dyn_smem);

cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
float occ = (float)(blocks_per_sm * 256 / prop.warpSize)
          / (prop.maxThreadsPerMultiProcessor / prop.warpSize);
printf("blocks/SM=%d theoretical occupancy=%.1f%%\n", blocks_per_sm, occ * 100);

// Let the runtime suggest a block size that maximises occupancy.
int min_grid = 0, best_block = 0;
cudaOccupancyMaxPotentialBlockSize(&min_grid, &best_block, tiled_kernel,
                                   dyn_smem, 0);

Two cautions. cudaOccupancyMaxPotentialBlockSize maximises occupancy, not speed; treat its answer as a candidate to benchmark. And when dynamic shared memory depends on block size, use the variant that takes a function mapping block size to bytes, or the suggestion will be computed against the wrong footprint.

To move the register limit, __launch_bounds__(256, 4) tells the compiler the kernel will be launched with at most 256 threads and that you want four blocks per SM, so it will budget registers to 64 per thread. -maxrregcount does the same for every kernel in a file, which is usually too blunt. Either way, re-read the ptxas output: if spill stores appear, local memory traffic can cost more than the extra occupancy buys.

Theoretical versus achieved

Theoretical occupancy is a ceiling. Nsight Compute's Occupancy section reports both the theoretical figure and achieved occupancy, the average number of warps actually resident while the kernel ran. A large gap usually means a tail effect (the grid is a few blocks more than a whole wave, so the last wave runs mostly empty SMs), imbalanced blocks finishing at different times, or too few blocks in total. Count the waves: with an H100 SXM's 132 SMs at 3 blocks each, one wave is 396 blocks, so a 400-block grid runs a second wave of four blocks on a nearly idle GPU. Changing occupancy changes the wave size, which is why a lower-occupancy configuration sometimes wins.

Failure modes

• Computing from source instead of from ptxas. Register counts change with compiler version, optimisation flags and template parameters. Feed the calculator from the build output for the exact binary.
• Ignoring the granularities. Dividing 65,536 by 72 x 32 gives 28.4 warps; ignoring per-warp rounding at, say, 70 registers would overstate capacity.
• Forgetting the per-block reservation and the 48 KB opt-in. The first makes the shared memory limit off by one block at the boundary; the second makes the launch fail.
• Assuming the maximum carveout. The driver may choose a smaller shared memory split for a kernel. If your calculation depends on the full carveout, set the preference.
• Maximising occupancy as a goal. Kernels that keep data in registers and exploit instruction-level parallelism, such as many GEMM and attention kernels, often run fastest at modest occupancy. See register pressure for the spill trade-off.

Trade-offs

Related reading: shared memory for tiling and bank conflicts, and warp scheduling for how resident warps are actually issued.

What to do next

1. Add -Xptxas -v to your build and save registers, shared memory and spills per kernel as a build artefact.
2. Encode the per-SM limits for the GPUs you deploy on, taken from NVIDIA's compute capability table.
3. Run the calculator for each hot kernel and record the binding limit, not just the percentage.
4. On a real device, check the result with cudaOccupancyMaxActiveBlocksPerMultiprocessor and fix any mismatch in your table.
5. Profile with Nsight Compute and compare achieved to theoretical occupancy; count waves for the grid size you launch.
6. Change one lever at a time, re-read ptxas for spills, and keep a change only if runtime improves.