是否可以保证WaveFront (OpenCL)中的所有线程总是同步的？

Question

众所周知，有翘曲(在CUDA中)和WaveFront (在OpenCL中)：http://courses.cs.washington.edu/courses/cse471/13sp/lectures/GPUsStudents.pdf

​

​

• WARP in CUDA：http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#simt-architecture

4.1。SIMT体系结构 ..。 翘曲一次执行一条公共指令，所以当一个翘曲的所有32个线程都同意它们的执行路径时，就能实现完全的效率。如果翘曲的线程通过依赖于数据的条件分支发散，则翘曲会依次执行所采取的每个分支路径，禁用不在该路径上的线程，并且当所有路径完成时，这些线程将收敛到相同的执行路径。分支发散只在翘曲中发生；不同的翘曲独立执行，不管它们是执行公共的还是不相交的代码路径。 SIMT结构类似于SIMD (单指令，多数据)向量组织，因为一个指令控制多个处理元素。一个关键的区别是SIMD向量组织向软件公开SIMD宽度，而SIMT指令指定单个线程的执行和分支行为。

• WaveFront in OpenCL：https://sites.google.com/site/csc8820/opencl-basics/opencl-terms-explained#TOC-Wavefront

在运行时，第一个波前被发送到计算单元运行，然后第二个波前被发送到计算单元，依此类推。一个波前内的工作项在并行和锁步骤中执行。但是不同的波前按顺序执行。

也就是说，我们知道：

• WARP (CUDA)中的线程是SIMT线程，每次执行相同的指令，并且始终保持同步的，即经纱线程与SIMD车道 (on )相同。
• WaveFront (OpenCL)中的线程是线程，它们总是并行执行，但不一定所有线程都执行完全相同的指令，也不一定所有线程都是同步的。

但是，是否可以保证WaveFront中的所有线程都是同步的，比如在翘曲中的线程或者在SIMD中作为通道的线程？

结论：

1. 波前线程(项)总是同步锁定步骤：“波前在锁步骤中执行许多相对于彼此的工作项。”
2. WaveFront映射在SIMD块上：“波前的所有工作项都流向流控制的两条路径”
3. 即：每个波前线(项目)映射到SIMD通道。

第20页：Guide-rev-2.7.pdf

第一章OpenCL体系结构与加速并行处理 1.1术语 ..。 波前和工作组是与提供数据并行粒度的计算核相关的两个概念.一个波前在锁步骤中执行许多相对于彼此的工作项。16个工作项在向量单元中并行执行，整个波前覆盖四个时钟周期。这是流量控制所能影响的最低水平。这意味着，如果一个波前内的两个工作项在流控制中走发散路径，则所有工作项都会到达流控制的两个路径。

这对于：Guide2.pdf是正确的

• (page-45)第2章GCN设备的OpenCL性能和优化
• (第81页)第3章OpenCL性能及对长青和北方群岛设备的优化

Answer 1

首先，您可以查询一些值：

CL_DEVICE_WAVEFRONT_WIDTH_AMD
CL_DEVICE_SIMD_WIDTH_AMD
CL_DEVICE_WARP_SIZE_NV
CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE

但据我所知只能从主人那一边。

让我们假设这些查询返回了64，并且您的问题重视线程的隐式同步。

如果有人选择本地范围= 4呢？

由于opencl从开发人员那里提取硬件时钟，所以您无法知道运行时内核执行中的实际SIMD或波前大小。

例如，AMD NCU有64个着色器，但它有16宽的SIMD，8宽的SIMD，4宽的SIMD，2宽的SIMD，甚至在同一个计算单元内有两个标量单元。

4个本地线程可以在两个标量和一个2宽的单元或任何其他SIMD组合上共享。内核代码不能知道这一点。即使知道某种计算方法，您也无法知道在运行时在随机计算单元(64个着色器)中，哪个SIMD组合将用于下一个内核执行(甚至下一个工作组)。

或者一个包含4x16SIMD的GCN可以为每个SIMD分配一个线程，使所有4个线程完全独立。如果他们都住在同一个SIMD里，你就很幸运了。不能保证在内核执行之前就知道这一点。即使在您知道之后，下一个内核也可能会有所不同，因为无法保证选择相同的SIMD分配(背景内核、三维可视化软件，甚至操作系统也可能在管道中放置气泡)。

没有保证在内核执行之前命令/暗示/查询N个线程以相同的SIMD或相同的翘曲运行。然后在内核中，没有获得线程的波前索引的命令，就像get_global_id(0)一样。然后，在内核之后，您不能依赖数组结果，因为下一个内核执行可能不会对完全相同的项使用相同的SIMD。即使是来自其他波前的一些项目也可以被替换为当前波前的一个项目，仅仅是通过驱动程序或硬件进行优化(nvidia最近有负载平衡器，并且可能正在这样做，NCU的amd也可能在将来使用类似的东西)。

即使你在硬件和驱动程序上猜对了SIMD上的线程组合，在另一台计算机上也可能完全不同。

如果是从性能角度来看，可以尝试：：

• 内核代码中的零分支
• 零内核在后台运行
• gpu不用于监视器输出。
• gpu不用于某些可视化软件。

只是为了确保%99概率，管道中没有气泡，所以所有线程都在同一周期退出一条指令(或者至少在最晚的退出时同步)。

或者，在每条指令之后添加一个栅栏，以便在全局或本地内存上同步，这是非常慢的。栅栏使工作项级同步，障碍使本地组级同步。没有波前同步命令。

然后，在同一个SIMD中运行的线程将同步运行。但是，您可能不知道它们是哪些线程以及哪个SIMD。

对于4线程示例，使用float16进行所有计算可以让驱动程序使用16宽的AMD CU的SIMD来计算，但它们不再是线程，而是变量。但是这应该比线程在数据上有更好的同步。

还有更复杂的情况，例如：

• 相同SIMD中的4个线程，但一个线程计算会生成一些NaN值，并在此基础上进行额外的规范化(可能需要1-2个周期)。3其他人应等待完成，但它的工作与数据相关的减速无关。
• 在同一个波前有4个线程处于一个循环中，其中一个永远被卡住。他们中的3个等待第四个永远完成，或者驱动检测并移动到另一个自由-空的SIMD？或者第四人在同一时间等待其他3人，因为他们没有移动！
• 执行原子操作的4个线程，一个接一个。
• Amd的HD5000系列gpu具有SIMD宽度4(或5)，但波前尺寸为64。

Answer 2

波前保证锁步。这就是为什么在旧的编译器上，如果本地组只包含一个波前，就可以省略同步。(您不能再在较新的编译器上这样做了，后者会错误地解释依赖项并给出错误的代码。但是另一方面，如果本地组只包含一个波前，则较新的编译器将省略同步。)

一个流处理器就像CPU的一个核心。它将重复运行一个16宽的矢量指令四次，以实现64所谓的“线程”在一个波前。实际上，一个波前更像是一个线程，而不是GPU上的线程。