我的理解是warp是在运行时通过任务调度程序定义的一组线程，CUDA的一个性能关键部分是warp中线程的分歧，有没有办法很好地猜测硬件将如何构建warp在线程块内？
例如，我在一个线程块中启动了一个具有 1024 个线程的内核，如何安排扭曲，我可以从线程索引中判断（或至少做出一个好的猜测）吗？
因为通过这样做，可以最大限度地减少给定warp中线程的分歧。
问问题
1333 次
2 回答
4
经线内的线程排列取决于实现，但 atm 我总是经历相同的行为：
一个warp由32个线程组成，但是warp调度器每次都会发出1条指令用于halp一个warp（16个线程）
如果您使用 1D 块（只有 threadIdx.x 维度有效），那么 warp 调度程序将为threadIdx.x = (0..15) (16..31) ... 等发出 1 条指令
如果您使用 2D 块（threadIdx.x 和 threadIdx.y 尺寸有效），那么 warp 调度程序将尝试按照以下方式发出：
threadIdx.y = 0 threadIdx.x = (0 ..15) (16..31) ... 等等
因此，具有连续 threadIdx.x 组件的线程将以 16 个为一组执行相同的指令。
于 2013-03-07T09:27:39.703 回答
2
一个 warp 由 32 个线程组成，它们将同时执行。在任何给定时间，一批 32 将在 GPU 上执行，这称为warp。
我还没有发现任何地方可以说明您可以控制接下来要执行的warp，您唯一知道的是它由32个线程组成，并且线程块应该始终是该数字的倍数。
单个块中的线程将在单个多处理器上执行，共享软件数据缓存，并可与同一块中的线程同步和共享数据；warp 将始终是来自单个块的线程子集。
关于内存操作和延迟，还有这个：
当 warp 中的线程发出设备内存操作时，由于长的内存延迟，该指令将花费很长时间，可能需要数百个时钟周期。主流架构会添加缓存层次结构以减少延迟，Fermi 确实包含一些硬件缓存，但大多数 GPU 是为流或吞吐量计算而设计的，其中缓存无效。相反，这些 GPU 通过使用高度的多线程来容忍内存延迟。Tesla 在每个多处理器上最多支持 32 个活动 warp，而 Fermi 最多支持 48个。当一个 warp 在内存操作上停止时，多处理器会选择另一个就绪的 warp 并切换到那个。通过这种方式，只要有足够的并行度让内核保持忙碌，内核就可以高效工作。
关于将线程块划分为经线，我发现了这一点：
如果块是 2D 或 3D，线程按第一个维度排序，然后是第二个维度，然后是第三个维度 - 然后分成 32 个经线
于 2013-03-07T09:18:35.633 回答