diff --git a/sycl/doc/extensions/proposed/SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES.asciidoc b/sycl/doc/extensions/proposed/SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES.asciidoc new file mode 100644 index 0000000000000..f095119b25c1c --- /dev/null +++ b/sycl/doc/extensions/proposed/SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES.asciidoc @@ -0,0 +1,209 @@ += 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 { + +size_t get_linear_id(); + +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++] +---- +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 +the SYCL 2020 specification. + +[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 the +global iteration space. + +[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 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 +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. +//--