GPU寄存器使用

  • 3 replies
  • 821 views
GPU寄存器使用
« 于: 九月 07, 2022, 03:52:21 pm »
手册上说若寄存器不足,会使用本地内存,而本地内存位于全局内存上,理论上不会有限制。
3090卡1个SM理论可以同时调度1536个线程,可是在实际使用中之恩那个调度1152个线程,通过ncu查看当前函数执行时一个线程使用寄存器56个【65536/1152约等于56】.这说明寄存器限制了SM可同时调度的线程数量。
请问为什么没有使用本地内存呢?

Re: GPU寄存器使用
« 回复 #1 于: 九月 11, 2022, 03:09:29 pm »
手册上说若寄存器不足,会使用本地内存,而本地内存位于全局内存上,理论上不会有限制。
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优化冷知识系列》)。

Re: GPU寄存器使用
« 回复 #2 于: 九月 15, 2022, 10:54:47 am »
好的,谢谢您的解答。我之前还以为如果寄存器被使用完,剩余的部分编译器会自动使用local memory来代替,从而不会出现限制occupancy的情况。

Re: GPU寄存器使用
« 回复 #3 于: 九月 16, 2022, 03:35:44 pm »
好的,谢谢您的解答。我之前还以为如果寄存器被使用完,剩余的部分编译器会自动使用local memory来代替,从而不会出现限制occupancy的情况。

如果用到了255个寄存器的话,的确只能使用local了,这个是必然的。但是你之前的问题,显然你并没有使用到255个,所以你最后的回复的结论并不成立。建议重新看一下。