use*_*128 2 gpu gpgpu gpu-warp
我的理解是，warp 是通过任务调度程序在运行时定义的一组线程，CUDA 的一个性能关键部分是 warp 内线程的分歧，有没有办法很好地猜测硬件将如何构造 warp在线程块内？
例如，我启动了一个线程块中包含 1024 个线程的内核，扭曲是如何排列的，我可以从线程索引中看出（或至少做出一个很好的猜测）吗？
因为通过这样做，可以最大限度地减少给定经纱内线程的发散。
warp 内的线程排列取决于实现，但 atm 我总是遇到相同的行为：
一个warp由32个线程组成，但warp调度程序每次都会发出1条指令来停止一个warp（16个线程）
如果您使用 1D 块（仅 threadIdx.x 维度有效），那么 warp 调度程序将发出 1 条指令threadIdx.x = (0..15) (16..31) ... 等
如果您使用 2D 块（threadIdx.x 和 threadIdx.y 维度有效），那么 warp 调度程序将尝试按照以下方式发出：
threadIdx.y = 0 threadIdx.x = (0 ..15) (16..31) ... 等等
因此，具有连续 threadIdx.x 组件的线程将以 16 个为一组执行相同的指令。