GPU 如何将线程分组为扭曲/波前？

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

例如，我启动了一个线程块中包含 1024 个线程的内核，扭曲是如何排列的，我可以从线程索引中看出（或至少做出一个很好的猜测）吗？

因为通过这样做，可以最大限度地减少给定经纱内线程的发散。

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

一个warp由32个线程组成，但warp调度程序每次会为半个warp发出1条指令（16个线程）

• 如果您使用 1D 块（仅 threadIdx.x 维度有效），那么 warp 调度程序将发出 1 条指令threadIdx.x = (0..15) (16..31) ... etc

• 如果您使用 2D 块（threadIdx.x 和 threadIdx.y 维度有效），那么 warp 调度程序将尝试按照以下方式发出：

线程Idx.y = 0 线程Idx.x = (0 ..15) (16..31) ... etc

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