1 parallel-processing synchronization cuda gpu
我有点困惑Warps发散并需要通过__syncthreads()函数同步.块中的所有元素都以SIMT方式处理相同的代码.他们怎么可能不同步？它与调度程序有关吗？不同的warp会有不同的计算时间吗？为什么使用时有开销__syncthreads()？
让我们说我们在一个区块中有12个不同的Warps已经完成了他们的工作.所以现在有空转,其他经线得到他们的计算时间.或者他们还有计算时间来完成这项__syncthreads()功能吗？
首先让我们小心术语.由于代码中的控制结构(if,while等),Warp divergence指的是单个warp中的线程采用不同的执行路径.您的问题实际上与warp和warp调度有关.
尽管SIMT模型可能表明所有线程都是锁步执行的,但事实并非如此.首先,不同块中的线程是完全独立的.它们可以相对于彼此以任何顺序执行.关于同一块中线程的问题,让我们首先观察一个块最多可以有1024个(或者更多)线程,但今天的SM(SM或SMX是处理线程块的GPU内部的"引擎")don'有1024个cuda核心,所以理论上,SM在锁步中执行线程块的所有线程甚至都不可能.请注意,单个线程块在单个SM上执行,而不是同时跨越所有(或多个)SM.因此,即使一台机器具有512或更多的总cuda核心,它们也不能全部用于处理单个线程块的线程,因为单个线程块在单个SM上执行.(这样做的一个原因是,特定于SM的资源,如共享内存,可以被线程块中的所有线程访问.)
那会发生什么？事实证明每个SM都有一个warp调度程序.一经无非是那个被组合在一起,安排在一起,一起执行32个线程一个集合.如果一个线程块有1024个线程,那么它每个warp有32个32个线程的warp.现在,例如,在Fermi上,一个SM有32个CUDA核心,所以考虑一个SM以锁步方式执行warp是合理的(那个是会发生什么,费米).通过锁步,我的意思是(忽略经线发散的情况,以及指令级并行的某些方面,我试图在这里保持解释简单......)在前一条指令之前,没有执行warp中的指令已经由warp中的所有线程执行.因此,Fermi SM实际上只能在任何给定时刻执行线程块中的一个warp.该线程块中的所有其他warp排队等待,准备好等待.
现在,当一个warp的执行由于任何原因遇到停顿时,warp调度程序可以自由地移动那个warp out并带来另一个准备好的warp(这个新的warp甚至可能不是来自同一个threadblock,但是我希望到现在你可以看到,如果一个threadblock中有超过32个线程,并不是所有的线程实际上都是以锁步方式执行的.一些经线在其他经线之前正在进行.
这种行为通常是可取的,除非它不是.有时您不希望线程块中的任何线程超出某个点,直到满足条件.这__syncthreads()是为了什么.例如,您可能正在将数据从全局复制到共享内存,并且您不希望在共享内存已正确填充之前开始任何线程块数据处理. __syncthreads()确保所有线程都有机会在任何线程超出障碍之前复制其数据元素,并且可能开始计算现在驻留在共享内存中的数据.
开销有__syncthreads()两种风格.首先,处理与此内置函数相关的机器级指令的成本非常低.其次,__syncthreads()通常会产生强制warp调度程序和SM对线程块中的所有warp进行洗牌的效果,直到每个warp都遇到障碍.如果这很有用,那很好.但如果不需要,那么你就花时间去做一些不需要的事情.因此,建议不要只是自由地__syncthreads()通过你的代码.在需要的地方谨慎使用它.如果你可以制作一个不像另一个那样使用它的算法,那么该算法可能会更好(更快).