求问几个有关优化方面的问题

  • 2 replies
  • 444 views
求问几个有关优化方面的问题
« 于: 十月 12, 2022, 05:04:09 pm »
问题1(图一所示):
 如Nsight Compute提示所示,如何有效的提高计算吞吐量和内存带宽?
问腿2 :
在这个项目中,我一个block内有128个线程,每个线程往全局内存中每次写一个ulonglong3(24字节),求问,这算是满足对齐的要求了吗?(每个线程写的地址都是连续的,相当于写了24*128字节)
问题3:(图二所示)
如Nsight Compute提示所示,如何提高活跃的线程束数量呢?我已经使用__launch_bounds__来限制寄存器的数量来提高了,请问还有其他方法来优化此问题吗?
问题4:(图三和图四)
如Nsight Compute提示所示,如何解决线程束发散导致的问题呢?(算法以及无法减少if分支)

感谢各位[名词6]的指导

Re: 求问几个有关优化方面的问题
« 回复 #1 于: 十月 13, 2022, 12:26:36 pm »
问题1(图一所示):
 如Nsight Compute提示所示,如何有效的提高计算吞吐量和内存带宽?
问腿2 :
在这个项目中,我一个block内有128个线程,每个线程往全局内存中每次写一个ulonglong3(24字节),求问,这算是满足对齐的要求了吗?(每个线程写的地址都是连续的,相当于写了24*128字节)
问题3:(图二所示)
如Nsight Compute提示所示,如何提高活跃的线程束数量呢?我已经使用__launch_bounds__来限制寄存器的数量来提高了,请问还有其他方法来优化此问题吗?
问题4:(图三和图四)
如Nsight Compute提示所示,如何解决线程束发散导致的问题呢?(算法以及无法减少if分支)

感谢各位[名词6]的指导

看上去你的SM和Memory没有充分忙碌,你的图综合说明了:
occupancy较低,理论occupancy才能到大约12个warps左右(看图估计的)/ SM, 也就是如果是8.6的卡才20%~30%,achieved occupancy更是大约10个warps/SM(~20%)。如果你认为你已经用launch bounds限制了,则可能你的launch bounds设定有误(或者你因为没有说明的苦衷而不能更激进的设定)。这里可以考虑尝试提升一下,看看是否有效果。

此外,可以尝试消除blocks内部的长尾效应(部分warps提前结束,部分warps拖后结束,导致提前结束的warps依然占用部分资源,从而achieved occupancy比theoretical的occupancy低一些。这里也可以考虑修一下。不过这个不是主要因素。

你的图看来主要矛盾点在访存,而不是SM使用率。这样综上尝试提升了occupancy后,看看能否让SM和MEM的忙碌程度提升一下(主要是访存),这样看看能否提速。

N卡只支持从每个线程的角度的1B, 2B, 4B, 8B, 16B粒度的访存的。你的"一次性"24B每个线程的写入,因为粒度和对齐的原因,可能会被拆分成3次8B的写入的,因此不是一次。从这个拆分的角度看,每个warp的这个写入,不是合并的(warp的32个8B写入,中间有很多16B的空洞)。不过因为有cache的存在,在8.6上这种写入其实还好,不太要紧。不过如果能提升,可以尝试修改下。例如改成32B每个线程的写入量,这样1个warp会拆分成两条无空洞的合并的16B写入,不一定有正面效果(因为写入量变大了),不过你可以试试看,特别是在有ECC的卡上,能减少空洞导致的读取--修改--写入的访存放大。(普通卡一般情况下可能无此放大问题)。

你说的因为分支,每个warp实际有效的计算量是13.50/32 = 42%,这的确是个问题。如果算法的确会导致warp内分支的话,根据每个分支体的代码量不同,可能的确没办法。也可能你1个warp内部,通过任务重排,例如将40%的分支1,60%的分支2,重新排列成前40%的线程计算分支1,后60%的线程计算分支2,可能会提升效率。如果你假设block有8个warps,256个线程:原本每个warp都是40%效率分支1 ---> 60%效率分支2。变成了前40%的线程(也就是warp0,1,2无分支的全效率执行分支1,和warp4,5,6,7全效率的执行分支2,和warp 3依然部分效率的两个分支都执行),这样的任务重拍可以有效的减少divergent branch, 让数字非常好看。但是不一定总是有正面效果。(因为太小的分支不够block里头任务重排的的成本)。

此外,目前等待MEM的延迟依然是主要矛盾(Long SB,倒数第二个图),看看能重写代码的结果或者算法的逻辑,减少这一点。





Re: 求问几个有关优化方面的问题
« 回复 #2 于: 十月 13, 2022, 02:11:13 pm »
看上去你的SM和Memory没有充分忙碌,你的图综合说明了:
occupancy较低,理论occupancy才能到大约12个warps左右(看图估计的)/ SM, 也就是如果是8.6的卡才20%~30%,achieved occupancy更是大约10个warps/SM(~20%)。如果你认为你已经用launch bounds限制了,则可能你的launch bounds设定有误(或者你因为没有说明的苦衷而不能更激进的设定)。这里可以考虑尝试提升一下,看看是否有效果。

此外,可以尝试消除blocks内部的长尾效应(部分warps提前结束,部分warps拖后结束,导致提前结束的warps依然占用部分资源,从而achieved occupancy比theoretical的occupancy低一些。这里也可以考虑修一下。不过这个不是主要因素。

你的图看来主要矛盾点在访存,而不是SM使用率。这样综上尝试提升了occupancy后,看看能否让SM和MEM的忙碌程度提升一下(主要是访存),这样看看能否提速。

N卡只支持从每个线程的角度的1B, 2B, 4B, 8B, 16B粒度的访存的。你的"一次性"24B每个线程的写入,因为粒度和对齐的原因,可能会被拆分成3次8B的写入的,因此不是一次。从这个拆分的角度看,每个warp的这个写入,不是合并的(warp的32个8B写入,中间有很多16B的空洞)。不过因为有cache的存在,在8.6上这种写入其实还好,不太要紧。不过如果能提升,可以尝试修改下。例如改成32B每个线程的写入量,这样1个warp会拆分成两条无空洞的合并的16B写入,不一定有正面效果(因为写入量变大了),不过你可以试试看,特别是在有ECC的卡上,能减少空洞导致的读取--修改--写入的访存放大。(普通卡一般情况下可能无此放大问题)。

你说的因为分支,每个warp实际有效的计算量是13.50/32 = 42%,这的确是个问题。如果算法的确会导致warp内分支的话,根据每个分支体的代码量不同,可能的确没办法。也可能你1个warp内部,通过任务重排,例如将40%的分支1,60%的分支2,重新排列成前40%的线程计算分支1,后60%的线程计算分支2,可能会提升效率。如果你假设block有8个warps,256个线程:原本每个warp都是40%效率分支1 ---> 60%效率分支2。变成了前40%的线程(也就是warp0,1,2无分支的全效率执行分支1,和warp4,5,6,7全效率的执行分支2,和warp 3依然部分效率的两个分支都执行),这样的任务重拍可以有效的减少divergent branch, 让数字非常好看。但是不一定总是有正面效果。(因为太小的分支不够block里头任务重排的的成本)。

此外,目前等待MEM的延迟依然是主要矛盾(Long SB,倒数第二个图),看看能重写代码的结果或者算法的逻辑,减少这一点。





十分感谢,看了您的建议受益匪浅