寄存器与共享内存限制

Each SM owns a fixed pool of registers and shared memory. Every resident block must carve its share from these pools.

Your kernel uses some number of registers per thread. Multiply by threads per block and you see one blocks register cost.

If each thread needs many registers, fewer threads fit, so fewer blocks stay resident. Register-heavy kernels lose occupancy.

Compile with this flag and nvcc prints registers per thread, letting you spot when a kernel is too register hungry.

You can force a ceiling with a launch bound so the compiler spends fewer registers and lets more warps stay resident.

Each block can request shared memory. The SM only fits as many blocks as its shared pool, often 48 to 100 KB, allows.

Ask for a large __shared__ tile and only one or two blocks fit per SM. Big tiles trade occupancy for data reuse.

The SM computes blocks allowed by registers, by shared memory, and by the warp cap. The smallest of these decides occupancy.

When a thread needs more registers than allowed, extras spill to slow local memory. Spills can hurt more than low occupancy.

Cutting registers or shared memory raises occupancy, but too aggressive a cut causes spills. The sweet spot is a balance.

Always read the actual register and shared usage from the compiler before tuning. Guessing usually picks the wrong knob.

Recall how the SM decides how many blocks can be resident.

You saw that registers and shared memory are fixed SM budgets, and the tightest limit caps occupancy. Watch for spills. 🧮