官方的OpenCL 2.2标准支持WaveFront吗？

Question

众所周知，AMD支持WaveFront (2015年8月)：Guide2.pdf

例如，and 7770 GPU支持超过25000个在飞工作项，并且可以在一个周期内切换到一个新的波前(包含多达64个工作项)。

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

WaveFront：https://www.khronos.org/registry/OpenCL/specs/一词都没有。

我还发现：

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

OpenCL是一个开放的标准。它仍然不支持这种摇摆不定的概念。它甚至还不支持波前/翘曲。

• 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？

结论

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

• 页面-100： 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. 

因此，即使它不是所谓的波前，它现在在运行时和

在没有同步功能(例如障碍)的情况下，子组中的工作项可以序列化。在子-group函数存在的情况下，子-group中的工作项可以在给定的子-group函数之前序列化，在动态遇到的子组函数之间以及在工作组函数和内核结束之间进行序列化。

即使是锁步方式，有时也可能会丢失。

再加上这些，

 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);
    }
    );
}

您应该能够生成您需要的任意大小的(升级的？)波前，并且它们与父内核同时工作(并且可以与子组内线程通信)，但它们不称为波前，因为它们不是硬件imho硬编码的。

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宽模拟和nvidia的32宽模拟，而不是一些虚构的fpga 95宽的计算核。可能是假波前？