有哪些因素导致了achieved occupancy小于theoretical occupancy?

  • 4 replies
  • 159 views
*

sisiy

  • *****
  • 180
    • 查看个人资料
问个问题:kernel通常有theoretical occupancy和achieved occupancy之分,theoretical occupancy代表理论上SM能够容纳的warp数目(SM资源受限)与SM能容纳的warp数目的最大值的比值,而achieved occupancy代表实际运行时每个时刻各SM常驻warp数目与SM能容纳的warp数目的最大值的比值的平均值。但是实际上achieved occupancy要小于theoretical occupancy,有哪些因素导致了achieved occupancy小于theoretical occupancy?是block级别的调度或warp级别的调度产生的调度开销导致的吗?

Re: 有哪些因素导致了achieved occupancy小于theoretical occupancy?
« 回复 #1 于: 十一月 20, 2019, 09:05:51 pm »
问个问题:kernel通常有theoretical occupancy和achieved occupancy之分,theoretical occupancy代表理论上SM能够容纳的warp数目(SM资源受限)与SM能容纳的warp数目的最大值的比值,而achieved occupancy代表实际运行时每个时刻各SM常驻warp数目与SM能容纳的warp数目的最大值的比值的平均值。但是实际上achieved occupancy要小于theoretical occupancy,有哪些因素导致了achieved occupancy小于theoretical occupancy?是block级别的调度或warp级别的调度产生的调度开销导致的吗?

我们要考虑到,导致achieved和theoretical occupancy不同的原因,主要在于,前者除了考虑后者需要考虑的各种影响occupancy的情况下(各种资源),实际运行期间,一个block上了SM后,如果其他有warps中途死掉,则这些死掉的warps占用的资源可能并不会释放,等到block整体结束了,才能释放资源,然后下一个block上去。而achieved occupancy只对每个周期的存活warps数量,在时间上,和跨SM上进行积分,然后按照一个SM有64个warps,或者32个warps(和计算能力有关),将积分结果和这个最大值相除,从而得到achieved occupancy.

这样的话,一旦有的warp执行的特别长,有的执行的特别短(例如被你提前用if判断,给return结束掉的),那么就会显著的降低achieved occupancy, 因为死掉的warps还占着坑不拉屎。影响其他人上来执行的机会(而理论occupancy并不考虑这点,也无法考虑实际执行中的这个情况)。

例如如下代码:
以<<<N, 512>>>启动的某kernel:
if (threadIdx.x > 384) return;
//其他人继续埋头苦干

这种代码就有可能会显著降低实际的occupancy.

建议的解决方案:
有用的:想法让工作量尽量平衡的分布在各个warps间,不要有人拖尾很久才能完成,而有人很快的就能完成。平衡比较重要。以及,如果warp内部还有人快速的结束了,有人还在苦干,则你需要将achieved occupancy和divergent branch的情况结合看。平衡+不要有人特别拖沓是关键。注意有的算法本身无法避免这点。

给老板看的,1秒见效:
不要让有人提前结束,在所有线程结束前来一次achieved occupancy, 来维持一种虚假的“活性”,这种情况下achieved occupancy立刻就上来了,然而并没有实际效果。但是可以立刻完成本次报告内容。

Re: 有哪些因素导致了achieved occupancy小于theoretical occupancy?
« 回复 #2 于: 十一月 20, 2019, 09:09:51 pm »
我们要考虑到,导致achieved和theoretical occupancy不同的原因,主要在于,前者除了考虑后者需要考虑的各种影响occupancy的情况下(各种资源),实际运行期间,一个block上了SM后,如果其他有warps中途死掉,则这些死掉的warps占用的资源可能并不会释放,等到block整体结束了,才能释放资源,然后下一个block上去。而achieved occupancy只对每个周期的存活warps数量,在时间上,和跨SM上进行积分,然后按照一个SM有64个warps,或者32个warps(和计算能力有关),将积分结果和这个最大值相除,从而得到achieved occupancy.

这样的话,一旦有的warp执行的特别长,有的执行的特别短(例如被你提前用if判断,给return结束掉的),那么就会显著的降低achieved occupancy, 因为死掉的warps还占着坑不拉屎。影响其他人上来执行的机会(而理论occupancy并不考虑这点,也无法考虑实际执行中的这个情况)。

例如如下代码:
以<<<N, 512>>>启动的某kernel:
if (threadIdx.x > 384) return;
//其他人继续埋头苦干

这种代码就有可能会显著降低实际的occupancy.

建议的解决方案:
有用的:想法让工作量尽量平衡的分布在各个warps间,不要有人拖尾很久才能完成,而有人很快的就能完成。平衡比较重要。以及,如果warp内部还有人快速的结束了,有人还在苦干,则你需要将achieved occupancy和divergent branch的情况结合看。平衡+不要有人特别拖沓是关键。注意有的算法本身无法避免这点。

给老板看的,1秒见效:
不要让有人提前结束,在所有线程结束前来一次achieved occupancy, 来维持一种虚假的“活性”,这种情况下achieved occupancy立刻就上来了,然而并没有实际效果。但是可以立刻完成本次报告内容。

此外,如果你没有好办法能手工平衡任务,和降低拖尾的某些线程,则最快的方案可能是在你的算法要求和卡实际情况允许的情况下,采用更小的block大小,这往往会导致任务会分布的更加离散,和资源分配的单位更小。这样的情况下,当一个kernel要处理的任务(假设以warp为单位好了),是长短不一,但是比较符合随机分布的情况下,往往可能效果会好一些(你想想为何)。

Re: 有哪些因素导致了achieved occupancy小于theoretical occupancy?
« 回复 #3 于: 十一月 20, 2019, 10:31:04 pm »
所以,用较小的线程块就是为了降低提前完成任务的线程束等待最后完成任务的线程束的时间?

Re: 有哪些因素导致了achieved occupancy小于theoretical occupancy?
« 回复 #4 于: 十一月 21, 2019, 12:47:37 pm »
程序代码: [选择]
所以,用较小的线程块就是为了降低提前完成任务的线程束等待最后完成任务的线程束的时间?应该是让执行时间不同的warp分散在不同的block上,这样执行时间特别短的warp都在同一个block中,执行时间特别长的warp也都在同一个block中,使得执行时间特别短的warp和执行时间特别长的warp不在同一个block中,也就不存在死掉的warps还占着坑。