GPU上的分支预测

Zk1*_*001 11 cuda gpu gpgpu opencl

我有一个关于GPU中分支预测的问题.据我所知,在GPU中,它们通过分支进行预测.

例如,我有一个这样的代码:

if (C)
 A
else
 B
Run Code Online (Sandbox Code Playgroud)

因此,如果A需要40个周期而B需要50个周期来完成执行,如果假设一个warp,A和B都被执行,那么完成这个分支需要总共90个周期吗?或者它们是否重叠A和B,即,当执行A的某些指令时,等待内存请求,然后执行B的某些指令,然后等待内存,依此类推?谢谢

tal*_*ies 12

到目前为止,所有发布的支持CUDA的架构都像SIMD机器一样运行.当warp中存在分支差异时,两个代码路径都由warp中的所有线程执行,而不跟随活动路径的线程执行NOP的功能等价物(我想我记得有条件执行)标志附加到warp中的每个线程,允许非执行线程被屏蔽掉).

所以在你的例子中,90个周期的答案可能是比替代方案更接近实际发生的事情.

  • 这是不正确的,每个warp处理条件执行,而不是每半warp处理.值得指出的是,如果分支条件在warp中不发散(例如`if(threadIdx.x> 64)`,则没有发散执行). (3认同)
  • 正如ptx_isa.pdf:"如果warp的线程通过依赖于数据的条件分支发散,则warp会串行执行所采用的每个分支路径,禁用不在该路径上的线程,并且当所有路径完成时,线程会聚回到相同的执行路径." 因此,在PTX中有一个条件分支,但是Warp的所有线程必须同时采用或不采用该分支才能统一(以获得性能) (2认同)
  • 在Fermi架构(当前架构)上,每个warp也处理内存事务.旧的GT200(又名特斯拉架构)每半个经线处理内存事务.费米的任何地方都没有使用半经线. (2认同)