官方的OpenCL 2.2標準是否支持WaveFront？

Question

衆所周知，AMD-OpenCL的支持波前（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 飛行工作項目，也可以切換到一個新的波前（包含從 到64個工作項目）。

---

但是，爲什麼在標準的OpenCL 1.0/2.0/2.2有關於波前沒有提及？

無PDF的還沒有一個字波前：https://www.khronos.org/registry/OpenCL/specs/

而且我發現：

• 下一篇：https://community.amd.com/thread/160658

OpenCL是一個開放的標準。它仍然不支持這種混亂的概念。它甚至不支持波前/扭曲。

• 下一篇：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中的標準，但在OpenCL中-2.0 有附屬基團與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. 

因此，即使它不是所謂的波陣面，它現在是queryab le在運行時和

在沒有同步功能（例如，一個障礙）， 子組內的工作項可能被序列化。在存在 子組函數的情況下，子組內的工作項可以在任何給定的子組函數之前，動態地遇到 對子組函數之間以及在工作組函數和 之間被串行化 內核的結束。

甚至鎖步的方式有時可能會丟失。

在這些之上，

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寬simds和NVIDIA公司的32寬simds與一些假想FPGA的95-寬計算核心。僞波前可能？