CUDA Occupancy: what to measure and why

The mechanism

What cuda occupancy means

A kernel launch creates thread blocks, and the GPU assigns those blocks to streaming multiprocessors (SMs). Registers, shared memory, threads, blocks, and architectural limits determine how many blocks can reside on an SM at once. Resident blocks contribute active warps. Achieved occupancy reports the active-warps ratio observed during execution; theoretical occupancy predicts the limit from the launch and resource requirements.

• Low occupancy can reduce the scheduler’s ability to hide memory and instruction latency.
• Registers per thread and shared memory per block can limit the number of resident blocks.
• A larger block does not guarantee higher occupancy; allocation granularity and resource use matter.
• After enough warps are available to hide the relevant latency, additional occupancy may not improve runtime.

A minimal pattern

Change one thing at a time

The snippets isolate the address or launch pattern being discussed. A real optimization still has to preserve output and be measured in the complete kernel.

constexpr int blockSize = 32;
const int gridSize = (n + blockSize - 1) / blockSize;
vectorAdd<<<gridSize, blockSize>>>(a, b, c, n);

constexpr int blockSize = 256;
const int gridSize = (n + blockSize - 1) / blockSize;
vectorAdd<<<gridSize, blockSize>>>(a, b, c, n);

Profiler evidence

What to measure

Compare correctness first, then kernel duration, achieved occupancy, active warps, eligible warps, and the resource identified as the occupancy limiter. Keep the input size and GPU fixed. If you change block size, verify that the grid still covers exactly the required elements.

Reason about the result

How to interpret the measurements

If occupancy increases and runtime falls, the original launch may not have exposed enough ready work. If occupancy increases but runtime stays flat, the kernel may already have had enough warps to hide its dominant latency. If runtime becomes worse, the new launch may have changed instruction, cache, scheduling, or tail behavior. Occupancy is evidence about available parallel work—not a standalone performance score.

Avoid false conclusions

Common mistakes

• Treating 100% occupancy as the optimization target for every kernel.
• Comparing different inputs, GPUs, compiler options, or profiling passes as if only block size changed.
• Using theoretical occupancy without checking achieved occupancy and runtime.
• Reducing register use solely to raise occupancy when spilling creates slower local-memory traffic.