OpenCL 2.2官方标准支持WaveFront吗？

Question

众所周知，AMD-OpenCL 支持 WaveFront（2015 年 8 月）：http://amd-dev.wpengine.netdna-cdn.com/wordpress/media/2013/12/AMD_OpenCL_Programming_Optimization_Guide2.pdf

以 AMD Radeon HD 7770 GPU 为例，支持超过 25,000 个 飞行中的工作项目，并且可以切换到新的wavefront（包含 到 64 个工作项目）在一个周期内。

---

但是为什么在 OpenCL 标准 1.0/2.0/2.2 中没有提到 WaveFront？

所有 PDF 都没有一个字 WaveFront： https://www.khronos.org/registry/OpenCL/specs/

我还发现：

• 2013：https://community.amd.com/thread/160658

OpenCL 是一个开放标准。它仍然不支持这种 swizzling 概念。它甚至还不支持波前/扭曲。

• 2013：https://stackoverflow.com/a/19874984/1558037

这就是为什么这个概念不在 OpenCL 规范本身中的原因。

• 2011：https://forums.khronos.org/showthread.php/7211-How-can-i-split-my-work-load-in-a-GPU-with-OpenCL

标准 OpenCL 没有“波前”的概念

• 2011：https://www.cvg.ethz.ch/teaching/2011spring/gpgpu/GPU-Optimization.pdf

确实官方的OpenCL 2.2标准仍然不支持WaveFront？

---

结论：

OpenCL 标准中没有 WaveFront，但 在 OpenCL-2.0 中存在类似于 WaveFronts 的 SIMD 执行模型子组。

• 第 100 页： http://amd-dev.wpengine.netdna-cdn.com/wordpress/media/2013/12/AMD_OpenCL_Programming_User_Guide2.pdf

6.4.2 工作组/子组级职能

OpenCL 2.0 引入了 Khronos

子组扩展。子组是 硬件 SIMD 执行模型的逻辑抽象类似于 波前、扭曲或矢量，并允许编程更接近 以独立于供应商的方式硬件。该扩展包括一组 与集合匹配的跨子组内置函数 上面指定的跨工作组内置函数。

Answer 1

他们一定采用了一种名为

sub-group

 的更加动态的方法：

https://www.khronos.org/registry/OpenCL/specs/opencl-2.2.pdf

Sub-group: Sub-groups are an implementation-dependent grouping of work-items within a
work-group. The size and number of sub-groups is implementation-defined.

和

Work-groups are further divided into sub-groups,
which provide an additional level of control over execution.

和

The mapping of work-items to
sub-groups is implementation-defined and may be queried at runtime. 

所以即使它不被称为波前，它现在可以在运行时查询并且

在没有同步功能（例如屏障）的情况下， 子组内的工作项可以被序列化。在......的存在下 子组功能、子组内的工作项可以序列化 在任何给定的子组函数之前，在动态遇到的之间 子组职能对以及工作组职能和 内核结束。

即使是步调一致的方式有时也可能会丢失。

除此之外，

 sub_group_all() and
sub_group_broadcast() and are described in OpenCL C++ kernel language and IL specifications.
The use of these sub-group functions implies sequenced-before relationships between statements
within the execution of a single work-item in order to satisfy data dependencies.

表示存在某种子组内沟通。因为现在 opencl 有子内核定义：

Device-side enqueue: A mechanism whereby a kernel-instance is enqueued by a kernel-instance
running on a device without direct involvement by the host program. This produces nested
parallelism; i.e. additional levels of concurrency are nested inside a running kernel-instance.
The kernel-instance executing on a device (the parent kernel) enqueues a kernel-instance (the
child kernel) to a device-side command queue. Child and parent kernels execute asynchronously
though a parent kernel does not complete until all of its child-kernels have completed. 

---

最终，像

kernel void launcher()
{
    ndrange_t ndrange = ndrange_1D(1);
    enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange,
    ^{
    size_t id = get_global_id(0);
    }
    );
}

您应该能够生成您自己的（升级的？）波前，具有您需要的任何大小，并且它们与父内核同时工作（并且可以通信子组内线程），但它们不被称为波前，因为它们不是由硬件硬编码的恕我直言。

---

2.0 API 规范说：

Extreme care should be exercised when writing code that uses
subgroups if the goal is to write portable OpenCL applications.

这让人想起 AMD 的 16 宽 simd 和 nvidia 的 32 宽 simd 与一些虚构的 FPGA 的 95 宽计算核心。也许是伪波前？