手册上说若寄存器不足，会使用本地内存，而本地内存位于全局内存上，理论上不会有限制。
3090卡1个SM理论可以同时调度1536个线程，可是在实际使用中之恩那个调度1152个线程，通过ncu查看当前函数执行时一个线程使用寄存器56个【65536/1152约等于56】.这说明寄存器限制了SM可同时调度的线程数量。
请问为什么没有使用本地内存呢？
正常现象，编译器选择使用较多的寄存器，而拒绝使用local，从而导致了较低的occupancy（如同你说的，1152/1536 = 75%）的现象是正常的。因为默认情况不会自动追求100% occupancy（注意100%的occupancy也不一定会有更好的性能）。
不过你总是可以手工来，强制寄存器使用更少的寄存器，和达到更大的occupancy的。编译的时候选择-maxrregcount或者源代码里使用__launch_bounds__修饰kernel，都可以手工做到这点。
如果需要手工要求编译器这样，请搜索本论坛或者网上的这2个关键字，获取正确的使用方法。（也可以看看我们最近的《CUDA优化冷知识系列》）。
