我在一本书中读到，在波前或扭曲中，所有线程共享一个公共程序计数器。那么它的后果是什么？为什么这很重要？
问问题
637 次
4 回答
4
NVIDIA GPU 一次执行 32 个线程（扭曲），而 AMD GPU 一次执行 64 个线程（波前）。控制逻辑、获取和数据路径的共享减少了面积并增加了性能/面积和性能/瓦特。
为了利用设计编程语言，开发人员需要了解如何合并内存访问以及如何管理控制流分歧。如果扭曲/波前中的每个线程采用不同的执行路径，或者如果每个线程访问显着不同的内存，那么设计的好处就会丧失并且性能将显着降低。
于 2014-08-25T04:07:31.683 回答
2
这意味着所有线程同时运行相同的命令。这对于确保所有线程在处理当前行时都完成了上一行非常重要。例如，如果您需要将数据从一个线程传递到另一个线程，则需要确保数据已由第一个线程写入。因为程序计数器是共享的，所以您知道一旦写入数据行完成，数据就存在于所有线程中。
于 2014-08-24T16:07:03.103 回答
2
像往常一样，了解幕后工作的原理有助于您提高性能。从 OCL 开发人员的角度来看，我们只知道
给定工作组中的工作项在单个计算单元的处理元素上同时执行。（OCL 规范 1.2 - 第 3.2 节）。
这和 SIMT 架构如今的工作方式在谈论分支时会导致这种考虑（来自这篇文章）：
仅当条件在本地工作组中的线程之间不一致时才会执行两个分支，这意味着如果条件在本地工作组中的工作项之间评估为不同的值，则当前一代 GPU 将执行两个分支，但仅执行正确的分支将写入值并产生副作用。
这是完全正确的，但没有告诉您如何避免分歧（请注意，这里我们仍处于工作组级别）。
但是知道一个工作组由 1 个或多个 Warp 组成，其中工作项共享一台 PC（不是在工作组级别）有时可以帮助您避免分歧。只有当扭曲中的某些工作项采用不同的路径时，您才会有分歧（两个分支都被执行）。考虑一下（来源）：
if (threadIdx.x > 2) {...} else {...}
和这个：
if (threadIdx.x / WARP_SIZE > 2) {...} else {...}
在第一种情况下，第一个扭曲（NVIDIA 的 32 个线程）内会有分歧。但不是在第二种情况下，无论工作组的大小如何，它始终是扭曲大小的倍数。显然这两个例子做的不是同一件事。但在某些情况下，您可能能够重新排列数据（或找到另一个技巧）以保持第二个示例的理念。
这似乎与现实相去甚远，但现实生活中的一个例子是减少。通过在“SIMD 友好结构”中命令您的操作，您可以在每个阶段放下一些经线（因此为来自另一个工作组的其他人留出空间）。有关完整说明和代码，请参阅本白皮书中的“利用交换性”部分。
于 2014-08-27T08:18:36.903 回答
2
正如其他一些答案所述，线程（扭曲/波前）在每个工作组的基础上彼此同步执行。对于开发人员来说，这意味着您需要特别注意任何分支/条件逻辑，因为如果组中的至少一个工作项达到“else”条件，则在执行该代码时所有其他工作项都会暂停。
那么GPU制造商为什么要这样做呢？缺少单独的程序计数器、分支预测和大型高速缓存存储器为芯片中的更多算术逻辑单元 (ALU) 节省了大量硅片。更多的 ALU 等于更多的工作组或并发线程。
于 2014-08-25T12:52:50.200 回答