Occupancy tuning

What bounds residency

Per SM there is a maximum of resident warps, resident blocks, registers, and shared memory; per block there is a maximum of threads and a rounded shared-memory allocation. Residency is the smallest of those divisions, and the binding one has a name.

int blocks = 0;
cudaOccupancyMaxActiveBlocksPerMultiprocessor(&blocks, mm_tiled, 256, 0);
int warps = blocks * 256 / 32;

ncu reports achieved and theoretical occupancy separately, plus the limiter. A large gap between achieved and theoretical is usually the tail of the grid, not a resource.

The occupancy APIs, the default shared-memory limit, and the per-SM maxima are architecture and release specific; read them from your device query and confirm the API in the guide for your toolkit.

Registers are usually the binding one

-Xptxas -v prints registers per thread. __launch_bounds__(256, 2) asks ptxas to keep the count low enough for two blocks of 256 per SM; -maxrregcount=64 clamps every kernel in the file. Both buy residency with spills, and a spilled accumulator is a DRAM access where a register used to be.

More warps is a means

Extra resident warps cover latency, but they also contend for shared memory, the load-store units, and L1. A kernel already at DRAM peak gains nothing from another resident block and can lose from cache pressure. The target is the resource the profiler names at the dominant stall, not a percentage on its own.

Things that break

A static __shared__ array above 48 KiB is rejected at compile time; the fix is dynamic shared memory plus an opt-in with cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, bytes). The occupancy calculator ignores the L1 and shared-memory carveout, so its answer is an upper bound. Register counts move with the toolkit and the architecture, so a launch bound tuned on one release can spill on the next.