[SYCL] Implement latest version of sycl_ext_oneapi_free_function_queries (#13257)

- Implement latest revision of the proposed extension
- Deprecate `sycl::ext::oneapi::experimental::` interfaces
- Remove old deprecated interfaces from `sycl::` namespace

Author: aelovikov-intel

llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp

@@ -153,7 +153,9 @@ class ESIMDVerifierImpl {
           if (!Name.starts_with("sycl::_V1::") ||
               Name.starts_with("sycl::_V1::detail::") ||
               Name.starts_with("sycl::_V1::ext::intel::esimd::") ||
-              Name.starts_with("sycl::_V1::ext::intel::experimental::esimd::"))
+              Name.starts_with(
+                  "sycl::_V1::ext::intel::experimental::esimd::") ||
+              Name.starts_with("sycl::_V1::ext::oneapi::this_work_item::"))
             continue;
 
           // Check if function name matches any allowed SYCL function name.

sycl/doc/extensions/experimental/sycl_ext_oneapi_free_function_queries.asciidoc

@@ -45,14 +45,7 @@ SYCL specification refer to that revision.
 
 == Status
 
-This is a proposed update to an existing extension.  Interfaces defined in this
-specification may not be implemented yet or may be in a preliminary state.  The
-specification itself may also change in incompatible ways before it is
-finalized.  *Shipping software products should not rely on APIs defined in this
-specification.*  See
-link:../experimental/sycl_ext_oneapi_free_function_queries.asciidoc[here] for
-the existing extension, which is implemented.
-
+This extension is implemented and fully supported by {dpcpp}.
 
 == Overview
 

sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_queries.asciidoc renamed to sycl/doc/extensions/supported/sycl_ext_oneapi_free_function_queries.asciidoc

@@ -43,7 +43,7 @@ Specifically, this library depends on the following SYCL extensions:
 * [sycl_ext_oneapi_complex](
     ../extensions/experimental/sycl_ext_oneapi_complex.asciidoc)
 * [sycl_ext_oneapi_free_function_queries](
-    ../extensions/experimental/sycl_ext_oneapi_free_function_queries.asciidoc)
+    ../extensions/supported/sycl_ext_oneapi_free_function_queries.asciidoc)
 * [sycl_ext_oneapi_assert](
     ../extensions/supported/sycl_ext_oneapi_assert.asciidoc)
 * [sycl_ext_oneapi_enqueue_barrier](

sycl/doc/syclcompat/README.md

@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include <sycl/ext/oneapi/free_function_queries.hpp>
 #include <sycl/ext/oneapi/properties/properties.hpp>
 #include <sycl/group.hpp>
 #include <sycl/memory_enums.hpp>
@@ -91,7 +92,8 @@ template <int Dimensions> sycl::sub_group get_child_group(group<Dimensions> g) {
 
 namespace this_kernel {
 template <int Dimensions> root_group<Dimensions> get_root_group() {
-  return this_nd_item<Dimensions>().ext_oneapi_get_root_group();
+  return sycl::ext::oneapi::this_work_item::get_nd_item<Dimensions>()
+      .ext_oneapi_get_root_group();
 }
 } // namespace this_kernel
 

sycl/include/sycl/ext/oneapi/experimental/root_group.hpp

@@ -0,0 +1,90 @@
+//==---- free_function_queries.hpp -- SYCL_INTEL_free_function_queries ext -==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/group.hpp>
+#include <sycl/nd_item.hpp>
+#include <sycl/sub_group.hpp>
+
+// For deprecated queries:
+#include <sycl/id.hpp>
+#include <sycl/item.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+namespace ext::oneapi::this_work_item {
+template <int Dimensions> nd_item<Dimensions> get_nd_item() {
+#ifdef __SYCL_DEVICE_ONLY__
+  return sycl::detail::Builder::getElement(
+      sycl::detail::declptr<nd_item<Dimensions>>());
+#else
+  throw sycl::exception(
+      sycl::make_error_code(sycl::errc::feature_not_supported),
+      "Free function calls are not supported on host");
+#endif
+}
+
+template <int Dimensions> group<Dimensions> get_work_group() {
+  return get_nd_item<Dimensions>().get_group();
+}
+
+inline sycl::sub_group get_sub_group() {
+#ifdef __SYCL_DEVICE_ONLY__
+  return sycl::sub_group();
+#else
+  throw sycl::exception(
+      sycl::make_error_code(sycl::errc::feature_not_supported),
+      "Free function calls are not supported on host");
+#endif
+}
+} // namespace ext::oneapi::this_work_item
+
+namespace ext::oneapi::experimental {
+template <int Dims>
+__SYCL_DEPRECATED(
+    "use sycl::ext::oneapi::this_work_item::get_nd_item() instead")
+nd_item<Dims> this_nd_item() {
+  return ext::oneapi::this_work_item::get_nd_item<Dims>();
+}
+
+template <int Dims>
+__SYCL_DEPRECATED(
+    "use sycl::ext::oneapi::this_work_item::get_work_group() instead")
+group<Dims> this_group() {
+  return ext::oneapi::this_work_item::get_work_group<Dims>();
+}
+
+__SYCL_DEPRECATED(
+    "use sycl::ext::oneapi::this_work_item::get_sub_group() instead")
+inline sycl::sub_group this_sub_group() {
+  return ext::oneapi::this_work_item::get_sub_group();
+}
+
+template <int Dims>
+__SYCL_DEPRECATED("use nd_range kernel and "
+                  "sycl::ext::oneapi::this_work_item::get_nd_item() instead")
+item<Dims> this_item() {
+#ifdef __SYCL_DEVICE_ONLY__
+  return sycl::detail::Builder::getElement(sycl::detail::declptr<item<Dims>>());
+#else
+  throw sycl::exception(
+      sycl::make_error_code(sycl::errc::feature_not_supported),
+      "Free function calls are not supported on host");
+#endif
+}
+
+template <int Dims>
+__SYCL_DEPRECATED("use nd_range kernel and "
+                  "sycl::ext::oneapi::this_work_item::get_nd_item() instead")
+id<Dims> this_id() {
+  return this_item<Dims>().get_id();
+}
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl