Pick an architecture, set block size, registers per thread, and shared memory per block. Get theoretical occupancy and the limiting resource.
Configuration
Architecture
sm_80 · Ampere A100
256
32
Read the real count with nvcc -Xptxas -v
8 KB
Theoretical occupancy
100.0%
Resident blocks
8
64 / 64 warp
Bottleneck resource:Thread / block-count cap and Registers. is already the best blockDim under these resources.
How many blocks each limit allows
Thread / block-count cap bottleneck8 blocks
min(32, floor(64 / 8))
Registers bottleneck8 blocks
floor(65536 / 8192) each block uses 8192
Shared memory20 blocks
floor(164KB / 8.00KB)
Occupancy vs blockDim
Registers 32 · shared memory 8KB
The sawtooth is allocation granularity: a few extra threads can cross a cliff.
322565127681024
Matching kernel constraints
CUDA C++
1// Target: blockDim = 256, at least 8 resident blocks per SM2__global__void__launch_bounds__(256,8)3myKernel(constfloat*__restrict__in,float*out,intn){4__shared__floattile[2048];// 8 KB5...6}78// Confirm real register count at compile time:9// nvcc -O3 -arch=sm_80 -Xptxas -v -c kernel.cu10// Let CUDA recommend blockDim at runtime:11// cudaOccupancyMaxPotentialBlockSize(&minGrid, &blockSize, myKernel, 0, 0);
This is theoretical occupancy. The model matches the CUDA Occupancy Calculator, including register allocation at warp granularity and shared memory aligned to the allocation unit. Achieved occupancy at runtime is usually lower because of tail effects and load imbalance. Nsight Compute is the final word.