From 0ef41cd8afd219d53585e23b62c03fe304fdbe13 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Wed, 8 Dec 2021 07:32:05 -0800 Subject: [PATCH 1/8] [SYCL][Doc] Update FreeFunctionQueries In preparation for moving the FreeFunctionQueries extension out of the experimental namespace, this commit makes the following changes: - Expands upon the motivation for the extension, and provides some non-normative recommendations for when it should be used. - Reformats the description of the functions to match the style used by ISO C++ and the newest sections of SYCL 2020. - Moves all functionality into the this_kernel namespace, to more closely align with the this_thread namespace in ISO C++. - Documents outstanding issues regarding undefined behavior. Signed-off-by: John Pennycook --- .../FreeFunctionQueries.asciidoc | 188 +++++++++++++++ .../SYCL_INTEL_free_function_queries.asciidoc | 226 ------------------ sycl/doc/extensions/README.md | 2 +- 3 files changed, 189 insertions(+), 227 deletions(-) create mode 100644 sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc delete mode 100644 sycl/doc/extensions/FreeFunctionQueries/SYCL_INTEL_free_function_queries.asciidoc diff --git a/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc b/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc new file mode 100644 index 0000000000000..2fd23851caec8 --- /dev/null +++ b/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc @@ -0,0 +1,188 @@ += SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES +:source-highlighter: coderay +:coderay-linenums-mode: table + +// This section needs to be after the document title. +:doctype: book +:toc2: +:toc: left +:encoding: utf-8 +:lang: en + +:blank: pass:[ +] + +// Set the default source code type in this document to C++, +// for syntax highlighting purposes. This is needed because +// docbook uses c++ and html5 uses cpp. +:language: {basebackend@docbook:c++:cpp} + +== Introduction +IMPORTANT: This specification is a draft. + +NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. used by permission by Khronos. + +This document describes an extension that adds features for SYCL work items and groups to be available globally. + +== Notice + +Copyright (c) 2020-2021 Intel Corporation. All rights reserved. + +== Status + +Working Draft + +This is a preview extension specification, intended to provide early access to a feature for review and community feedback. When the feature matures, this specification may be released as a formal extension. + +Because the interfaces defined by this specification are not final and are subject to change they are not intended to be used by shipping software products. + +== Version + +Revision: 1 + +== Contact +Ruslan Arutyunyan, Intel (ruslan 'dot' arutyunyan 'at' intel 'dot' com) + +John Pennycook, Intel (john 'dot' pennycook 'at' intel 'dot' com) + +== Dependencies + +This extension is written against the SYCL 2020 specification, Revision 4. + +== Feature Test Macro + +This extension provides a feature-test macro as described in the core SYCL +specification section 6.3.3 "Feature test macros". Therefore, an implementation +supporting this extension must predefine the macro `SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES` +to one of the values defined in the table below. Applications can test for the +existence of this macro to determine if the implementation supports this +feature, or applications can test the macro's value to determine which of the +extension's APIs the implementation supports. + +[%header,cols="1,5"] +|=== +|Value |Description +|1 |Initial extension version. Base features are supported. +|=== + +== Overview + +The extension allows to get `sycl::id`, `sycl::item`, `sycl::nd_item`, +`sycl::group` and `sycl::sub_group` instances globally, without having to +explicitly pass them as arguments to each function used on the device. + +NOTE: Passing such instances as arguments can result in a clearer interface +that is less error-prone to use. For example, if a function accepts a +`sycl::group`, the caller must assume that function may call a +`sycl::group_barrier` and ensure that associated control flow requirements are +satisfied. It is recommended that this extension is used only when modifying +existing interfaces is not feasible. + +== Accessing Instances of Special SYCL Classes + +The `sycl::ext::oneapi::this_kernel` namespace contains functionality related +to the currently executing kernel. + +It is the user's responsibility to ensure that that these functions are called +in a manner that is compatible with the kernel's launch parameters, as detailed +in the definition of each function. Calling these functions from an incompatible +kernel results in undefined behavior. + +[source,c++] +---- +namespace sycl { +namespace ext { +namespace oneapi { +namespace this_kernel { + +template +id get_id(); + +template +item get_item(); + +template +nd_item get_nd_item(); + +template +group get_group(); + +sub_group get_sub_group(); + +} +} +} +} +---- + +[source,c++] +---- +template +id get_id(); +---- +_Preconditions_: `Dimensions` must match the dimensionality of the currently +executing kernel. The currently executing kernel must have been launched with a +`sycl::range` or `sycl::nd_range` argument. + +_Returns_: A `sycl::id` instance representing the current work-item in a +`sycl::range`. + +[source,c++] +---- +template +item get_item(); +---- +_Preconditions_: `Dimensions` must match the dimensionality of the currently +executing kernel. The currently executing kernel must have been launched with a +`sycl::range` or `sycl::nd_range` argument. + +_Returns_: A `sycl::item` instance representing the current work-item in a +`sycl::range`. + +[source,c++] +---- +template +nd_item get_nd_item(); +---- +_Preconditions_: `Dimensions` must match the dimensionality of the currently +executing kernel. The currently executing kernel must have been launched with a +`sycl::nd_range` argument. + +_Returns_: A `sycl::nd_item` instance representing the current work-item in a +`sycl::nd_range`. + +[source,c++] +---- +template +group get_group(); +---- +_Preconditions_: `Dimensions` must match the dimensionality of the currently +executing kernel. The currently executing kernel must have been launched with a +`sycl::nd_range` argument. + +_Returns_: A `sycl::group` instance representing the work-group to which the +current work-item belongs. + +[source,c++] +---- +sub_group get_sub_group(); +---- +_Preconditions_: The currently executing kernel must have been launched with a +`sycl::nd_range` argument. + +_Returns_: A `sycl::sub_group` instance representing the sub-group to which the +current work-item belongs. + +== Issues + +. Can undefined behavior be avoided or detected? +-- +*UNRESOLVED*: Good run-time errors would likely require support for device-side +assertions or exceptions, while good compile-time errors would likely require +some additional compiler modifications and/or kernel properties. +-- + +//. asd +//+ +//-- +//*RESOLUTION*: Not resolved. +//-- diff --git a/sycl/doc/extensions/FreeFunctionQueries/SYCL_INTEL_free_function_queries.asciidoc b/sycl/doc/extensions/FreeFunctionQueries/SYCL_INTEL_free_function_queries.asciidoc deleted file mode 100644 index ff5da3a417d91..0000000000000 --- a/sycl/doc/extensions/FreeFunctionQueries/SYCL_INTEL_free_function_queries.asciidoc +++ /dev/null @@ -1,226 +0,0 @@ -= SYCL_INTEL_free_function_queries -:source-highlighter: coderay -:coderay-linenums-mode: table - -// This section needs to be after the document title. -:doctype: book -:toc2: -:toc: left -:encoding: utf-8 -:lang: en - - -:blank: pass:[ +] - -// Set the default source code type in this document to C++, -// for syntax highlighting purposes. This is needed because -// docbook uses c++ and html5 uses cpp. -:language: {basebackend@docbook:c++:cpp} - -== Introduction -IMPORTANT: This specification is a draft. - -NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. used by permission by Khronos. - -NOTE: This document is better viewed when rendered as html with asciidoctor. GitHub does not render image icons. - -This document describes an extension that adds features for SYCL work items and groups to be available globally. - - -== Name Strings - -+SYCL_INTEL_free_function_queries+ - -== Notice - -Copyright (c) 2020 Intel Corporation. All rights reserved. - -== Status - -Working Draft - -This is a preview extension specification, intended to provide early access to a feature for review and community feedback. When the feature matures, this specification may be released as a formal extension. - -Because the interfaces defined by this specification are not final and are subject to change they are not intended to be used by shipping software products. - -== Version - -Built On: {docdate} + -Revision: 1 - -== Contact -Ruslan Arutyunyan, Intel (ruslan 'dot' arutyunyan 'at' intel 'dot' com) - -John Pennycook, Intel (john 'dot' pennycook 'at' intel 'dot' com) - -== Dependencies - -This extension is written against the SYCL 2020 Provisional specification, Revision 1. - -== Overview - -The extension allows to get `id`, `item`, `nd_item`, `group`, `sub_group` instances globally - -== Modifications of SYCL 2020 provisional specification - -=== Changes in 4.10.1.3 (id class) - -==== Changes in synopsis - -===== Add the following free function - -[source,c++,multiptr,linenums] ----- -template -id this_id(); ----- - -===== Add the table with free functions description - -|=== -a| -[source,c++,multiptr,linenums] ----- -template -id this_id(); ----- a| -Returns the `id` instance representing the current work item in a `range`. The prerequisites: - -* `dimensions` value passed to `this_id` function shall match `dimensions` value passed to the `range` when the `parallel_for` is called - -* shall only be invoked within the `parallel_for` with `range` or `nd_range` as the argument. - -If prerequisites are not fulfilled the behavior is undefined -|=== - -=== Changes in 4.10.1.4 (item class) - -==== Changes in synopsis - -===== Add the following free function - -[source,c++,multiptr,linenums] ----- -template -item this_item(); ----- - -===== Add the table with free functions description - -|=== -a| -[source,c++,multiptr,linenums] ----- -template -item this_item(); ----- a| -Returns the `item` instance representing the current work item in a `range`. The prerequisites: - -* `dimensions` value passed to `this_item` function shall match `dimensions` value passed to the `range` when the `parallel_for` is called - -* shall only be invoked within the `parallel_for` with `range` or `nd_range` as the argument - -If prerequisites are not fulfilled the behavior is undefined -|=== - -=== Changes in 4.10.1.5 (nd_item class) - -==== Changes in synopsis - -===== Add the following free function - -[source,c++,multiptr,linenums] ----- -template -nd_item this_nd_item(); ----- - -===== Add the table with free functions description - -|=== -a| -[source,c++,multiptr,linenums] ----- -template -nd_item this_nd_item(); ----- a| -Returns the `nd_item` instance representing the current work item in a `nd_range`. The prerequisites: - -* `dimensions` value passed to `this_nd_item` function shall match `dimensions` value passed to the `nd_range` when the `parallel_for` is called - -* shall only be invoked within the `parallel_for` with `nd_range` as the argument - -If prerequisites are not fulfilled the behavior is undefined -|=== - -=== Changes in 4.10.1.7 (group class) - -==== Changes in synopsis - -===== Add the following free function - -[source,c++,multiptr,linenums] ----- -template -group this_group(); ----- - -===== Add the table with free functions description - -|=== -a| -[source,c++,multiptr,linenums] ----- -template -group this_group(); ----- a| -Returns the `group` instance representing the particular work-group within the parallel execution. The prerequisites: - -* `dimensions` value passed to `this_group` function shall match `dimensions` value passed to the `range` when the `parallel_for` is called - -* shall only be called from within the `parallel_for` with `nd_range` as the argument - -If prerequisites are not fulfilled the behavior is undefined -|=== - -=== Changes in 4.10.1.8 (sub_group class) - -==== Changes in synopsis - -===== Add the following free function - -[source,c++,multiptr,linenums] ----- -sub_group this_sub_group(); ----- - -===== Add the table with free functions description - -|=== -a| -[source,c++,multiptr,linenums] ----- -sub_group this_sub_group(); ----- a| -Returns the `sub_group` instance representing the particular `sub-group` within the parallel execution. The prerequisites: - -* shall only be called from within the `parallel_for` with `nd_range` as the argument - -If prerequisites are not fulfilled the behavior is undefined -|=== - -== Feature test macro - -This extension provides a feature-test macro as described in the core SYCL -specification section 6.3.3 "Feature test macros". Therefore, an implementation -supporting this extension must predefine the macro `SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES` -to one of the values defined in the table below. Applications can test for the -existence of this macro to determine if the implementation supports this -feature, or applications can test the macro's value to determine which of the -extension's APIs the implementation supports. - -[%header,cols="1,5"] -|=== -|Value |Description -|1 |Initial extension version. Base features are supported. -|=== diff --git a/sycl/doc/extensions/README.md b/sycl/doc/extensions/README.md index 9f6eb9c864ed2..cff6bddd5e795 100755 --- a/sycl/doc/extensions/README.md +++ b/sycl/doc/extensions/README.md @@ -42,7 +42,7 @@ DPC++ extensions status: | [Uniform](Uniform/Uniform.asciidoc) | Proposal | | | [Assert](Assert/SYCL_ONEAPI_ASSERT.asciidoc) | Proposal | | | [Matrix](Matrix/dpcpp-joint-matrix.asciidoc) | Partially supported(AMX AOT) | Not supported: dynamic-extent, wg and wi scopes, layouts other than packed| -| [SYCL_INTEL_free_function_queries](FreeFunctionQueries/SYCL_INTEL_free_function_queries.asciidoc) | Supported (experimental) | | +| [Free Function Queries](FreeFunctionQueries/FreeFunctionQueries.asciidoc) | Supported (experimental) | | | [EXT_ONEAPI_max_work_groups](MaxWorkGroupQueries/max_work_group_query.md) | Supported | | | [SYCL_EXT_ONEAPI_DEVICE_GLOBAL](DeviceGlobal/SYCL_INTEL_device_global.asciidoc) | Proposal | | | [SYCL_INTEL_bf16_conversion](Bf16Conversion/SYCL_INTEL_bf16_conversion.asciidoc) | Partially supported (Level Zero: GPU) | Currently available only on Xe HP GPU. ext_intel_bf16_conversion aspect is not supported. | From b657e080465b177323fa53a6e72b05883169eecf Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Wed, 15 Dec 2021 14:42:01 -0800 Subject: [PATCH 2/8] Add get_linear_id() Signed-off-by: John Pennycook --- .../FreeFunctionQueries/FreeFunctionQueries.asciidoc | 11 +++++++++++ 1 file changed, 11 insertions(+) diff --git a/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc b/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc index 2fd23851caec8..df3823d61436f 100644 --- a/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc +++ b/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc @@ -94,6 +94,8 @@ namespace ext { namespace oneapi { namespace this_kernel { +size_t get_linear_id(); + template id get_id(); @@ -114,6 +116,15 @@ sub_group get_sub_group(); } ---- +[source,c++] +---- +size_t get_linear_id(); +---- +_Returns_: A `size_t` representing the current work-item's linear position in +the global iteration space. Multi-dimensional indices are converted into a +linear index following the linearization equations defined in Section 3.11.1 of +the SYCL 2020 specification. + [source,c++] ---- template From b21d20eca1ce7b7c9d6e898ccc8a96cc6328fbfa Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Wed, 26 Jan 2022 14:12:26 -0800 Subject: [PATCH 3/8] Generalize description of global iteration space Co-authored-by: Greg Lueck --- .../FreeFunctionQueries/FreeFunctionQueries.asciidoc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc b/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc index df3823d61436f..c58a9ac418d8c 100644 --- a/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc +++ b/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc @@ -134,8 +134,8 @@ _Preconditions_: `Dimensions` must match the dimensionality of the currently executing kernel. The currently executing kernel must have been launched with a `sycl::range` or `sycl::nd_range` argument. -_Returns_: A `sycl::id` instance representing the current work-item in a -`sycl::range`. +_Returns_: A `sycl::id` instance representing the current work-item in the +global iteration space. [source,c++] ---- From ff5f5136f2f426ca12ebfd28a28f1eb841b46845 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Wed, 26 Jan 2022 14:13:59 -0800 Subject: [PATCH 4/8] Generalize description of global iteration space for item Co-authored-by: Greg Lueck --- .../FreeFunctionQueries/FreeFunctionQueries.asciidoc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc b/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc index c58a9ac418d8c..770d20e610e79 100644 --- a/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc +++ b/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc @@ -146,8 +146,8 @@ _Preconditions_: `Dimensions` must match the dimensionality of the currently executing kernel. The currently executing kernel must have been launched with a `sycl::range` or `sycl::nd_range` argument. -_Returns_: A `sycl::item` instance representing the current work-item in a -`sycl::range`. +_Returns_: A `sycl::item` instance representing the current work-item in the +global iteration space. [source,c++] ---- From d58457f8260b7b82649c65e64703919cb5bfe65f Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Wed, 26 Jan 2022 14:21:15 -0800 Subject: [PATCH 5/8] Restrict get_linear_id to parallel_for Signed-off-by: John Pennycook --- .../FreeFunctionQueries/FreeFunctionQueries.asciidoc | 3 +++ 1 file changed, 3 insertions(+) diff --git a/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc b/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc index 770d20e610e79..a0de7c8a27d52 100644 --- a/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc +++ b/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc @@ -120,6 +120,9 @@ sub_group get_sub_group(); ---- size_t get_linear_id(); ---- +_Preconditions_: The currently executing kernel must have been launched with a +`sycl::range` or `sycl::nd_range` argument. + _Returns_: A `size_t` representing the current work-item's linear position in the global iteration space. Multi-dimensional indices are converted into a linear index following the linearization equations defined in Section 3.11.1 of From 055e1578a2845bf8092071f172c320aae835bf69 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Fri, 28 Jan 2022 09:19:08 -0800 Subject: [PATCH 6/8] Make get_item() return item without an offset --- .../FreeFunctionQueries/FreeFunctionQueries.asciidoc | 11 +++++++++-- 1 file changed, 9 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc b/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc index a0de7c8a27d52..f095119b25c1c 100644 --- a/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc +++ b/sycl/doc/extensions/FreeFunctionQueries/FreeFunctionQueries.asciidoc @@ -100,7 +100,7 @@ template id get_id(); template -item get_item(); +item get_item(); template nd_item get_nd_item(); @@ -143,7 +143,7 @@ global iteration space. [source,c++] ---- template -item get_item(); +item get_item(); ---- _Preconditions_: `Dimensions` must match the dimensionality of the currently executing kernel. The currently executing kernel must have been launched with a @@ -152,6 +152,13 @@ executing kernel. The currently executing kernel must have been launched with a _Returns_: A `sycl::item` instance representing the current work-item in the global iteration space. +NOTE: The `offset` parameter to `parallel_for` is deprecated in SYCL 2020, as +is the ability of an `item` to carry an offset. This extension returns an +`item` where the `WithOffset` template parameter is set to `false` to prevent +usage of the new queries in conjunction with deprecated functionality. The +return type of `get_item()` is expected to change when the `offset` parameter +is removed in a future version of the SYCL specification. + [source,c++] ---- template From 05f1f7773e452d50cc6f37edd72d3d7cb26c9229 Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Wed, 2 Feb 2022 07:44:29 -0800 Subject: [PATCH 7/8] Move documentation from experimental/ to proposed/ --- .../SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES.asciidoc | 0 1 file changed, 0 insertions(+), 0 deletions(-) rename sycl/doc/extensions/{experimental => proposed}/SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES.asciidoc (100%) diff --git a/sycl/doc/extensions/experimental/SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES.asciidoc b/sycl/doc/extensions/proposed/SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES.asciidoc similarity index 100% rename from sycl/doc/extensions/experimental/SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES.asciidoc rename to sycl/doc/extensions/proposed/SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES.asciidoc From 07f0cbaf9a7e59d9be32f189a320efd60733b5cb Mon Sep 17 00:00:00 2001 From: John Pennycook Date: Wed, 2 Feb 2022 12:04:56 -0800 Subject: [PATCH 8/8] Restore experimental version of the extension --- ..._EXT_ONEAPI_FREE_FUNCTION_QUERIES.asciidoc | 226 ++++++++++++++++++ 1 file changed, 226 insertions(+) create mode 100644 sycl/doc/extensions/experimental/SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES.asciidoc diff --git a/sycl/doc/extensions/experimental/SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES.asciidoc b/sycl/doc/extensions/experimental/SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES.asciidoc new file mode 100644 index 0000000000000..ff5da3a417d91 --- /dev/null +++ b/sycl/doc/extensions/experimental/SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES.asciidoc @@ -0,0 +1,226 @@ += SYCL_INTEL_free_function_queries +:source-highlighter: coderay +:coderay-linenums-mode: table + +// This section needs to be after the document title. +:doctype: book +:toc2: +:toc: left +:encoding: utf-8 +:lang: en + + +:blank: pass:[ +] + +// Set the default source code type in this document to C++, +// for syntax highlighting purposes. This is needed because +// docbook uses c++ and html5 uses cpp. +:language: {basebackend@docbook:c++:cpp} + +== Introduction +IMPORTANT: This specification is a draft. + +NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. used by permission by Khronos. + +NOTE: This document is better viewed when rendered as html with asciidoctor. GitHub does not render image icons. + +This document describes an extension that adds features for SYCL work items and groups to be available globally. + + +== Name Strings + ++SYCL_INTEL_free_function_queries+ + +== Notice + +Copyright (c) 2020 Intel Corporation. All rights reserved. + +== Status + +Working Draft + +This is a preview extension specification, intended to provide early access to a feature for review and community feedback. When the feature matures, this specification may be released as a formal extension. + +Because the interfaces defined by this specification are not final and are subject to change they are not intended to be used by shipping software products. + +== Version + +Built On: {docdate} + +Revision: 1 + +== Contact +Ruslan Arutyunyan, Intel (ruslan 'dot' arutyunyan 'at' intel 'dot' com) + +John Pennycook, Intel (john 'dot' pennycook 'at' intel 'dot' com) + +== Dependencies + +This extension is written against the SYCL 2020 Provisional specification, Revision 1. + +== Overview + +The extension allows to get `id`, `item`, `nd_item`, `group`, `sub_group` instances globally + +== Modifications of SYCL 2020 provisional specification + +=== Changes in 4.10.1.3 (id class) + +==== Changes in synopsis + +===== Add the following free function + +[source,c++,multiptr,linenums] +---- +template +id this_id(); +---- + +===== Add the table with free functions description + +|=== +a| +[source,c++,multiptr,linenums] +---- +template +id this_id(); +---- a| +Returns the `id` instance representing the current work item in a `range`. The prerequisites: + +* `dimensions` value passed to `this_id` function shall match `dimensions` value passed to the `range` when the `parallel_for` is called + +* shall only be invoked within the `parallel_for` with `range` or `nd_range` as the argument. + +If prerequisites are not fulfilled the behavior is undefined +|=== + +=== Changes in 4.10.1.4 (item class) + +==== Changes in synopsis + +===== Add the following free function + +[source,c++,multiptr,linenums] +---- +template +item this_item(); +---- + +===== Add the table with free functions description + +|=== +a| +[source,c++,multiptr,linenums] +---- +template +item this_item(); +---- a| +Returns the `item` instance representing the current work item in a `range`. The prerequisites: + +* `dimensions` value passed to `this_item` function shall match `dimensions` value passed to the `range` when the `parallel_for` is called + +* shall only be invoked within the `parallel_for` with `range` or `nd_range` as the argument + +If prerequisites are not fulfilled the behavior is undefined +|=== + +=== Changes in 4.10.1.5 (nd_item class) + +==== Changes in synopsis + +===== Add the following free function + +[source,c++,multiptr,linenums] +---- +template +nd_item this_nd_item(); +---- + +===== Add the table with free functions description + +|=== +a| +[source,c++,multiptr,linenums] +---- +template +nd_item this_nd_item(); +---- a| +Returns the `nd_item` instance representing the current work item in a `nd_range`. The prerequisites: + +* `dimensions` value passed to `this_nd_item` function shall match `dimensions` value passed to the `nd_range` when the `parallel_for` is called + +* shall only be invoked within the `parallel_for` with `nd_range` as the argument + +If prerequisites are not fulfilled the behavior is undefined +|=== + +=== Changes in 4.10.1.7 (group class) + +==== Changes in synopsis + +===== Add the following free function + +[source,c++,multiptr,linenums] +---- +template +group this_group(); +---- + +===== Add the table with free functions description + +|=== +a| +[source,c++,multiptr,linenums] +---- +template +group this_group(); +---- a| +Returns the `group` instance representing the particular work-group within the parallel execution. The prerequisites: + +* `dimensions` value passed to `this_group` function shall match `dimensions` value passed to the `range` when the `parallel_for` is called + +* shall only be called from within the `parallel_for` with `nd_range` as the argument + +If prerequisites are not fulfilled the behavior is undefined +|=== + +=== Changes in 4.10.1.8 (sub_group class) + +==== Changes in synopsis + +===== Add the following free function + +[source,c++,multiptr,linenums] +---- +sub_group this_sub_group(); +---- + +===== Add the table with free functions description + +|=== +a| +[source,c++,multiptr,linenums] +---- +sub_group this_sub_group(); +---- a| +Returns the `sub_group` instance representing the particular `sub-group` within the parallel execution. The prerequisites: + +* shall only be called from within the `parallel_for` with `nd_range` as the argument + +If prerequisites are not fulfilled the behavior is undefined +|=== + +== Feature test macro + +This extension provides a feature-test macro as described in the core SYCL +specification section 6.3.3 "Feature test macros". Therefore, an implementation +supporting this extension must predefine the macro `SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES` +to one of the values defined in the table below. Applications can test for the +existence of this macro to determine if the implementation supports this +feature, or applications can test the macro's value to determine which of the +extension's APIs the implementation supports. + +[%header,cols="1,5"] +|=== +|Value |Description +|1 |Initial extension version. Base features are supported. +|===