[SYCL][Doc] Add new free function queries proposal (#5106)

The new proposal 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 <[email protected]>

Author: Pennycook

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 <int Dimensions>
+id<Dimensions> get_id();
+
+template <int Dimensions>
+item<Dimensions, false> get_item();
+
+template <int Dimensions>
+nd_item<Dimensions> get_nd_item();
+
+template <int Dimensions>
+group<Dimensions> 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 <int Dimensions>
+id<Dimensions> 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 <int Dimensions>
+item<Dimensions, false> 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 <int Dimensions>
+nd_item<Dimensions> 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 <int Dimensions>
+group<Dimensions> 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.
+//--