GPGPU：使用普通 PC 的后果答案

【问题标题】：GPGPU: Consequence of having a common PC in a warpGPGPU：使用普通 PC 的后果
【发布时间】：2014-10-17 21:02:12
【问题描述】：

我在一本书中读到，在波前或扭曲中，所有线程共享一个公共程序计数器。那么它的后果是什么？为什么这很重要？

【问题讨论】：

有人可以提供参考来证明“波前/扭曲具有公共程序计数器”这一说法吗？
这本书声称相反“每个GPU线程都有自己的标量寄存器、线程私有内存、线程执行状态、线程ID、独立执行和分支路径、有效程序计数器，并且可以独立寻址内存. 虽然当线程的 PC 相同时，一组线程（例如，32 个线程的扭曲）执行效率更高，但这不是必需的。（见books.google.de/…）
我自己找到了解决方案。重要的是上面引文中的“有效”一词。事实上，每个扭曲/波前只有 1 个程序计数器，因为硬件组织为 SIMD，您只能为所有 SIMD 通道发出 1 条指令（禁用这些通道，其控制路径发散并在下一个时钟周期执行它们）。但是从逻辑上讲，每个 CUDA 线程/OpenCL 工作项似乎有 1 台 PC。也就是说，为什么 SIMT 在 SIMD 硬件上有时称为 SPMD。（计算机架构：定量方法和haifux.org/lectures/267/Introduction-to-GPUs.pdf）

【解决方案1】：

NVIDIA GPU 一次执行 32 个线程（扭曲），而 AMD GPU 一次执行 64 个线程（波前）。控制逻辑、获取和数据路径的共享减少了面积并增加了性能/面积和性能/瓦特。

为了利用设计编程语言，开发人员需要了解如何合并内存访问以及如何管理控制流分歧。如果 warp/wavefront 中的每个线程采用不同的执行路径，或者如果每个线程访问显着不同的内存，那么设计的好处就会丧失并且性能会显着下降。

【讨论】：

【解决方案2】：

这意味着所有线程同时运行相同的命令。这对于确保在处理当前行时所有线程都完成上一行非常重要。例如，如果您需要将数据从一个线程传递到另一个线程，则需要确保数据已由第一个线程写入。因为程序计数器是共享的，所以你知道一旦写数据行完成，数据就存在于所有线程中。

【讨论】：

Warp 同步编程是一种特定于实现的优化。开发人员可以通过利用 warp 同步编程实现显着收益。但是，大多数 GPGPU 计算语言都没有很好地定义 warp 同步编程，因此鼓励开发人员使用所有必需的线程围栏和屏障，就像所有线程独立执行一样。

【解决方案3】：

正如其他一些答案所述，线程（扭曲/波前）在每个工作组的基础上彼此同步执行。对于开发人员来说，这意味着您需要特别注意任何分支/条件逻辑，因为如果组中的至少一个工作项达到“else”条件，则在执行该代码时所有其他工作项都会暂停。

那么 GPU 制造商为什么要这样做呢？缺少单独的程序计数器、分支预测和大型高速缓存存储器为芯片中的更多算术逻辑单元 (ALU) 节省了大量硅片。更多的 ALU 等于更多的工作组或并发线程。

【讨论】：

【解决方案4】：

像往常一样，了解幕后工作的原理有助于您提高性能。从 OCL 开发人员的角度来看，我们只知道

给定工作组中的工作项同时在单个计算单元的处理元素。（OCL 规范 1.2 - 第 3.2 节）。

这和如今 SIMT 架构的工作方式在谈论分支时会导致这种考虑（来自post）：

仅当条件不一致时才会执行两个分支在本地工作组中的线程之间，这意味着如果条件评估本地工作中工作项之间的不同值组，当前一代 GPU 将执行两个分支，但仅正确的分支会写入值并产生副作用。

这是完全正确的，但没有告诉您如何避免分歧（请注意，这里我们仍处于工作组级别）。

但知道一个工作组由 1 个或多个 Warp 组成，其中工作项共享一台 PC（不是在工作组级别）有时可以帮助您避免分歧。只有当一些工作项在一个warp中采用不同的路径时，你才会有分歧（两个分支都被执行）。考虑一下（source）：

if (threadIdx.x > 2) {...} else {...}

还有这个：

if (threadIdx.x / WARP_SIZE > 2) {...} else {...}

在第一种情况下，第一个扭曲（NVIDIA 的 32 个线程）内会有分歧。但不是在第二种情况下，无论工作组的大小如何，它始终是扭曲大小的倍数。显然这两个例子做的不是同一件事。但在某些情况下，您可能能够重新排列数据（或找到另一个技巧）以保持第二个示例的理念。

这似乎与现实相去甚远，但现实生活中的一个例子就是还原。通过在“SIMD 友好结构”中命令您的操作，您可以在每个阶段放下一些经线（因此为来自另一个工作组的其他人留出空间）。有关完整说明和代码，请参阅此 whitepaper 中的“利用交换性”部分。

【讨论】：