是否可以保证 WaveFront (OpenCL) 中的所有线程始终同步？
Is there any guarantee that all of threads in WaveFront (OpenCL) always synchronized?
众所周知，有 WARP（在 CUDA 中）和 WaveFront（在 OpenCL 中）：http://courses.cs.washington.edu/courses/cse471/13sp/lectures/GPUsStudents.pdf
4.1. SIMT Architecture
...
A warp executes one common instruction at a time, so full efficiency
is realized when all 32 threads of a warp agree on their execution
path. If threads of a warp diverge via a data-dependent conditional
branch, the warp serially executes each branch path taken, disabling
threads that are not on that path, and when all paths complete, the
threads converge back to the same execution path. Branch divergence
occurs only within a warp; different warps execute independently
regardless of whether they are executing common or disjoint code
paths.
The SIMT architecture is akin to SIMD (Single Instruction, Multiple
Data) vector organizations in that a single instruction controls
multiple processing elements. A key difference is that SIMD vector
organizations expose the SIMD width to the software, whereas SIMT
instructions specify the execution and branching behavior of a single
thread.
- WaveFront 在 OpenCL 中：https://sites.google.com/site/csc8820/opencl-basics/opencl-terms-explained#TOC-Wavefront
During runtime, the first wavefront is sent to the compute unit to
run, then the second wavefront is sent to the compute unit, and so on.
Work items within one wavefront are executed in parallel and in lock
steps. But different wavefronts are executed sequentially.
即我们知道：
WARP (CUDA) 中的线程 - 是 SIMT 线程，每次总是执行相同的指令，并且始终保持不变synchronized - 即 WARP 线程与 lanes of SIMD（在 CPU 上）
WaveFront (OpenCL) 中的线程 - 是始终并行执行的线程，但不一定所有线程都执行完全相同的指令，而不是必然所有线程都是同步的
但是是否可以保证 WaveFront 中的所有线程始终同步，例如 WARP 中的线程或 SIMD 中的通道？
结论：
- WaveFront-线程（项）始终同步 - 锁步："wavefront executes a number of work-items in lock step relative to each other."
- WaveFront 映射到 SIMD 块："all work-items in the wavefront go to both paths of flow control"
- 即映射到 SIMD 通道的每个 WaveFront 线程（项目）
Chapter 1 OpenCL Architecture and AMD Accelerated Parallel Processing
1.1 Terminology
...
Wavefronts and work-groups are two concepts relating to compute
kernels that provide data-parallel granularity. A wavefront
executes a number of work-items in lock step relative to each
other. Sixteen workitems are execute in parallel across the vector
unit, and the whole wavefront is covered over four clock cycles. It
is the lowest level that flow control can affect. This means that if
two work-items inside of a wavefront go divergent paths of flow
control, all work-items in the wavefront go to both paths of flow
control.
- （第 45 页） 第 2 章 GCN 设备的 OpenCL 性能和优化
- （第 81 页） 第 3 章 Evergreen 和 Northern Islands 设备的 OpenCL 性能和优化
首先，您可以查询一些值：
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 或 WAVEFRONT 大小。
例如，AMD NCU 有 64 个着色器，但它在同一计算单元内有 16 宽 SIMD、8 宽 SIMD、4 宽 SIMD、2 宽 SIMD 甚至两个标量单元。
4 个本地线程可以在两个标量和一个 2 宽单元或任何其他 SIMD 组合上共享。内核代码无法知道这一点。即使它知道以某种方式计算东西，你也无法知道在 运行 时间随机 compute-unit( 64 个着色器）。
或者一个 GCN CU，其中有 4x16 个 SIMD，可以为每个 SIMD 分配 1 个线程，使所有 4 个线程完全独立。如果它们都位于同一个 SIMD 中，那你就幸运了。无法保证知道 "before" 内核执行。即使在您知道之后，下一个内核也可能会有所不同，因为无法保证选择相同的 SIMD 分配（后台内核、3d 可视化软件，甚至 OS 可能会在管道中放置气泡）
无法保证 commanding/hinting/querying 的 N 个线程 运行 与内核执行前相同的 SIMD 或相同的 WARP 。那么在内核中，并没有像get_global_id(0)那样获取线程波前索引的命令。然后在内核之后，您不能依赖数组结果，因为下一次内核执行可能不会对完全相同的项目使用相同的 SIMD。甚至其他 wavefronts 的一些项目也可以与当前 wavefront 的项目交换，只是为了通过驱动程序或硬件进行优化（nvidia 最近有负载平衡器并且可能已经这样做了，amd 的 NCU 将来也可能使用类似的东西）
即使您猜对了硬件和驱动程序上 SIMD 上的线程组合，在另一台计算机上也可能完全不同。
如果从性能角度考虑，可以试试：
- zero-branch 内核代码
- 零内核 运行在后台运行
- gpu 未用于监视器输出
- gpu 未用于某些可视化软件
只是为了确保 %99 的概率，管道中没有气泡，因此所有线程都在同一周期退出指令（或者至少同步最新退出的指令）。
或者，在每条指令后添加栅栏以同步全局或本地内存，这非常慢。栅栏使工作项级别同步，屏障使本地组级别同步。没有波前同步命令。
然后，那些 运行 在同一个 SIMD 中的线程将表现同步。但你可能不知道这些是哪些线程以及哪些 SIMD。
对于4线程的例子，所有计算都使用float16可能会让驱动程序使用AMD GCN CU的16宽SIMD来计算，但它们不再是线程，只是变量。但这应该比线程有更好的数据同步。
还有更复杂的情况如：
4 个线程在同一个 SIMD 中，但一个线程计算生成一些 NaN 值并对其进行额外的标准化（可能需要 1-2 个周期）。其他 3 个应该等待完成，但它独立于数据相关的减速。
同一波前的 4 个线程处于循环中，其中一个永远卡住了。其中 3 个等待第 4 个永远完成，或者驱动程序检测到并将其移动到另一个 free-empty SIMD？或第4个同时等待其他3个，因为他们也没有移动！
4个线程逐一进行原子操作。
AMD 的 HD5000 系列 gpu 的 SIMD 宽度为 4（或 5），但波前尺寸为 64。
波前保证步调一致。这就是为什么在较旧的编译器上，如果您的本地组仅包含一个波前，您可以省略同步。 （你不能再在较新的编译器上这样做，它们会错误地解释依赖关系并给你错误的代码。但另一方面，如果你的本地组只包含一个波前，较新的编译器会为你忽略同步。）
一个流处理器相当于CPU的一个核心。它会重复 运行 一条 16 宽矢量指令四次，以在波前完成 64 所谓的 "threads" 。实际上一个波阵面更像是一个线程，而不是我们在 GPU 上所说的线程。
众所周知，有 WARP（在 CUDA 中）和 WaveFront（在 OpenCL 中）：http://courses.cs.washington.edu/courses/cse471/13sp/lectures/GPUsStudents.pdf
4.1. SIMT Architecture
...
A warp executes one common instruction at a time, so full efficiency is realized when all 32 threads of a warp agree on their execution path. If threads of a warp diverge via a data-dependent conditional branch, the warp serially executes each branch path taken, disabling threads that are not on that path, and when all paths complete, the threads converge back to the same execution path. Branch divergence occurs only within a warp; different warps execute independently regardless of whether they are executing common or disjoint code paths.
The SIMT architecture is akin to SIMD (Single Instruction, Multiple Data) vector organizations in that a single instruction controls multiple processing elements. A key difference is that SIMD vector organizations expose the SIMD width to the software, whereas SIMT instructions specify the execution and branching behavior of a single thread.
- WaveFront 在 OpenCL 中：https://sites.google.com/site/csc8820/opencl-basics/opencl-terms-explained#TOC-Wavefront
During runtime, the first wavefront is sent to the compute unit to run, then the second wavefront is sent to the compute unit, and so on. Work items within one wavefront are executed in parallel and in lock steps. But different wavefronts are executed sequentially.
即我们知道：
WARP (CUDA) 中的线程 - 是 SIMT 线程，每次总是执行相同的指令，并且始终保持不变synchronized - 即 WARP 线程与 lanes of SIMD（在 CPU 上）
相同
WaveFront (OpenCL) 中的线程 - 是始终并行执行的线程，但不一定所有线程都执行完全相同的指令，而不是必然所有线程都是同步的
但是是否可以保证 WaveFront 中的所有线程始终同步，例如 WARP 中的线程或 SIMD 中的通道？
结论：
- WaveFront-线程（项）始终同步 - 锁步："wavefront executes a number of work-items in lock step relative to each other."
- WaveFront 映射到 SIMD 块："all work-items in the wavefront go to both paths of flow control"
- 即映射到 SIMD 通道的每个 WaveFront 线程（项目）
Chapter 1 OpenCL Architecture and AMD Accelerated Parallel Processing
1.1 Terminology
...
Wavefronts and work-groups are two concepts relating to compute kernels that provide data-parallel granularity. A wavefront executes a number of work-items in lock step relative to each other. Sixteen workitems are execute in parallel across the vector unit, and the whole wavefront is covered over four clock cycles. It is the lowest level that flow control can affect. This means that if two work-items inside of a wavefront go divergent paths of flow control, all work-items in the wavefront go to both paths of flow control.
- （第 45 页） 第 2 章 GCN 设备的 OpenCL 性能和优化
- （第 81 页） 第 3 章 Evergreen 和 Northern Islands 设备的 OpenCL 性能和优化
首先，您可以查询一些值：
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 或 WAVEFRONT 大小。
例如，AMD NCU 有 64 个着色器，但它在同一计算单元内有 16 宽 SIMD、8 宽 SIMD、4 宽 SIMD、2 宽 SIMD 甚至两个标量单元。
4 个本地线程可以在两个标量和一个 2 宽单元或任何其他 SIMD 组合上共享。内核代码无法知道这一点。即使它知道以某种方式计算东西，你也无法知道在 运行 时间随机 compute-unit( 64 个着色器）。
或者一个 GCN CU，其中有 4x16 个 SIMD，可以为每个 SIMD 分配 1 个线程，使所有 4 个线程完全独立。如果它们都位于同一个 SIMD 中，那你就幸运了。无法保证知道 "before" 内核执行。即使在您知道之后，下一个内核也可能会有所不同，因为无法保证选择相同的 SIMD 分配（后台内核、3d 可视化软件，甚至 OS 可能会在管道中放置气泡）
无法保证 commanding/hinting/querying 的 N 个线程 运行 与内核执行前相同的 SIMD 或相同的 WARP 。那么在内核中，并没有像get_global_id(0)那样获取线程波前索引的命令。然后在内核之后，您不能依赖数组结果，因为下一次内核执行可能不会对完全相同的项目使用相同的 SIMD。甚至其他 wavefronts 的一些项目也可以与当前 wavefront 的项目交换，只是为了通过驱动程序或硬件进行优化（nvidia 最近有负载平衡器并且可能已经这样做了，amd 的 NCU 将来也可能使用类似的东西）
即使您猜对了硬件和驱动程序上 SIMD 上的线程组合，在另一台计算机上也可能完全不同。
如果从性能角度考虑，可以试试：
- zero-branch 内核代码
- 零内核 运行在后台运行
- gpu 未用于监视器输出
- gpu 未用于某些可视化软件
只是为了确保 %99 的概率，管道中没有气泡，因此所有线程都在同一周期退出指令（或者至少同步最新退出的指令）。
或者，在每条指令后添加栅栏以同步全局或本地内存，这非常慢。栅栏使工作项级别同步，屏障使本地组级别同步。没有波前同步命令。
然后，那些 运行 在同一个 SIMD 中的线程将表现同步。但你可能不知道这些是哪些线程以及哪些 SIMD。
对于4线程的例子，所有计算都使用float16可能会让驱动程序使用AMD GCN CU的16宽SIMD来计算，但它们不再是线程，只是变量。但这应该比线程有更好的数据同步。
还有更复杂的情况如：
4 个线程在同一个 SIMD 中，但一个线程计算生成一些 NaN 值并对其进行额外的标准化（可能需要 1-2 个周期）。其他 3 个应该等待完成，但它独立于数据相关的减速。
同一波前的 4 个线程处于循环中，其中一个永远卡住了。其中 3 个等待第 4 个永远完成，或者驱动程序检测到并将其移动到另一个 free-empty SIMD？或第4个同时等待其他3个，因为他们也没有移动！
4个线程逐一进行原子操作。
AMD 的 HD5000 系列 gpu 的 SIMD 宽度为 4（或 5），但波前尺寸为 64。
波前保证步调一致。这就是为什么在较旧的编译器上，如果您的本地组仅包含一个波前，您可以省略同步。 （你不能再在较新的编译器上这样做，它们会错误地解释依赖关系并给你错误的代码。但另一方面，如果你的本地组只包含一个波前，较新的编译器会为你忽略同步。）
一个流处理器相当于CPU的一个核心。它会重复 运行 一条 16 宽矢量指令四次，以在波前完成 64 所谓的 "threads" 。实际上一个波阵面更像是一个线程，而不是我们在 GPU 上所说的线程。