The CUDA Handbook ▸ Chapter 7. Kernel Execution

7.3 Occupancy

Occupancy is a ratio that measures the number of threads/SM that will run in a given kernel launch, as opposed to the maximum number of threads that potentially could be running on that SM:

\[\frac{Warps\ per\ SM}{Max.\ \ Warps\ per\ SM}\]

The denominator (maximum warps per SM) is a constant that depends only on the compute capability of the device. The numerator of this expression, which determines the occupancy, is a function of the following:

• Compute Capability,

• Threads per block,

• Registers per thread,

• Shared memory configuration, and

• Shared memory per block.

Rather than work the occupancy out by hand, an application can query it for a specific kernel at runtime. cudaOccupancyMaxActiveBlocksPerMultiprocessor() reports how many thread blocks of a given size (and dynamic shared-memory allocation) can be simultaneously resident on an SM for a particular kernel, from which the occupancy follows directly; cudaOccupancyMaxPotentialBlockSize() goes a step further and suggests a block size that maximizes it. The driver API exposes the same pair as cuOccupancyMaxActiveBlocksPerMultiprocessor() and cuOccupancyMaxPotentialBlockSize(). These functions also pair naturally with the cooperative launch APIs (Section 7.6): when an application requires every warp in the launch to be simultaneously active, so the grid can synchronize as a whole, it sizes the grid at cudaOccupancyMaxActiveBlocksPerMultiprocessor() blocks per SM times the number of multiprocessors, guaranteeing the entire launch is co-resident.

To help developers assess the tradeoffs between these parameters, the CUDA Toolkit once included an occupancy calculator in the form of an Excel spreadsheet¹⁰. A more recent, GitHub-hosted tool covers CUDA 11-era hardware. Given the inputs above, these tools calculate the following:

• Active thread count,

• Active warp count,

• Active block count, and

• Occupancy (active warp count divided by the hardware’s maximum number of active warps).

These tools were intended to help developers identify whichever parameter is limiting the occupancy:

• Registers per multiprocessor,

• Maximum number of warps or blocks per multiprocessor, or

• Shared memory per multiprocessor.

Part of the reason NVIDIA no longer maintains its occupancy calculator may be that, contrary to the developer education guidance of early CUDA years, occupancy is not the be-all and end-all of CUDA performance¹¹; often it is better to use more registers per thread and rely on instruction-level parallelism (ILP) to deliver performance. NVIDIA has a good presentation on warps and occupancy that discusses the tradeoffs.

An example of a low-occupancy kernel that can achieve near-maximum global memory bandwidth is given in Section 5.2.8 (Listing 5-6). The inner loop of the GlobalReads() kernel can be unrolled according to a template parameter; as the number of unrolled iterations increases, the number of needed registers increases and the occupancy goes down. For a GeForce RTX 3060, for example, the peak read bandwidth reported is 318GiB/s, with occupancy of 67%. Volkov reports achieving near-peak memory bandwidth when running kernels whose occupancy is in the single digits.

A complementary example appears in Section 15.5, where the normalized cross-correlation kernel already runs at 100% occupancy yet is limited by the latency of its shared-memory reads rather than by arithmetic: each thread carries a single dependent chain of accumulations. Having each thread compute two independent output columns exposes enough instruction-level parallelism to hide that latency and more than doubles the throughput, with no change in occupancy at all—and it unlocks a further speedup from cheaper arithmetic (the __dp4a instruction) that the latency-bound version could not exploit. The sample illustrates that high occupancy is desirable but not sufficient: the schedulers still need independent instructions to issue while earlier ones are in flight.

The occupancy a kernel reaches is not entirely the compiler’s to choose. The __launch_bounds__ qualifier, attached to a __global__ function, tells ptxas the launch configuration the kernel must support – a maximum thread-block size, and optionally a minimum number of blocks to keep resident per multiprocessor:

__global__ void __launch_bounds__(256, 4) myKernel( ... ) { ... }

From those two numbers, ptxas derives a register budget: to fit the requested blocks of the requested size on one SM, each thread may use at most the SM’s register file divided by (threads per block × blocks per SM). The compiler caps the kernel’s register allocation to that budget, spilling to local memory if the code would otherwise want more. Left to itself, ptxas allocates registers to make each thread fast in isolation, which for a register-hungry kernel can pin occupancy far below what the hardware allows; __launch_bounds__ gives back some per-thread speed in exchange for keeping more warps resident. It is the in-source, per-kernel counterpart to -maxrregcount, which caps registers across a whole compilation unit, and to the runtime occupancy queries above: the flag and the queries bound or measure occupancy, while __launch_bounds__ compiles the kernel to hit a target.

Whether the trade is worth making depends on the kernel. The GlobalReads() kernel of Section 5.2.8 shows one side: its unrolling raises the register count and lowers occupancy, yet it still saturates memory bandwidth, so a forced register cap would only cost bandwidth. The cross-correlation kernel of Section 15.5 shows the other: there the extra registers buy the instruction-level parallelism that hides latency, and taking them away would remove the parallelism the kernel depends on. __launch_bounds__ helps only in the remaining case: a kernel whose occupancy is limited by registers, running a workload latency-bound enough that more resident warps would cover the stalls, and not so register-hungry that the enforced cap spills heavily. Cap the registers there and the added warps raise throughput; apply it anywhere else and it has no effect, or costs more in spills than the added occupancy recovers. The qualifier also makes a hard promise: a launch whose block exceeds the declared maximum fails rather than running, so the kernel must never be launched with a block larger than the maximum it declares. It belongs after profiling has named registers as the limiter, not on every kernel by reflex.