gpu - GPU组如何进入扭曲/波阵面？

Question

我的理解是，warp是在运行时通过任务调度程序定义的一组线程，CUDA的性能关键部分是warp中的线程差异，有没有办法很好地猜测硬件将如何构造warp在线程块中？

例如，我已经在一个线程块中启动了一个具有1024个线程的内核，那么线程的排列方式如何，我能否从线程索引中得知(或至少可以作一个好猜测)？

由于这样做，可以使给定经线内的线散度最小。

Answer 1

经纱内部的线程排列取决于实现，但是atm我始终遇到相同的行为:

一个warp由32个线程组成，但是warp scheduller每次将发出1条指令以暂停warp(16个线程)

如果使用1D块(仅threadIdx.x尺寸有效)，则扭曲调度程序将针对 threadIdx.x =(0..15)(16..31)发出1条指令 ...等

如果您使用2D块(threadIdx.x和threadIdx.y尺寸有效)，那么扭曲调度程序将尝试按照以下方式发出:

threadIdx.y = 0 threadIdx.x =(0 ..15)(16..31) ...等

因此，具有连续threadIdx.x组件的线程将以16为一组执行同一条指令。