OpenCL 3.0 Specifications v3.0.7 (#97)
diff --git a/index.php b/index.php
index 759e353..c0d1e50 100644
--- a/index.php
+++ b/index.php
@@ -31,19 +31,19 @@
<li> OpenCL API Specification
(<a href="specs/3.0-unified/html/OpenCL_API.html">HTML</a>)
(<a href="specs/3.0-unified/pdf/OpenCL_API.pdf">PDF</a>)
- (December 18, 2020). </li>
+ (April 23, 2021). </li>
<li> OpenCL C Language Specification
(<a href="specs/3.0-unified/html/OpenCL_C.html">HTML</a>)
(<a href="specs/3.0-unified/pdf/OpenCL_C.pdf">PDF</a>)
- (December 18, 2020). </li>
+ (April 23, 2021). </li>
<li> OpenCL Extension Specification
(<a href="specs/3.0-unified/html/OpenCL_Ext.html">HTML</a>)
(<a href="specs/3.0-unified/pdf/OpenCL_Ext.pdf">PDF</a>)
- (December 18, 2020. </li>
+ (April 23, 2021). </li>
<li> OpenCL SPIR-V Environment Specification
(<a href="specs/3.0-unified/html/OpenCL_Env.html">HTML</a>)
(<a href="specs/3.0-unified/pdf/OpenCL_Env.pdf">PDF</a>)
- (December 18, 2020). </li>
+ (April 23, 2021). </li>
<li> <a href="https://www.khronos.org/files/opencl30-reference-guide.pdf">
OpenCL 3.0 Reference Guide.</a> </li>
</ul>
diff --git a/specs/3.0-unified/html/OpenCL_API.html b/specs/3.0-unified/html/OpenCL_API.html
index 69315ee..5b31e52 100644
--- a/specs/3.0-unified/html/OpenCL_API.html
+++ b/specs/3.0-unified/html/OpenCL_API.html
@@ -841,9 +841,9 @@
<h1>The OpenCL<sup>™</sup> Specification</h1>
<div class="details">
<span id="author" class="author">Khronos<sup>®</sup> OpenCL Working Group</span><br>
-<span id="revnumber">version v3.0.6,</span>
-<span id="revdate">Fri, 18 Dec 2020 12:00:00 +0000</span>
-<br><span id="revremark">from git branch: master commit: e9a4d468b1a0a38c1e10b8af484bb2bbb495e2b7</span>
+<span id="revnumber">version v3.0.7,</span>
+<span id="revdate">Fri, 23 Apr 2021 20:00:00 +0000</span>
+<br><span id="revremark">from git branch: master commit: 133503b85911ca8fd6642721d440e89460d491a3</span>
</div>
<div id="toc" class="toc2">
<div id="toctitle">Table of Contents</div>
@@ -18118,7 +18118,8 @@
in the unsafe math optimization section of the OpenCL C or OpenCL SPIR-V
Environment specifications, and (d) may violate edge case behavior in the
OpenCL C or OpenCL SPIR-V Environment specifications.
-This option includes the <code>-cl-no-signed-zeros</code> and <code>-cl-mad-enable</code> options.</p>
+This option includes the <code>-cl-no-signed-zeros</code>, <code>-cl-mad-enable</code>, and
+<code>-cl-denorms-are-zero</code> <sup class="footnote">[<a id="_footnoteref_26" class="footnote" href="#_footnotedef_26" title="View footnote.">26</a>]</sup> options.</p>
</dd>
<dt class="hdlist1"><code>-cl-finite-math-only</code> </dt>
<dd>
@@ -18355,7 +18356,7 @@
</div>
<div class="paragraph">
<p><code>-cl-denorms-are-zero</code><br>
-<code>-cl-no-signed-zeroes</code><br>
+<code>-cl-no-signed-zeros</code><br>
<code>-cl-unsafe-math-optimizations</code><br>
<code>-cl-finite-math-only</code><br>
<code>-cl-fast-relaxed-math</code><br>
@@ -18528,7 +18529,7 @@
</thead>
<tbody>
<tr>
-<td class="tableblock halign-left valign-top"><p class="tableblock"><a id="CL_PROGRAM_REFERENCE_COUNT"></a><a href="#CL_PROGRAM_REFERENCE_COUNT"><code>CL_PROGRAM_<wbr>REFERENCE_<wbr>COUNT</code></a> <sup class="footnote">[<a id="_footnoteref_26" class="footnote" href="#_footnotedef_26" title="View footnote.">26</a>]</sup></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><a id="CL_PROGRAM_REFERENCE_COUNT"></a><a href="#CL_PROGRAM_REFERENCE_COUNT"><code>CL_PROGRAM_<wbr>REFERENCE_<wbr>COUNT</code></a> <sup class="footnote">[<a id="_footnoteref_27" class="footnote" href="#_footnotedef_27" title="View footnote.">27</a>]</sup></p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock"><code>cl_uint</code></p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">Return the <em>program</em> reference count.</p></td>
</tr>
@@ -19856,7 +19857,7 @@
<td class="tableblock halign-left valign-top"><p class="tableblock">Return the number of arguments to kernel.</p></td>
</tr>
<tr>
-<td class="tableblock halign-left valign-top"><p class="tableblock"><a id="CL_KERNEL_REFERENCE_COUNT"></a><a href="#CL_KERNEL_REFERENCE_COUNT"><code>CL_KERNEL_<wbr>REFERENCE_<wbr>COUNT</code></a> <sup class="footnote">[<a id="_footnoteref_27" class="footnote" href="#_footnotedef_27" title="View footnote.">27</a>]</sup></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><a id="CL_KERNEL_REFERENCE_COUNT"></a><a href="#CL_KERNEL_REFERENCE_COUNT"><code>CL_KERNEL_<wbr>REFERENCE_<wbr>COUNT</code></a> <sup class="footnote">[<a id="_footnoteref_28" class="footnote" href="#_footnotedef_28" title="View footnote.">28</a>]</sup></p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock"><code>cl_uint</code></p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">Return the <em>kernel</em> reference count.</p></td>
</tr>
@@ -20448,8 +20449,8 @@
for the argument given by <em>arg_index</em>.
The returned values can be:</p>
<p class="tableblock"> <a id="CL_KERNEL_ARG_TYPE_CONST"></a><a href="#CL_KERNEL_ARG_TYPE_CONST"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>CONST</code></a>
- <sup class="footnote">[<a id="_footnoteref_28" class="footnote" href="#_footnotedef_28" title="View footnote.">28</a>]</sup>
- <sup class="footnote">[<a id="_footnoteref_29" class="footnote" href="#_footnotedef_29" title="View footnote.">29</a>]</sup><br>
+ <sup class="footnote">[<a id="_footnoteref_29" class="footnote" href="#_footnotedef_29" title="View footnote.">29</a>]</sup>
+ <sup class="footnote">[<a id="_footnoteref_30" class="footnote" href="#_footnotedef_30" title="View footnote.">30</a>]</sup><br>
<a id="CL_KERNEL_ARG_TYPE_RESTRICT"></a><a href="#CL_KERNEL_ARG_TYPE_RESTRICT"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>RESTRICT</code></a><br>
<a id="CL_KERNEL_ARG_TYPE_VOLATILE"></a><a href="#CL_KERNEL_ARG_TYPE_VOLATILE"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>VOLATILE</code></a><br>
<a id="CL_KERNEL_ARG_TYPE_PIPE"></a><a href="#CL_KERNEL_ARG_TYPE_PIPE"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>PIPE</code></a>, or<br>
@@ -20693,8 +20694,7 @@
</li>
<li>
<p><a href="#CL_INVALID_KERNEL_ARGS"><code>CL_INVALID_<wbr>KERNEL_<wbr>ARGS</code></a> if the kernel argument values have not been
-specified or if a kernel argument declared to be a pointer to a type
-does not point to a named address space.</p>
+specified.</p>
</li>
<li>
<p><a href="#CL_INVALID_WORK_DIMENSION"><code>CL_INVALID_<wbr>WORK_<wbr>DIMENSION</code></a> if <em>work_dim</em> is not a valid value (i.e. a
@@ -20907,8 +20907,7 @@
</li>
<li>
<p><a href="#CL_INVALID_KERNEL_ARGS"><code>CL_INVALID_<wbr>KERNEL_<wbr>ARGS</code></a> if the kernel argument values have not been
-specified or if a kernel argument declared to be a pointer to a type
-does not point to a named address space.</p>
+specified.</p>
</li>
<li>
<p><a href="#CL_INVALID_WORK_GROUP_SIZE"><code>CL_INVALID_<wbr>WORK_<wbr>GROUP_<wbr>SIZE</code></a> if a work-group size is specified for <em>kernel</em>
@@ -21533,7 +21532,7 @@
<a href="#event-command-type-table">Event Command Types</a> table.</p></td>
</tr>
<tr>
-<td class="tableblock halign-left valign-top"><p class="tableblock"><a id="CL_EVENT_COMMAND_EXECUTION_STATUS"></a><a href="#CL_EVENT_COMMAND_EXECUTION_STATUS"><code>CL_EVENT_<wbr>COMMAND_<wbr>EXECUTION_<wbr>STATUS</code></a> <sup class="footnote">[<a id="_footnoteref_30" class="footnote" href="#_footnotedef_30" title="View footnote.">30</a>]</sup></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><a id="CL_EVENT_COMMAND_EXECUTION_STATUS"></a><a href="#CL_EVENT_COMMAND_EXECUTION_STATUS"><code>CL_EVENT_<wbr>COMMAND_<wbr>EXECUTION_<wbr>STATUS</code></a> <sup class="footnote">[<a id="_footnoteref_31" class="footnote" href="#_footnotedef_31" title="View footnote.">31</a>]</sup></p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock"><code>cl_int</code></p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">Return the execution status of the command identified by event.
Valid values are:</p>
@@ -21550,7 +21549,7 @@
errcode_ret values.</p></td>
</tr>
<tr>
-<td class="tableblock halign-left valign-top"><p class="tableblock"><a id="CL_EVENT_REFERENCE_COUNT"></a><a href="#CL_EVENT_REFERENCE_COUNT"><code>CL_EVENT_<wbr>REFERENCE_<wbr>COUNT</code></a> <sup class="footnote">[<a id="_footnoteref_31" class="footnote" href="#_footnotedef_31" title="View footnote.">31</a>]</sup></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><a id="CL_EVENT_REFERENCE_COUNT"></a><a href="#CL_EVENT_REFERENCE_COUNT"><code>CL_EVENT_<wbr>REFERENCE_<wbr>COUNT</code></a> <sup class="footnote">[<a id="_footnoteref_32" class="footnote" href="#_footnotedef_32" title="View footnote.">32</a>]</sup></p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock"><code>cl_uint</code></p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">Return the <em>event</em> reference count.</p></td>
</tr>
@@ -21782,7 +21781,7 @@
which the callback is registered.
The command execution callback values for which a callback can be registered
are: <a href="#CL_SUBMITTED"><code>CL_SUBMITTED</code></a>, <a href="#CL_RUNNING"><code>CL_RUNNING</code></a>, or
-<a href="#CL_COMPLETE"><code>CL_COMPLETE</code></a> <sup class="footnote">[<a id="_footnoteref_32" class="footnote" href="#_footnotedef_32" title="View footnote.">32</a>]</sup>.
+<a href="#CL_COMPLETE"><code>CL_COMPLETE</code></a> <sup class="footnote">[<a id="_footnoteref_33" class="footnote" href="#_footnotedef_33" title="View footnote.">33</a>]</sup>.
There is no guarantee that the callback functions registered for various
execution status values for an event will be called in the exact order that
the execution status of a command changes.
@@ -22865,7 +22864,7 @@
<p>Image and image arrays created with an
<code>image_channel_data_type</code> value of <a href="#CL_FLOAT"><code>CL_FLOAT</code></a> or <a href="#CL_HALF_FLOAT"><code>CL_HALF_<wbr>FLOAT</code></a> can only be
used with samplers that use a filter mode of <a href="#CL_FILTER_NEAREST"><code>CL_FILTER_<wbr>NEAREST</code></a>.
-The values returned by <strong>read_imagef</strong> <sup class="footnote">[<a id="_footnoteref_33" class="footnote" href="#_footnotedef_33" title="View footnote.">33</a>]</sup> for 2D and 3D
+The values returned by <strong>read_imagef</strong> <sup class="footnote">[<a id="_footnoteref_34" class="footnote" href="#_footnotedef_34" title="View footnote.">34</a>]</sup> for 2D and 3D
images if <code>image_channel_data_type</code> value is <a href="#CL_FLOAT"><code>CL_FLOAT</code></a> or <a href="#CL_HALF_FLOAT"><code>CL_HALF_<wbr>FLOAT</code></a>
and sampler with filter_mode = <a href="#CL_FILTER_LINEAR"><code>CL_FILTER_<wbr>LINEAR</code></a> are undefined.</p>
</li>
@@ -22880,7 +22879,7 @@
<li>
<p>64 bit integers i.e. long, ulong including the appropriate vector data
types and operations on 64-bit integers are optional.
-The <strong>cles_khr_int64</strong> <sup class="footnote">[<a id="_footnoteref_34" class="footnote" href="#_footnotedef_34" title="View footnote.">34</a>]</sup> extension string will
+The <strong>cles_khr_int64</strong> <sup class="footnote">[<a id="_footnoteref_35" class="footnote" href="#_footnotedef_35" title="View footnote.">35</a>]</sup> extension string will
be reported if the embedded profile implementation supports 64-bit integers.
If double precision is supported i.e. <a href="#CL_DEVICE_DOUBLE_FP_CONFIG"><code>CL_DEVICE_<wbr>DOUBLE_<wbr>FP_<wbr>CONFIG</code></a> is not
zero, then <strong>cles_khr_int64</strong> must also be supported.</p>
@@ -23360,7 +23359,7 @@
<div class="sect2">
<h3 id="_multiple_host_threads"><a class="anchor" href="#_multiple_host_threads"></a>Multiple Host Threads</h3>
<div class="paragraph">
-<p>All OpenCL API calls are thread-safe <sup class="footnote">[<a id="_footnoteref_35" class="footnote" href="#_footnotedef_35" title="View footnote.">35</a>]</sup> except those
+<p>All OpenCL API calls are thread-safe <sup class="footnote">[<a id="_footnoteref_36" class="footnote" href="#_footnotedef_36" title="View footnote.">36</a>]</sup> except those
that modify the state of <code>cl_kernel</code> objects: <a href="#clSetKernelArg"><strong>clSetKernelArg</strong></a>,
<a href="#clSetKernelArgSVMPointer"><strong>clSetKernelArgSVMPointer</strong></a>, <a href="#clSetKernelExecInfo"><strong>clSetKernelExecInfo</strong></a> and <a href="#clCloneKernel"><strong>clCloneKernel</strong></a>.
<a href="#clSetKernelArg"><strong>clSetKernelArg</strong></a>, <a href="#clSetKernelArgSVMPointer"><strong>clSetKernelArgSVMPointer</strong></a>, <a href="#clSetKernelExecInfo"><strong>clSetKernelExecInfo</strong></a> and
@@ -23370,7 +23369,7 @@
The state of the <code>cl_kernel</code> object is undefined if <a href="#clSetKernelArg"><strong>clSetKernelArg</strong></a>,
<a href="#clSetKernelArgSVMPointer"><strong>clSetKernelArgSVMPointer</strong></a>, <a href="#clSetKernelExecInfo"><strong>clSetKernelExecInfo</strong></a> or <a href="#clCloneKernel"><strong>clCloneKernel</strong></a> are
called from multiple host threads on the same <code>cl_kernel</code> object at the same
-time <sup class="footnote">[<a id="_footnoteref_36" class="footnote" href="#_footnotedef_36" title="View footnote.">36</a>]</sup>.
+time <sup class="footnote">[<a id="_footnoteref_37" class="footnote" href="#_footnotedef_37" title="View footnote.">37</a>]</sup>.
Please note that there are additional limitations as to which OpenCL APIs
may be called from <a href="#event-objects">OpenCL callback functions</a>.</p>
</div>
@@ -23542,7 +23541,7 @@
That is the transformation is equally valid whether that vector was a
<code>uchar16</code> or a <code>ulong2</code>.
Of course, as is well known, little-endian machines
-actually <sup class="footnote">[<a id="_footnoteref_37" class="footnote" href="#_footnotedef_37" title="View footnote.">37</a>]</sup> store their data in reverse byte order to
+actually <sup class="footnote">[<a id="_footnoteref_38" class="footnote" href="#_footnotedef_38" title="View footnote.">38</a>]</sup> store their data in reverse byte order to
compensate for the little-endian storage format of the array elements:</p>
</div>
<div class="paragraph">
@@ -24913,6 +24912,9 @@
<li>
<p>The <code>CL_VERSION_2_0</code> and <code>NULL</code> macros.</p>
</li>
+<li>
+<p>The <code>opencl_unroll_hint</code> attribute.</p>
+</li>
</ul>
</div>
<div class="paragraph">
@@ -25649,6 +25651,18 @@
<p>OpenCL 3.0 breaks backwards compatibility with earlier versions of OpenCL by making some features that were previously required for FULL_PROFILE or EMBEDDED_PROFILE devices optional.
This appendix describes the features that were previously required that are now optional, how to detect whether an optional feature is supported, and expected behavior when an optional feature is not supported.</p>
</div>
+<div class="admonitionblock note">
+<table>
+<tr>
+<td class="icon">
+<i class="fa icon-note" title="Note"></i>
+</td>
+<td class="content">
+Informally, in the tables below the first row usually describes a feature detection mechanism ("May return this value indicating that the feature is not supported") and subsequent rows usually describe behavior when a feature is not supported ("Returns this value if the feature is not supported").
+</td>
+</tr>
+</table>
+</div>
<div class="sect2">
<h3 id="_shared_virtual_memory_2"><a class="anchor" href="#_shared_virtual_memory_2"></a>Shared Virtual Memory</h3>
<div class="paragraph">
@@ -26645,46 +26659,49 @@
<a href="#_footnoteref_25">25</a>. The reference count returned should be considered immediately stale. It is unsuitable for general use in applications. This feature is provided for identifying memory leaks.
</div>
<div class="footnote" id="_footnotedef_26">
-<a href="#_footnoteref_26">26</a>. The reference count returned should be considered immediately stale. It is unsuitable for general use in applications. This feature is provided for identifying memory leaks.
+<a href="#_footnoteref_26">26</a>. As per the definition of <code>-cl-denorms-are-zero</code>, the inclusion of this option with <code>-cl-unsafe-math-optimizations</code> means that the implementation may flush denormal numbers to zero but is not required to.
</div>
<div class="footnote" id="_footnotedef_27">
<a href="#_footnoteref_27">27</a>. The reference count returned should be considered immediately stale. It is unsuitable for general use in applications. This feature is provided for identifying memory leaks.
</div>
<div class="footnote" id="_footnotedef_28">
-<a href="#_footnoteref_28">28</a>. <a href="#CL_KERNEL_ARG_TYPE_CONST"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>CONST</code></a> is returned if the argument is a pointer and the referenced type is declared with the const qualifier. For example, a kernel argument declared as <code>global int const *x</code> returns <a href="#CL_KERNEL_ARG_TYPE_CONST"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>CONST</code></a> but a kernel argument declared as <code>global int * const x</code> does not.<br> Similarly, <a href="#CL_KERNEL_ARG_TYPE_RESTRICT"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>RESTRICT</code></a> will be returned if the pointer type is marked <code>restrict</code>. For example, <code>global int * restrict x</code> returns <a href="#CL_KERNEL_ARG_TYPE_RESTRICT"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>RESTRICT</code></a>.<br> <a href="#CL_KERNEL_ARG_TYPE_VOLATILE"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>VOLATILE</code></a> is returned for <a href="#CL_KERNEL_ARG_TYPE_QUALIFIER"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>QUALIFIER</code></a> if the argument is a pointer and the referenced type is declared with the volatile qualifier. For example, a kernel argument declared as <code>global int volatile *x</code> returns <a href="#CL_KERNEL_ARG_TYPE_VOLATILE"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>VOLATILE</code></a> but a kernel argument declared as <code>global int * volatile x</code> does not.
+<a href="#_footnoteref_28">28</a>. The reference count returned should be considered immediately stale. It is unsuitable for general use in applications. This feature is provided for identifying memory leaks.
</div>
<div class="footnote" id="_footnotedef_29">
-<a href="#_footnoteref_29">29</a>. <a href="#CL_KERNEL_ARG_TYPE_CONST"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>CONST</code></a> is returned for <a href="#CL_KERNEL_ARG_TYPE_QUALIFIER"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>QUALIFIER</code></a> if the argument is declared with the <code>constant</code> address space qualifier.
+<a href="#_footnoteref_29">29</a>. <a href="#CL_KERNEL_ARG_TYPE_CONST"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>CONST</code></a> is returned if the argument is a pointer and the referenced type is declared with the const qualifier. For example, a kernel argument declared as <code>global int const *x</code> returns <a href="#CL_KERNEL_ARG_TYPE_CONST"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>CONST</code></a> but a kernel argument declared as <code>global int * const x</code> does not.<br> Similarly, <a href="#CL_KERNEL_ARG_TYPE_RESTRICT"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>RESTRICT</code></a> will be returned if the pointer type is marked <code>restrict</code>. For example, <code>global int * restrict x</code> returns <a href="#CL_KERNEL_ARG_TYPE_RESTRICT"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>RESTRICT</code></a>.<br> <a href="#CL_KERNEL_ARG_TYPE_VOLATILE"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>VOLATILE</code></a> is returned for <a href="#CL_KERNEL_ARG_TYPE_QUALIFIER"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>QUALIFIER</code></a> if the argument is a pointer and the referenced type is declared with the volatile qualifier. For example, a kernel argument declared as <code>global int volatile *x</code> returns <a href="#CL_KERNEL_ARG_TYPE_VOLATILE"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>VOLATILE</code></a> but a kernel argument declared as <code>global int * volatile x</code> does not.
</div>
<div class="footnote" id="_footnotedef_30">
-<a href="#_footnoteref_30">30</a>. The error code values are negative, and event state values are positive. The event state values are ordered from the largest value <a href="#CL_QUEUED"><code>CL_QUEUED</code></a> for the first or initial state to the smallest value (<a href="#CL_COMPLETE"><code>CL_COMPLETE</code></a> or negative integer value) for the last or complete state. The value of <a href="#CL_COMPLETE"><code>CL_COMPLETE</code></a> and <a href="#CL_SUCCESS"><code>CL_SUCCESS</code></a> are the same.
+<a href="#_footnoteref_30">30</a>. <a href="#CL_KERNEL_ARG_TYPE_CONST"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>CONST</code></a> is returned for <a href="#CL_KERNEL_ARG_TYPE_QUALIFIER"><code>CL_KERNEL_<wbr>ARG_<wbr>TYPE_<wbr>QUALIFIER</code></a> if the argument is declared with the <code>constant</code> address space qualifier.
</div>
<div class="footnote" id="_footnotedef_31">
-<a href="#_footnoteref_31">31</a>. The reference count returned should be considered immediately stale. It is unsuitable for general use in applications. This feature is provided for identifying memory leaks.
+<a href="#_footnoteref_31">31</a>. The error code values are negative, and event state values are positive. The event state values are ordered from the largest value <a href="#CL_QUEUED"><code>CL_QUEUED</code></a> for the first or initial state to the smallest value (<a href="#CL_COMPLETE"><code>CL_COMPLETE</code></a> or negative integer value) for the last or complete state. The value of <a href="#CL_COMPLETE"><code>CL_COMPLETE</code></a> and <a href="#CL_SUCCESS"><code>CL_SUCCESS</code></a> are the same.
</div>
<div class="footnote" id="_footnotedef_32">
-<a href="#_footnoteref_32">32</a>. The callback function registered for a <em>command_exec_callback_type</em> value of <a href="#CL_COMPLETE"><code>CL_COMPLETE</code></a> will be called when the command has completed successfully or is abnormally terminated.
+<a href="#_footnoteref_32">32</a>. The reference count returned should be considered immediately stale. It is unsuitable for general use in applications. This feature is provided for identifying memory leaks.
</div>
<div class="footnote" id="_footnotedef_33">
-<a href="#_footnoteref_33">33</a>. And <strong>read_imageh</strong>, if the <strong>cl_khr_fp16</strong> extension is supported.
+<a href="#_footnoteref_33">33</a>. The callback function registered for a <em>command_exec_callback_type</em> value of <a href="#CL_COMPLETE"><code>CL_COMPLETE</code></a> will be called when the command has completed successfully or is abnormally terminated.
</div>
<div class="footnote" id="_footnotedef_34">
-<a href="#_footnoteref_34">34</a>. Note that the performance of 64-bit integer arithmetic can vary significantly between embedded devices.
+<a href="#_footnoteref_34">34</a>. And <strong>read_imageh</strong>, if the <strong>cl_khr_fp16</strong> extension is supported.
</div>
<div class="footnote" id="_footnotedef_35">
-<a href="#_footnoteref_35">35</a>. Please refer to the OpenCL glossary for the OpenCL definition of thread-safe. This definition may be different from usage of the term in other contexts.
+<a href="#_footnoteref_35">35</a>. Note that the performance of 64-bit integer arithmetic can vary significantly between embedded devices.
</div>
<div class="footnote" id="_footnotedef_36">
-<a href="#_footnoteref_36">36</a>. There is an inherent race condition in the design of OpenCL that occurs between setting a kernel argument and using the kernel with <a href="#clEnqueueNDRangeKernel"><strong>clEnqueueNDRangeKernel</strong></a>. Another host thread might change the kernel arguments between when a host thread sets the kernel arguments and then enqueues the kernel, causing the wrong kernel arguments to be enqueued. Rather than attempt to share <code>cl_kernel</code> objects among multiple host threads, applications are strongly encouraged to make additional <code>cl_kernel</code> objects for kernel functions for each host thread.
+<a href="#_footnoteref_36">36</a>. Please refer to the OpenCL glossary for the OpenCL definition of thread-safe. This definition may be different from usage of the term in other contexts.
</div>
<div class="footnote" id="_footnotedef_37">
-<a href="#_footnoteref_37">37</a>. Note that we are talking about the programming model here. In reality, little endian systems might choose to simply address their bytes from "the right" or reverse the "order" of the bits in the byte. Either of these choices would mean that no big swap would need to occur in hardware.
+<a href="#_footnoteref_37">37</a>. There is an inherent race condition in the design of OpenCL that occurs between setting a kernel argument and using the kernel with <a href="#clEnqueueNDRangeKernel"><strong>clEnqueueNDRangeKernel</strong></a>. Another host thread might change the kernel arguments between when a host thread sets the kernel arguments and then enqueues the kernel, causing the wrong kernel arguments to be enqueued. Rather than attempt to share <code>cl_kernel</code> objects among multiple host threads, applications are strongly encouraged to make additional <code>cl_kernel</code> objects for kernel functions for each host thread.
+</div>
+<div class="footnote" id="_footnotedef_38">
+<a href="#_footnoteref_38">38</a>. Note that we are talking about the programming model here. In reality, little endian systems might choose to simply address their bytes from "the right" or reverse the "order" of the bits in the byte. Either of these choices would mean that no big swap would need to occur in hardware.
</div>
</div>
<div id="footer">
<div id="footer-text">
-Version v3.0.6<br>
-Last updated 2020-12-13 22:09:28 -0800
+Version v3.0.7<br>
+Last updated 2021-04-19 17:02:25 -0700
</div>
</div>
diff --git a/specs/3.0-unified/html/OpenCL_C.html b/specs/3.0-unified/html/OpenCL_C.html
index 4fad35c..b201b59 100644
--- a/specs/3.0-unified/html/OpenCL_C.html
+++ b/specs/3.0-unified/html/OpenCL_C.html
@@ -841,9 +841,9 @@
<h1>The OpenCL<sup>™</sup> C Specification</h1>
<div class="details">
<span id="author" class="author">Khronos<sup>®</sup> OpenCL Working Group</span><br>
-<span id="revnumber">version v3.0.6,</span>
-<span id="revdate">Fri, 18 Dec 2020 12:00:00 +0000</span>
-<br><span id="revremark">from git branch: master commit: e9a4d468b1a0a38c1e10b8af484bb2bbb495e2b7</span>
+<span id="revnumber">version v3.0.7,</span>
+<span id="revdate">Fri, 23 Apr 2021 20:00:00 +0000</span>
+<br><span id="revremark">from git branch: master commit: 133503b85911ca8fd6642721d440e89460d491a3</span>
</div>
<div id="toc" class="toc2">
<div id="toctitle">Table of Contents</div>
@@ -1117,7 +1117,7 @@
</li>
<li>
<p>Requires support for OpenCL C 2.0, or OpenCL C 3.0 or newer and the
-<code>__opencl_c_feature_name</code> feature:
+<code>__opencl_c_<wbr><feature_<wbr>name></em></code> feature:
Features that were introduced in OpenCL C 2.0 as mandatory, but made
<a href="#optional-functionality">optional</a> in OpenCL C 3.0.
Compilers for versions of OpenCL C 1.2 or below will not provide these
@@ -1126,7 +1126,7 @@
</li>
<li>
<p>Requires support for OpenCL C 3.0 or newer and the
-<code>__opencl_c_feature_name</code> feature: <a href="#optional-functionality">Optional</a> features that were introduced in OpenCL C 3.0.
+<code>__opencl_c_<wbr><feature_<wbr>name></em></code> feature: <a href="#optional-functionality">Optional</a> features that were introduced in OpenCL C 3.0.
Compilers for an earlier version of OpenCL C will not provide these
features, compilers for OpenCL C 3.0 or newer may provide these features.</p>
</li>
@@ -1174,7 +1174,7 @@
<div class="paragraph">
<p>The following table describes OpenCL C 3.0 or newer features and their
meaning. The naming convention for the feature macros is
-<code>__opencl_c_<name></code>.</p>
+<code>__opencl_c_<wbr><feature_<wbr>name></em></code>.</p>
</div>
<div class="paragraph">
<p>Feature macro identifiers are used as names of features in this document.</p>
@@ -1249,7 +1249,7 @@
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><code>__opencl_c_<wbr>pipes</code></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">The OpenCL C compiler supports the pipe modifier and built-in functions
+<td class="tableblock halign-left valign-top"><p class="tableblock">The OpenCL C compiler supports the pipe specifier and built-in functions
to read and write from a pipe.</p>
<p class="tableblock">OpenCL C compilers that define the feature macro <code>__opencl_c_<wbr>pipes</code> must
also define the feature macro <code>__opencl_c_<wbr>generic_<wbr>address_<wbr>space</code>.</p></td>
@@ -3411,7 +3411,9 @@
<code>image3d_t</code>, <code>image2d_array_t</code>, <code>image1d_t</code>, <code>image1d_buffer_t</code>,
<code>image1d_array_t</code>, <code>image2d_depth_t</code>, <code>image2d_array_depth_t</code>,
<code>sampler_t</code>, <code>queue_t</code>, <code>ndrange_t</code>, <code>clk_event_t</code>, <code>reserve_id_t</code>, and
-<code>event_t</code> types is implementation-defined.</p>
+<code>event_t</code> types is implementation-defined. Additionally, the behavior of
+applying the <code>sizeof</code> operator to a pipe object (a type with the <code>pipe</code> type
+specifier keyword) is implementation-defined.</p>
</div>
</div>
</div>
@@ -3956,7 +3958,7 @@
without <code>__opencl_c_<wbr>program_<wbr>scope_<wbr>global_<wbr>variables</code> feature.</p>
</li>
<li>
-<p>Can be qualified by either <code><em>constant</code> or <code></em>global</code> for OpenCL C 2.0 or
+<p>Can be qualified by either <code>__constant</code> or <code>__global</code> for OpenCL C 2.0 or
OpenCL C 3.0 with <code>__opencl_c_<wbr>program_<wbr>scope_<wbr>global_<wbr>variables</code> feature.</p>
</li>
</ul>
@@ -4026,7 +4028,7 @@
cannot be initialized.</p>
</div>
<div class="paragraph">
-<p>Variables allocated in the __constant address space are required to be initialized
+<p>Variables allocated in the <code>__constant</code> address space are required to be initialized
and the values used to initialize these variables must be a compile time constant.</p>
</div>
<div class="paragraph">
@@ -5839,6 +5841,19 @@
<h4 id="specifying-attribute-for-unrolling-loops"><a class="anchor" href="#specifying-attribute-for-unrolling-loops"></a>6.13.5. Specifying Attribute For Unrolling Loops</h4>
<div class="openblock">
<div class="content">
+<div class="admonitionblock note">
+<table>
+<tr>
+<td class="icon">
+<i class="fa icon-note" title="Note"></i>
+</td>
+<td class="content">
+The functionality described in this section <a href="#unified-spec">requires</a>
+support for OpenCL C 2.0 or newer.
+</td>
+</tr>
+</table>
+</div>
<div class="paragraph">
<p>The <code>__attribute__((opencl_unroll_hint))</code> and
<code>__attribute__((opencl_unroll_hint(n)))</code> attribute qualifiers can be used
@@ -11038,11 +11053,12 @@
int <strong>vec_step</strong>(<em>type</em>)</p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">The <strong>vec_step</strong> built-in function takes a built-in scalar or vector
data type argument and returns an integer value representing the
- number of elements in the scalar or vector.</p>
+ number of elements in the scalar or vector. The argument is not
+ evaluated.</p>
<p class="tableblock"> For all scalar types, <strong>vec_step</strong> returns 1.</p>
<p class="tableblock"> The <strong>vec_step</strong> built-in functions that take a 3-component vector
return 4.</p>
-<p class="tableblock"> <strong>vec_step</strong> may also take a pure type as an argument, e.g.
+<p class="tableblock"> <strong>vec_step</strong> may also take a type name as an argument, e.g.
<strong>vec_step</strong>(float2)</p>
<p class="tableblock"> <a href="#unified-spec">Requires</a> support for OpenCL C 1.1 or newer.</p></td>
</tr>
@@ -11145,7 +11161,7 @@
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>Description</strong></p></td>
</tr>
<tr>
-<td class="tableblock halign-left valign-top"><p class="tableblock">int <strong>printf</strong>(constant char restrict <em>format</em>, …​)</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">int <strong>printf</strong>(constant char *restrict <em>format</em>, …​)</p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">The <strong>printf</strong> built-in function writes output to an
implementation-defined stream such as stdout under control of the
string pointed to by <em>format</em> that specifies how subsequent arguments
@@ -13250,7 +13266,7 @@
</table>
</div>
<div class="paragraph">
-<p>This section decribes built-in functions that perform collective options
+<p>This section describes built-in functions that perform collective options
across a work-group.
These built-in functions must be encountered by all work-items in a
work-group executing the kernel.
@@ -13288,10 +13304,12 @@
size_t <em>local_id_y</em>)<br>
gentype <strong>work_group_broadcast</strong>(gentype <em>a</em>, size_t <em>local_id_x</em>,
size_t <em>local_id_y</em>, size_t <em>local_id_z</em>)</p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Broadcast the value of <em>x</em> for work-item identified by <em>local_id</em> to
+<td class="tableblock halign-left valign-top"><p class="tableblock">Broadcast the value of <em>a</em> for work-item identified by <em>local_id</em> to
all work-items in the work-group.</p>
-<p class="tableblock"> <em>local_id</em> must be the same value for all work-items in the
- work-group.</p></td>
+<p class="tableblock"> Behavior is undefined when the value of <em>local_id</em> is not equivalent for
+ all work-items in the work-group.</p>
+<p class="tableblock"> Behavior is undefined when <em>local_id</em> is greater or equal to the
+ work-group size in the corresponding dimension.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock">gentype <strong>work_group_reduce_<op></strong>(gentype <em>x</em>)</p></td>
@@ -13322,16 +13340,9 @@
<strong>min</strong> or <strong>max</strong>.</p>
</div>
<div class="paragraph">
-<p>The inclusive scan operation takes a binary operator <strong>op</strong> with an identity I
-and <em>n</em> (where <em>n</em> is the size of the work-group) elements [a<sub>0</sub>, a<sub>1</sub>, …​
-a<sub>n-1</sub>] and returns [a<sub>0</sub>, (a<sub>0</sub> <strong>op</strong> a<sub>1</sub>), …​ (a<sub>0</sub> <strong>op</strong> a<sub>1</sub> <strong>op</strong> …​
-<strong>op</strong> a<sub>n-1</sub>)].
-If <strong>op</strong> = add, the identity I is 0.
-If <strong>op</strong> = min, the identity I is <code>INT_MAX</code>, <code>UINT_MAX</code>, <code>LONG_MAX</code>,
-<code>ULONG_MAX</code>, for <code>int</code>, <code>uint</code>, <code>long</code>, <code>ulong</code> types and is <code>+INF</code> for
-floating-point types.
-Similarly if <strong>op</strong> = max, the identity I is <code>INT_MIN</code>, 0, <code>LONG_MIN</code>, 0 and
-<code>-INF</code>.</p>
+<p>The inclusive scan operation takes a binary operator <strong>op</strong> with <em>n</em> (where <em>n</em>
+is the size of the work-group) elements [a<sub>0</sub>, a<sub>1</sub>, …​ a<sub>n-1</sub>] and returns
+[a<sub>0</sub>, (a<sub>0</sub> <strong>op</strong> a<sub>1</sub>), …​ (a<sub>0</sub> <strong>op</strong> a<sub>1</sub> <strong>op</strong> …​ <strong>op</strong> a<sub>n-1</sub>)].</p>
</div>
<div class="paragraph">
<p>Consider the following example:</p>
@@ -13359,6 +13370,12 @@
an identity I and n (where n is the size of the work-group) elements [a<sub>0</sub>,
a<sub>1</sub>, …​ a<sub>n-1</sub>] and returns [I, a<sub>0</sub>, (a<sub>0</sub> <strong>op</strong> a<sub>1</sub>), …​ (a<sub>0</sub> <strong>op</strong>
a<sub>1</sub> <strong>op</strong> …​ <strong>op</strong> a<sub>n-2</sub>)].
+If <strong>op</strong> = add, the identity I is 0.
+If <strong>op</strong> = min, the identity I is <code>INT_MAX</code>, <code>UINT_MAX</code>, <code>LONG_MAX</code>,
+<code>ULONG_MAX</code>, for <code>int</code>, <code>uint</code>, <code>long</code>, <code>ulong</code> types and is <code>+INF</code> for
+floating-point types.
+Similarly if <strong>op</strong> = max, the identity I is <code>INT_MIN</code>, 0, <code>LONG_MIN</code>, 0 and
+<code>-INF</code>.
For the example above, the exclusive scan add operation on the ordered set
[3 1 7 0 4 1 6 3] would return [0 3 4 11 11 15 16 22].</p>
</div>
@@ -13401,8 +13418,8 @@
</div>
<div class="paragraph">
<p>A pipe is identified by specifying the <code>pipe</code> keyword with a type.
-The data type specifies the size of each packet in the pipe.
-The <code>pipe</code> keyword is a type modifier.
+The data type specifies the type of each element in the pipe.
+The <code>pipe</code> keyword is a type specifier.
When it is applied to another type <strong>T</strong>, the result is a pipe type whose
elements (or packets) are of type <strong>T</strong>.
The packet type <strong>T</strong> may be any supported OpenCL C scalar and vector integer
@@ -13421,7 +13438,7 @@
</div>
<div class="paragraph">
<p>The <code>read_only</code> (or <code>__read_only</code>) and <code>write_only</code> (or <code>__write_only</code>)
-qualifiers must be used with the <code>pipe</code> qualifier when a pipe is a parameter
+qualifiers must be used with the <code>pipe</code> specifier when a pipe is a parameter
of a kernel or of a user-defined function to identify if a pipe can be read
from or written to by a kernel and its callees and enqueued child kernels.
If no qualifier is specified, <code>read_only</code> is assumed.</p>
@@ -13429,7 +13446,7 @@
<div class="paragraph">
<p>A kernel cannot read from and write to the same pipe object.
Using the <code>read_write</code> (or <code>__read_write</code>) qualifier with the <code>pipe</code>
-qualifier is a compilation error.</p>
+specifier is a compilation error.</p>
</div>
<div class="paragraph">
<p>In the following example</p>
@@ -13458,10 +13475,10 @@
<p>Pipes can only be passed as arguments to a function (including kernel
functions).
The <a href="#operators">C operators</a> cannot be used with variables declared
-with the pipe qualifier.</p>
+with the pipe specifier.</p>
</li>
<li>
-<p>The <code>pipe</code> qualifier cannot be used with variables declared inside a
+<p>The <code>pipe</code> specifier cannot be used with variables declared inside a
kernel, a structure or union field, a pointer type, an array, global
variables declared in program scope or the return type of a function.</p>
</li>
@@ -14638,8 +14655,10 @@
<td class="tableblock halign-left valign-top"><p class="tableblock">Broadcast the value of <em>x</em> for work-item identified by
<em>sub_group_local_id</em> (value returned by <strong>get_sub_group_local_id</strong>) to all
work-items in the subgroup.</p>
-<p class="tableblock"> <em>sub_group_local_id</em> must be the same value for all work-items in the
- subgroup.</p></td>
+<p class="tableblock"> Behavior is undefined when the value of <em>sub_group_local_id</em> is not
+ equivalent for all work-items in the subgroup.</p>
+<p class="tableblock"> Behavior is undefined when <em>sub_group_local_id</em> is greater or equal to the
+ subgroup size.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock">gentype <strong>sub_group_reduce_<op></strong> (<br>
@@ -15001,7 +15020,7 @@
<td class="tableblock halign-left valign-top"><p class="tableblock">Correctly rounded</p></td>
</tr>
<tr>
-<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>1.0 / <em>x</em></strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">1.0 / <em>x</em></p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">≤ 2.5 ulp</p></td>
</tr>
<tr>
@@ -15508,7 +15527,7 @@
<td class="tableblock halign-left valign-top"><p class="tableblock">Correctly rounded</p></td>
</tr>
<tr>
-<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>1.0 / <em>x</em></strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">1.0 / <em>x</em></p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">≤ 3 ulp</p></td>
</tr>
<tr>
@@ -15995,16 +16014,16 @@
<table id="table-float-ulp-relaxed" class="tableblock frame-all grid-all stretch">
<caption class="title">Table 47. ULP values for single precision built-in math functions with unsafe math optimizations in the full and embedded profiles</caption>
<colgroup>
-<col style="width: 50%;">
-<col style="width: 50%;">
+<col style="width: 30%;">
+<col style="width: 70%;">
</colgroup>
<tbody>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>Function</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>Min Accuracy - ULP values</strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>Minimum Accuracy</strong></p></td>
</tr>
<tr>
-<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>1.0 / <em>x</em></strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">1.0 / <em>x</em></p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">≤ 2.5 ulp for <em>x</em> in the domain of 2<sup>-126</sup> to 2<sup>126</sup> for the full
profile, and ≤ 3 ulp for the embedded profile.</p></td>
</tr>
@@ -16019,8 +16038,13 @@
<td class="tableblock halign-left valign-top"><p class="tableblock">≤ 4096 ulp</p></td>
</tr>
<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>acosh</strong>(<em>x</em>)</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>log</strong>(<em>x</em> + <strong>sqrt</strong>(<em>x</em> * <em>x</em> - 1)).
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
+</tr>
+<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>acospi</strong>(<em>x</em>)</p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Implemented as <strong>acos</strong>(<em>x</em>) * <code>M_PI_F</code>.
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>acos</strong>(<em>x</em>) * <code>M_PI_F</code>.
For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
@@ -16028,8 +16052,13 @@
<td class="tableblock halign-left valign-top"><p class="tableblock">≤ 4096 ulp</p></td>
</tr>
<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>asinh</strong>(<em>x</em>)</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>log</strong>(<em>x</em> + <strong>sqrt</strong>(<em>x</em> * <em>x</em> + 1)).
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
+</tr>
+<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>asinpi</strong>(<em>x</em>)</p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Implemented as <strong>asin</strong>(<em>x</em>) * <code>M_PI_F</code>.
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>asin</strong>(<em>x</em>) * <code>M_PI_F</code>.
For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
@@ -16037,32 +16066,32 @@
<td class="tableblock halign-left valign-top"><p class="tableblock">≤ 4096 ulp</p></td>
</tr>
<tr>
-<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>atan2</strong>(<em>y</em>, <em>x</em>)</p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Implemented as <strong>atan</strong>(<em>y</em> / <em>x</em>) for <em>x</em> > 0, <strong>atan</strong>(<em>y</em> / <em>x</em>)<br>
- <code>M_1_PI_F</code> for <em>x</em> < 0 and <em>y</em> > 0 and <strong>atan</strong>(<em>y</em> / <em>x</em>) -
- <code>M_1_PI_F</code> for <em>x</em> < 0 and <em>y</em> < 0.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>atanh</strong>(<em>x</em>)</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Defined for <em>x</em> in the domain (-1, 1).
+ For <em>x</em> in [-2<sup>-10</sup>, 2<sup>-10</sup>], derived implementations may implement as <em>x</em>.
+ For <em>x</em> outside of [-2<sup>-10</sup>, 2<sup>-10</sup>], derived implementations may implement as
+ 0.5f * <strong>log</strong>1.0f + <em>x</em>) / (1.0f - <em>x</em>.
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>atanpi</strong>(<em>x</em>)</p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Implemented as <strong>atan</strong>(<em>x</em>) * <code>M_1_PI_F</code>.
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>atan</strong>(<em>x</em>) * <code>M_1_PI_F</code>.
For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>atan2</strong>(<em>y</em>, <em>x</em>)</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>atan</strong>(<em>y</em> / <em>x</em>) for <em>x</em> > 0,
+ <strong>atan</strong>(<em>y</em> / <em>x</em>) + <code>M_PI_F</code> for <em>x</em> < 0 and <em>y</em> > 0, and
+ <strong>atan</strong>(<em>y</em> / <em>x</em>) - <code>M_PI_F</code> for <em>x</em> < 0 and <em>y</em> < 0.</p></td>
+</tr>
+<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>atan2pi</strong>(<em>y</em>, <em>x</em>)</p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Implemented as <strong>atan2</strong>(<em>y</em>, <em>x</em>) * <code>M_PI_F</code>.
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>atan2</strong>(<em>y</em>, <em>x</em>) * <code>M_1_PI_F</code>.
For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
-<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>acosh</strong>(<em>x</em>)</p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Implemented as <strong>log</strong>(<em>x</em> + <strong>sqrt</strong>(<em>x</em> * <em>x</em> - 1)).</p></td>
-</tr>
-<tr>
-<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>asinh</strong>(<em>x</em>)</p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Implemented as <strong>log</strong>(<em>x</em> + <strong>sqrt</strong>(<em>x</em> * <em>x</em> + 1)).</p></td>
-</tr>
-<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>cbrt</strong>(<em>x</em>)</p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Implemented as <strong>rootn</strong>(<em>x</em>, 3).
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>rootn</strong>(<em>x</em>, 3).
For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
@@ -16072,8 +16101,8 @@
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>cosh</strong>(<em>x</em>)</p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Defined for <em>x</em> in the domain [-88,88] and implemented as 0.5f *
- (<strong>exp</strong>(<em>x</em>) + <strong>exp</strong>(-<em>x</em>)).
+<td class="tableblock halign-left valign-top"><p class="tableblock">Defined for <em>x</em> in the domain [-88, 88].
+ Derived implementations may implement as 0.5f * (<strong>exp</strong>(<em>x</em>) + <strong>exp</strong>(-<em>x</em>)).
For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
@@ -16093,74 +16122,74 @@
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>exp10</strong>(<em>x</em>)</p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations implement this as <strong>exp2</strong>(<em>x</em> * <strong>log2</strong>(10)).
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>exp2</strong>(<em>x</em> * <strong>log2</strong>(10)).
For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>expm1</strong>(<em>x</em>)</p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations implement this as <strong>exp</strong>(<em>x</em>) - 1.
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>exp</strong>(<em>x</em>) - 1.
For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>log</strong>(<em>x</em>)</p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">For <em>x</em> in the domain [0.5, 2] the maximum absolute error is ≤
- 2<sup>-21</sup>; otherwise the maximum error is ≤3 ulp for the full profile
- and ≤ 4 ulp for the embedded profile</p></td>
+ 2<sup>-21</sup>; otherwise the maximum error is ≤ 3 ulp for the full profile
+ and ≤ 4 ulp for the embedded profile.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>log2</strong>(<em>x</em>)</p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">For <em>x</em> in the domain [0.5, 2] the maximum absolute error is ≤
- 2<sup>-21</sup>; otherwise the maximum error is ≤3 ulp for the full profile
- and ≤ 4 ulp for the embedded profile</p></td>
+ 2<sup>-21</sup>; otherwise the maximum error is ≤ 3 ulp for the full profile
+ and ≤ 4 ulp for the embedded profile.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>log10</strong>(<em>x</em>)</p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">For <em>x</em> in the domain [0.5, 2] the maximum absolute error is ≤
- 2<sup>-21</sup>; otherwise the maximum error is ≤3 ulp for the full profile
- and ≤ 4 ulp for the embedded profile</p></td>
+ 2<sup>-21</sup>; otherwise the maximum error is ≤ 3 ulp for the full profile
+ and ≤ 4 ulp for the embedded profile.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>log1p</strong>(<em>x</em>)</p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations implement this as <strong>log</strong>(<em>x</em> + 1).
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>log</strong>(<em>x</em> + 1).
For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>pow</strong>(<em>x</em>, <em>y</em>)</p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">Undefined for <em>x</em> = 0 and <em>y</em> = 0.
- Undefined for <em>x</em> < 0 and non-integer y.
- Undefined for <em>x</em> < 0 and <em>y</em> outside the domain [-2^24, 2^24].
- For <em>x</em> > 0 or <em>x</em> < 0 and even <em>y</em>, derived implementations implement
- this as <strong>exp2</strong>(<em>y</em> * <strong>log2</strong>(<strong>fabs</strong>(<em>x</em>))).
- For <em>x</em> < 0 and odd <em>y</em>, derived implementations implement this as
+ Undefined for <em>x</em> < 0 and non-integer <em>y</em>.
+ Undefined for <em>x</em> < 0 and <em>y</em> outside the domain [-2<sup>24</sup>, 2<sup>24</sup>].
+ For <em>x</em> > 0 or <em>x</em> < 0 and even <em>y</em>, derived implementations may implement as
+ <strong>exp2</strong>(<em>y</em> * <strong>log2</strong>(<strong>fabs</strong>(<em>x</em>))).
+ For <em>x</em> < 0 and odd <em>y</em>, derived implementations may implement as
-<strong>exp2</strong>(<em>y</em> * <strong>log2</strong>(<strong>fabs</strong>(<em>x</em>)).
- For <em>x</em> == 0 and nonzero <em>y</em>, derived implementations return zero.
+ For <em>x</em> == 0 and non-zero <em>y</em>, for derived implementations may return zero.
For non-derived implementations, the error is ≤ 8192 ulp.
<sup class="footnote">[<a id="_footnoteref_82" class="footnote" href="#_footnotedef_82" title="View footnote.">82</a>]</sup></p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>pown</strong>(<em>x</em>, <em>y</em>)</p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Defined only for integer values of y.
+<td class="tableblock halign-left valign-top"><p class="tableblock">Defined only for integer values of <em>y</em>.
Undefined for <em>x</em> = 0 and <em>y</em> = 0.
- For <em>x</em> >= 0 or <em>x</em> < 0 and even <em>y</em>, derived implementations
- implement this as <strong>exp2</strong>(<em>y</em> * <strong>log2</strong>(<strong>fabs</strong>(<em>x</em>))).
- For <em>x</em> < 0 and odd <em>y</em>, derived implementations implement this as
- -<strong>exp2</strong>(<em>y</em> * <strong>log2</strong>(<strong>fabs</strong>(<em>x</em>)).
+ For <em>x</em> >= 0 or <em>x</em> < 0 and even <em>y</em>, derived implementations may implement as
+ <strong>exp2</strong>(<em>y</em> * <strong>log2</strong>(<strong>fabs</strong>(<em>x</em>))).
+ For <em>x</em> < 0 and odd <em>y</em>, derived implementations may implement as
+ -<strong>exp2</strong>(<em>y</em> * <strong>log2</strong>(<strong>fabs</strong>(<em>x</em>))).
For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>powr</strong>(<em>x</em>, <em>y</em>)</p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">Defined only for <em>x</em> >= 0.
Undefined for <em>x</em> = 0 and <em>y</em> = 0.
- Derived implementations implement this as <strong>exp2</strong>(<em>y</em> * <strong>log2</strong>(<em>x</em>)).
+ Derived implementations may implement as <strong>exp2</strong>(<em>y</em> * <strong>log2</strong>(<em>x</em>)).
For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>rootn</strong>(<em>x</em>, <em>y</em>)</p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Defined for <em>x</em> > 0 when <em>y</em> is nonzero, derived implementations
- implement this case as <strong>exp2</strong>(log2(<em>x</em>) / <em>y</em>).
- Defined for <em>x</em> < 0 when <em>y</em> is odd, derived implementations implement
- this case as -<strong>exp2</strong>(<strong>log2</strong>(-<em>x</em>) / <em>y</em>).
- Defined for <em>x</em> = +/-0 when <em>y</em> > 0, derived implementations will
+<td class="tableblock halign-left valign-top"><p class="tableblock">Defined for <em>x</em> > 0 when <em>y</em> is non-zero, derived implementations
+ may implement this case as <strong>exp2</strong>(<strong>log2</strong>(<em>x</em>) / <em>y</em>).
+ Defined for <em>x</em> < 0 when <em>y</em> is odd, derived implementations
+ may implement this case as -<strong>exp2</strong>(<strong>log2</strong>(-<em>x</em>) / <em>y</em>).
+ Defined for <em>x</em> = +/-0 when <em>y</em> > 0, derived implementations may
return +0 in this case.
For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
@@ -16171,14 +16200,15 @@
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>sincos</strong>(<em>x</em>)</p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">ulp values as defined for <strong>sin</strong>(<em>x</em>) and <strong>cos</strong>(<em>x</em>)</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">ulp values as defined for <strong>sin</strong>(<em>x</em>) and <strong>cos</strong>(<em>x</em>).</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>sinh</strong>(<em>x</em>)</p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">Defined for <em>x</em> in the domain [-88,88].
- For <em>x</em> in [-2<sup>-10,2</sup>-10], derived implementations implement as <em>x</em>.
- For <em>x</em> outside of [-2<sup>10,2</sup>10], derived implement as <strong>0.5f *
- (*exp</strong>(<em>x</em>) - <strong>exp</strong>(-<em>x</em>)).
+ For <em>x</em> in [-2<sup>-10</sup>, 2<sup>-10</sup>], derived implementations
+ may implement as <em>x</em>.
+ For <em>x</em> outside of [-2<sup>-10</sup>, 2<sup>-10</sup>], derived implementations
+ may implement as 0.5f * (<strong>exp</strong>(<em>x</em>) - <strong>exp</strong>(-<em>x</em>)).
For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
@@ -16188,13 +16218,22 @@
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>tan</strong>(<em>x</em>)</p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations implement this as <strong>sin</strong>(<em>x</em>) * (<code>1.0f</code> /
- <strong>cos</strong>(<em>x</em>)).
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as
+ <strong>sin</strong>(<em>x</em>) * (1.0f / <strong>cos</strong>(<em>x</em>)).
For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>tanh</strong>(<em>x</em>)</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Defined for <em>x</em> in the domain [-∞, ∞].
+ For <em>x</em> in [-2<sup>-10</sup>, 2<sup>-10</sup>], derived implementations
+ may implement as <em>x</em>.
+ For <em>x</em> outside of [-2<sup>-10</sup>, 2<sup>-10</sup>], derived implementations
+ may implement as (<strong>exp</strong>(<em>x</em>) - <strong>exp</strong>(-<em>x</em>)) / (<strong>exp</strong>(<em>x</em>) + <strong>exp</strong>(-<em>x</em>)).
+ For non-derived implementations, the error is ≤ 8192 ULP.</p></td>
+</tr>
+<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>tanpi</strong>(<em>x</em>)</p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations implement this as <strong>tan</strong>(<em>x</em> * <code>M_PI_F</code>).
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>tan</strong>(<em>x</em> * <code>M_PI_F</code>).
For non-derived implementations, the error is ≤ 8192 ulp for <em>x</em>
in the domain [-1, 1].</p></td>
</tr>
@@ -18340,8 +18379,8 @@
</div>
<div id="footer">
<div id="footer-text">
-Version v3.0.6<br>
-Last updated 2020-12-18 09:20:25 -0800
+Version v3.0.7<br>
+Last updated 2021-04-23 12:57:31 -0700
</div>
</div>
diff --git a/specs/3.0-unified/html/OpenCL_Env.html b/specs/3.0-unified/html/OpenCL_Env.html
index 3ad4e0d..d61cb3b 100644
--- a/specs/3.0-unified/html/OpenCL_Env.html
+++ b/specs/3.0-unified/html/OpenCL_Env.html
@@ -841,9 +841,9 @@
<h1>The OpenCL<sup>™</sup> SPIR-V Environment Specification</h1>
<div class="details">
<span id="author" class="author">Khronos<sup>®</sup> OpenCL Working Group</span><br>
-<span id="revnumber">version v3.0.6,</span>
-<span id="revdate">Fri, 18 Dec 2020 12:00:00 +0000</span>
-<br><span id="revremark">from git branch: master commit: e9a4d468b1a0a38c1e10b8af484bb2bbb495e2b7</span>
+<span id="revnumber">version v3.0.7,</span>
+<span id="revdate">Fri, 23 Apr 2021 20:00:00 +0000</span>
+<br><span id="revremark">from git branch: master commit: 133503b85911ca8fd6642721d440e89460d491a3</span>
</div>
<div id="toc" class="toc2">
<div id="toctitle">Table of Contents</div>
@@ -1178,6 +1178,9 @@
<div class="paragraph">
<p><strong>OpTypeSampler</strong> may be used to declare sampler types in OpenCL environments.</p>
</div>
+<div class="paragraph">
+<p><strong>OpTypeSampledImage</strong> may be used to declare combined image and sampler types in OpenCL environments.</p>
+</div>
</div>
<div class="sect3">
<h4 id="_other_data_types"><a class="anchor" href="#_other_data_types"></a>2.5.3. Other Data Types</h4>
@@ -1526,6 +1529,124 @@
a built-in variable.
All built-in variables must be in the <strong>Input</strong> storage class.</p>
</div>
+<div class="paragraph">
+<p>The following table describes the required SPIR-V type for built-in variables.
+In this table, <code>size_t</code> is used as a generic type to represent:</p>
+</div>
+<div class="ulist">
+<ul>
+<li>
+<p><strong>OpTypeInt</strong> with <em>Width</em> equal to 32 if the <em>Addressing Model</em> declared in <strong>OpMemoryModel</strong> is <strong>Physical32</strong>.</p>
+</li>
+<li>
+<p><strong>OpTypeInt</strong> with <em>Width</em> equal to 64 if the <em>Addressing Model</em> declared in <strong>OpMemoryModel</strong> is <strong>Physical64</strong>.</p>
+</li>
+</ul>
+</div>
+<div class="paragraph">
+<p>The mapping from an OpenCL C built-in function to the SPIR-V <strong>BuiltIn</strong> is informational and non-normative.</p>
+</div>
+<table class="tableblock frame-all grid-all stretch">
+<colgroup>
+<col style="width: 28.5714%;">
+<col style="width: 28.5714%;">
+<col style="width: 42.8572%;">
+</colgroup>
+<thead>
+<tr>
+<th class="tableblock halign-left valign-top"><strong>OpenCL C Function</strong></th>
+<th class="tableblock halign-left valign-top"><strong>SPIR-V BuiltIn</strong></th>
+<th class="tableblock halign-left valign-top"><strong>Required SPIR-V Type</strong></th>
+</tr>
+</thead>
+<tbody>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>get_work_dim</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>WorkDim</strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpTypeInt</strong> with <em>Width</em> equal to 32</p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>get_global_size</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>GlobalSize</strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpTypeVector</strong> of 3 components of <code>size_t</code></p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>get_global_id</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>GlobalInvocationId</strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpTypeVector</strong> of 3 components of <code>size_t</code></p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>get_local_size</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>WorkgroupSize</strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpTypeVector</strong> of 3 components of <code>size_t</code></p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>get_enqueued_local_size</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>EnqueuedWorkgroupSize</strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpTypeVector</strong> of 3 components of <code>size_t</code></p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>get_local_id</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>LocalInvocationId</strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpTypeVector</strong> of 3 components of <code>size_t</code></p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>get_num_groups</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>NumWorkgroups</strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpTypeVector</strong> of 3 components of <code>size_t</code></p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>get_group_id</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>WorkgroupId</strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpTypeVector</strong> of 3 components of <code>size_t</code></p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>get_global_offset</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>GlobalOffset</strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpTypeVector</strong> of 3 components of <code>size_t</code></p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>get_global_linear_id</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>GlobalLinearId</strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>size_t</code></p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>get_local_linear_id</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>LocalInvocationIndex</strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>size_t</code></p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>get_sub_group_size</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>SubgroupSize</strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpTypeInt</strong> with <em>Width</em> equal to 32</p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>get_max_sub_group_size</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>SubgroupMaxSize</strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpTypeInt</strong> with <em>Width</em> equal to 32</p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>get_num_sub_groups</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>NumSubgroups</strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpTypeInt</strong> with <em>Width</em> equal to 32</p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>get_enqueued_num_sub_groups</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>NumEnqueuedSubgroups</strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpTypeInt</strong> with <em>Width</em> equal to 32</p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>get_sub_group_id</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>SubgroupId</strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpTypeInt</strong> with <em>Width</em> equal to 32</p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>get_sub_group_local_id</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>SubgroupLocalInvocationId</strong></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpTypeInt</strong> with <em>Width</em> equal to 32</p></td>
+</tr>
+</tbody>
+</table>
</div>
<div class="sect2">
<h3 id="_alignment_of_types"><a class="anchor" href="#_alignment_of_types"></a>2.10. Alignment of Types</h3>
@@ -1821,7 +1942,7 @@
</li>
<li>
<p>The <em>Pointer</em> operand must be a pointer to the <strong>Function</strong>, <strong>Workgroup</strong>,
-or <strong>CrossWorkGroup</strong> <em>Storage Classes</em>. Note that an <strong>Atomic Instruction</strong>
+or <strong>CrossWorkgroup</strong> <em>Storage Classes</em>. Note that an <strong>Atomic Instruction</strong>
on a pointer to the <strong>Function</strong> <em>Storage Class</em> is valid, but does not
have defined behavior.</p>
</li>
@@ -2765,6 +2886,22 @@
<p>When the <strong>GroupNonUniformClustered</strong> capability is declared, the <strong>GroupNonUniformArithmetic</strong> scan and reduction instructions may include the optional <em>ClusterSize</em> operand.</p>
</div>
</div>
+<div class="sect3">
+<h4 id="_cl_khr_spirv_extended_debug_info"><a class="anchor" href="#_cl_khr_spirv_extended_debug_info"></a>5.2.21. <code>cl_khr_spirv_extended_debug_info</code></h4>
+<div class="paragraph">
+<p>If the OpenCL environment supports the extension <code>cl_khr_spirv_extended_debug_info</code>, then the environment must accept modules
+that import the <code>OpenCL.DebugInfo.100</code> extended instruction set via <strong>OpExtInstImport</strong>.</p>
+</div>
+</div>
+<div class="sect3">
+<h4 id="_cl_khr_spirv_linkonce_odr"><a class="anchor" href="#_cl_khr_spirv_linkonce_odr"></a>5.2.22. <code>cl_khr_spirv_linkonce_odr</code></h4>
+<div class="paragraph">
+<p>If the OpenCL environment supports the extension <code>cl_khr_spirv_linkonce_odr</code>, then the environment must accept modules that declare use of the extension <code>SPV_KHR_linkonce_odr</code> via <strong>OpExtension</strong>.</p>
+</div>
+<div class="paragraph">
+<p>If the OpenCL environment supports the extension <code>cl_khr_spirv_linkonce_odr</code> and use of the SPIR-V extension <code>SPV_KHR_linkonce_odr</code> is declared in the module via <strong>OpExtension</strong>, then the environment must accept modules that include the <strong>LinkOnceODR</strong> linkage type.</p>
+</div>
+</div>
</div>
<div class="sect2">
<h3 id="_embedded_profile_extensions"><a class="anchor" href="#_embedded_profile_extensions"></a>5.3. Embedded Profile Extensions</h3>
@@ -4442,229 +4579,235 @@
</colgroup>
<thead>
<tr>
-<th class="tableblock halign-left valign-top"><strong>SPIR-V Instruction</strong></th>
+<th class="tableblock halign-left valign-top"><strong>Function</strong></th>
<th class="tableblock halign-left valign-top"><strong>Minimum Accuracy</strong></th>
</tr>
</thead>
<tbody>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpFDiv</strong> for 1.0 / <em>x</em></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock"><= 2.5 ulp for x in the domain of 2<sup>-126</sup> to 2<sup>126</sup> for the full profile,
- and <= 3 ulp for the embedded profile.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">≤ 2.5 ulp for <em>x</em> in the domain of 2<sup>-126</sup> to 2<sup>126</sup> for the full
+ profile, and ≤ 3 ulp for the embedded profile.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpFDiv</strong> for <em>x</em> / <em>y</em></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock"><= 2.5 ulp for x in the domain of 2<sup>-62</sup> to 2<sup>62</sup> and <em>y</em> in the domain
- of 2<sup>-62</sup> to 2<sup>62</sup> for the full profile, and <= 3 ulp for the embedded
- profile.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">≤ 2.5 ulp for <em>x</em> in the domain of 2<sup>-62</sup> to 2<sup>62</sup> and <em>y</em> in the
+ domain of 2<sup>-62</sup> to 2<sup>62</sup> for the full profile, and ≤ 3 ulp for
+ the embedded profile.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>acos</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock"><= 4096 ulp</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">≤ 4096 ulp</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>acosh</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Implemented as log( x + sqrt(x*x - 1) ).</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>log</strong>(<em>x</em> + <strong>sqrt</strong>(<em>x</em> * <em>x</em> - 1)).
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>acospi</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Implemented as acos(x) * M_PI_F.
- For non-derived implementations, the error is <= 8192 ulp.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>acos</strong>(<em>x</em>) * <code>M_PI_F</code>.
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>asin</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock"><= 4096 ulp</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">≤ 4096 ulp</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>asinh</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Implemented as log( x + sqrt(x*x + 1) ).</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>log</strong>(<em>x</em> + <strong>sqrt</strong>(<em>x</em> * <em>x</em> + 1)).
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>asinpi</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Implemented as asin(x) * M_PI_F.
- For non-derived implementations, the error is <= 8192 ulp.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>asin</strong>(<em>x</em>) * <code>M_PI_F</code>.
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>atan</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock"><= 4096 ulp</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">≤ 4096 ulp</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>atanh</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Defined for x in the domain (-1, 1).
- For x in [-2<sup>-10</sup>, 2<sup>-10</sup>], implemented as x.
- For x outside of [-2<sup>-10</sup>, 2<sup>-10</sup>], implemented as 0.5f * log( (1.0f + x)
- / (1.0f - x) ).
- For non-derived implementations, the error is <= 8192 ulp.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Defined for <em>x</em> in the domain (-1, 1).
+ For <em>x</em> in [-2<sup>-10</sup>, 2<sup>-10</sup>], derived implementations may implement as <em>x</em>.
+ For <em>x</em> outside of [-2<sup>-10</sup>, 2<sup>-10</sup>], derived implementations may implement as
+ 0.5f * <strong>log</strong>1.0f + <em>x</em>) / (1.0f - <em>x</em>.
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>atanpi</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Implemented as atan(x) * M_1_PI_F.
- For non-derived implementations, the error is <= 8192 ulp.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>atan</strong>(<em>x</em>) * <code>M_1_PI_F</code>.
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>atan2</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Implemented as atan(y/x) for x > 0, atan(y/x) + M_PI_F for x < 0 and y >
- 0, and atan(y/x) - M_PI_F for x < 0 and y < 0.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>atan</strong>(<em>y</em> / <em>x</em>) for <em>x</em> > 0,
+ <strong>atan</strong>(<em>y</em> / <em>x</em>) + <code>M_PI_F</code> for <em>x</em> < 0 and <em>y</em> > 0, and
+ <strong>atan</strong>(<em>y</em> / <em>x</em>) - <code>M_PI_F</code> for <em>x</em> < 0 and <em>y</em> < 0.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>atan2pi</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Implemented as atan2(y, x) * M_1_PI_F.
- For non-derived implementations, the error is <= 8192 ulp.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>atan2</strong>(<em>y</em>, <em>x</em>) * <code>M_1_PI_F</code>.
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>cbrt</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Implemented as rootn(x, 3).
- For non-derived implementations, the error is <= 8192 ulp.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>rootn</strong>(<em>x</em>, 3).
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>cos</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">For x in the domain [-π, π], the maximum absolute error is <=
- 2<sup>-11</sup> and larger otherwise.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">For <em>x</em> in the domain [-π, π], the maximum absolute error
+ is ≤ 2<sup>-11</sup> and larger otherwise.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>cosh</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Defined for x in the domain [-∞, ∞] and implemented as 0.5f * (
- exp(x) + exp(-x) ).
- For non-derived implementations, the error is <= 8192 ULP.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Defined for <em>x</em> in the domain [-88, 88].
+ Derived implementations may implement as 0.5f * (<strong>exp</strong>(<em>x</em>) + <strong>exp</strong>(-<em>x</em>)).
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>cospi</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">For x in the domain [-1, 1], the maximum absolute error is <= 2<sup>-11</sup> and
- larger otherwise.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">For <em>x</em> in the domain [-1, 1], the maximum absolute error is ≤
+ 2<sup>-11</sup> and larger otherwise.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>exp</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock"><= 3 + floor( fabs(2 * x) ) ulp for the full profile, and <= 4 ulp for
- the embedded profile.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">≤ 3 + <strong>floor</strong>(<strong>fabs</strong>(2 * <em>x</em>)) ulp for the full profile, and ≤
+ 4 ulp for the embedded profile.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>exp2</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock"><= 3 + floor( fabs(2 * x) ) ulp for the full profile, and <= 4 ulp for
- the embedded profile.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">≤ 3 + <strong>floor</strong>(<strong>fabs</strong>(2 * <em>x</em>)) ulp for the full profile, and ≤
+ 4 ulp for the embedded profile.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>exp10</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations implement this as exp2( x * log2(10) ).
- For non-derived implementations, the error is <= 8192 ulp.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>exp2</strong>(<em>x</em> * <strong>log2</strong>(10)).
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>expm1</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations implement this as exp(x) - 1.
- For non-derived implementations, the error is <= 8192 ulp.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>exp</strong>(<em>x</em>) - 1.
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>log</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">For x in the domain [0.5, 2] the maximum absolute error is <= 2<sup>-21</sup>;
- otherwise the maximum error is <=3 ulp for the full profile and <= 4 ulp
- for the embedded profile</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">For <em>x</em> in the domain [0.5, 2] the maximum absolute error is ≤
+ 2<sup>-21</sup>; otherwise the maximum error is ≤ 3 ulp for the full profile
+ and ≤ 4 ulp for the embedded profile.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>log2</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">For x in the domain [0.5, 2] the maximum absolute error is <= 2<sup>-21</sup>;
- otherwise the maximum error is <=3 ulp for the full profile and <= 4 ulp
- for the embedded profile</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">For <em>x</em> in the domain [0.5, 2] the maximum absolute error is ≤
+ 2<sup>-21</sup>; otherwise the maximum error is ≤ 3 ulp for the full profile
+ and ≤ 4 ulp for the embedded profile.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>log10</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">For x in the domain [0.5, 2] the maximum absolute error is <= 2<sup>-21</sup>;
- otherwise the maximum error is <=3 ulp for the full profile and <= 4 ulp
- for the embedded profile</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">For <em>x</em> in the domain [0.5, 2] the maximum absolute error is ≤
+ 2<sup>-21</sup>; otherwise the maximum error is ≤ 3 ulp for the full profile
+ and ≤ 4 ulp for the embedded profile.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>log1p</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations implement this as log(x + 1).
- For non-derived implementations, the error is <= 8192 ulp.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>log</strong>(<em>x</em> + 1).
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>pow</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Undefined for x = 0 and y = 0.
- Undefined for x < 0 and non-integer y.
- Undefined for x < 0 and y outside the domain [-2<sup>24</sup>, 2<sup>24</sup>].
- For x > 0 or x < 0 and even y, derived implementations implement this as
- exp2( y * log2( fabs(x) ) ).
- For x < 0 and odd y, derived implementations implement this as -exp2( y *
- log2( fabs(x) ).
- For x == 0 and nonzero y, derived implementations return zero.
- For non-derived implementations, the error is <= 8192 ULP.</p>
-<p class="tableblock"> On some implementations, powr() or pown() may perform faster
- than pow().
- If x is known to be >= 0, consider using powr() in place of pow(),
- or if y is known to be an integer, consider using pown() in place of
- pow().</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Undefined for <em>x</em> = 0 and <em>y</em> = 0.
+ Undefined for <em>x</em> < 0 and non-integer <em>y</em>.
+ Undefined for <em>x</em> < 0 and <em>y</em> outside the domain [-2<sup>24</sup>, 2<sup>24</sup>].
+ For <em>x</em> > 0 or <em>x</em> < 0 and even <em>y</em>, derived implementations may implement as
+ <strong>exp2</strong>(<em>y</em> * <strong>log2</strong>(<strong>fabs</strong>(<em>x</em>))).
+ For <em>x</em> < 0 and odd <em>y</em>, derived implementations may implement as
+ -<strong>exp2</strong>(<em>y</em> * <strong>log2</strong>(<strong>fabs</strong>(<em>x</em>)).
+ For <em>x</em> == 0 and non-zero <em>y</em>, for derived implementations may return zero.
+ For non-derived implementations, the error is ≤ 8192 ulp.</p>
+<p class="tableblock"> On some implementations, powr() or pown() may perform faster
+ than pow().
+ If x is known to be >= 0, consider using powr() in place of pow(),
+ or if y is known to be an integer, consider using pown() in place of
+ pow().</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>pown</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Defined only for integer values of y.
- Undefined for x = 0 and y = 0.
- For x >= 0 or x < 0 and even y, derived implementations implement this as
- exp2( y * log2( fabs(x) ) ).
- For x < 0 and odd y, derived implementations implement this as -exp2( y *
- log2( fabs(x) ) ).
- For non-derived implementations, the error is <= 8192 ulp.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Defined only for integer values of <em>y</em>.
+ Undefined for <em>x</em> = 0 and <em>y</em> = 0.
+ For <em>x</em> >= 0 or <em>x</em> < 0 and even <em>y</em>, derived implementations may implement as
+ <strong>exp2</strong>(<em>y</em> * <strong>log2</strong>(<strong>fabs</strong>(<em>x</em>))).
+ For <em>x</em> < 0 and odd <em>y</em>, derived implementations may implement as
+ -<strong>exp2</strong>(<em>y</em> * <strong>log2</strong>(<strong>fabs</strong>(<em>x</em>))).
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>powr</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Defined only for x >= 0.
- Undefined for x = 0 and y = 0.
- Derived implementations implement this as exp2( y * log2(x) ).
- For non-derived implementations, the error is <= 8192 ulp.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Defined only for <em>x</em> >= 0.
+ Undefined for <em>x</em> = 0 and <em>y</em> = 0.
+ Derived implementations may implement as <strong>exp2</strong>(<em>y</em> * <strong>log2</strong>(<em>x</em>)).
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>rootn</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Defined for x > 0 when y is non-zero, derived implementations implement
- this case as exp2( log2(x) / y ).
- Defined for x < 0 when y is odd, derived implementations implement this
- case as -exp2( log2(-x) / y ).
- Defined for x = +/-0 when y > 0, derived implementations will return +0 in
- this case.
- For non-derived implementations, the error is <= 8192 ULP.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Defined for <em>x</em> > 0 when <em>y</em> is non-zero, derived implementations
+ may implement this case as <strong>exp2</strong>(<strong>log2</strong>(<em>x</em>) / <em>y</em>).
+ Defined for <em>x</em> < 0 when <em>y</em> is odd, derived implementations
+ may implement this case as -<strong>exp2</strong>(<strong>log2</strong>(-<em>x</em>) / <em>y</em>).
+ Defined for <em>x</em> = +/-0 when <em>y</em> > 0, derived implementations may
+ return +0 in this case.
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>sin</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">For x in the domain [-π, π], the maximum absolute error is <=
- 2<sup>-11</sup> and larger otherwise.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">For <em>x</em> in the domain [-π, π], the maximum absolute error is
+ ≤ 2<sup>-11</sup> and larger otherwise.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>sincos</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">ulp values as defined for sin(x) and cos(x).</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">ulp values as defined for <strong>sin</strong>(<em>x</em>) and <strong>cos</strong>(<em>x</em>).</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>sinh</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Defined for x in the domain [-∞, ∞].
- For x in [-2<sup>-10</sup>, 2<sup>-10</sup>], derived implementations implement as x.
- For x outside of [-2<sup>-10</sup>, 2<sup>-10</sup>], derived implement as 0.5f * ( exp(x) -
- exp(-x) ).
- For non-derived implementations, the error is <= 8192 ULP.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Defined for <em>x</em> in the domain [-88, 88].
+ For <em>x</em> in [-2<sup>-10</sup>, 2<sup>-10</sup>], derived implementations
+ may implement as <em>x</em>.
+ For <em>x</em> outside of [-2<sup>-10</sup>, 2<sup>-10</sup>], derived implementations
+ may implement as 0.5f * (<strong>exp</strong>(<em>x</em>) - <strong>exp</strong>(-<em>x</em>)).
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>sinpi</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">For x in the domain [-1, 1], the maximum absolute error is <= 2<sup>-11</sup> and
- larger otherwise.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">For <em>x</em> in the domain [-1, 1], the maximum absolute error is ≤
+ 2<sup>-11</sup> and larger otherwise.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>tan</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations implement this as sin(x) * ( 1.0f / cos(x) ).
- For non-derived implementations, the error is <= 8192 ulp.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as
+ <strong>sin</strong>(<em>x</em>) * (1.0f / <strong>cos</strong>(<em>x</em>)).
+ For non-derived implementations, the error is ≤ 8192 ulp.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>tanh</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Defined for x in the domain [-∞, ∞].
- For x in [-2<sup>-10</sup>, 2<sup>-10</sup>], derived implementations implement as x.
- For x outside of [-2<sup>-10</sup>, 2<sup>-10</sup>], derived implementations implement as
- (exp(x) - exp(-x)) / (exp(x) + exp(-x)).
- For non-derived implementations, the error is <= 8192 ULP.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Defined for <em>x</em> in the domain [-∞, ∞].
+ For <em>x</em> in [-2<sup>-10</sup>, 2<sup>-10</sup>], derived implementations
+ may implement as <em>x</em>.
+ For <em>x</em> outside of [-2<sup>-10</sup>, 2<sup>-10</sup>], derived implementations
+ may implement as (<strong>exp</strong>(<em>x</em>) - <strong>exp</strong>(-<em>x</em>)) / (<strong>exp</strong>(<em>x</em>) + <strong>exp</strong>(-<em>x</em>)).
+ For non-derived implementations, the error is ≤ 8192 ULP.</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpExtInst</strong> <strong>tanpi</strong></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations implement this as tan(x * M_PI_F).
- For non-derived implementations, the error is <= 8192 ulp for x in the
- domain [-1, 1].</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Derived implementations may implement as <strong>tan</strong>(<em>x</em> * <code>M_PI_F</code>).
+ For non-derived implementations, the error is ≤ 8192 ulp for <em>x</em>
+ in the domain [-1, 1].</p></td>
</tr>
<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><strong>OpFMul</strong> and <strong>OpFAdd</strong>,<br>
for <em>x</em> * <em>y</em> + <em>z</em></p></td>
-<td class="tableblock halign-left valign-top"><p class="tableblock">Implemented either as a correctly rounded fma or as a multiply and an add
- both of which are correctly rounded.</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Implemented either as a correctly rounded <strong>fma</strong> or as a multiply and
+ an add both of which are correctly rounded.</p></td>
</tr>
</tbody>
</table>
@@ -6658,8 +6801,8 @@
</div>
<div id="footer">
<div id="footer-text">
-Version v3.0.6<br>
-Last updated 2020-12-08 14:02:39 -0800
+Version v3.0.7<br>
+Last updated 2021-04-19 17:02:25 -0700
</div>
</div>
diff --git a/specs/3.0-unified/html/OpenCL_Ext.html b/specs/3.0-unified/html/OpenCL_Ext.html
index 8a4c4e0..45f58c6 100644
--- a/specs/3.0-unified/html/OpenCL_Ext.html
+++ b/specs/3.0-unified/html/OpenCL_Ext.html
@@ -841,9 +841,9 @@
<h1>The OpenCL<sup>™</sup> Extension Specification</h1>
<div class="details">
<span id="author" class="author">Khronos<sup>®</sup> OpenCL Working Group</span><br>
-<span id="revnumber">version v3.0.6,</span>
-<span id="revdate">Fri, 18 Dec 2020 12:00:00 +0000</span>
-<br><span id="revremark">from git branch: master commit: e9a4d468b1a0a38c1e10b8af484bb2bbb495e2b7</span>
+<span id="revnumber">version v3.0.7,</span>
+<span id="revdate">Fri, 23 Apr 2021 20:00:00 +0000</span>
+<br><span id="revremark">from git branch: master commit: 133503b85911ca8fd6642721d440e89460d491a3</span>
</div>
<div id="toc" class="toc2">
<div id="toctitle">Table of Contents</div>
@@ -886,7 +886,10 @@
<li><a href="#cl_khr_device_uuid">36. Unique Device Identifiers</a></li>
<li><a href="#cl_khr_extended_versioning">37. Extended versioning</a></li>
<li><a href="#_extended_subgroup_functions">38. Extended Subgroup Functions</a></li>
-<li><a href="#spirv_extensions">39. Extensions to the OpenCL SPIR-V Environment</a></li>
+<li><a href="#cl_khr_pci_bus_info">39. PCI Bus Information Query</a></li>
+<li><a href="#cl_khr_extended_bit_ops">40. Extended Bit Operations</a></li>
+<li><a href="#cl_khr_suggested_local_work_size">41. Suggested Local Work Size Query</a></li>
+<li><a href="#spirv_extensions">42. Extensions to the OpenCL SPIR-V Environment</a></li>
<li><a href="#_extensions_promoted_to_core_features">Appendix A: Extensions Promoted to Core Features</a></li>
<li><a href="#_deprecated_extensions">Appendix B: Deprecated Extensions</a></li>
<li><a href="#_quick_reference">Appendix C: Quick Reference</a></li>
@@ -1188,7 +1191,7 @@
<span class="comment">// function pointer typedefs must use the</span>
<span class="comment">// following naming convention</span>
-<span class="keyword">typedef</span> CL_API_ENTRY return_type
+<span class="keyword">typedef</span> return_type
(CL_API_CALL *clExtensionFunctionNameTAG_fn)(...);
<span class="preprocessor">#endif</span> <span class="comment">// _extension_name_</span></code></pre>
@@ -1219,7 +1222,7 @@
<span class="comment">// function pointer typedefs must use the</span>
<span class="comment">// following naming convention</span>
-<span class="keyword">typedef</span> CL_API_ENTRY cl_int
+<span class="keyword">typedef</span> cl_int
(CL_API_CALL *clGetGLContextInfoKHR_fn)(
<span class="directive">const</span> cl_context_properties * <span class="comment">/* properties */</span>,
cl_gl_context_info <span class="comment">/* param_name */</span>,
@@ -15957,7 +15960,7 @@
<p>The exclusive scan operation takes a binary operator <strong>op</strong> with an identity I and <em>n</em> (where <em>n</em> is the size of the sub-group) elements [a<sub>0</sub>, a<sub>1</sub>, …​ a<sub>n-1</sub>] and returns [I, a<sub>0</sub>, (a<sub>0</sub> <strong>op</strong> a<sub>1</sub>), …​ (a<sub>0</sub> <strong>op</strong> a<sub>1</sub> <strong>op</strong> …​ <strong>op</strong> a<sub>n-2</sub>)].</p>
</div>
<div class="paragraph">
-<p>The inclusive scan operation takes a binary operator <strong>op</strong> with an identity I and <em>n</em> (where <em>n</em> is the size of the sub-group) elements [a<sub>0</sub>, a<sub>1</sub>, …​ a<sub>n-1</sub>] and returns [a<sub>0</sub>, (a<sub>0</sub> <strong>op</strong> a<sub>1</sub>), …​ (a<sub>0</sub> <strong>op</strong> a<sub>1</sub> <strong>op</strong> …​ <strong>op</strong> a<sub>n-1</sub>)].</p>
+<p>The inclusive scan operation takes a binary operator <strong>op</strong> with <em>n</em> (where <em>n</em> is the size of the sub-group) elements [a<sub>0</sub>, a<sub>1</sub>, …​ a<sub>n-1</sub>] and returns [a<sub>0</sub>, (a<sub>0</sub> <strong>op</strong> a<sub>1</sub>), …​ (a<sub>0</sub> <strong>op</strong> a<sub>1</sub> <strong>op</strong> …​ <strong>op</strong> a<sub>n-1</sub>)].</p>
</div>
<div class="paragraph">
<p>If <strong>op</strong> = <strong>add</strong>, the identity I is 0.
@@ -19398,7 +19401,512 @@
</div>
</div>
<div class="sect1">
-<h2 id="spirv_extensions"><a class="anchor" href="#spirv_extensions"></a>39. Extensions to the OpenCL SPIR-V Environment</h2>
+<h2 id="cl_khr_pci_bus_info"><a class="anchor" href="#cl_khr_pci_bus_info"></a>39. PCI Bus Information Query</h2>
+<div class="sectionbody">
+<div class="paragraph">
+<p>This extension adds a new query to obtain PCI bus information about an OpenCL
+device.</p>
+</div>
+<div class="paragraph">
+<p>Not all OpenCL devices have PCI bus information, either due to the device not
+being connected to the system through a PCI interface or due to platform
+specific restrictions and policies. Thus this extension is only expected to be
+supported by OpenCL devices which can provide the information.</p>
+</div>
+<div class="paragraph">
+<p>As a consequence, applications should always check for the presence of the
+extension string for each individual OpenCL device for which they intend to
+issue the new query for and should not have any assumptions about the
+availability of the extension on any given platform.</p>
+</div>
+<div class="sect2">
+<h3 id="_general_information_38"><a class="anchor" href="#_general_information_38"></a>39.1. General information</h3>
+<div class="sect3">
+<h4 id="_name_strings_2"><a class="anchor" href="#_name_strings_2"></a>39.1.1. Name Strings</h4>
+<div class="paragraph">
+<p><code>cl_khr_pci_bus_info</code></p>
+</div>
+</div>
+<div class="sect3">
+<h4 id="_version_history_38"><a class="anchor" href="#_version_history_38"></a>39.1.2. Version History</h4>
+<table class="tableblock frame-all grid-all stretch">
+<colgroup>
+<col style="width: 20%;">
+<col style="width: 20%;">
+<col style="width: 60%;">
+</colgroup>
+<thead>
+<tr>
+<th class="tableblock halign-left valign-top"><strong>Date</strong></th>
+<th class="tableblock halign-left valign-top"><strong>Version</strong></th>
+<th class="tableblock halign-left valign-top"><strong>Description</strong></th>
+</tr>
+</thead>
+<tbody>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock">2021-04-19</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">1.0.0</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Initial version.</p></td>
+</tr>
+</tbody>
+</table>
+</div>
+<div class="sect3">
+<h4 id="_dependencies_2"><a class="anchor" href="#_dependencies_2"></a>39.1.3. Dependencies</h4>
+<div class="paragraph">
+<p>This extension is written against the OpenCL API Specification Version V3.0.6.</p>
+</div>
+<div class="paragraph">
+<p>This extension requires OpenCL 1.0.</p>
+</div>
+</div>
+</div>
+<div class="sect2">
+<h3 id="_new_api_types_3"><a class="anchor" href="#_new_api_types_3"></a>39.2. New API Types</h3>
+<div class="paragraph">
+<p>Structure returned by the device info query for <code>CL_DEVICE_PCI_BUS_INFO_KHR</code>:</p>
+</div>
+<div class="listingblock">
+<div class="content">
+<pre class="CodeRay highlight"><code data-lang="c"><span class="keyword">typedef</span> <span class="keyword">struct</span> _cl_device_pci_bus_info_khr {
+ cl_uint pci_domain;
+ cl_uint pci_bus;
+ cl_uint pci_device;
+ cl_uint pci_function;
+} cl_device_pci_bus_info_khr;</code></pre>
+</div>
+</div>
+</div>
+<div class="sect2">
+<h3 id="_new_api_enums_2"><a class="anchor" href="#_new_api_enums_2"></a>39.3. New API Enums</h3>
+<div class="paragraph">
+<p>Accepted value for the <em>param_name</em> parameter to <strong>clGetDeviceInfo</strong>:</p>
+</div>
+<div class="listingblock">
+<div class="content">
+<pre class="CodeRay highlight"><code data-lang="c"><span class="preprocessor">#define</span> CL_DEVICE_PCI_BUS_INFO_KHR <span class="hex">0x410F</span></code></pre>
+</div>
+</div>
+</div>
+<div class="sect2">
+<h3 id="_modifications_to_the_opencl_api_specification_2"><a class="anchor" href="#_modifications_to_the_opencl_api_specification_2"></a>39.4. Modifications to the OpenCL API Specification</h3>
+<div class="sect3">
+<h4 id="_section_4_2_querying_devices"><a class="anchor" href="#_section_4_2_querying_devices"></a>39.4.1. Section 4.2 - Querying Devices:</h4>
+<div class="paragraph">
+<p>Add to Table 5 - OpenCL Device Queries:</p>
+</div>
+<table class="tableblock frame-all grid-all stretch">
+<caption class="title">Table 5. OpenCL Device Queries</caption>
+<colgroup>
+<col style="width: 30%;">
+<col style="width: 20%;">
+<col style="width: 50%;">
+</colgroup>
+<thead>
+<tr>
+<th class="tableblock halign-left valign-top"><strong>cl_device_info</strong></th>
+<th class="tableblock halign-left valign-top">Return Type</th>
+<th class="tableblock halign-left valign-top">Description</th>
+</tr>
+</thead>
+<tbody>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>CL_DEVICE_PCI_BUS_INFO_KHR</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><code>cl_device_pci_bus_info_khr</code></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Returns PCI bus information for the device.</p>
+<p class="tableblock"> The PCI bus information is returned as a single structure that includes
+ the PCI bus domain, the PCI bus identifier, the PCI device identifier, and
+ the PCI device function identifier.</p></td>
+</tr>
+</tbody>
+</table>
+</div>
+</div>
+</div>
+</div>
+<div class="sect1">
+<h2 id="cl_khr_extended_bit_ops"><a class="anchor" href="#cl_khr_extended_bit_ops"></a>40. Extended Bit Operations</h2>
+<div class="sectionbody">
+<div class="paragraph">
+<p>This extension adds OpenCL C functions for performing extended bit operations.
+Specifically, the following functions are added:</p>
+</div>
+<div class="ulist">
+<ul>
+<li>
+<p>bitfield insert: insert bits from one source operand into another source operand.</p>
+</li>
+<li>
+<p>bitfield extract: extract bits from a source operand, with sign- or zero-extension.</p>
+</li>
+<li>
+<p>bit reverse: reverse the bits of a source operand.</p>
+</li>
+</ul>
+</div>
+<div class="sect2">
+<h3 id="_general_information_39"><a class="anchor" href="#_general_information_39"></a>40.1. General Information</h3>
+<div class="sect3">
+<h4 id="_name_strings_3"><a class="anchor" href="#_name_strings_3"></a>40.1.1. Name Strings</h4>
+<div class="paragraph">
+<p><code>cl_khr_extended_bit_ops</code></p>
+</div>
+</div>
+<div class="sect3">
+<h4 id="_version_history_39"><a class="anchor" href="#_version_history_39"></a>40.1.2. Version History</h4>
+<table class="tableblock frame-all grid-all stretch">
+<colgroup>
+<col style="width: 20%;">
+<col style="width: 20%;">
+<col style="width: 60%;">
+</colgroup>
+<thead>
+<tr>
+<th class="tableblock halign-left valign-top"><strong>Date</strong></th>
+<th class="tableblock halign-left valign-top"><strong>Version</strong></th>
+<th class="tableblock halign-left valign-top"><strong>Description</strong></th>
+</tr>
+</thead>
+<tbody>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock">2021-04-22</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">1.0.0</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Initial version.</p></td>
+</tr>
+</tbody>
+</table>
+</div>
+<div class="sect3">
+<h4 id="_dependencies_3"><a class="anchor" href="#_dependencies_3"></a>40.1.3. Dependencies</h4>
+<div class="paragraph">
+<p>This extension is written against the OpenCL 3.0 C Language Specification and the OpenCL SPIR-V Environment Specification Version V3.0.6.</p>
+</div>
+<div class="paragraph">
+<p>This extension requires OpenCL 1.0.</p>
+</div>
+</div>
+</div>
+<div class="sect2">
+<h3 id="_new_opencl_c_functions"><a class="anchor" href="#_new_opencl_c_functions"></a>40.2. New OpenCL C Functions</h3>
+<div class="listingblock">
+<div class="content">
+<pre class="CodeRay highlight"><code>gentype bitfield_insert( gentype base, gentype insert, uint offset, uint count )
+igentype bitfield_extract_signed( gentype base, uint offset, uint count )
+ugentype bitfield_extract_unsigned( gentype base, uint offset, uint count )
+gentype bit_reverse( gentype base )</code></pre>
+</div>
+</div>
+</div>
+<div class="sect2">
+<h3 id="_modifications_to_the_opencl_c_specification"><a class="anchor" href="#_modifications_to_the_opencl_c_specification"></a>40.3. Modifications to the OpenCL C Specification</h3>
+<div class="sect3">
+<h4 id="_modify_section_6_15_3_integer_functions"><a class="anchor" href="#_modify_section_6_15_3_integer_functions"></a>40.3.1. Modify Section 6.15.3. Integer Functions:</h4>
+<div class="dlist">
+<dl>
+<dt class="hdlist1">Add a new Section 6.15.3.X. Extended Bit Operations: </dt>
+<dd>
+<div class="openblock">
+<div class="content">
+<div class="paragraph">
+<p>The functions described in the following table can be used with built-in scalar or vector integer types to perform extended bit operations.
+The functions that operate on vector types operate component-wise.
+The description is per-component.</p>
+</div>
+<div class="paragraph">
+<p>In the table below, the generic type name <code>gentype</code> refers to the built-in integer types <code>char</code>, <code>char<em>n</em></code>, <code>uchar</code>, <code>uchar<em>n</em></code>, <code>short</code>, <code>short<em>n</em></code>, <code>ushort</code>, <code>ushort<em>n</em></code>, <code>int</code>, <code>int<em>n</em></code>, <code>uint</code>, <code>uint<em>n</em></code>, <code>long</code>, <code>long<em>n</em></code>, <code>ulong</code>, and <code>ulong<em>n</em></code>.
+The generic type name <code>igentype</code> refers to the built-in signed integer types <code>char</code>, <code>char<em>n</em></code>, <code>short</code>, <code>short<em>n</em></code>, <code>int</code>, <code>int<em>n</em></code>, <code>long</code>, and <code>long<em>n</em></code>.
+The generic type name <code>ugentype</code> refers to the built-in unsigned integer types <code>uchar</code>, <code>uchar<em>n</em></code>, <code>ushort</code>, <code>ushort<em>n</em></code>, <code>uint</code>, <code>uint<em>n</em></code>, <code>ulong</code>, and <code>ulong<em>n</em></code>.
+<em>n</em> is 2, 3, 4, 8, or 16.</p>
+</div>
+<table class="tableblock frame-all grid-all stretch">
+<caption class="title">Table 47. Built-in Scalar and Vector Extended Bit Operations</caption>
+<colgroup>
+<col style="width: 50%;">
+<col style="width: 50%;">
+</colgroup>
+<thead>
+<tr>
+<th class="tableblock halign-left valign-top"><strong>Function</strong></th>
+<th class="tableblock halign-left valign-top"><strong>Description</strong></th>
+</tr>
+</thead>
+<tbody>
+<tr>
+<td class="tableblock halign-left valign-top"><div class="content"><div class="listingblock">
+<div class="content">
+<pre class="CodeRay highlight"><code data-lang="c">gentype bitfield_insert(
+ gentype base, gentype insert,
+ uint offset, uint count)</code></pre>
+</div>
+</div></div></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Returns a copy of <em>base</em>, with a modified bitfield that comes from <em>insert</em>.</p>
+<p class="tableblock">Any bits of the result value numbered outside [<em>offset</em>, <em>offset</em> + <em>count</em> - 1] (inclusive) will come from the corresponding bits in <em>base</em>.</p>
+<p class="tableblock">Any bits of the result value numbered inside [<em>offset</em>, <em>offset</em> + <em>count</em> - 1] (inclusive) will come from the bits numbered [0, <em>count</em> - 1] (inclusive) of <em>insert</em>.</p>
+<p class="tableblock"><em>count</em> is the number of bits to be modified.
+If <em>count</em> equals 0, the return value will be equal to <em>base</em>.</p>
+<p class="tableblock">If <em>count</em> or <em>offset</em> or <em>offset</em> + <em>count</em> is greater than number of bits in <code>gentype</code> (for scalar types) or components of <code>gentype</code> (for vector types), the result is undefined.</p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><div class="content"><div class="listingblock">
+<div class="content">
+<pre class="CodeRay highlight"><code data-lang="c">igentype bitfield_extract_signed(
+ gentype base,
+ uint offset, uint count)</code></pre>
+</div>
+</div></div></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Returns an extracted bitfield from <em>base</em> with sign extension.
+The type of the return value is always a signed type.</p>
+<p class="tableblock">The bits of <em>base</em> numbered in [<em>offset</em>, <em>offset</em> + <em>count</em> - 1] (inclusive) are returned as the bits numbered in [0, <em>count</em> - 1] (inclusive) of the result.
+The remaining bits in the result will be sign extended by replicating the bit numbered <em>offset</em> + <em>count</em> - 1 of <em>base</em>.</p>
+<p class="tableblock"><em>count</em> is the number of bits to be extracted.
+If <em>count</em> equals 0, the result is 0.</p>
+<p class="tableblock">If the <em>count</em> or <em>offset</em> or <em>offset</em> + <em>count</em> is greater than number of bits in <code>gentype</code> (for scalar types) or components of <code>gentype</code> (for vector types), the result is undefined.</p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><div class="content"><div class="listingblock">
+<div class="content">
+<pre class="CodeRay highlight"><code data-lang="c">ugentype bitfield_extract_unsigned(
+ gentype base,
+ uint offset, uint count)</code></pre>
+</div>
+</div></div></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Returns an extracted bitfield from <em>base</em> with zero extension.
+The type of the return value is always an unsigned type.</p>
+<p class="tableblock">The bits of <em>base</em> numbered in [<em>offset</em>, <em>offset</em> + <em>count</em> - 1] (inclusive) are returned as the bits numbered in [0, <em>count</em> - 1] (inclusive) of the result.
+The remaining bits in the result will be zero.</p>
+<p class="tableblock"><em>count</em> is the number of bits to be extracted.
+If <em>count</em> equals 0, the result is 0.</p>
+<p class="tableblock">If the <em>count</em> or <em>offset</em> or <em>offset</em> + <em>count</em> is greater than number of bits in <code>gentype</code> (for scalar types) or components of <code>gentype</code> (for vector types), the result is undefined.</p></td>
+</tr>
+<tr>
+<td class="tableblock halign-left valign-top"><div class="content"><div class="listingblock">
+<div class="content">
+<pre class="CodeRay highlight"><code data-lang="c">gentype bit_reverse(
+ gentype base)</code></pre>
+</div>
+</div></div></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Returns the value of <em>base</em> with reversed bits.
+That is, the bit numbered <em>n</em> of the result value will be taken from the bit numbered <em>width</em> - <em>n</em> - 1 of <em>base</em> (for scalar types) or a component of <em>base</em> (for vector types), where <em>width</em> is number of bits of <code>gentype</code> (for scalar types) or components of <code>gentype</code> (for vector types).</p></td>
+</tr>
+</tbody>
+</table>
+</div>
+</div>
+</dd>
+</dl>
+</div>
+</div>
+</div>
+<div class="sect2">
+<h3 id="_modifications_to_the_opencl_spir_v_environment_specification"><a class="anchor" href="#_modifications_to_the_opencl_spir_v_environment_specification"></a>40.4. Modifications to the OpenCL SPIR-V Environment Specification</h3>
+<div class="sect3">
+<h4 id="_add_to_section_5_opencl_extensions"><a class="anchor" href="#_add_to_section_5_opencl_extensions"></a>40.4.1. Add to Section 5 - OpenCL Extensions</h4>
+<div class="dlist">
+<dl>
+<dt class="hdlist1">Add a new Section 5.2.X - <code>cl_khr_extended_bit_ops</code>: </dt>
+<dd>
+<div class="openblock">
+<div class="content">
+<div class="paragraph">
+<p>If the OpenCL environment supports the extension <code>cl_khr_extended_bit_ops</code>, then the environment must accept modules that declare use of the extension <code>SPV_KHR_bit_instructions</code> via <strong>OpExtension</strong>.</p>
+</div>
+<div class="paragraph">
+<p>If the OpenCL environment supports the extension <code>cl_khr_extended_bit_ops</code> and use of the SPIR-V extension <code>SPV_KHR_bit_instructions</code> is declared in the module via <strong>OpExtension</strong>, then the environment must accept modules that declare the <strong>BitInstructions</strong> capability.</p>
+</div>
+</div>
+</div>
+</dd>
+</dl>
+</div>
+</div>
+</div>
+</div>
+</div>
+<div class="sect1">
+<h2 id="cl_khr_suggested_local_work_size"><a class="anchor" href="#cl_khr_suggested_local_work_size"></a>41. Suggested Local Work Size Query</h2>
+<div class="sectionbody">
+<div class="paragraph">
+<p>This extension adds the ability to query a suggested local work group size for a kernel running on a device for a specified global work size and global work offset.
+The suggested local work group size will match the work group size that would be chosen if the kernel were enqueued with the specified global work size and global work offset and a <code>NULL</code> local work size.</p>
+</div>
+<div class="paragraph">
+<p>By using the suggested local work group size query an application has greater insight into the local work group size chosen by the OpenCL implementation, and the OpenCL implementation need not re-compute the local work group size if the same kernel is enqueued multiple times with the same parameters.</p>
+</div>
+<div class="sect2">
+<h3 id="_general_information_40"><a class="anchor" href="#_general_information_40"></a>41.1. General Information</h3>
+<div class="sect3">
+<h4 id="_name_strings_4"><a class="anchor" href="#_name_strings_4"></a>41.1.1. Name Strings</h4>
+<div class="paragraph">
+<p><code>cl_khr_suggested_local_work_size</code></p>
+</div>
+</div>
+<div class="sect3">
+<h4 id="_version_history_40"><a class="anchor" href="#_version_history_40"></a>41.1.2. Version History</h4>
+<table class="tableblock frame-all grid-all stretch">
+<colgroup>
+<col style="width: 20%;">
+<col style="width: 20%;">
+<col style="width: 60%;">
+</colgroup>
+<thead>
+<tr>
+<th class="tableblock halign-left valign-top"><strong>Date</strong></th>
+<th class="tableblock halign-left valign-top"><strong>Version</strong></th>
+<th class="tableblock halign-left valign-top"><strong>Description</strong></th>
+</tr>
+</thead>
+<tbody>
+<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock">2021-04-22</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">1.0.0</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Initial version.</p></td>
+</tr>
+</tbody>
+</table>
+</div>
+<div class="sect3">
+<h4 id="_dependencies_4"><a class="anchor" href="#_dependencies_4"></a>41.1.3. Dependencies</h4>
+<div class="paragraph">
+<p>This extension is written against the OpenCL API Specification Version V3.0.6.</p>
+</div>
+<div class="paragraph">
+<p>This extension requires OpenCL 1.0.</p>
+</div>
+</div>
+</div>
+<div class="sect2">
+<h3 id="_new_api_functions_2"><a class="anchor" href="#_new_api_functions_2"></a>41.2. New API Functions</h3>
+<div class="listingblock">
+<div class="content">
+<pre class="CodeRay highlight"><code>cl_int clGetKernelSuggestedLocalWorkSizeKHR(
+ cl_command_queue command_queue,
+ cl_kernel kernel,
+ cl_uint work_dim,
+ const size_t *global_work_offset,
+ const size_t *global_work_size,
+ size_t *suggested_local_work_size);</code></pre>
+</div>
+</div>
+</div>
+<div class="sect2">
+<h3 id="_modifications_to_the_opencl_api_specification_3"><a class="anchor" href="#_modifications_to_the_opencl_api_specification_3"></a>41.3. Modifications to the OpenCL API Specification</h3>
+<div class="sect3">
+<h4 id="_section_5_9_kernel_objects"><a class="anchor" href="#_section_5_9_kernel_objects"></a>41.3.1. Section 5.9 - Kernel Objects:</h4>
+<div class="sect4">
+<h5 id="_new_section_5_9_4_x_suggested_local_work_size_query"><a class="anchor" href="#_new_section_5_9_4_x_suggested_local_work_size_query"></a>41.3.1.1. New Section 5.9.4.X - Suggested Local Work Size Query</h5>
+<div class="paragraph">
+<p>To query a suggested local work size for a kernel object, call the function</p>
+</div>
+<div class="listingblock">
+<div class="content">
+<pre class="CodeRay highlight"><code>cl_int clGetKernelSuggestedLocalWorkSizeKHR(
+ cl_command_queue command_queue,
+ cl_kernel kernel,
+ cl_uint work_dim,
+ const size_t *global_work_offset,
+ const size_t *global_work_size,
+ size_t *suggested_local_work_size);</code></pre>
+</div>
+</div>
+<div class="paragraph">
+<p>The returned suggested local work size is expected to match the local work size that would be chosen if the specified kernel object, with the same kernel arguments, were enqueued into the specified command queue with the specified global work size, specified global work offset, and with a <code>NULL</code> local work size.</p>
+</div>
+<div class="ulist">
+<ul>
+<li>
+<p><em>command_queue</em> specifies the command queue and device for the query.</p>
+</li>
+<li>
+<p><em>kernel</em> specifies the kernel object and kernel arguments for the query.
+The OpenCL context associated with <em>kernel</em> and <em>command_queue</em> must the same.</p>
+</li>
+<li>
+<p><em>work_dim</em> specifies the number of work dimensions in the input global work offset and global work size, and the output suggested local work size.</p>
+</li>
+<li>
+<p><em>global_work_offset</em> can be used to specify an array of at least <em>work_dim</em> global ID offset values for the query.
+This is optional and may be <code>NULL</code> to indicate there is no global ID offset.</p>
+</li>
+<li>
+<p><em>global_work_size</em> is an array of at least <em>work_dim</em> values describing the global work size for the query.</p>
+</li>
+<li>
+<p><em>suggested_local_work_size</em> is an output array of at least <em>work_dim</em> values that will contain the result of the query.</p>
+</li>
+</ul>
+</div>
+<div class="paragraph">
+<p><strong>clGetKernelSuggestedLocalWorkSizeKHR</strong> returns <code>CL_SUCCESS</code> if the query executed successfully.
+Otherwise, it returns one of the following errors:</p>
+</div>
+<div class="ulist">
+<ul>
+<li>
+<p><code>CL_INVALID_COMMAND_QUEUE</code> if <em>command_queue</em> is not a valid host command queue.</p>
+</li>
+<li>
+<p><code>CL_INVALID_KERNEL</code> if <em>kernel</em> is not a valid kernel object.</p>
+</li>
+<li>
+<p><code>CL_INVALID_CONTEXT</code> if the context associated with <em>kernel</em> is not the same as the context associated with <em>command_queue</em>.</p>
+</li>
+<li>
+<p><code>CL_INVALID_PROGRAM_EXECUTABLE</code> if there is no successfully built program executable available for <em>kernel</em> for the device associated with <em>command_queue</em>.</p>
+</li>
+<li>
+<p><code>CL_INVALID_KERNEL_ARGS</code> if all argument values for <em>kernel</em> have not been set.</p>
+</li>
+<li>
+<p><code>CL_MISALIGNED_SUB_BUFFER_OFFSET</code> if a sub-buffer object is set as an argument to <em>kernel</em> and the offset specified when the sub-buffer object was created is not aligned to <code>CL_DEVICE_MEM_BASE_ADDR_ALIGN</code> for the device associated with <em>command_queue</em>.</p>
+</li>
+<li>
+<p><code>CL_INVALID_IMAGE_SIZE</code> if an image object is set as an argument to <em>kernel</em> and the image dimensions are not supported by device associated with <em>command_queue</em>.</p>
+</li>
+<li>
+<p><code>CL_IMAGE_FORMAT_NOT_SUPPORTED</code> if an image object is set as an argument to <em>kernel</em> and the image format is not supported by the device associated with <em>command_queue</em>.</p>
+</li>
+<li>
+<p><code>CL_INVALID_OPERATION</code> if an SVM pointer is set as an argument to <em>kernel</em> and the device associated with <em>command_queue</em> does not support SVM or the required SVM capabilities for the SVM pointer.</p>
+</li>
+<li>
+<p><code>CL_INVALID_WORK_DIMENSION</code> if <em>work_dim</em> is not a valid value (i.e. a value between 1 and <code>CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS</code>).</p>
+</li>
+<li>
+<p><code>CL_INVALID_GLOBAL_WORK_SIZE</code> if <em>global_work_size</em> is NULL or if any of the values specified in <em>global_work_size</em> are 0.</p>
+</li>
+<li>
+<p><code>CL_INVALID_GLOBAL_WORK_SIZE</code> if any of the values specified in <em>global_work_size</em> exceed the maximum value representable by <code>size_t</code> on the device associated with <em>command_queue</em>.</p>
+</li>
+<li>
+<p><code>CL_INVALID_GLOBAL_OFFSET</code> if the value specified in <em>global_work_size</em> plus the corresponding value in <em>global_work_offset</em> for dimension exceeds the maximum value representable by <code>size_t</code> on the device associated with <em>command_queue</em>.</p>
+</li>
+<li>
+<p><code>CL_OUT_OF_RESOURCES</code> if there is a failure to allocate resources required by the OpenCL implementation on the device.</p>
+</li>
+<li>
+<p><code>CL_OUT_OF_HOST_RESOURCES</code> if there is a failure to allocate resources required by the OpenCL implementation on the host.</p>
+</li>
+</ul>
+</div>
+<div class="admonitionblock note">
+<table>
+<tr>
+<td class="icon">
+<i class="fa icon-note" title="Note"></i>
+</td>
+<td class="content">
+These error conditions are consistent with error conditions for <strong>clEnqueueNDRangeKernel</strong>.
+</td>
+</tr>
+</table>
+</div>
+</div>
+</div>
+</div>
+</div>
+</div>
+<div class="sect1">
+<h2 id="spirv_extensions"><a class="anchor" href="#spirv_extensions"></a>42. Extensions to the OpenCL SPIR-V Environment</h2>
<div class="sectionbody">
<div class="paragraph">
<p>An OpenCL SPIR-V environment may be modified by OpenCL extensions.
@@ -19410,6 +19918,12 @@
<li>
<p><code>cl_khr_spirv_no_integer_wrap_decoration</code></p>
</li>
+<li>
+<p><code>cl_khr_spirv_extended_debug_info</code></p>
+</li>
+<li>
+<p><code>cl_khr_spirv_linkonce_odr</code></p>
+</li>
</ul>
</div>
</div>
@@ -19583,6 +20097,11 @@
<td class="tableblock halign-left valign-top"><p class="tableblock">Provisional Extension</p></td>
</tr>
<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><a href="#cl_khr_extended_bit_ops">cl_khr_extended_bit_ops</a></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Bit Insert, Extract, and Reverse Operations</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Extension</p></td>
+</tr>
+<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><a href="#cl_khr_extended_versioning">cl_khr_extended_versioning</a></p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">Extend versioning of platform, devices, extensions, etc.</p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">Extension</p></td>
@@ -19678,6 +20197,11 @@
<td class="tableblock halign-left valign-top"><p class="tableblock">Extension</p></td>
</tr>
<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><a href="#cl_khr_pci_bus_info">cl_khr_pci_bus_info</a></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Query PCI Bus Information for an OpenCL Device</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Extension</p></td>
+</tr>
+<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><a href="#cl_khr_priority_hints">cl_khr_priority_hints</a></p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">Create Command Queues with Different Priorities</p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">Extension</p></td>
@@ -19743,6 +20267,11 @@
<td class="tableblock halign-left valign-top"><p class="tableblock">Extension</p></td>
</tr>
<tr>
+<td class="tableblock halign-left valign-top"><p class="tableblock"><a href="#cl_khr_suggested_local_work_size">cl_khr_suggested_local_work_size</a></p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Query a Suggested Local Work Size</p></td>
+<td class="tableblock halign-left valign-top"><p class="tableblock">Extension</p></td>
+</tr>
+<tr>
<td class="tableblock halign-left valign-top"><p class="tableblock"><a href="#cl_khr_terminate_context">cl_khr_terminate_context</a></p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">Terminate an OpenCL Context</p></td>
<td class="tableblock halign-left valign-top"><p class="tableblock">Extension</p></td>
@@ -19759,8 +20288,8 @@
</div>
<div id="footer">
<div id="footer-text">
-Version v3.0.6<br>
-Last updated 2020-12-18 09:20:25 -0800
+Version v3.0.7<br>
+Last updated 2021-04-23 12:57:31 -0700
</div>
</div>
diff --git a/specs/3.0-unified/pdf/OpenCL_API.pdf b/specs/3.0-unified/pdf/OpenCL_API.pdf
index a3cbfb5..dcdf69e 100644
--- a/specs/3.0-unified/pdf/OpenCL_API.pdf
+++ b/specs/3.0-unified/pdf/OpenCL_API.pdf
Binary files differ
diff --git a/specs/3.0-unified/pdf/OpenCL_C.pdf b/specs/3.0-unified/pdf/OpenCL_C.pdf
index 70b71fd..0009161 100644
--- a/specs/3.0-unified/pdf/OpenCL_C.pdf
+++ b/specs/3.0-unified/pdf/OpenCL_C.pdf
Binary files differ
diff --git a/specs/3.0-unified/pdf/OpenCL_Env.pdf b/specs/3.0-unified/pdf/OpenCL_Env.pdf
index a07337e..bfef7ec 100644
--- a/specs/3.0-unified/pdf/OpenCL_Env.pdf
+++ b/specs/3.0-unified/pdf/OpenCL_Env.pdf
Binary files differ
diff --git a/specs/3.0-unified/pdf/OpenCL_Ext.pdf b/specs/3.0-unified/pdf/OpenCL_Ext.pdf
index ab5b09e..7b4576f 100644
--- a/specs/3.0-unified/pdf/OpenCL_Ext.pdf
+++ b/specs/3.0-unified/pdf/OpenCL_Ext.pdf
Binary files differ
